OpenCL 本地内存声明答案

【问题标题】：OpenCL Local Memory DeclarationOpenCL 本地内存声明
【发布时间】：2015-02-11 09:28:48
【问题描述】：

如下声明本地内存有什么区别：

__kernel void mmul(const int Ndim, const int Mdim, const int Pdim,
                      const __global int* A,
                      const __global int* B,
                     __global char* C,
                     __local int* restrict block_a,
                     __local int* restrict block_b)

并在内核中声明本地内存

#define a_size 1024
#define b_size 1024 * 1024
__kernel void mmul(const int Ndim, const int Mdim, const int Pdim,
                      const __global int* A,
                      const __global int* B,
                     __global char* C) {

__local int block_a[a_size]
__local int block_b[b_size]

... 
}

在这两种情况下，所有线程都将更新共享 A 和 B 数组中的单个单元格

我知道内核中不可能有“可变”长度数组（因此第二个内核顶部的#define），但还有其他区别吗？释放内存的时间有什么区别吗？

【问题讨论】：

在内核中定义本地内存意味着它将为每个工作组分配。将本地内存作为参数意味着，在分配内存之前的任何地方，对于每个工作组来说，它都是相同的指针。不要考虑何时以及是否释放内存。 OpenCL 就像 1. 编译 2. 调用具有多少组和工作人员的内核 3. 现在隐式知道将使用多少内存。
优化提示：“如果每个线程都只是读取和写入本地内存的一个位置。为什么还需要本地内存？”

标签： opencl gpu

【解决方案1】：

在这两种情况下，本地内存都存在于工作组的生命周期中。正如您所指出的，唯一的区别是，将本地内存指针作为参数传递允许动态指定缓冲区的大小，而不是作为编译时常量。不同的工作组将始终使用不同的本地内存分配。

【讨论】：

Nit 采摘，但如果您使用 CL_KERNEL_LOCAL_MEM_SIZE 标志调用 clGetKernelWorkGroupInfo 也会有所不同。在第一种情况下，大小将被假定为 0。在第二种情况下，它将返回正确的大小...

【解决方案2】：

如果要将代码移植到CUDA，第二种方法更好，因为CUDA中的__shared__内存（相当于OpenCL中的__local）不支持像第一种情况那样声明。

【讨论】：