如何加快 OpenCL 内核中的随机写入速度

Question

我有这个 OpenCL 内核用于计算 Buddhabrot-Set：

// ... initializing things
while (true)
    {
        Z2 = Z * Z;
        Z = (float2)(Z2.x - Z2.y + Point.x, 2 * Z.x * Z.y + Point.y);

        X = ((Z.x - MinView.x) * iScale.x);
        Y = ((Z.y - MinView.y) * iScale.y);

        if (X >= 0 && X < Res.x && Y >= 0 && Y < Res.y)
        {
            Result[Offset + Y * Res.x + X]++; // <--- Commenting out this line gives a 10x-speedup
            count++;
        }

        if (dot(Z, Z) > 4)
        {
            PointsOutPos = atomic_add(&atomicCounter[2], 1);
            if (PointsOutPos >= PointsOutSize)
                return;
            Point = convert_float2(PointsOut[PointsOutPos]); // Load next point to minimize thread divergence
            Z = 0;
            i = 0;
        }
        i++;
    }

注释掉

Result[Offset + Y * Res.x + X]++

行可以使速度提高大约 10 倍（当然没有结果）。所以随机写入似乎是一个巨大的瓶颈。此写入不仅会发生在输出数组中的随机位置，而且还会发生在不同的时间/迭代（尽管在大多数迭代中）。

有人知道如何加快速度吗？使用本地数组作为缓冲区是行不通的，因为它不会改变写入——只是改变写入完成的时间（如果一次全部完成，可能会导致内存带宽饱和）。还是我的想法错误，其他线程可以在写入内存时进行计算？

Answer 1

您可以使用基于图块的方法，通过使用本地内存作为缓冲区来最小化全局内存写入：

__kernel void buddhabrot(__global float2* PointsOut, __global int* Result, __global int* atomicCounter, 
                         int PointsOutSize, float2 MinView, float2 MaxView, int2 Res) {
    int id = get_global_id(0);
    if (id >= PointsOutSize) return;

    float2 Point = PointsOut[id];
    float2 Z = (0, 0);
    int2 iScale = (Res.x / (MaxView.x - MinView.x), Res.y / (MaxView.y - MinView.y));
    
    __local int localBuffer[LOCAL_SIZE]; // Define as per your tile size
    for (int i = 0; i < LOCAL_SIZE; i++) localBuffer[i] = 0;
    barrier(CLK_LOCAL_MEM_FENCE);

    for (int i = 0; i < MAX_ITERATIONS; i++) {
        float2 Z2 = Z * Z;
        Z = (float2)(Z2.x - Z2.y + Point.x, 2 * Z.x * Z.y + Point.y);
        int2 pos = (int2)((Z.x - MinView.x) * iScale.x, (Z.y - MinView.y) * iScale.y);

        if (all(pos >= 0) && all(pos < Res)) {
            int localIdx = (pos.y % TILE_SIZE) * TILE_SIZE + (pos.x % TILE_SIZE);
            atomic_add(&localBuffer[localIdx], 1);
        }

        if (dot(Z, Z) > 4) break;
    }

    barrier(CLK_LOCAL_MEM_FENCE);
    // Write back to global memory
    for (int i = 0; i < LOCAL_SIZE; i++) {
        if (localBuffer[i] > 0) {
            int globalIdx = calculateGlobalIndex(i, ...); // Implement this
            atomic_add(&Result[globalIdx], localBuffer[i]);
        }
    }
}