【发布时间】:2013-12-01 03:17:27
【问题描述】:
我一直致力于创建基于 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 函数使空间成为环形。我该如何修复此代码以使其按预期工作。为什么全局内存不会随着任务的每次更改而更新?不应该是共享内存吗?
【问题讨论】:
-
答案很简单。在单个内核调用中从不同线程读取和写入同一内存位置是未定义的。让它们发挥作用的唯一方法是设置障碍,即使这些障碍也只能在一个工作组中发挥作用。
-
那么这不会在与工作组接壤的单元格中失败吗?同样,if 中的障碍将如何发挥作用?因为不是所有的任务都会遇到它。
-
@HunterLarco 是的,确实它会在与工作组接壤的单元格中失败,但前提是您设置了适当的障碍。因为你没有,它在任何单元格中都会失败。你不能在 if 中设置障碍,因为所有工作项都会遇到障碍 HAS。
标签: java c++ parallel-processing opencl jocl