我目前在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是可取的。它是否正确?如果不是,为什么?
当每个工作项目读取相邻的32位数量时,会发生最大合并。所以我会让每个工作项目读取4个字节(最小值),这可以使用char4向量来完成。 – Dithermaster 2014-11-09 23:28:50
列主要命令拧紧内存合并 - 最重要的考虑!见下面的答案。 – wcochran 2017-02-15 18:44:14