优化内存访问OpenCL [英] Optimising Memory Access OpenCL

问题描述

我目前正在使用挡路矩阵乘法在OpenCL内核中乘以字节矩阵：我将矩阵细分为瓦片(32x32)，将这些瓦片加载到本地内存中，然后将其写回全局内存。

目前，内存访问是瓶颈。我正在尝试对其进行优化。

假设我要乘以C=A x B，其中A、B、C是字符*

A(NDIM，PDIM)、B(PDIM，MDim)、C(Ndim，MDim)。

我当前将A设置为行主格式，B设置为列主格式，以确保每个矩阵的内存访问在工作组内是按顺序进行的。

每个工作项将单个字节加载到本地内存中，并负责处理该字节。我的内核的维数是全局工作项的{ndim，mdim}和本地工作项的{挡路_SIZE，挡路_SIZE}。

代码与http://www.nvidia.com/content/cudazone/download/OpenCL/NVIDIA_OpenCL_ProgrammingGuide.pdf几乎相同(只是A以列主格式存储)

我的问题：如何优化内存访问？我听说过很多关于合并的事情，但是我很难理解合并和并行之间的权衡。

选项0：保持原样，即使每个线程访问一个字节，它也会合并，以便工作组内的每个线程都获得已经访问的数据。 ->不太可能，因为我的访问不是字节对齐的。我怀疑我每次都要装货 4字节+x，其中x是线程的偏移量。

选项1：使用整数矩阵减少并行度 如果我将矩阵作为整数，我将能够一次加载更多，但会显著降低并行度(降低4倍)，其中每个字节乘法都必须按顺序执行。

选项2：使用整数矩阵，但保持相同的并行度 这基本上意味着内存中的数据将由每个 直观地说，这对应于int foo=get_global_id(0)，然后假设 我将foo转换为具有字节x=foo[get_local_id(0))的char[]foo_bytes； 我的理解是，第一个线程将使用get_global_id(0)将数据加载到内存中，而工作组中的其余线程将看到它已经加载

选项3：使用整数矩阵，减少并行度，但使用向量类型 在工作项中处理数据 我知道OpenCL支持向量类型，如果我加载一个32位整数，我可以将 将其转换为向量类型，这样工作项将并行处理这4个字节。 我的理解是，这只是语法问题，在OpenCL中使用这样的向量类型不会带来任何性能提升。

据我所知，选项2更可取。这是对的吗？如果没有，为什么？

推荐答案

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

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

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

选项3-这似乎是选项1的自然扩展。如果您打算尝试选项1，将值映射到向量是下一个要测试的合乎逻辑的事情。不过，可能不会对每种架构都有性能提升--GPU喜欢浮点数、双精度数和整数，而不一定是字节。

更多想法/评论：

我认为对您的全局访问性能进行的最大优化是您已经实现的以列为主的排序。

您有没有想过使用Halfn和Halfn类型？对于支持一半的设备，您应该能够获得比浮点数/浮点数高出一倍的数据密度。这不如打包为int或char4的4字节，但是任何支持Half类型的设备都可能支持点(Halfn，Halfn)，这可能会让您一次计算4、8或16个MAD。

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

如果您喜欢使用方形块：3*64x64字节=12KB，3*96x96=27KB

如果您希望处理输出矩阵‘C’的32x32：

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工作项并循环访问数据。

这篇关于优化内存访问OpenCL的文章就介绍到这了，希望我们推荐的答案对大家有所帮助，也希望大家多多支持IT屋！