OpenCL

Question

我确实了解全局内存和本地内存之间的一般区别。但是我在使用本地内存时遇到问题。

1) 将全局内存变量转换为局部内存变量需要考虑什么？

2) 如何使用局部屏障？

也许有人可以帮我举个例子。

我尝试使用本地内存进行 jacobi 计算，但结果只得到 0。也许有人可以给我一个建议。

工作解决方案：

#define IDX(_M,_i,_j) (_M)[(_i) * N + (_j)]
#define U(_i, _j)     IDX(uL, _i, _j)

__kernel void jacobi(__global VALUE* u, __global VALUE* f, __global VALUE* tmp, VALUE factor) {

int i = get_global_id(0);
int j = get_global_id(1);

int iL = get_local_id(0);
int jL = get_local_id(1);

__local VALUE uL[(N+2)*(N+2)];
__local VALUE fL[(N+2)*(N+2)];

IDX(uL, iL, jL) = IDX(u, i, j);
IDX(fL, iL, jL) = IDX(f, i, j);

barrier(CLK_LOCAL_MEM_FENCE);

IDX(tmp, i, j) = (VALUE)0.25 * ( U(iL-1, jL) + U(iL, jL-1) + U(iL, jL+1) + U(iL+1, jL) - factor * IDX(fL, iL, jL));

}

谢谢。

Answer 1

1) 查询CL_DEVICE_LOCAL_MEM_SIZE值，最小16kB，不同硬件增加。如果你的局部变量可以适应这个并且如果它们被多次重复使用，你应该在使用前将它们放在本地内存中。即使您不这样做，在访问 gpu 的全局内存时自动使用 L2 缓存对于内核的利用仍然有效。

如果全局-本地复制占用了大量时间，您可以在核心计算的同时进行异步工作组复制。

另一个重要的部分是，更多可用的本地内存 space 意味着每个内核有更多的并发线程。如果 gpu 每个计算单元有 64 个核心，那么当使用所有本地内存时，只有 64 个线程可以运行。当它有更多space时，如果没有其他限制，128,192,...2560个线程可以同时运行。

探查器可以显示瓶颈，因此您可以考虑是否值得一试。

例如，使用嵌套循环的简单矩阵-矩阵乘法依赖于缓存 l1 l2，但子矩阵可以放入本地内存。也许 48x48 的浮点子矩阵可以装进一个中档显卡计算单元，并且可以在被下一个子矩阵替换之前用于整个计算 N 次。

CL_DEVICE_LOCAL_MEM_TYPE 查询可以 return LOCAL 或 GLOBAL 这也说如果是 GLOBAL 不建议使用本地内存。

最后，任何内存 space 分配（__private 除外）大小必须在编译时知道（对于设备，而不是主机），因为它必须知道可以发出多少波阵面以达到最大值性能（and/or 可能是其他编译器优化）。这就是 opencl 1.2 不允许递归函数的原因。但是你可以复制一个函数并重命名n次来实现伪递归。
2) 屏障是工作组中所有工作组线程的交汇点。类似于循环障碍，它们都停在那里，等待所有直到继续。如果它是本地屏障，则所有工作组线程在离开该点之前完成所有本地内存操作。如果你想将一些数字 1,2,3,4.. 赋给一个本地数组，你无法确定是否所有线程都写入了这些数字或已经写入，直到通过一个本地屏障，然后才能确定该数组将具有已写入的最终值。

所有工作组线程必须遇到相同的障碍。如果无法访问它，内核就会卡住或出现错误。

__local int localArray[64]; // not each thread. For all threads. 
                            // per compute unit.

if(localThreadId!=0)               
    localArray[localThreadId]=localThreadId; // 64 values written in O(1)
// not sure if 2nd thread done writing, just like last thread

if(localThreadId==0) // 1st core of each compute unit loads from VRAM
    localArray[localThreadId]=globalArray[globalThreadId];

barrier(CLK_LOCAL_MEM_FENCE); // probably all threads wait 1st thread
                              // (maybe even 1st SIMD or 
                              // could be even whole 1st wavefront!)
// here all threads written their own id to local array. safe to read.
// except first element which is a variable from global memory
// lets add that value to all other values
if(localThreadId!=0)
   localArrray[localThreadId]+=localArray[0];

工作示例（本地工作组大小=64）：

输入：0,1,2,3,4,0,0,0,0,0,0,..

    __kernel void vecAdd(__global float* x )
    {
       int id = get_global_id(0);
       int idL = get_local_id(0);
       __local float loc[64];
       loc[idL]=x[id];
       barrier (CLK_LOCAL_MEM_FENCE);
       float distance_square_sum=0;
       for(int i=0;i<64;i++)
       { 
            float diff=loc[idL]-loc[i];
            float diff_squared=diff*diff;
            distance_square_sum+=diff_squared;
       }       
       x[id]=distance_square_sum;

    }

输出：30、74、246、546、974、30、30、30...

OpenCL - 本地内存

OpenCL - Local Memory

c

memory

kernel