我一直在努力创建一个基于 GPU 的 conway 游戏生活程序。如果你不熟悉它,这里是维基百科页面。我创建了一个版本,该版本通过保留一个值数组来工作,其中 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 数组。
此外,包装功能使空间呈环形。我如何修复此代码,使其按预期工作。为什么全局内存不会随着任务的每次更改而更新?它不应该是共享内存吗?
正如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上创建了一个"Live"游戏,我发现更简单,更快捷的方法是创建一个大的2D位数组(位寻址)。然后编写一段没有任何分支的代码,这些分支只是分析neibours并获取该单元格的更新值。由于使用了位寻址,因此每个线程的读/写内存量远低于寻址字符/整数/其他。我在非常旧的OpenCL HW(nVIDIA 9100M G)中实现了33Mcells/sec。只是为了让你知道你的if/else方法可能不是最有效的方法。
作为参考,我在这里让你实现我的生活游戏(OpenCL 内核):
//Each work-item processess one 4x2 block of cells, but needs to access to the (3x3)x(4x2) block of cells surrounding it
// . . . . . .
// . * * * * .
// . * * * * .
// . . . . . .
__kernel void life (__global unsigned char * input, __global unsigned char * output){
int x_length = get_global_size(0);
int x_id = get_global_id(0);
int y_length = get_global_size(1);
int y_id = get_global_id(1);
//int lx_length = get_local_size(0);
//int ly_length = get_local_size(1);
int x_n = (x_length+x_id-1)%x_length; //Negative X
int x_p = (x_length+x_id+1)%x_length; //Positive X
int y_n = (y_length+y_id-1)%y_length; //Negative Y
int y_p = (y_length+y_id+1)%y_length; //Positive X
//Get the data of the surrounding blocks (TODO: Make this shared across the local group)
unsigned char block[3][3];
block[0][0] = input[x_n + y_n*x_length];
block[1][0] = input[x_id + y_n*x_length];
block[2][0] = input[x_p + y_n*x_length];
block[0][1] = input[x_n + y_id*x_length];
block[1][1] = input[x_id + y_id*x_length];
block[2][1] = input[x_p + y_id*x_length];
block[0][2] = input[x_n + y_p*x_length];
block[1][2] = input[x_id + y_p*x_length];
block[2][2] = input[x_p + y_p*x_length];
//Expand the block to points (bool array)
bool point[6][4];
point[0][0] = (bool)(block[0][0] & 1);
point[1][0] = (bool)(block[1][0] & 8);
point[2][0] = (bool)(block[1][0] & 4);
point[3][0] = (bool)(block[1][0] & 2);
point[4][0] = (bool)(block[1][0] & 1);
point[5][0] = (bool)(block[2][0] & 8);
point[0][1] = (bool)(block[0][1] & 16);
point[1][1] = (bool)(block[1][1] & 128);
point[2][1] = (bool)(block[1][1] & 64);
point[3][1] = (bool)(block[1][1] & 32);
point[4][1] = (bool)(block[1][1] & 16);
point[5][1] = (bool)(block[2][1] & 128);
point[0][2] = (bool)(block[0][1] & 1);
point[1][2] = (bool)(block[1][1] & 8);
point[2][2] = (bool)(block[1][1] & 4);
point[3][2] = (bool)(block[1][1] & 2);
point[4][2] = (bool)(block[1][1] & 1);
point[5][2] = (bool)(block[2][1] & 8);
point[0][3] = (bool)(block[0][2] & 16);
point[1][3] = (bool)(block[1][2] & 128);
point[2][3] = (bool)(block[1][2] & 64);
point[3][3] = (bool)(block[1][2] & 32);
point[4][3] = (bool)(block[1][2] & 16);
point[5][3] = (bool)(block[2][2] & 128);
//Process one point of the game of life!
unsigned char out = (unsigned char)0;
for(int j=0; j<2; j++){
for(int i=0; i<4; i++){
char num = point[i][j] + point[i+1][j] + point[i+2][j] + point[i][j+1] + point[i+2][j+1] + point[i][j+2] + point[i+1][j+2] + point[i+2][j+2];
if(num == 3 || num == 2 && point[i+1][j+1] ){
out |= (128>>(i+4*j));
}
}
}
output[x_id + y_id*x_length] = out; //Assign to the output the new cells value
};
在这里,您不保存任何中间状态,只保存最后的单元格状态(活/死)。它没有分支,因此在此过程中非常快。