在 OpenCL 中使用 clEnqueueNDRangeKernel

Question

我需要有关 OpenCL 中一项功能的帮助。当我开始使用clEnqueueNDRangeKernel 而不是clEnqueueTask 时，程序需要更多时间才能成功。为什么这样？据我了解，程序应该使用数据并行模型，它会运行得更快，我错了吗？如果我是，如何更改代码以查看数据并行模型的实际工作？

__kernel void black_white_img(__global unsigned char *pDataIn, __global unsigned char *pDataOut, unsigned int InSize, unsigned int OutSize)
{
    for (int i = 0, j = 0; i < InSize; i+=4, j++)
    {
        unsigned char Value = (pDataIn[i] + pDataIn[i + 1] + pDataIn[i + 2]) / 3;
        pDataOut[j] = Value;
    }
}

int iWidth, iHeight, iBpp;
vector<unsigned char> pDataIn;
vector<unsigned char> pDataOut;


int err = LoadBmpFile(L"3840x2160.bmp", iWidth, iHeight, iBpp, pDataIn);

if (err != 0 || pDataIn.size() == 0 || iBpp != 32)
{
    std::cout << "error load input file!\n";
}

pDataOut.resize(pDataIn.size()/4);


cl_device_id device_id = NULL;
cl_context context = NULL;
cl_command_queue command_queue = NULL;
cl_mem memobj = NULL;
cl_mem memobj1 = NULL;
cl_program program = NULL;
cl_kernel kernel = NULL;
cl_platform_id platform_id = NULL;
cl_uint ret_num_devices;
cl_uint ret_num_platforms;
cl_int ret;

unsigned int SizeIn, SizeOut;

SizeIn = pDataIn.size();
SizeOut = pDataOut.size();

FILE *fp;
char fileName[] = "./kernel.cl";
char *source_str;
size_t source_size;

//Loading kernel
fp = fopen(fileName, "r");
if (!fp) {
    fprintf(stderr, "Failed to load kernel.\n");
    system("PAUSE");
    exit(1);
}
source_str = (char*)malloc(MAX_SOURCE_SIZE);
source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp);
fclose(fp);

//Getting Platform and Device
ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);

ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices);


//Create context
context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);


//create kernel program
program = clCreateProgramWithSource(context, 1, (const char **)&source_str,
(const size_t *)&source_size, &ret);

//build it
ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);

//create queue
command_queue = clCreateCommandQueue(context, device_id, 0, &ret);

//create bufer
memobj = clCreateBuffer(context, CL_MEM_READ_WRITE, pDataIn.size(), NULL, &ret);

memobj1 = clCreateBuffer(context, CL_MEM_READ_WRITE,pDataOut.size(), NULL, &ret);
//copy buffer to kernel

ret = clEnqueueWriteBuffer(command_queue, memobj, CL_TRUE, 0, pDataIn.size(), pDataIn.data(), 0, NULL, NULL);


//create opencl kernel
kernel = clCreateKernel(program, "red_to_green", &ret);


//set kernel args
ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobj);
ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&memobj1);
ret = clSetKernelArg(kernel, 2, sizeof(unsigned int), (void *)&SizeIn);
ret = clSetKernelArg(kernel, 3, sizeof(unsigned int), (void *)&SizeOut);

const size_t cycles_max = 10;
clock_t t0 = clock();
for (int i = 0; i<cycles_max; i++){

    float start_time =  clock();
    float search_time = 0;
    //float last_time = 0;

    //execute opencl kernel
    //ret = clEnqueueTask(command_queue, kernel, 0, NULL, NULL);

    size_t global_item_size = 8;
    size_t local_item_size = 4;

    ret = clEnqueueNDRangeKernel(command_queue,kernel, 1, NULL, &global_item_size, &local_item_size, 0, NULL, NULL);

    //copy from buffer
    ret = clEnqueueReadBuffer(command_queue, memobj1, CL_TRUE, 0, pDataOut.size(), pDataOut.data(), 0, NULL, NULL);

    ret = clFinish(command_queue);

    float end_time = clock();
    search_time = end_time - start_time;
    //float last_time = last_time + search_time;
    cout << search_time << endl;

}

clock_t t1 = clock();
double time_seconds = (t1-t0)*CLOCKS_PER_SEC/cycles_max;
cout << time_seconds/1000 <<endl;
WriteBmpFile(L"3840x2160_wb.bmp", iWidth, iHeight, 8, pDataOut.size(), pDataOut.data(), false);
system("PAUSE");

Answer 1

from the docs page:

内核使用单个工作项执行。

clEnqueueTask 相当于调用 clEnqueueNDRangeKernel work_dim = 1，global_work_offset = NULL，global_work_size[0] 设置为 1， 并且 local_work_size[0] 设置为 1。

当您使用 clEnqueueNDRangeKernel 时，您使用的是 4 个工作项的 2 个工作组，但它们都在做同样的工作。它们都从同一个全局内存中读取，但更重要的是，它们都尝试写入全局内存中的相同位置。

在进行计算时，您需要考虑工作人员的全局 ID。

__kernel void black_white_img(__global unsigned char *pDataIn, __global unsigned char *pDataOut, unsigned int InSize, unsigned int OutSize)
{
    int gid = get_global_id(0);
    int gsize = get_global_size(0);

    for (int j = gid; j < (InSize >> 2); j+= gsize)
    {
        unsigned char Value = (pDataIn[j*4] + pDataIn[j*4 + 1] + pDataIn[j*4 + 2]) / 3;
        pDataOut[j] = Value;
    }
}

Answer 2

看起来您正在遍历内核中输入图像的所有像素。这将导致所有线程计算所有像素的图像强度。尝试为每个像素启动一个线程。为此，请将您的内核源代码更改为仅计算一个像素的输出值：

__kernel void black_white_img(__global unsigned char *pDataIn, __global unsigned char *pDataOut) {
    int j = get_global_id(0);
    int i = j*4;
    pDataOut[i] = (pDataIn[j] + pDataIn[j + 1] + pDataIn[j + 2]) / 3;
}

此代码现在将对位置 i 处的单个像素的 RGBA 输入图像的 RGB 值执行平均。现在您需要做的就是启动与图像像素一样多的线程。相关变化：

//create opencl kernel
kernel = clCreateKernel(program, "black_white_img", &ret);


//set kernel args
ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobj);
ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&memobj1);

const size_t cycles_max = 10;
clock_t t0 = clock();
for (int i = 0; i<cycles_max; i++){

float start_time =  clock();
float search_time = 0;
//float last_time = 0;

//execute opencl kernel
//ret = clEnqueueTask(command_queue, kernel, 0, NULL, NULL);

size_t global_item_size = iWidth * iHeight;

ret = clEnqueueNDRangeKernel(command_queue,kernel, 1, NULL, &global_item_size, NULL, 0, NULL, NULL);

与您的代码相比，这应该会大大加快速度。