OpenCL - 本地内存

Question

OpenCL - 本地内存

cmemorykernelopencl

3

我理解全局内存和局部内存的区别，但使用局部内存时遇到了问题。

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));

}

谢谢。

- SteveOhio

需要为本地数组分配空间，例如__local VALUE uL[128]。然后，如果需要，可以从中获取指针。大小必须在编译时知道（至少对于OpenCL 1.2）。 - huseyin tugrul buyukisik

好的，它可行了 - 非常感谢。最后一个问题：当本地内存版本的结果比全局内存版本的结果大得多时，问题可能出在哪里？两个版本都使用相同的10x10矩阵和100个Jacobi迭代。 - 我已经纠正了我的矩阵索引i，j到iL和jL。 - SteveOhio

如果数字类似于1231434252或其负数，可能会访问超出范围（并获取垃圾值，甚至其他变量？）。这个错误是现在还是追溯到第一个版本？你是指单载入版本吗？ - huseyin tugrul buyukisik

让我们在聊天中[继续这个讨论]（http://chat.stackoverflow.com/rooms/128060/discussion-between-steveohio-and-huseyin-tugrul-buyukisik）。 - SteveOhio

对我而言，本地内存只有在与工作组中的其他工作项协调以减少全局内存带宽时才有用。例如，当重复使用从全局内存读取的值以供其他工作项使用，或者为了能够使用来自全局内存的协同读取或写入。例如，在矩阵乘法期间。 - Dithermaster

显示剩余5条评论

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- huseyin tugrul buyukisik · Accepted Answer

1) 查询CL_DEVICE_LOCAL_MEM_SIZE值，它至少为16kB并且会因硬件不同而增加。如果您的局部变量可以适合其中，并且它们被多次重复使用，那么在使用之前应将它们放入本地内存中。即使您不这样做，在访问gpu的全局内存时自动使用L2缓存仍然可以有效利用核心。

如果全局-局部复制占用了重要的时间片段，则可以在核心计算事物时进行异步工作组复制。

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

分析器可以显示瓶颈，因此您可以考虑尝试或不尝试。

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

CL_DEVICE_LOCAL_MEM_TYPE查询可以返回LOCAL或GLOBAL，这也说明如果是全局，则不建议使用本地内存。

最后，任何内存空间分配（除__private之外）大小必须在编译时（对于设备而不是主机）知道，因为它必须知道可以发出多少波前以实现最大性能（和/或可能其他编译器优化）。这就是为什么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……