▶ 一个完整的两向量加和的过程，包括查询平台、查询设备、创建山下文、创建命令队列、编译程序、创建内核、设置内核参数、执行内核、数据拷贝等。
● C 代码
#include <stdio.h>
#include <stdlib.h>
#include <cl.h> const int nElement = ; // 参与运算的矢量长度
// 内核函数代码，通常写在另一个 .cl 文件中
const char *programSource = " \
__kernel void vectorAdd(__global int *A, __global int *B, __global int *C) \
{ \
int idx = get_global_id(); \
C[idx] = A[idx] + B[idx]; \
return; \
} \
"; int main()
{
const size_t datasize = sizeof(int) * nElement;
int i, *A, *B, *C;
cl_int status; // 用于接收 OpenCL 函数状态返回值 A = (int*)malloc(datasize);
B = (int*)malloc(datasize);
C = (int*)malloc(datasize);
for (i = ; i < nElement; A[i] = B[i] = i, i++); // 初始化平台
cl_uint nPlatform;
status = clGetPlatformIDs(, NULL, &nPlatform); // 首次调用，获取平台数，注意返回值是显式赋值
cl_platform_id *listPlatform = (cl_platform_id*)malloc(nPlatform * sizeof(cl_platform_id)); // 用得到的平台数申请内存，保存平台数据结构
status = clGetPlatformIDs(nPlatform, listPlatform, NULL); // 再次调用 // 初始化设备
cl_uint nDevice = ;
status = clGetDeviceIDs(listPlatform[], CL_DEVICE_TYPE_ALL, , NULL, &nDevice); // 首次调用，获取设备数，需要传入平台数据结构
cl_device_id *listDevice = (cl_device_id*)malloc(nDevice * sizeof(cl_device_id)); // 用得到的设备数申请内存，保存设备数据结构
status = clGetDeviceIDs(listPlatform[], CL_DEVICE_TYPE_ALL, nDevice, listDevice, NULL);// 再次调用，初始化 // 创建上下文
cl_context context = clCreateContext(NULL, nDevice, listDevice, NULL, NULL, &status); // 需要传入设备数据结构 // 创建命令队列
cl_command_queue cmdQueue = clCreateCommandQueue(context, listDevice[], , &status); // 需要传入上下文和设备数据结构 // 创建数据缓冲区
cl_mem bufferA, bufferB, bufferC;
bufferA = clCreateBuffer(context, CL_MEM_READ_ONLY, datasize, NULL, &status); // 需要传入上下文
bufferB = clCreateBuffer(context, CL_MEM_READ_ONLY, datasize, NULL, &status);
bufferC = clCreateBuffer(context, CL_MEM_WRITE_ONLY, datasize, NULL, &status); // 主机数据写入设备（缓冲区）
status = clEnqueueWriteBuffer(cmdQueue, bufferA, CL_FALSE, , datasize, A, , NULL, NULL); // 需要传入命令队列
status = clEnqueueWriteBuffer(cmdQueue, bufferB, CL_FALSE, , datasize, B, , NULL, NULL); // 创建和编译程序
cl_program program = clCreateProgramWithSource(context, , (const char**)&programSource, NULL, &status);// 需要传入上下文
status = clBuildProgram(program, nDevice, listDevice, NULL, NULL, NULL); // 创建内核
cl_kernel kernel = clCreateKernel(program, "vectorAdd", &status); // 从 program 中抽取内核函数 vecadd // 声明内核调用时用到的参数
status = clSetKernelArg(kernel, , sizeof(cl_mem), &bufferA);
status = clSetKernelArg(kernel, , sizeof(cl_mem), &bufferB);
status = clSetKernelArg(kernel, , sizeof(cl_mem), &bufferC); // 设置工作项结构，即并行结构
size_t globalSize[] = { nElement }, localSize[] = {}; // 使用一维全局尺寸，不使用工作组尺寸 // 将内核入队，执行计算
status = clEnqueueNDRangeKernel(cmdQueue, kernel, , NULL, globalSize, localSize, , NULL, NULL); // 需要传入命令队列 // 计算结果返回主机
clEnqueueReadBuffer(cmdQueue, bufferC, CL_TRUE, , datasize, C, , NULL, NULL); // 检查结果
for (i = ; i < nElement; i++)
{
if (C[i] != i + i)
break;
}
printf("Output is %s.\n", (i==nElement) ? "correct" : "incorrect"); // 释放资源，包括主机资源和 OpenCL 资源
free(A);
free(B);
free(C);
free(listPlatform);
free(listDevice);
clReleaseContext(context);
clReleaseMemObject(bufferA);
clReleaseMemObject(bufferB);
clReleaseMemObject(bufferC);
clReleaseCommandQueue(cmdQueue);
clReleaseProgram(program);
clReleaseKernel(kernel);
getchar();
return ;
}
● 输出结果
Output is correct.
● 用到的宏和函数定义
// cl_platform.h
// 各种数据类型，没有更多信息
typedef struct _cl_platform_id * cl_platform_id;
typedef struct _cl_device_id * cl_device_id;
typedef struct _cl_context * cl_context;
typedef struct _cl_command_queue * cl_command_queue;
typedef struct _cl_mem * cl_mem;
typedef struct _cl_program * cl_program;
typedef struct _cl_kernel * cl_kernel;
typedef struct _cl_event * cl_event;
typedef struct _cl_sampler * cl_sampler; // 函数 clGetDeviceIDs 中的设备类型
#define CL_DEVICE_TYPE_DEFAULT (1 << 0)
#define CL_DEVICE_TYPE_CPU (1 << 1)
#define CL_DEVICE_TYPE_GPU (1 << 2)
#define CL_DEVICE_TYPE_ACCELERATOR (1 << 3)
#define CL_DEVICE_TYPE_CUSTOM (1 << 4)
#define CL_DEVICE_TYPE_ALL 0xFFFFFFFF // 上下文属性
#define CL_CONTEXT_PLATFORM 0x1084
#define CL_CONTEXT_INTEROP_USER_SYNC 0x1085 // 命令队列的属性
#define CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE (1 << 0)
#define CL_QUEUE_PROFILING_ENABLE (1 << 1) // 函数 clCreateBuffer 创建的缓冲区类型
#define CL_MEM_READ_WRITE (1 << 0)
#define CL_MEM_WRITE_ONLY (1 << 1)
#define CL_MEM_READ_ONLY (1 << 2)
#define CL_MEM_USE_HOST_PTR (1 << 3)
#define CL_MEM_ALLOC_HOST_PTR (1 << 4)
#define CL_MEM_COPY_HOST_PTR (1 << 5)
/* reserved (1 << 6) */
#define CL_MEM_HOST_WRITE_ONLY (1 << 7)
#define CL_MEM_HOST_READ_ONLY (1 << 8)
#define CL_MEM_HOST_NO_ACCESS (1 << 9) // 布尔变量宏
#define CL_FALSE 0
#define CL_TRUE 1
#define CL_BLOCKING CL_TRUE
#define CL_NON_BLOCKING CL_FALSE extern CL_API_ENTRY cl_int CL_API_CALL clGetPlatformIDs( // 查询平台数和初始化平台数据结构
cl_uint, // 需要初始化的平台数，首次调用时可以传入 0
cl_platform_id *, // 输入指针，首次调用传入 NULL 表查询平台数，再次调用传入平台据结构列表指针，表初始化
cl_uint * // 输出指针，首次调用要给一个保存平台数的整形变量的指针，再次调用传入 NULL 即可
) CL_API_SUFFIX__VERSION_1_0; // 版本号 extern CL_API_ENTRY cl_int CL_API_CALL clGetDeviceIDs( // 查询设备数和初始化设备数据结构
cl_platform_id, // 平台数据结构
cl_device_type, // 限定设备类型，CPU 或 GPU 或 全部
cl_uint, // 设备数量，首次调用传入 0 表查询设备数，再次调用传入已知的设备数量，表初始化
cl_device_id *, // 输入指针，首次调用传入 NULL 表查询设备数，再次调用传入设备据结构列表指针，表初始化
cl_uint *, // 输出指针，首次调用要给一个保存设备数的整形变量的指针，再次调用传入 NULL 即可
) CL_API_SUFFIX__VERSION_1_0; extern CL_API_ENTRY cl_context CL_API_CALL clCreateContext( // 创建上下文
const cl_context_properties *, // 上下文属性
cl_uint, // 使用该上下文的设备数量
const cl_device_id *, // 设备数据结构列表指针
void (CL_CALLBACK *)(const char *, const void *, size_t, void *),// 返回状态的回调函数
void *, // 用户数据
cl_int * // 返回结果状态的指针
) CL_API_SUFFIX__VERSION_1_0; extern CL_API_ENTRY cl_context CL_API_CALL clCreateContextFromType( // 批量创建平台上所有符合要求的设备的上下文，一并放在这里
const cl_context_properties *,
cl_device_type, // 设备类型
void (CL_CALLBACK *)(const char *, const void *, size_t, void *),
void *,
cl_int *
) CL_API_SUFFIX__VERSION_1_0; extern CL_API_ENTRY cl_command_queue CL_API_CALL clCreateCommandQueue( // 创建命令队列
cl_context, // 提供命令队列所处的上下文
cl_device_id, // 运行命令的设备数据类型
cl_command_queue_properties, // 命令队列属性，用宏定义
cl_int * // 返回结果状态的指针
) CL_API_SUFFIX__VERSION_1_0; extern CL_API_ENTRY cl_mem CL_API_CALL clCreateBuffer( // 创建数据缓冲区，显式示返回
cl_context, // 提供数据所处的上下文
cl_mem_flags, // 缓冲区属性，用宏定义
size_t, // 缓冲区大小（Byte）
void *, // 提供一个主机指针，在使用具有某些属性的缓冲区时要用到，否则传入 NULL 即可
cl_int * // 返回结果状态的指针
) CL_API_SUFFIX__VERSION_1_0; extern CL_API_ENTRY cl_int CL_API_CALL clEnqueueWriteBuffer( // 主机数据写入设备缓冲区
cl_command_queue, // 提供命令队列
cl_mem, // 目标缓冲区
cl_bool, // 是否阻塞写入，CL_TRUE / CL_FALSE
size_t, // 写入数据在缓冲区中的偏移量
size_t, // 写入数据尺寸（Byte）
const void *, // 主机数据源
cl_uint, // 等待列表中的事件数，不用时传入 0
const cl_event *, // 等待列表，不用时传入 NULL
cl_event * // 用于标记本次写入的事件的指针
) CL_API_SUFFIX__VERSION_1_0; extern CL_API_ENTRY cl_int CL_API_CALL clEnqueueReadBuffer( // 设备缓冲区数据写回主机
cl_command_queue, // 提供命令队列
cl_mem, // 数据源缓冲区
cl_bool, // 是否阻塞写入
size_t, // 写入数据在缓冲区中的偏移量
size_t, // 写入数据尺寸（Byte）
void *, // 主机数据源
cl_uint, // 等待列表中的事件个数，不用时传入 0
const cl_event *, // 等待列表，不用时传入 NULL
cl_event * // 用于标记本次写入的事件的指针
) CL_API_SUFFIX__VERSION_1_0; extern CL_API_ENTRY cl_program CL_API_CALL clCreateProgramWithSource( // 从文本编译程序
cl_context, // 提供上下文
cl_uint, // 程序数
const char **, // 程序代码
const size_t *, // 长度？实例代码中传入了 NULL
cl_int * // 返回结果状态的指针
) CL_API_SUFFIX__VERSION_1_0; extern CL_API_ENTRY cl_int CL_API_CALL clBuildProgram( // 在设备中新建编译好的程序
cl_program, // 程序名
cl_uint, // 要运行程序的设备数量
const cl_device_id *, // 要运行程序的设备数据结构列表指针
const char *, // 程序运行选项，没有时传入 NULL 即可
void (CL_CALLBACK *)(cl_program, void *), // 返回状态的回调函数
void * // 用户数据
) CL_API_SUFFIX__VERSION_1_0; extern CL_API_ENTRY cl_kernel CL_API_CALL clCreateKernel( // 创建程序内核
cl_program, // 程序名
const char *, // 要运行的函数名
cl_int * // 返回结果状态的指针
)CL_API_SUFFIX__VERSION_1_0; extern CL_API_ENTRY cl_int CL_API_CALL clSetKernelArg( // 声明内核调用时用到的参数
cl_kernel, // 使用的内核
cl_uint, // 参数编号，从 0 开始
size_t, // 参数大小（Byte）
const void * // 参数缓冲区的地址（类型为 (void*)cl_mem*）
) CL_API_SUFFIX__VERSION_1_0; extern CL_API_ENTRY cl_int CL_API_CALL clEnqueueNDRangeKernel( // 执行内核
cl_command_queue, // 提供命令队列
cl_kernel, // 使用的内核
cl_uint, // 工作项维度，1 或 2 或 3
const size_t *, // 工作项偏移，不用时传入 NULL
const size_t *, // 工作项尺寸，至多三维
const size_t *, // 工作组尺寸至多三维，不用时传入 NULL
cl_uint, // 等待列表中的事件编号，不用时传入 0 即可
const cl_event *, // 等待列表
cl_event * // 返回结果状态的指针
) CL_API_SUFFIX__VERSION_1_0; extern CL_API_ENTRY cl_int CL_API_CALL clReleaseContext(cl_context) CL_API_SUFFIX__VERSION_1_0; // 释放上下文资源 extern CL_API_ENTRY cl_int CL_API_CALL clReleaseMemObject(cl_mem) CL_API_SUFFIX__VERSION_1_0; // 释放内存对象资源 extern CL_API_ENTRY cl_int CL_API_CALL clReleaseCommandQueue(cl_command_queue) CL_API_SUFFIX__VERSION_1_0;// 释放命令队列资源 extern CL_API_ENTRY cl_int CL_API_CALL clReleaseProgram(cl_program) CL_API_SUFFIX__VERSION_1_0; // 释放程序资源 extern CL_API_ENTRY cl_int CL_API_CALL clReleaseKernel(cl_kernel) CL_API_SUFFIX__VERSION_1_0; // 释放内核资源
● C++ 代码，不知为何缺少 cl::Error 成员，无法返回错误信息
#include <iostream>
#include <fstream>
#include <string>
#include <cl.hpp> #define __NO_STD_VECTOR // 禁用 std：：vector，使用 cl::vector
#define __CL_ENABLE_EXCEPTIONS // 允许错误抛出 const int nElement = ;
const char *sourceProgram = " \
__kernel void vectorAdd(__global int *A, __global int *B, __global int *C) \
{ \
int idx = get_global_id(); \
C[idx] = A[idx] + B[idx]; \
return; \
} \
"; int main()
{
const size_t datasize = sizeof(int) * nElement;
int i;
int *A = new int[nElement], *B = new int[nElement], *C = new int[nElement];
for (i = ; i < nElement; A[i] = B[i] = i, i++); // 查询平台
std::vector<cl::Platform> listPlatform;
cl::Platform::get(&listPlatform); // 查询设备
std::vector<cl::Device> listDevice;
listPlatform[].getDevices(CL_DEVICE_TYPE_ALL, &listDevice); // 创建上下文，注意参数变少，后面几个函数也减少了输入参数
cl::Context context(listDevice); // 创建命令队列
cl::CommandQueue queue = cl::CommandQueue(context, listDevice[]); // 创建缓冲区
cl::Buffer bufferA = cl::Buffer(context, CL_MEM_READ_ONLY, datasize);
cl::Buffer bufferB = cl::Buffer(context, CL_MEM_READ_ONLY, datasize);
cl::Buffer bufferC = cl::Buffer(context, CL_MEM_WRITE_ONLY, datasize); // 主机数据写入缓冲区
queue.enqueueWriteBuffer(bufferA, CL_TRUE, , datasize, A);
queue.enqueueWriteBuffer(bufferB, CL_TRUE, , datasize, B); // 读取程序，如果是从文件读取的程序文件 XX.cl，则需要采用下面的三行代码，将文件内容转化为字符串，否则直接使用存储在字符串中的代码即可
//std::ifstream sourceFile("XX.cl");
//std::string sourceCode(std::istreambuf_iterator<char>(sourceFile), (std::istreambuf_iterator<char>()));
//cl::Program::Sources programSource(1, std::make_pair(sourceCode.c_str(), sourceCode.length() + 1));
cl::Program program = cl::Program(context, sourceProgram);
program.build(listDevice); // 创建内核
cl::Kernel kernel(program, "vectorAdd"); // 声明内核参数
kernel.setArg(, bufferA);
kernel.setArg(, bufferB);
kernel.setArg(, bufferC); // 执行内核
cl::NDRange globalSize(nElement), localSize();
queue.enqueueNDRangeKernel(kernel, cl::NullRange, globalSize, localSize); // 结果返回主机
queue.enqueueReadBuffer(bufferC, CL_TRUE, , datasize, C); // 检查结果
for (i = ; i < nElement; i++)
{
if (C[i] != i + i)
break;
}
printf("Output is %s.\n", (i == nElement) ? "correct" : "incorrect"); delete[] A, B, C;
getchar();
return ;
}
●输出结果同 C 代码
