OpenCL Kernel与C/C++传统循环的对应关系及编写疑问

你对执行逻辑的猜测有一半是对的，但二者的运行模型有本质区别：

• 传统C/C++里的for循环是单线程串行执行：同一个线程按顺序让循环变量从0递增到n-1，逐轮完成每一次迭代的计算，循环边界n是代码里显式写的判断条件，满足i < n就继续跑，不满足就停止。
• OpenCL Kernel是多线程并行执行模型：Kernel本身没有内置的循环终止判断，你在主机端调用OpenCL API启动Kernel时，会通过clEnqueueNDRangeKernel的global_work_size参数显式指定要启动的并行工作项（可以理解为轻量线程）总数，这个值通常就等于你要处理的数组长度n。
  启动后OpenCL运行时会自动创建对应数量的工作项，每个工作项会独立执行一次你写的Kernel函数，函数内调用get_global_id(0)拿到的就是当前工作项的全局唯一编号，取值范围刚好是0到你设置的全局工作大小减1，正好对应数组0到n-1的下标。每个工作项只需要处理自己编号对应位置的元素即可，不需要在Kernel里写循环遍历整个数组。

举个最常见的向量加法例子，传统串行代码是：

// 传统C串行实现
const int n = 1024;
float a[n], b[n], c[n];
for (int i = 0; i < n; i++) {
    c[i] = a[i] + b[i];
}

对应的Kernel代码不需要写循环，直接写单元素的处理逻辑即可：

// OpenCL Kernel实现
__kernel void vec_add(__global const float* a, __global const float* b, __global float* c) {
    int i = get_global_id(0);
    c[i] = a[i] + b[i];
}

要注意的是，Kernel本身不会自动做边界检查，如果你主机端设置的全局工作大小大于数组实际长度n，多出来的工作项访问数组就会出现越界问题，和串行循环写错边界导致越界是一样的。

无数据依赖的多层嵌套循环，直接对应OpenCL的多维工作项配置即可，OpenCL最多支持3维的工作项网格，足够覆盖绝大多数双层、三层循环的场景。
以最常见的二维矩阵加法为例，传统嵌套循环代码是：

// 传统C串行嵌套循环，矩阵大小为rows行 * cols列
const int rows = 512, cols = 1024;
float a[rows*cols], b[rows*cols], c[rows*cols];
for (int row = 0; row < rows; row++) {
    for (int col = 0; col < cols; col++) {
        int idx = row * cols + col;
        c[idx] = a[idx] + b[idx];
    }
}

写对应Kernel的时候，主机端启动Kernel时把global_work_size设为二维值：第一维长度为cols，第二维长度为rows。Kernel里分别取两个维度的全局ID，对应内外层循环的循环变量即可，不需要写嵌套for：

// 对应嵌套循环的OpenCL Kernel
__kernel void mat_add(__global const float* a, __global const float* b, __global float* c, const int cols) {
    int col = get_global_id(0); // 对应内层循环变量col
    int row = get_global_id(1); // 对应外层循环变量row
    int idx = row * cols + col;
    c[idx] = a[idx] + b[idx];
}

如果你的循环层数超过3层，或者迭代之间存在前后依赖（比如后一次计算需要用到前一次迭代的结果），才需要在Kernel内部手动写循环实现，这类场景没法直接通过多维工作项完全并行化。

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