使用2D内核工作组计算OpenCL数组偏移量



我很难使用2D工作组计算正确的数组偏移量。我用这些参数调用我的内核

ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&a_mem_obj);
ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&b_mem_obj);
ret = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&c_mem_obj);
// Initialize kernel work dimensions 
const size_t globalWorkSizeCalc[2] = { 32, 32 }; //rows and cols ie x and y
const size_t localWorkSizeCalc[2] =  { 4, 4 }; //rows and cols ie x and y
//Total WorkGroups = (32/4) * (32/4) = 64
//LocalWorkGroups = (4*4) * 64 = 1024
ret = clEnqueueNDRangeKernel(command_queue, kernel, 2, NULL, globalWorkSizeCalc, localWorkSizeCalc, 0, NULL, NULL);

内核内部如何计算以更正简单数组添加的索引偏移量?

//use vectors for conciseness
int2 globalId = (int2)(get_global_id(0), get_global_id(1));  // 0..31 x 0..31
int2 localId = (int2)(get_local_id(0), get_local_id(1)); //0..3 x 0..3
int2 groupId = (int2)(get_group_id(0), get_group_id(1)); //0..7 x 0..7
int2 globalSize = (int2)(get_global_size(0), get_global_size(1)); //32 x 32
int2 localSize = (int2)(get_local_size(0), get_local_size(1));  //4 x 4
int2 numberOfGrp = (int2)(get_num_groups(0), get_num_groups(1));  //8 x 8
int index = ???
C[index] = A[index] + B[index];

运行相同的代码,localworksize为32X32,如下所示,

const size_t localWorkSizeCalc[2] =  { 32, 32 }; 

我成功地计算了偏移量,但对于4X4,我有一个问题。

int index = (localSize.x *  localId.y) + localId.x;
C[index] = A[index] + B[index];

**EDIT(解决方案(:它似乎与下面的代码一起工作,但我想知道我是否可以使用localId和Groups来实现同样的功能。有什么建议吗?

int index = (globalId.y * globalSize.x) + globalId.x;

带有'int index=(globalId.y*globalSize.x(+globalId.x;'你已经找到了解决方案。这是一个简单的2D线性索引。您也可以用'groupId.x*localSize.x+localID.x'(与'y'相同(替换'globalId.x',用'numberOfGrp.x*localSize.x'替换'globalSize.x'(全局大小必须是本地大小的倍数才能起作用(,但在这种情况下,这没有提供额外的好处。'localId只有在使用本地内存时才有用。

相关内容

最新更新