使用 2D 内核工作组的 OpenCL 阵列偏移计算

Question

我很难使用 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];

**编辑（解决方案）：它似乎适用于下面的代码，但我想知道我是否可以使用 localId 和组来完成同样的工作。有什么推荐吗？

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

Answer 1

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

使用 2D 内核工作组的 OpenCL 阵列偏移计算

OpenCL array offset calculation using a 2D kernel workgroups

multidimensional-array

gpu

opencl