基本数组复制 OpenCL GPU答案

【问题标题】：Basic arrays copy OpenCL GPU基本数组复制 OpenCL GPU
【发布时间】：2022-06-25 03:52:16
【问题描述】：

有人可以帮我弄清楚如何将 CPU 的 C 代码转换为 GPU 的内核代码

int a[N], b[N];
b[0] = a[0];
b[N] = a[N];

for (i=1; i<N-1; i++) 
    b[i]= a[i-1] + a[i] + a[i+1];

我想过这样写，但我想找到一个性能更好的解决方案

__kernel void adjacentCopy(__global double *a, __global double *b, const unsigned int n) {

    int gid = get_global_id(0);

    if (gid < N) 
        b[gid] = a[gid-1]+a[gid]+a[gid+1];
    
}
    // and than execute the two base case into the host

任何人都可以提出一种方法来组织代码以使用本地内存，并将两种极端情况带回内核，而不会增加分支分歧

关于输入代码，看起来数组很小（因为堆栈大小有限）。请注意，这些值未初始化，因此存在未定义的行为，b[0] = b[0]; a[N] = a[N]; 完全没用。事实上，a[N] = a[N]; 导致了另一个未定义的行为......计算太便宜以至于 GPU 无法真正有用，数据传输的延迟和数据传输的速度肯定会导致 GPU 上的计算速度变慢。简而言之：您的输入代码是伪造的，在这里使用 GPU 毫无用处。
嗨，谢谢你的回复，这段代码只是理解如何在 gpu 上高效工作的概念
*我更正了两个基本情况

标签： arrays c optimization gpu opencl

【解决方案1】：

kernel 本质上是一个for 循环，其中每次迭代并行运行。确切的执行顺序是随机的，因此从一次迭代到下一次迭代不能有任何数据依赖关系；否则你必须使用双缓冲区（只从一个缓冲区读取，只写入另一个缓冲区）。

在您的情况下，内核将读取：

__kernel void adjacentCopy(const __global double *a, __global double *b, const unsigned int N) {
    int gid = get_global_id(0);
    if(gid==0||gid==N-1) return; // guard clause: do not execute the first and last element
    b[gid] = a[gid-1]+a[gid]+a[gid+1]; // double buffers to resolve data dependencies: only read from a and only write to b
}

对于极端情况gid==0||gid==N-1，在这样的计算网格上，您通常使用周期性边界条件。然后内核将变得无分支，如下所示：

__kernel void adjacentCopy(const __global double *a, __global double *b, const unsigned int N) {
    int gid = get_global_id(0);
    b[gid] = a[(gid+N-1)%N]+a[gid]+a[(gid+1)%N]; // periodic boundaries with modulo; in "(gid+N-1)" the "+N" ensures that the argument of the modulo operator always is positive
}

现在对于local 内存优化：没有它，对于每个线程，您从慢速global 内存中读取a 的3 个相邻值。理论上，每个线程只能从global 内存加载一个元素，并使用快速local 内存在工作组内共享数据。但是gid==0||gid==N-1 的两个线程必须从global 内存中加载2 个值，从而引入分支，这可能会扼杀任何潜在的性能提升。在这种情况下，增加的复杂性以及没有显着的性能提升使得local 内存优化成为不利的选择。这就是内核的样子：

#define def_workgroup_size 128 // set this to the size of the workgroup
__kernel void adjacentCopy(const __global double *a, __global double *b, const unsigned int N) {
    int gid = get_global_id(0);
    int lid = get_local_id(0);
    __local double cached_a[def_workgroup_size+2]; // as large as the workgroup, plus neighbors on the left and right sides of the workgroup
    if(lid==0) cached_a[lid] = a[(gid+N-1)%N]; // first thread in workgroup also has to load left neighbor
    cached_a[lid+1] = a[gid];
    if(lid==def_workgroup_size-1) cached_a[lid+1] = a[(gid+1)%N]; // last thread in workgroup also has to load right neighbor
    barrier(CLK_LOCAL_MEM_FENCE); // barrier to make sure cached_a is entirely filled up
    b[gid] = cached_a[lid]+cached_a[lid+1]+cached_a[lid+2]; // read 3 values from local memory
}

【讨论】：