我一直在努力创造生活程序的基于GPU康威的比赛。 如果你不熟悉它，这里是维基百科页面 。 我创建了一个版本，通过保持数值数组，其中0表示死细胞和1个活的工作。 然后，内核简单地写入到图像缓冲器数据阵列来绘制基于小区数据的图像，然后检查每个单元的邻居来更新用于下一个执行的渲染单元阵列。
然而，更快的方法，而不是代表为负数，如果死了，只要活着正数一个单元格的值。 该小区的数量表示它具有邻居的量加一（使零是不可能的值，因为我们无法通过-0区分0）。 然而，这意味着，产卵或杀死细胞的时候，我们必须更新其相应的八个邻居的值。 因此不像工作程序，其中只有从相邻内存插槽读取，这个过程必须写入的插槽。 这样做是不一致的和输出的阵列是无效的。 例如细胞包含数字，例如14，其指示13楼的邻居，一个不可能的值。 该代码是正确的，因为我写上的CPU相同的程序和它的作品如预期。 经过测试，我相信，当任务尝试写入，同时内存也导致某种类型的写入错误的延迟。 例如，也许有读取在该时间内数据被改变使得另一个任务的程序不正确的阵列数据和设置之间的延迟。 我已经使用semaphors和障碍试过，但刚刚得知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常数分别是图像缓冲器的宽度和高度。
- 颜色表数组包含黑白分别被用于适当地改变图像缓冲器的值，以使颜色的RGB整数值。
- 所述newCellMap阵列是一旦渲染被确定来计算更新的细胞图。
- 该historyBuffer是细胞在内核调用开始的老态。 每次执行内核的时候，这个数组被更新到newCellMap阵列。
此外，该包装功能，使空间环形。 我怎么能解决这个代码，这样它按预期工作。 为什么不与任务的每个变化的全球内存更新？ 是不是应该是共享内存？
作为sharpneli他回答说，你正在阅读从不同的线程写入同一个存储区域，并且给出了一个未定义的行为。
解决方案：你需要分割你的newCellMap在2个阵列，一个是以前的执行和一个地方的新值将被保存。 然后，你需要从主机端更改内核参数在每次调用，从而使oldvalues下一次迭代的是newvalues前一次迭代。 由于你如何structurize你的算法，你也将需要执行的copybuffer 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的不具有共享跨越HOST-设备内存 。
当您创建用于该设备的内存缓冲区，你首先必须初始化与该存储区域clEnqueueWriteBuffer()与阅读clEnqueueWriteBuffer()得到的结果。 即使你有一个指针的存储区，你的指针指向到该区域的主机端副本。 这很可能是不具有设备计算输出的最后一个版本。
PD：我创建了很久以前就OpenCL的一个“活”的游戏，我发现easyer和更快的方式做到这一点很简单，就是创造位（位寻址）的一个大的二维数组。 然后写一段代码，而不简单地analize的neibours并获取该小区的更新值的任何分支机构。 由于位用于寻址，存储器的量的读/写每个线程是相当低的，解决字符/整数/其它。 我在一个非常古老的OpenCL HW（nVIDIA的9100M G）实现33Mcells /秒。 只是为了让你知道，你的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
};
在这里，你不保存任何的中间状态，只是在最后（活/死亡）的细胞状态。 它没有分支机构，所以它是相当快的过程。
