类积分系统OpenCL同步问题及百万步矩阵仿真性能优化咨询

优化方案1：改进版单步Kernel方案（兼容所有OpenCL 1.x设备） #

这是实现成本最低、兼容性最好的可行方案，只需调整你现有方案2的两个高开销设计即可：

• 去掉每次循环的clFinish调用：只要你创建命令队列时没有指定CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE参数，队列默认是有序执行的，后入队的Kernel一定会等前一个Kernel完全执行完成后才会启动，天然保证每一步的全局数据一致性，不需要手动加同步。你只需要在所有Kernel都入队完成后，调用一次clFinish等待全部执行结束即可。
• Kernel内部使用双缓冲避免读写冲突：不要原地更新矩阵，申请两块全局内存空间matrix_prev和matrix_curr，每一步的Kernel只读matrix_prev的上一步结果，写入matrix_curr，单步执行完成后交换两个指针的角色即可，不需要在Kernel内加任何barrier，也不需要设置localWork参数。

改进后的主机端代码示例：

// 提前申请两块大小匹配的全局内存
cl_mem buf[2] = {create_buffer(matrix_size), create_buffer(matrix_size)};
int curr_buf_idx = 0;
// 循环仅执行入队操作，无同步开销
for(int t = 0; t < nSteps; t++)
{
    clSetKernelArg(kernel, 0, sizeof(cl_mem), &buf[curr_buf_idx ^ 1]); // 上一步数据为只读输入
    clSetKernelArg(kernel, 1, sizeof(cl_mem), &buf[curr_buf_idx]);   // 当前步数据为输出
    clSetKernelArg(kernel, 2, sizeof(int), &t);
    clEnqueueNDRangeKernel(queue, kernel, 2, NULL, globalWork, NULL, 0, NULL, NULL);
    curr_buf_idx ^= 1;
}
// 全部入队完成后仅做一次全局同步
clFinish(queue);

该方案的实际开销极低，主流OpenCL驱动的单次Kernel入队开销仅为几微秒，就算执行1e6次调用总开销也只有几秒，远低于GPU计算本身的耗时，完全满足性能要求。

优化方案2：步长折叠进一步降低调用开销 #

如果你的旧显卡驱动入队开销较高，可以将K个时间步的逻辑折叠到一个Kernel内执行，K可以取128/256/512等不会触发旧显卡超长循环崩溃的数值，主机端的循环次数直接降低为原来的1/K，开销可以忽略不计。

折叠后的Kernel代码示例：

kernel void FoldStep(global Vector* matrix_prev, global Vector* matrix_curr, const int start_t, const int step_num)
{
    int i = get_global_id(0);
    int j = get_global_id(1);
    // 用私有变量缓存当前点的数值，减少全局内存读写次数
    Vector curr_val = matrix_prev[i * WIDTH + j];
    global Vector* tmp;
    for(int k = 0; k < step_num; k++)
    {
        curr_val = DoSomething(i, j, matrix_prev, start_t + k, curr_val);
        // 每步结束后同步全局内存，保证依赖的其他点数据是上一步的结果
        barrier(CLK_GLOBAL_MEM_FENCE);
        // 双缓冲指针交换
        tmp = matrix_prev;
        matrix_prev = matrix_curr;
        matrix_curr = tmp;
    }
    matrix_curr[i * WIDTH + j] = curr_val;
}

优化方案3：设备端入队（仅支持OpenCL 2.0及以上设备） #

如果你的显卡支持OpenCL 2.0标准，可以直接在设备端创建命令队列，循环入队单步Kernel，完全不需要主机参与调用，开销几乎为0，完美避开主机调用的所有开销，同时不受单个工作组大小的限制，也不会触发超长循环崩溃问题。

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