OpenCL C 内核(Kernel)

1. 内核（Kernel）的基本概念

• 内核函数：内核是一个特殊的函数，它在 OpenCL 设备上并行执行。当主机（CPU）程序发起执行请求时，这个函数会被大量的工作项（Work-Items）同时执行。

• 并行模型：OpenCL 使用 NDRange（N-Dimensional Range）模型来定义并行性。你可以把它想象成一个一维、二维或三维的网格，网格中的每一个点都是一个工作项，每个工作项都执行相同的内核代码，但通过不同的全局ID来区分彼此，从而处理不同的数据。

2. 内核函数的编写规则

一个标准的 OpenCL C 内核函数需要遵循以下规则：

1. 使用 __kernel 关键字：这个限定符声明一个函数是内核函数，可以从主机端调用。

2. 返回值必须是 void：内核函数不能有返回值。

3. 参数限制：所有参数都必须位于特定的地址空间（__global, __constant, __local, __private）。指针必须明确指定地址空间。

3. 内核函数的参数：地址空间限定符

这是 OpenCL C 中最关键的概念之一，它指明了数据存储在设备内存的哪个部分。

示例：

c

__kernel void myKernel(
    __global const float* inputA, // 输入缓冲区 A，只读
    __global const float* inputB, // 输入缓冲区 B，只读
    __global float* output,       // 输出缓冲区
    __constant float* coefficients, // 常量参数（如滤波器系数）
    __local float* sharedTemp,    // 局部共享内存
    int arraySize                 // 一个标量值，默认为 private
) {
    // ... 内核代码 ...
}

4. 获取工作项标识

在内核内部，你需要知道当前是哪个工作项在执行，以便处理正确的数据。OpenCL 提供了内置函数来获取这些信息：

• get_global_id(dim): 返回在 全局NDRange 中指定维度（0, 1, 2）的ID。

• get_local_id(dim): 返回在 工作组内部 指定维度的ID。

• get_group_id(dim): 返回当前工作组的 工作组ID。

• get_global_size(dim): 返回 全局NDRange 在指定维度的大小。

• get_local_size(dim): 返回 工作组 在指定维度的大小。

一维示例：向量加法
这是最经典的入门示例，演示了如何通过全局ID来映射数据。

c

// Kernel: Vector Addition (VecAdd)
// 每个工作项计算一个输出元素 output[i] = inputA[i] + inputB[i]

__kernel void vecAdd(
    __global const float* a,
    __global const float* b,
    __global float* c)
{
    // 获取当前工作项的全局一维ID
    int gid = get_global_id(0);

    // 执行加法操作
    c[gid] = a[gid] + b[gid];
}

主机端需要确保启动的全局工作项数量（Global Work Size）至少等于向量的长度。

二维示例：图像处理（旋转、模糊等）

c

// Kernel: 图像处理（例如，每个工作项处理一个像素）

__kernel void imageFilter(
    __global const uchar4* inputImage,
    __global uchar4* outputImage,
    int width,
    int height)
{
    // 获取当前工作项在2D网格中的坐标
    int x = get_global_id(0);
    int y = get_global_id(1);

    // 检查边界，防止越界
    if (x < width && y < height) {
        // 计算一维索引
        int idx = y * width + x;

        // 读取输入像素（例如，一个包含RGBA的4分量向量）
        uchar4 pixel = inputImage[idx];

        // 进行处理（例如，简单的颜色反转）
        uchar4 outputPixel;
        outputPixel.x = 255 - pixel.x; // R
        outputPixel.y = 255 - pixel.y; // G
        outputPixel.z = 255 - pixel.z; // B
        outputPixel.w = pixel.w;       // A (Alpha通道保持不变)

        // 写入输出像素
        outputImage[idx] = outputPixel;
    }
}

5. 使用局部内存（__local）的工作组同步

局部内存允许一个工作组内的所有工作项高效地共享和协作。这通常需要配合屏障（Barrier） 来同步工作组内所有工作项的执行。

经典示例：并行归约（求和）

c

__kernel void sumReduction(
    __global const float* input,
    __global float* partialSums,
    __local float* localSums) // 局部内存，大小由主机在运行时指定
{
    int gid = get_global_id(0);
    int lid = get_local_id(0); // 工作组内的本地ID
    int groupId = get_group_id(0);

    // 将全局数据拷贝到局部内存
    localSums[lid] = input[gid];

    // 等待工作组内所有工作项都完成拷贝
    barrier(CLK_LOCAL_MEM_FENCE);

    // 在工作组内部进行归约求和
    for (int stride = get_local_size(0) / 2; stride > 0; stride >&