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 中最关键的概念之一，它指明了数据存储在设备内存的哪个部分。

限定符 描述 用途 类比（以 NVIDIA GPU 为例） __global 全局内存。所有工作项都可读写，但访问速度较慢。 用于输入和输出的大型数据缓冲区。 显存（DRAM） __constant 常量内存。只读，所有工作项可访问，缓存速度快。 用于在内核执行期间不会改变的只读数据（如查找表、配置参数）。 常量缓存 __local 局部内存。工作组（Work-Group）内的工作项可共享的内存。速度比全局内存快。 用于工作组内部的通信和共享数据的临时存储。 共享内存 __private 私有内存。每个工作项的私有内存（默认）。 用于函数内部的局部变量和寄存器溢出。 寄存器和线程私有内存

示例：

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 >&