在 OpenCL 中写入全局内存

Posted 2023-02-16

【中文标题】在 OpenCL 中写入全局内存

【英文标题】：Writing into global memory in OpenCL

【发布时间】：2014-09-26 18:09:51

【问题描述】：

我正在尝试优化一些最初用 Fortran 编写的代码。

该算法涉及在多个迭代中对大型数组（约 2700 万个单元）进行操作。每个单元格可以在一次迭代期间独立评估。但是，迭代不能并行化，因为在 t+1 完成的计算取决于在 t 完成的计算结果。

一个粗略的、简化的非并行伪代码示例：

for (t=0; t<tmax; t++)

A = A + B;

B  = B + A /2;

其中 A 和 B 是大数组。

目前，我通过在主机 C++ 代码的循环中调用 EnqueueNDRangeKernel 来实现这一点。因为需要上一次迭代的结果，所以每次都写入全局内存。

每次迭代执行 2700 万次全局内存写入会影响我的性能。我有两个正在使用的内核版本；与 Fortran 相比，版本 1 快了约 2.5 倍；版本 2 快 4 倍。

我试过摆弄算法，以及用指针做一些事情（版本 2）。

我的问题如下：有什么办法可以避免这个全局内存写入阻塞点？

谢谢！

---

请求的代码：

在 C++ 中调用：

NDRange global(nxp1*ny*nz);
NDRange local(nz);

    for (w=0; w<100; w++)
    
       queue.enqueueNDRangeKernel(kernA, NullRange, global, local);
        queue.enqueueBarrierWithWaitList();
       queue.enqueueNDRangeKernel(kernB, NullRange, global, local);
    

queue.finish();

内核：

__kernel void kernA(__global double *A, __global double *B)
    
    int i = get_global_id(0);

    double A_l;
    A_l = A[i];
    double B_l;
    B_l = B[i];

    A_l = A_l + B_l;

    A[i] = A_l; //if this line is removed, everything goes much faster.

    

   __kernel void kernB(__global double *A, __global double *B)
    
    int i = get_global_id(0);

    double A_l;
    A_l = A[i];
    double B_l;
    B_l = B[i];

    B_l = B_l + A_l/2;

    B[i] = B_l; //if this line is removed, everything goes much faster.
    

为了阐明算法，我已经简化了内核代码。但我的想法是我根据 B 更新 A；然后我根据 A 更新 B。这会重复几次迭代。

【问题讨论】：

为了计算新 A 或 B 的 (i, j) 条目，除了旧 A 和 B 的 (i, j) 条目之外，您还需要引用任何其他内容吗？如果没有，那么除了第一次迭代的输入和最后一次迭代的输出之外，您不需要将全局内存用于任何其他内容，并且您可以在单个内核调用中计算多次迭代。

我只需要知道上一次迭代中的旧 A(i) 和 B(i)。您是否建议在内核中运行迭代循环？这会保留迭代的顺序吗？

细胞由什么组成？一个单一的价值？ int，float，double，还是别的什么？您可能可以在同一个内核调用中计算不止 1 次迭代。

你能发布你的内核和你的 C/C++ 代码吗？您是否以任何方式使用本地内存？向量？

A、B 是双精度数，包含约 27 M 个细胞。稍后将发布代码的相关部分。

【参考方案1】：

没有办法完全避免全局写入问题。你只写一次值，你的速度是受硬件限制的。不过，您可以减少全局读取的数量，只要您不介意一次计算多个步骤。这仍然可以节省沿途的每一步。

__kernel void myKernel(__global double *A, __global double *B, __global uint outDataMultiple)              
                                                                                      
    const uint gid = get_global_id(0);
    const uint inDataSize = get_global_size(0);

    double2 nextValue;
    nextValue.x = A[gid];
    nextValue.y = B[gid];
    for(uint i=0; i<outDataMultiple; i++)
        nextValue.x = nextValue.x + nextValue.y;
        nextValue.y = nextValue.y + nextValue.x /2;
        A[gid+i+1] = nextValue.x;
        B[gid+i+1] = nextValue.y;
    

使用上面的内核，一个工作项将处理单个单元的多次迭代。你需要分配 outDataMultiple 倍的内存，剩下的会被内核填满。全局工作项计数决定了初始输入的大小。 outDataMultiple 仅受全局内存分配的限制，并且可能会受到每次迭代的数学复杂性的限制。

所需的全局内存总量： 27M * sizeof(double2) * (1+outDataMultiple)

__kernel void myKernel(__global double2 *data, __global uint outDataMultiple)              
                                                                                      
    const uint gid = get_global_id(0);
    const uint inDataSize = get_global_size(0);

    double2 nextValue = data[gid];
    for(uint i=0; i<outDataMultiple; i++)
        nextValue.x = nextValue.x + nextValue.y;
        nextValue.y = nextValue.y + nextValue.x /2;
        data[gid+i+1] = nextValue;
    

只要您可以交错 A 和 B 向量，相同内核的 double2 版本可能是可能的。这将结合读取和写入以保证 8 字节块，并可能进一步提高性能。

【讨论】：

严格从概念上讲，步骤 1. 从全局内存中加载内容 2. 开始循环 3. 在私有上运行 A = f(B), B = f(A) 4. 结束循环 5. 复制回全局？我认为这些失败是因为寄存器中没有足够的可用内存，并且因为 opencl 产生了一堆线程，所以循环将在每个单元上独立运行（这很糟糕，因为 a[i] 取决于例如 b[i]但是 b[i+1] 也是）。

27M 线程不会同时启动。驱动程序/硬件将为您管理所有这些。您仍应使用 enqueueNDrange 指定工作组大小。在您的示例代码中计算 A[i+1] 时，我看不到哪里需要 B[i+1] 。我列出的代码将重现原始循环，但在不同的内存位置，因此每次 outDataMultiple 写入只需一次全局读取即可。

【参考方案2】：

减少 OpenCL 设备从全局内存中获取时间的简单方法是将全局内存批量缓冲到本地内存，对本地内存进行操作，然后将本地内存批量写入全局内存。

本地内存与线程内存具有基本相同的延迟，并且可以从全局内存中以块的形式读取。本地内存可以在主机上声明并传递给内核（参见下面的示例）或在内核中分配并使用（参见下面列出的 AMD 优化指南中的示例）。例如：

__kernel void kernA(__global double *A, 
                    __global double *B, 
                    __local double *BufferA,
                    __local double *BufferB)
    
    BufferA[get_local_id(0)] = A[get_global_id(0)];
    BufferB[get_local_id(0)] = B[get_global_id(0)];
    mem_fence(CLK_LOCAL_MEM_FENCE);

    double tmp = BufferA[get_local_id(0)] + BufferB[get_local_id(0)];

    A[get_global_id(0)] = BufferA[get_local_id(0)];
    mem_fence(CLK_GLOBAL_MEM_FENCE);
    

还有更多可以做的事情，包括：

内存平铺 - 使用 2d 全局 ID 访问内存 优化设备的本地内存大小 尽可能使用浮点数 尽可能多地填入内核 Read this guide by AMD for optmizing OpenCL kernels