OpenCL 学习step by step (5) 使用二维NDRange workgroup

在本教程中，我们使用二维NDRange来设置workgroup，这样在opencl中，workitme的组织形式是二维的，Kernel中的代码也要做相应的改变，我们先看一下clEnqueueNDRangeKernel函数的变化。首先我们指定了workgroup size为localx*localy,通常这个值为64的倍数，但最好不要超过256。

//执行kernel,Range用2维，work itmes size为width*height,
cl_event ev;
size_t globalThreads[] = {width, height};
size_t localx, localy;
if(width/8 > 4)
localx = 16;
else if(width < 8)
localx = width;
else localx = 8;

if(height/8 > 4)
localy = 16;
else if (height < 8)
localy = height;
else localy = 8;

size_t localThreads[] = {localx, localy}; // localx*localy应该是64的倍数
printf("global_work_size =(%d,%d), local_work_size=(%d, %d)\n",width,height,localx,localy);

clTimer.Reset();
clTimer.Start();
clEnqueueNDRangeKernel( queue,
    kernel,
    2,
    NULL,
    globalThreads,
    localThreads, 0, NULL, &ev);

注意：在上面代码中，定义global threads以及local threads数量，都是通过二维数组的方式进行的。

新的Kernel代码如下：

#pragma OPENCL EXTENSION cl_amd_printf : enable

__kernel void vecadd(__global const float* a, __global const float* b, __global float* c)
{
    int x = get_global_id(0);
    int y = get_global_id(1);
    int width = get_global_size(0);
    int height = get_global_size(1);
    if(x == 1 && y ==1)
         printf("%d, %d,%d,%d,%d,%d\n",get_local_size(0),get_local_size(1),get_local_id(0),get_local_id(1),get_group_id(0),get_group_id(1));
    
    c[x + y * width] = a[x + y * width] + b[x + y * width];
    
}