OpenCL 设备队列

▶ 按书上写的设备队列的代码，需要 OpenCL2.0 的平台和设备，先把代码堆上来

● 程序主要功能：用主机上的数组 Ahost 和 Bhost 创建设备缓冲区 Adevice 和 Bdevice，调用核函数 foo 及其子核函数 fooChild 计算 factor * Adevice .* Bdevice，结果写入 Cdevice，最后拷贝回主机数组 Chost 检查结果。

● 代码

//deviceQueue.cl
__kernel void fooChild(const int nElement, const float factor,
    __global const float *A, __global const float *B, __global float *C)
{
    uint gid = get_global_id(0);
    if (gid < nElement)
        C[gid] = factor * A[gid] * B[gid];
}

__kernel void foo(const int nElement, const float factor,
    __global const float *A, __global const float *B, __global float *C)
{
    uint gid = get_global_id(0), gsize = get_global_size(0);
    uint childGsize = nElement / gsize, childOffset = gid * childGsize;

    __global const float *Achild = &A[childOffset];
    __global const float *Bchild = &B[childOffset];
    __global const float *Cchild = &C[childOffset];

    queue_t defQ = get_default_queue();
    ndrange_t ndrange = ndrange_1D(childGsize);
    void(^fooChildWrapper)(void) = ^{ fooChild(childGsize, factor, Achild, Bchild, Cchild); };
    enqueue_kernel(defQ, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, saxpyDpChildWrapper);
}

//main.c
#include <stdio.h>  
#include <stdlib.h>  
#include <cl.h>

const char *sourceCode = "D:/Code/deviceQueue.cl";

char* readSource(const char* kernelPath)// 读取文本文件，存储为 char *
{
    FILE *fp;
    char *source;
    long int size;
    //printf("readSource, Program file: %s
", kernelPath);
    fopen_s(&fp, kernelPath, "rb");
    if (!fp)
    {
        printf("Open kernel file failed
");
        exit(-1);
    }
    if (fseek(fp, 0, SEEK_END) != 0)
    {
        printf("Seek end of file faildd
");
        exit(-1);
    }
    if ((size = ftell(fp)) < 0)
    {
        printf("Get file position failed
");
        exit(-1);
    }
    rewind(fp);
    if ((source = (char *)malloc(size + 1)) == NULL)
    {
        printf("Allocate space failed
");
        exit(-1);
    }
    fread(source, 1, size, fp);
    fclose(fp);
    source[size] = '';
    return source;
}

int main()
{
    const int nElement = 8196, nChildElement = 128, dataSize = nElement * sizeof(float);
    float factor = 2.3f;    
    char info[1024] = { 0 };
    int i;

    // 初始化平台
    cl_int status;
    cl_platform_id platform;
    status = clGetPlatformIDs(1, &platform, NULL);
    cl_device_id device;
    status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
    cl_context_properties contextProp[] = { CL_CONTEXT_PLATFORM,(cl_context_properties)(platform), 0 };
    cl_context context = clCreateContext(contextProp, 1, &device, NULL, contextProp, &status);
    cl_queue_properties queueProp[3] = { CL_QUEUE_PROPERTIES,CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE | CL_QUEUE_ON_DEVICE | CL_QUEUE_ON_DEVICE_DEFAULT, 0 };
    cl_command_queue queue = clCreateCommandQueueWithProperties(context, device, queueProp, &status);
    cl_event eventProducer, eventConsumer;

    const char* source = readSource(sourceCode);
    cl_program program = clCreateProgramWithSource(context, 1, &source, NULL, &status);
    status = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
    if (status)
    {
        clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 1024, info, NULL);
        printf("Build log:
%s
", info);
    }

    cl_kernel foo = clCreateKernel(program, "foo", &status);    
    size_t globalSize = nElement / nChildElement, localSize = 1;// 每个父工作项调度 nChildElement 个子工作项

    float *Ahost = (float *)malloc(dataSize);
    float *Bhost = (float *)malloc(dataSize);
    float *Chost = (float *)malloc(dataSize);
    for (i = 0; i < nElement; Ahost[i] = i, Bhost[i] = i + 1, Chost[i] = 0.f, i++);

    cl_mem Adevice, Bdevice, Cdevice;
    Adevice = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, dataSize, Ahost, &status);
    Bdevice = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, dataSize, Bhost, &status);
    Cdevice = clCreateBuffer(context, CL_MEM_WRITE_ONLY, dataSize, NULL, &status);
    
    clSetKernelArg(foo, 0, sizeof(int), (void*)&nElement);
    clSetKernelArg(foo, 1, sizeof(float), (void*)&factor);
    clSetKernelArg(foo, 2, sizeof(cl_mem), Adevice);
    clSetKernelArg(foo, 3, sizeof(cl_mem), Bdevice);
    clSetKernelArg(foo, 4, sizeof(cl_mem), Cdevice);
    
    clEnqueueNDRangeKernel(queue, foo, 1, NULL, &globalSize, &localSize, 0, NULL, &eventProducer);
    clFinish(queue);

    clEnqueueReadBuffer(queue, Cdevice, CL_TRUE, dataSize, dataSize, Chost, 1, &eventConsumer, NULL);
    clFinish(queue);

    for (i = 0; i < nElement; i++)
    {
        if (Chost[i] != factor*i*(i + 1))
            break;
    }
    printf("Output is %s.
", (i == nElement) ? "correct" : "incorrect");

    free(Ahost);
    free(Bhost);
    free(Chost);
    clReleaseContext(context);
    clReleaseCommandQueue(queue);
    clReleaseProgram(program);
    clReleaseKernel(foo);
    clReleaseMemObject(Adevice);
    clReleaseMemObject(Bdevice);
    clReleaseMemObject(Cdevice);
    getchar();
    return 0;
}

● 输出结果

■ 一直卡在函数 clCreateCommandQueueWithProperties 的调用上，返回值 -6（CL_OUT_OF_HOST_MEMORY），原因不明，stackExchange 上有人说换了显卡驱动就好了（https://stackoverflow.com/questions/39864947/opencl-cl-out-of-host-memory-on-clcreatecommandqueuewithproperties-with-minima），还有人说是设备位数的问题（https://stackoverflow.com/questions/45231329/opencl-clcreatecommandqueue-cl-out-of-host-memory-error），但是我更新了显卡驱动，工程改成 32 位（才发现显卡是 32 位的）还是不行。

■ 强行忽略上面的问题（clCreateCommandQueueWithProperties 第四参数用 NULL）仍然程序编译失败，返回 -11（CL_BUILD_PROGRAM_FAILURE），原因是不能支持和函数中的 queue_t 和 ndrange_t 数据类型，后面的块语法就更别想了，应该是平台和设备不能完全支持 OpenCL2.0 所致。

■ 更新心啊卡驱动和 AMD OpenCL2.0 驱动，编译失败，报错 -11，报错信息变成了【Error: Inserting openCl Source to binary】和【Error: Compiling CL to IR】