2014-11-09 101 views
0

我目前在openCL内核中使用块矩阵乘法算法乘以字节矩阵:我将矩阵细分为块(32 x 32),将这些块加载到本地内存中,并将其写回到全球记忆。优化内存访问OpenCL

目前,内存访问是瓶颈。我试图看看我可以对其进行优化。

比方说,我乘C =,其中A,B,C是字符一个X B *

A(NDIM,PDIM),B(PDIM,MDim),C(NDIM,MDim)。

我目前有A行的主要格式和B列的主要格式,以确保内存访问在每个矩阵的工作组内连续。

每个工作项都将一个字节加载到本地内存中,并负责处理该字节。我的内核的dimensiosn是用于全局工作项目的{Ndim,Mdim},以及用于本地工作项目的{block_size,block_size}。

代码几乎是相同的http://www.nvidia.com/content/cudazone/download/OpenCL/NVIDIA_OpenCL_ProgrammingGuide.pdf(前提是A存储在列主要格式除外)

我的问题:如何优化存储器存取?我听到很多关于合并的内容,但我很难理解合并和并行之间的折衷。

选项0:保持原样,即使每个线程访问一个字节,它都会合并,因此工作组中的每个线程都会获取已经访问过的数据。 - >不太可能,因为我的访问不是字节对齐的。我怀疑我每次都会加载 4字节+ x,其中x是线程的偏移量。

选项1:使用整数矩阵减少并行 如果我是有矩阵作为整体,我将能够加载在时间多得多,但将显著降低平行度(通过因子4) ,每个字节乘法都必须按顺序执行。

选项2:使用整数矩阵但保持平行度相同 这基本上意味着,在存储器中的数据将通过每个 直观地被加载多次,这相当于到int富= get_global_id(0),然后,假设 我将foo转换为字节x = foo [get_local_id(0))的char [] foo_bytes; 我的理解是,第一个线程将使用get_global_id(0)加载数据到内存中,而工作组中剩余的线程会看到它已经加载

选项3:使用整数矩阵,减少并行,但在工作项目中使用矢量类型 来处理数据 我知道opencl支持矢量类型,如果我加载一个32位整数,我可以将 转换为矢量类型,这样工作项将处理4个字节平行。 我的理解是,这只是句法,我不会从使用OpenCL中的矢量类型获得任何性能改进。

从我的理解,选项2是可取的。它是否正确?如果不是,为什么?

回答

1

选项0 - 如果它保持代码简单并且您的当前性能足够好,那么这并不是那么糟糕。

选项1 - 我认为这值得一试。您希望将4个字节作为单个int加载,并使用单个线程处理它。 ALU饱和度正是您的调度程序需要隐藏您正在经历的全局内存延迟的原因。我认为这对选项#2来说是非常接近的第二位。

选项2 - 可能是您提到的最好的一个,因为它会利用许多现代设备上提供的内存广播优势。每个int值将被每4个线程读取一次。我认为在每4个线程处理超过1个int时(也许每4个线程4个,总共16个字节),性能测试是值得的。

选项3 - 这似乎是选项#1的自然延伸。如果你打算给选项1一个镜头,将值映射到矢量是测试的下一个合理的事情。尽管每个体系结构可能都没有性能提升 - GPUs喜欢漂浮,双打和整数,而不一定是字节。

更多的想法/评论:

我认为您的全球访问性能的最大优化是你已经实现了列主顺序。

你有没有使用半和半?对于支持一半的设备,您应该能够通过float/floatn获得两倍的数据密度。这不如4个字节打包为int或char4,但任何支持half类型的设备都可能支持dot(halfn,halfn),这可以让您一次计算4个,8个或16个MAD。

选项4 -我强烈建议将更大的块读入本地内存。当从本地存储器中乘以32x32矩阵时,每个元素被读取32次,但只从全局存储器读取一次。当您对64x64块执行相同操作时,每个元素从本地内存读取64次。 OpenCL设备具有32KB的共享内存,并且当您有三个32x32字节的矩阵时,您只能使用3KB。

如果你喜欢使用正方形块:3个* 64×64字节= 12KB,3 * 96×96 = 27KB

如果您希望在输出矩阵 'C' 的32×32的工作:

blockDim = ((32768 - 32*32) /2)/32 = 496 
1) read 496x32 block from A, store locally 
2) read 496x32 block from B, store locally 
3) read or initialize 32x32 block of C in local memory 
4) do the math 
5) write the 32x32 block to global memory C 

496比大多数工作组维度所允许的大,但我个人更喜欢使用32x1工作项目并循环访问数据。

+0

当每个工作项目读取相邻的32位数量时,会发生最大合并。所以我会让每个工作项目读取4个字节(最小值),这可以使用char4向量来完成。 – Dithermaster 2014-11-09 23:28:50

+0

列主要命令拧紧内存合并 - 最重要的考虑!见下面的答案。 – wcochran 2017-02-15 18:44:14

1

Memory coalescing is the single most important performance consideration用于编程nVidia GPU。如果线程i正在从存储位置n读取,则具有线程i + 1从位置读取n + 1。如果这些线程处于同一个warp中,那么这些读取将“合并”为一个事务。

请注意,在将每个子矩阵加载到共享内存的nVidia示例中,矩阵均为行主要订单。这意味着,该线程用于(行,列)将读取存储器单元行X步幅+山口和线程(行,列+ 1)将读取存储器单元行X步幅+ COL + 1这确实在记忆中彼此相邻。这很可能是since the threads are ordered in row-major order.

如果基质为列优先顺序这螺丝一切行动 - 如果线程在同一经线这将是coelesced! (row,col + 1)的线程将在内存中读取存储单元(col + 1)x stride + row这不是col x stride + row

因此,你对列主要命令的小改动打破了在nVidia GPU中优化的最重要的事情!