java - 任务间的 OpenCL 共享内存

Question

我一直致力于创建一个基于 GPU 的康威生命游戏程序。如果您不熟悉它，这里是 Wikipedia Page .我创建了一个版本，它通过保留一组值来工作，其中 0 代表死细胞，1 代表活细胞。内核然后简单地写入图像缓冲区数据数组以根据单元格数据绘制图像，然后检查每个单元格的邻居以更新单元格数组以供下一次执行渲染。

但是，一种更快的方法将细胞的值表示为负数(如果死亡)和正数(如果存活)。该单元格的编号表示它所拥有的邻居的数量加一(使零成为不可能的值，因为我们无法将 0 与 -0 区分开来)。然而，这意味着当产生或杀死一个细胞时，我们必须相应地更新它的八个邻居的值。因此，与只需要从相邻内存槽中读取的工作过程不同，该过程必须写入这些槽。这样做是不一致的，输出的数组无效。例如，单元格包含诸如 14 之类的数字，表示 13 个邻居，这是一个不可能的值。代码是正确的，因为我在 cpu 上编写了相同的程序，并且按预期工作。经过测试，我相信当任务尝试同时写入内存时，会有延迟导致某种写入错误。例如，在读取数组数据和设置数据更改时间之间可能存在延迟，从而导致另一个任务的过程不正确。我试过使用信号量和障碍，但刚刚学习了 OpenCL 和并行处理，还没有完全掌握它们。内核如下。

int wrap(int val, int limit){
    int response = val;
    if(response<0){response+=limit;}
    if(response>=limit){response-=limit;}
    return response;
}

__kernel void optimizedModel(
        __global uint *output,
        int sizeX, int sizeY,
        __global uint *colorMap,
        __global uint *newCellMap,
        __global uint *historyBuffer
)
{
    // the x and y coordinates that currently being computed
    unsigned int x = get_global_id(0);
    unsigned int y = get_global_id(1);

    int cellValue = historyBuffer[sizeX*y+x];
    int neighborCount = abs(cellValue)-1;
    output[y*sizeX+x] = colorMap[cellValue > 0 ? 1 : 0];

    if(cellValue > 0){// if alive
        if(neighborCount < 2 || neighborCount > 3){
            // kill

            for(int i=-1; i<2; i++){
                for(int j=-1; j<2; j++){
                    if(i!=0 || j!=0){
                        int wxc = wrap(x+i, sizeX);
                        int wyc = wrap(y+j, sizeY);
                        newCellMap[sizeX*wyc+wxc] -= newCellMap[sizeX*wyc+wxc] > 0 ? 1 : -1;
                    }
                }
            }
            newCellMap[sizeX*y+x] *= -1;

            // end kill
        }
    }else{
        if(neighborCount==3){
            // spawn

            for(int i=-1; i<2; i++){
                for(int j=-1; j<2; j++){
                    if(i!=0 || j!=0){
                        int wxc = wrap(x+i, sizeX);
                        int wyc = wrap(y+j, sizeY);
                        newCellMap[sizeX*wyc+wxc] += newCellMap[sizeX*wyc+wxc] > 0 ? 1 : -1;
                    }
                }
            }
            newCellMap[sizeX*y+x] *= -1;

            // end spawn
        }
    }
}

1. 数组output 是用于渲染的图像缓冲区数据 内核的计算。
2. sizeX 和sizeY 常量分别是图像缓冲区的宽度和高度。
3. colorMap 数组包含分别用于黑色和白色的 rgb 整数值，用于正确更改图像缓冲区的值以呈现颜色。
4. newCellMap 数组是在确定渲染后计算的更新单元格映射。
5. historyBuffer 是内核调用开始时细胞的旧状态。每次执行内核时，这个数组都会更新为 newCellMap 数组。

此外，wrap 功能使空间呈环形。我怎样才能修复此代码以使其按预期工作。为什么全局内存不会随着任务的每次更改而更新？不应该是共享内存吗？

Answer 1

正如 sharpneli 在他的回答中所说，您正在从不同的线程读取和写入相同的内存区域，这会产生未定义的行为。

解决方案: 您需要将 newCellMap 分成 2 个数组，一个用于上一次执行，另一个用于存储新值。然后，您需要在每次调用中从主机端更改内核参数，以便下一次迭代的 oldvalues 是上一次迭代的 newvalues。由于您构建算法的方式，您还需要在运行之前将 oldvalues 复制到 newvalues。

__kernel void optimizedModel(
        __global uint *output,
        int sizeX, int sizeY,
        __global uint *colorMap,
        __global uint *oldCellMap,
        __global uint *newCellMap,
        __global uint *historyBuffer
)
{
    // the x and y coordinates that currently being computed
    unsigned int x = get_global_id(0);
    unsigned int y = get_global_id(1);

    int cellValue = historyBuffer[sizeX*y+x];
    int neighborCount = abs(cellValue)-1;
    output[y*sizeX+x] = colorMap[cellValue > 0 ? 1 : 0];

    if(cellValue > 0){// if alive
        if(neighborCount < 2 || neighborCount > 3){
            // kill

            for(int i=-1; i<2; i++){
                for(int j=-1; j<2; j++){
                    if(i!=0 || j!=0){
                        int wxc = wrap(x+i, sizeX);
                        int wyc = wrap(y+j, sizeY);
                        newCellMap[sizeX*wyc+wxc] -= oldCellMap[sizeX*wyc+wxc] > 0 ? 1 : -1;
                    }
                }
            }
            newCellMap[sizeX*y+x] *= -1;

            // end kill
        }
    }else{
        if(neighborCount==3){
            // spawn

            for(int i=-1; i<2; i++){
                for(int j=-1; j<2; j++){
                    if(i!=0 || j!=0){
                        int wxc = wrap(x+i, sizeX);
                        int wyc = wrap(y+j, sizeY);
                        newCellMap[sizeX*wyc+wxc] += oldCellMap[sizeX*wyc+wxc] > 0 ? 1 : -1;
                    }
                }
            }
            newCellMap[sizeX*y+x] *= -1;

            // end spawn
        }
    }
}

关于您关于共享内存的问题有一个简单的答案。 OpenCL 没有跨主机设备的共享内存。

当您为设备创建内存缓冲区时，您首先必须使用 clEnqueueWriteBuffer() 初始化该内存区域，然后使用 clEnqueueWriteBuffer() 读取它以获得结果.即使您确实有指向内存区域的指针，您的指针也是指向该区域的主机端拷贝的指针。这很可能没有最新版本的设备计算输出。

PD:很久以前，我在 OpenCL 上创建了一个“实时”游戏，我发现更简单快捷的方法就是创建一个大的二维位数组(位寻址)。然后编写一段没有任何分支的代码，简单地分析邻居并获取该单元格的更新值。由于使用位寻址，每个线程读取/写入的内存量比寻址字符/整数/其他要低得多。我在一个非常老的 OpenCL HW (nVIDIA 9100M G) 上达到了 33Mcells/sec。只是想让您知道您的 if/else 方法可能不是最有效的方法。