OpenCL核函数入队后未执行 GPU空闲 全流程错误码返回CL_SUCCESS

需求背景 #

需对不同元胞自动机的规则集合逐一验证双射属性，规则总规模精确为2^32，因此借助GPU并行算力完成批量校验。本人为OpenCL开发新手，已调试一周仍未定位到异常根因，请求相关经验开发者协助。

异常现象 #

• 核函数入队后，系统监测GPU处于空闲状态
• 在核函数中添加测试埋点语句验证执行状态，未观察到任何语句触发执行的痕迹
• 全流程所有OpenCL接口返回的错误码均为CL_SUCCESS
• 核函数执行完成后，从设备缓冲区读回的bijective_aut数组与主机端初始化状态完全一致，核函数中的所有数组修改语句从未被触发执行

代码示例 #

主机端代码 #

#define CL_USE_DEPRECATED_OPENCL_2_0_APIS

// 头文件引用部分

#define GLOBAL_SIZE 4096
#define LOCAL_SIZE 256
#define GLOBAL_SCOPE 0xffffffff
int main()
{
    // 假设GLOBAL_SCOPE可被GLOBAL_SIZE整除，无需处理余数
    long rules_per_thread = GLOBAL_SCOPE / GLOBAL_SIZE;
    
    int * starts = new int[GLOBAL_SIZE];
    int * stops = new int[GLOBAL_SIZE];

    int count = 0;
    for (int i = 0; i < GLOBAL_SIZE; i++) {
        starts[i] = count;
        count += rules_per_thread;
        stops[i] = count;
        count++;
    }
    
    ...
    // 获取平台、设备、构建程序部分
    ...

    /*====创建缓冲区====*/
    // 存储符合要求的元胞自动机的缓冲区
    const int bufSize = 10000;  // 缓冲区大小为随机估值，后续可能需要扩充
    uint32_t* bijective_aut = new uint32_t[bufSize];
    std::fill(&bijective_aut[0], &bijective_aut[bufSize - 1], 0);
    // 数组第一个元素为全局迭代器，指向下一个空闲写入位置，初始设为基础偏移量
    bijective_aut[0] = 3;
    // 数组第二个元素为数组长度标记
    bijective_aut[1] = bufSize;

    cl::Buffer buf(context, CL_MEM_READ_WRITE, sizeof(uint32_t) * bufSize);
    cl::Buffer starts_buf(context, CL_MEM_READ_ONLY, sizeof(int) * GLOBAL_SIZE);
    cl::Buffer stops_buf(context, CL_MEM_READ_ONLY, sizeof(int) * GLOBAL_SIZE);

    /*====设置命令队列====*/
    cl::CommandQueue queue(context, device);

    err = queue.enqueueWriteBuffer(buf, CL_FALSE, 0, sizeof(uint32_t) * bufSize, bijective_aut);
    err = queue.enqueueWriteBuffer(starts_buf, CL_FALSE, 0, sizeof(int) * GLOBAL_SIZE, starts);
    err = queue.enqueueWriteBuffer(stops_buf, CL_FALSE, 0, sizeof(int) * GLOBAL_SIZE, stops);

    /*====创建内核、设置参数====*/
    cl::Kernel bc_kernel(program, "bijection_check", &err);
    err = bc_kernel.setArg(0, buf);
    err = bc_kernel.setArg(1, starts_buf);
    err = bc_kernel.setArg(2, stops_buf);

    /*====执行内核====*/
    cl::Event event;
    err = queue.enqueueNDRangeKernel(bc_kernel, cl::NullRange, cl::NDRange(GLOBAL_SIZE), cl::NDRange(LOCAL_SIZE), nullptr, &event);
    event.wait();
    err = queue.enqueueReadBuffer(buf, CL_FALSE, 0, sizeof(uint32_t) * bufSize, bijective_aut);
    cl::finish();
}

设备端核函数代码 #

__kernel void bijection_check (
                                __global uint * bijective_rules, // 存储符合双射性要求的规则
                                __global const uint * starts, 
                                __global const uint * stops 
                              ) 
{
    __private int idx = get_global_id(0);
    int iterator = bijective_rules[0];  // 指向下一个待写入的空闲位置
    int start = starts[idx];
    int stop = stops[idx];
    bool check = true;
    // 测试所需的其他变量

    // 遍历[start, stop)区间内的所有规则
    for (uint rule = start; rule < stop; rule++)
    {
        ...
        // 规则双射性校验逻辑
        ...

        // 若规则符合双射性要求，写入全局列表
        if ((check == true) && (iterator < 10000))
        {
            bijective_rules[iterator] = rule;
            bijective_rules[0]++;
        }
        else
        {
            bijective_rules[2]++;
        }
    }
    bijective_rules[3]++;
}

内容的提问来源于stack exchange，提问作者absqua