我一直在努力创建一个基于 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
}
}
}
- 数组输出是用于渲染内核计算的图像缓冲区数据。
- sizeX和sizeY常量分别是图像缓冲区的宽度和高度。
- colorMap数组分别包含黑色和白色的 rgb 整数值，用于正确更改图像缓冲区的值以呈现颜色。
- newCellMap数组是确定渲染后计算的更新单元格图。
- historyBuffer是内核调用开始时单元的旧状态。每次执行内核时，这个数组都会更新为 newCellMap 数组。
此外，wrap功能使空间呈环形。我该如何修复此代码以使其按预期工作。为什么全局内存不会随着任务的每次更改而更新？不应该是共享内存吗？