OpenCL 基础篇

OpenCL 是由 Khronos Group 针对异质性计算装置（heterogeneous device）进行平行化运算所设计的标准 API 以及程式语言。本文将全面介绍OpenCL技术，从基本介绍到环境搭建等，很适合初学者学习阅读。

内容目录： OpenCL 简介 OpenCL 的架构 OpenCL 环境设定 开始撰写 OpenCL 程式 建立 Command Queue 产生资料  配置记忆体并复制资料 编译 OpenCL kernel 程式 执行 OpenCL kernel OpenCL 简介 OpenCL 是由 Khronos Group针对异质性计算装置（heterogeneous device）进行平行化运算所设计的标准 API 以及程式语言。所谓的「异质性计算装置」，是指在同一个电脑系统中，有两种以上架构差异很大的计算装置，例如一般的 CPU 以及显示晶片，或是像 CELL 的 PPE 以及 SPE。目前，最为常见的就是所谓的 GPGPU 应用，也就是利用一般的显示晶片（即 GPU）进行 3D 绘图以外的计算工作。 过去 GPGPU 的应用，有各种不同的使用方式。最早的 GPGPU，多半是直接透过 3D 绘图的 API 进行，例如 OpenGL 或 D3D 的 HLSL（High Level Shading Language）。但是，这样做有很多缺点，主要是即使想要进行的运算和 3D 绘图无关，仍需要处理很多 3D 绘图方面的动作（例如建立 texture，安排 render-to-texture 动作等等）。这让 GPGPU 变得十分复杂。后来开始有些尝试把这些 3D 绘图部份隐藏起來的想法，例如由 Stanford 大学设计的BrookGPU，可以透过不同的 backend 将 Brook 程式转换成由 CPU、Direct3D、或 OpenGL 来执行。另外，也有各家显示卡厂商自行开发的系统，包括 ATI 针对其产品设计的 Close to Metal（以及后来的 AMD Stream），以及 NVIDIA 的 CUDA。Microsoft 也在 DirectX 11 中加入了特別为 GPGPU 设计的 DirectCompute。 由于各家厂商的 GPGPU 方案都是互不相容的（例如 AMD Stream 的程式无法在 NVIDIA 的显示晶片上执行，而 CUDA 的程式也不能在 AMD 的显示晶片上执行），这对 GPGPU 的发展是不利的，因为程式开发者必须为不同厂商的显示晶片分別撰写程式，或是选择只支援某个显示晶片厂商。由于显示晶片的发展愈来愈弹性化，GPGPU 的应用范围也增加，因此 Apple 决定提出一个统一的 GPGPU 方案。这个方案得到包括 AMD、IBM、Intel、NVIDIA 等相关厂商的支持，并很快就交由 Khronos Group 进行标准化。整个计划只花了五个月的时间，并在 2008 年十二月时正式公开。第一个正式支援 OpenCL 的作业系统是 Apple 的 MacOS X 10.6 "Snow Leopard"。AMD 和 NVIDIA 也随后推出了在 Windows 及 Linux 上的 OpenCL 实作。IBM 也推出了支援 CELL 的 OpenCL 实作。 OpenCL 的主要设计目的，是要提供一个容易使用、且适用于各种不同装置的平行化计算平台。因此，它提供了两种平行化的模式，包括 task parallel 以及 data parallel。目前 GPGPU 的应用，主要是以 data parallel 为主，这里也是以这个部份为主要重点。所谓的 data parallel，指的是有大量的资料，都进行同样的处理。这种形式的平行化，在很多工作上都可以见到。例如，影像处理的程式，经常要对一个影像的每个 pixel 进行同样的动作（例如 Gaussian blur）。因此，这类工作很适合 data parallel 的模式。 OpenCL 的架构 OpenCL 包括一组 API 和一个程式语言。基本上，程式透过 OpenCL API 取得 OpenCL 装置（例如显示晶片）的相关资料，并将要在装置上执行的程式（使用 OpenCL 程式语言撰写）编译成适合的格式，在装置上执行。OpenCL API 也提供许多装置控制方面的动作，例如在 OpenCL 装置上取得一块记忆体、把资料从主记忆体复制到 OpenCL 装置上（或从 OpenCL 装置上复制到主记忆体中）、取得装置动作的资讯（例如上一个程式执行所花费的时间）等等。 例如，我们先考虑一个简单的工作：把一群数字相加。在一般的 C 程式中，可能是如下： float a[DATA_SIZE]; float b[DATA_SIZE]; float result[DATA_SIZE]; // ... for(int i = 0; i < DATA_SIZE; i++) { result[i] = a[i] + b[i]; } 在 OpenCL 中，则大致的流程是： 把 OpenCL 装置初始化。 在 OpenCL 装置上配置三块记忆体，以存放 a、b、c 三个阵列的资料。 把 a 阵列和 b 阵列的內容，复制到 OpenCL 装置上。 编译要执行的 OpenCL 程式（称为 kernel）。 执行编译好的 kernel。 把计算結果从 OpenCL 装置上，复制到 result 阵列中。 透过 data parallel 的模式，这里的 OpenCL 程式非常简单，如下所示： __kernel void adder(__global const float* a, __global const float* b, __global float* result) { int idx = get_global_id(0); result[idx] = a[idx] + b[idx]; } 在一般的版本中，是透过一个回圈，执行 DATA_SIZE 次数的加法动作。而在 OpenCL 中，则是建立DATA_SIZE 个 work item，每个 work item 都执行上面所示的 kernel。可以看到，OpenCL 程式语言和一般的 C 语言非常类似。__kernel 表示这个函式是在 OpenCL 装置上执行的。__global 则表示这个指标是在 global memory 中（即 OpenCL 装置上的主要記记忆体）。而 get_global_id(0) 会传回 work item 的编号，例如，如果有 1024 个 work item，则编号会分別是 0 ~ 1023（实际上编号可以是二维或三维，但在这里先只考虑一维的情形）。 要如何让上面这个简单的 OpenCL kernel 实际在 OpenCL 装置上执行呢？这就需要透过 OpenCL API 的帮助了。以下会一步一步说明使用 OpenCL API 的方法。 OpenCL 环境设定 在使用 OpenCL API 之前，不免要进行一些环境的设定。相关的动作可以参考下列的文章： 在 Windows 下使用 OpenCL 在 Xcode 中使用 OpenCL 开始撰写 OpenCL 程式 在使用 OpenCL API 之前，和绝大部份所有其它的 API 一样，都需要 include 相关的 header 档案。由于在 MacOS X 10.6 下 OpenCL 的 header 档案命名方式和在其它作业系统下不同，因此，通常要使用一个#ifdef 来进行区分。如下所示： #ifdef __APPLE__ #include <OpenCL/opencl.h> #else #include <CL/cl.h> #endif 这样就可以在 MacOS X 10.6 下，以及其它的作业系统下，都可以 include 正确的 OpenCL header 档。 接着，要先取得系统上所有的 OpenCL platform。在 MacOS X 10.6 下，目前只有一个由 Apple 提供的 OpenCL platform，但是在其它系统上，可能会有不同厂商提供的多个不同的 OpenCL platform，因此需要先取得 platform 的数目： cl_int err; cl_uint num; err = clGetPlatformIDs(0, 0, &num); if(err != CL_SUCCESS) { std::cerr << "Unable to get platforms\n"; return 0; } 大部份的 OpenCL API 会传错误值。如果传回值是 CL_SUCCESS 则表示执行成功，否则会传回某个错误值，表示失败的原因。 接着，再取得 platform 的 ID，这在建立 OpenCL context 时会用到： std::vector<cl_platform_id> platforms(num); err = clGetPlatformIDs(num, &platforms[0], &num); if(err != CL_SUCCESS) { std::cerr << "Unable to get platform ID\n"; return 0; } 在 OpenCL 中，类似这样的模式很常出现：先呼叫第一次以取得数目，以便配置足够的记忆体量。接着，再呼叫第二次，取得实际的资料。 接下来，要建立一个 OpenCL context。如下： cl_context_properties prop[] = { CL_CONTEXT_PLATFORM, reinterpret_cast<cl_context_properties>(platforms[0]), 0 };   cl_context context = clCreateContextFromType(prop, CL_DEVICE_TYPE_DEFAULT, NULL, NULL, NULL); if(context == 0) { std::cerr << "Can't create OpenCL context\n"; return 0; } clReleaseContext(context); return 0; 在上面的程式中，clCreateContextFromType 是一个 OpenCL 的 API，它可以从指定的装置类別中，建立一个 OpenCL context。第一个参数是指定 context 的 property。在 OpenCL 中，是透过一个 property 的阵列，以「property 种类」及「property 內容」成对出现，并以 0 做为结束。例如，以上面的例子来说，要指定的 property 种类是 CL_CONTEXT_PLATFORM，即要使用的 platform ID，而 property 內容则是由之前取得的 platform ID 中的第一个（即 platforms[0]）。由于 property 的內容可能是不同的资料型态，因此需要使用 reinterpret_cast 来进行强制转型。 第二个参数可以指定要使用的装置类別。目前可以使用的类別包括： CL_DEVICE_TYPE_CPU：使用 CPU 装置 CL_DEVICE_TYPE_GPU：使用显示晶片装置 CL_DEVICE_TYPE_ACCELERATOR：特定的 OpenCL 加速装置，例如 CELL CL_DEVICE_TYPE_DEFAULT：系统预设的 OpenCL 装置 CL_DEVICE_TYPE_ALL：所有系统计中的 OpenCL 装置 这里使用的是 CL_DEVICE_TYPE_DEFAULT，也就是指定使用预设的装置。另外，在这里，直接使用了之前取得的 OpenCL platform ID 中的第一个 ID（实际的程式中，可能会需要让使用者可以指定要使用哪一个 platform）。 如果建立 OpenCL context 失败，会传回 0。因此，要进行检查，并显示错误讯息。如果建立成功的话，在使用完后，要记得将 context 释放。这可以透过呼叫 clReleaseContext 来达成。 这个程式基本上已经可以编译执行了，但是当然它并沒有真的做什么事情。 一个 OpenCL context 中可以包括一个或多个装置，所以接下来的工作是要取得装置的列表。要取得任何和 OpenCL context 相光的资料，可以使用 clGetContextInfo 函式。以下是取得装置列表的方式： size_t cb; clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &cb); std::vector<cl_device_id> devices(cb / sizeof(cl_device_id)); clGetContextInfo(context, CL_CONTEXT_DEVICES, cb, &devices[0], 0); CL_CONTEXT_DEVICES 表示要取得装置的列表。和前面取得 platform ID 的情形相同，clGetContextInfo 被呼叫了两次：第一次是要取得需要存放装置列表所需的记忆体空间大小（也就是传入 &cb），然后第二次呼叫才真正取得所有装置的列表。 接下来，可能会想要确定倒底找到的 OpenCL 装置是什么。所以，可以透过 OpenCL API 取得装置的名称，并将它印出來。取得和装置相关的资料，是使用 clGetDeviceInfo 函式，和前面的 clGetContextInfo 函式相当类似。以下是取得装置名称的方式： clGetDeviceInfo(devices[0], CL_DEVICE_NAME, 0, NULL, &cb); std::string devname; devname.resize(cb); clGetDeviceInfo(devices[0], CL_DEVICE_NAME, cb, &devname[0], 0); std::cout << "Device: " << devname.c_str() << "\n"; 到目前为止，完整的程式应该如下所示：   // OpenCL tutorial 1 #include <iostream> #include <string> #include <vector> #ifdef __APPLE__ #include <OpenCL/opencl.h> #else #include <CL/cl.h> #endif int main() { cl_context context = clCreateContextFromType(prop, CL_DEVICE_TYPE_DEFAULT, NULL, NULL, NULL); if(context == 0) { std::cerr << "Can't create OpenCL context\n"; return 0; } size_t cb; clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &cb); std::vector<cl_device_id> devices(cb / sizeof(cl_device_id)); clGetContextInfo(context, CL_CONTEXT_DEVICES, cb, &devices[0], 0); clGetDeviceInfo(devices[0], CL_DEVICE_NAME, 0, NULL, &cb); std::string devname; devname.resize(cb); clGetDeviceInfo(devices[0], CL_DEVICE_NAME, cb, &devname[0], 0); std::cout << "Device: " << devname.c_str() << "\n"; clReleaseContext(context); return 0; } 执行这个程式，如果建立 OpenCL context 成功的话，应该会显示出找到的 OpenCL 装置的名称，例如 Device: GeForce GTX 285 建立 Command Queue 大部份 OpenCL 的操作，都要透过 command queue。Command queue 可以接收对一个 OpenCL 装置的各种操作，并按照顺序执行（OpenCL 也容许把一个 command queue 指定成不照顺序执行，即 out-of-order execution，但是这里先不讨论这个使用方式）。所以，下一步是建立一个 command queue： cl_command_queue queue = clCreateCommandQueue(context, devices[0], 0, 0); if(queue == 0) { std::cerr << "Can't create command queue\n"; clReleaseContext(context); return 0; } 和 context 一样，在程式结束前，要把 command queue 释放，即： clReleaseCommandQueue(queue); 上面的程式中，是把装置列表中的第一个装置（即 devices[0]）建立 command queue。如果想要同时使用多个 OpenCL 装置，则每个装置都要有自己的 command queue。 产生资料 由于这个程式的目的是要把一大堆数字进行相加，所以需要产生一些「测试资料」：   const int DATA_SIZE = 1048576; std::vector<float> a(DATA_SIZE), b(DATA_SIZE), res(DATA_SIZE); for(int i = 0; i < DATA_SIZE; i++) { a[i] = std::rand(); b[i] = std::rand(); } 配置记忆体并复制资料 要使用 OpenCL 装置进行运算时，通常会需要在 OpenCL 装置上配置记忆体，并把资料从主记忆体中复制到装置上。有些 OpenCL 装置可以直接从主记忆体存取资料，但是速度通常会比较慢，因为 OpenCL 装置（例如显示卡）通常会有专用的高速记忆体。以下的程式配置三块记忆体： cl_mem cl_a = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(cl_float) * DATA_SIZE, &a[0], NULL); cl_mem cl_b = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(cl_float) * DATA_SIZE, &b[0], NULL); cl_mem cl_res = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(cl_float) * DATA_SIZE, NULL, NULL); if(cl_a == 0 || cl_b == 0 || cl_res == 0) { std::cerr << "Can't create OpenCL buffer\n"; clReleaseMemObject(cl_a); clReleaseMemObject(cl_b); clReleaseMemObject(cl_res); clReleaseCommandQueue(queue); clReleaseContext(context); return 0; } clCreateBuffer 函式可以用来配置记忆体。它的第二个参数可以指定记忆体的使用方式，包括： CL_MEM_READ_ONLY：表示 OpenCL kernel 只会对这块记忆体进行读取的动作 CL_MEM_WRITE_ONLY：表示 OpenCL kernel 只会对这块记忆体进行写入的动作 CL_MEM_READ_WRITE：表示 OpenCL kernel 会对这块记忆体进行读取和写入的动作 CL_MEM_USE_HOST_PTR：表示希望 OpenCL 装置直接使用指定的主记忆体位址。要注意的是，如果 OpenCL 装置无法直接存取主记忆体，它可能会将指定的主记忆体位址的资料复制到 OpenCL 装置上。 CL_MEM_ALLOC_HOST_PTR：表示希望配置的记忆体是在主记忆体中，而不是在 OpenCL 装置上。不能和 CL_MEM_USE_HOST_PTR 同时使用。 CL_MEM_COPY_HOST_PTR：将指定的主记忆体位址的资料，复制到配置好的记忆体中。不能和 CL_MEM_USE_HOST_PTR 同时使用。 第三个参数是指定要配置的记忆体大小，以 bytes 为单位。在上面的程式中，指定的大小是sizeof(cl_float) * DATA_SIZE。 第四个参数是指定主记忆体的位置。因为对 cl_a 和 cl_b 来说，在第二个参数中，指定了 CL_MEM_COPY_HOST_PTR，因此要指定想要复制的资料的位址。cl_res 则不需要指定。 第五个参数是指定错误码的传回位址。在这里并沒有使用到。 如果 clCreateBuffer 因为某些原因无法配置记忆体（例如 OpenCL 装置上的记忆体不夠），则会传回 0。要释放配置的记忆体，可以使用 clReleaseMemObject 函式。 編译 OpenCL kernel 程式 现在执行 OpenCL kernel 的准备工作已经大致完成了。所以，现在剩下的工作，就是把 OpenCL kernel 程式编译并执行。首先，先把前面提过的 OpenCL kernel 程式，存放在一个文字档中，命名为 shader.cl： __kernel void adder(__global const float* a, __global const float* b, __global float* result) { int idx = get_global_id(0); result[idx] = a[idx] + b[idx]; } 要编译这个 kernel 程式，首先要把档案內容读进来，再使用 clCreateProgramWithSource 这个函式，然后再使用 clBuildProgram 编译。如下所示： cl_program load_program(cl_context context, const char* filename) { std::ifstream in(filename, std::ios_base::binary); if(!in.good()) { return 0; } // get file length in.seekg(0, std::ios_base::end); size_t length = in.tellg(); in.seekg(0, std::ios_base::beg); // read program source std::vector<char> data(length + 1); in.read(&data[0], length); data[length] = 0; // create and build program  const char* source = &data[0]; cl_program program = clCreateProgramWithSource(context, 1, &source, 0, 0); if(program == 0) { return 0; } if(clBuildProgram(program, 0, 0, 0, 0, 0) != CL_SUCCESS) { return 0; } return program; } 上面的程式，就是直接将档案读到记忆体中，再呼叫 clCreateProgramWithSource 建立一个 program object。建立成功后，再呼叫 clBuildProgram 函式编译程式。clBuildProgram 函式可以指定很多参数，不过在这里暂时沒有使用到。 有了这个函式，在 main 函式中，直接呼叫： cl_program program = load_program(context, "shader.cl"); if(program == 0) { std::cerr << "Can't load or build program\n"; clReleaseMemObject(cl_a); clReleaseMemObject(cl_b); clReleaseMemObject(cl_res); clReleaseCommandQueue(queue); clReleaseContext(context); return 0; } 同样的，在程式结束前，要记得将 program object 释放： clReleaseProgram(program); 一个 OpenCL kernel 程式里面可以有很多个函式。因此，还要取得程式中函式的进入点： cl_kernel adder = clCreateKernel(program, "adder", 0); if(adder == 0) { std::cerr << "Can't load kernel\n"; clReleaseProgram(program); clReleaseMemObject(cl_a); clReleaseMemObject(cl_b); clReleaseMemObject(cl_res); clReleaseCommandQueue(queue); clReleaseContext(context); return 0; } 和 program object 一样，取得的 kernel object 也需要在程式结束前释放： clReleaseKernel(adder); 执行 OpenCL kernel 弄了这么多，总算可以执行 OpenCL kernel 程式了。要执行 kernel 程式，只需要先设定好函式的参数。adder 函式有三个参数要设定： clSetKernelArg(adder, 0, sizeof(cl_mem), &cl_a); clSetKernelArg(adder, 1, sizeof(cl_mem), &cl_b); clSetKernelArg(adder, 2, sizeof(cl_mem), &cl_res); 设定参数是使用 clSetKernelArg 函式。它的参数很简单：第一个参数是要设定的 kernel object，第二个是参数的编号（从 0 开始），第三个参数是要设定的参数的大小，第四个参数则是实际上要设定的参数內部。以这里的 adder 函式来说，三个参数都是指向 memory object 的指标。 设定好参数后，就可以开始执行了。如下： size_t work_size = DATA_SIZE; err = clEnqueueNDRangeKernel(queue, adder, 1, 0, &work_size, 0, 0, 0, 0); clEnqueueNDRangeKernel 会把执行一个 kernel 的动作加到 command queue 里面。第三个参数（1）是指定 work item 数目的维度，在这里就是一维。第五个参数是指定 work item 的总数目，也就是 DATA_SIZE。后面的参数现在暂时先不用管。如果成功加入的话，会传回 CL_SUCCESS。否则会传错误值。 在执行 kernel 被加到 command queue 之后，就可能会开始执行（如果 command queue 现在沒有別的工作的话）。但是 clEnqueueNDRangeKernel 是非同步的，也就是说，它并不会等待 OpenCL 装置执行完毕才传回。这样可以让 CPU 在 OpenCL 装置在进行运算的同时，进行其它的动作。 由于执行的结果是在 OpenCL 裝置的记忆体中，所以要取得结果，需要把它的內容复制到 CPU 能存取的主记忆体中。这可以透过下面的程式完成： if(err == CL_SUCCESS) { err = clEnqueueReadBuffer(queue, cl_res, CL_TRUE, 0, sizeof(float) * DATA_SIZE, &res[0], 0, 0, 0); } clEnqueueReadBuffer 函式会把「将记忆体资料从 OpenCL 装置复制到主记忆体」的动作加到 command queue 中。第三个参数表示是否要等待复制的动作完成才传回，CL_TRUE 表示要等待。第五个参数是要复制的资料大小，第六个参数则是目标的位址。 由于这里指定要等待复制动作完成，所以当函式传回时，资料已经完全复制完成了。最后是进行验证，确定资料正确： if(err == CL_SUCCESS) { bool correct = true; for(int i = 0; i < DATA_SIZE; i++) { if(a[i] + b[i] != res[i]) { correct = false; break; } } if(correct) { std::cout << "Data is correct\n"; } else { std::cout << "Data is incorrect\n"; } } else { std::cerr << "Can't run kernel or read back data\n"; } 到这里，整个程式就算是完成了。编译后执行，如果顺利的话，应该改会印出 Data is correct 的讯息。 以下是整个程式的全貌：   // OpenCL tutorial 1 #include <iostream> #include <fstream> #include <string> #include <vector> #include <cstdlib> #ifdef __APPLE__ #include <OpenCL/opencl.h> #else #include <CL/cl.h> #endif cl_program load_program(cl_context context, const char* filename) { std::ifstream in(filename, std::ios_base::binary); if(!in.good()) { return 0; } // get file length in.seekg(0, std::ios_base::end); size_t length = in.tellg(); in.seekg(0, std::ios_base::beg); // read program source std::vector<char> data(length + 1); in.read(&data[0], length); data[length] = 0; // create and build program  const char* source = &data[0]; cl_program program = clCreateProgramWithSource(context, 1, &source, 0, 0); if(program == 0) { return 0; } if(clBuildProgram(program, 0, 0, 0, 0, 0) != CL_SUCCESS) { return 0; } return program; } int main() { cl_context context = clCreateContextFromType(prop, CL_DEVICE_TYPE_DEFAULT, NULL, NULL, NULL); if(context == 0) { std::cerr << "Can't create OpenCL context\n"; return 0; } size_t cb; clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &cb); std::vector<cl_device_id> devices(cb / sizeof(cl_device_id)); clGetContextInfo(context, CL_CONTEXT_DEVICES, cb, &devices[0], 0); clGetDeviceInfo(devices[0], CL_DEVICE_NAME, 0, NULL, &cb); std::string devname; devname.resize(cb); clGetDeviceInfo(devices[0], CL_DEVICE_NAME, cb, &devname[0], 0); std::cout << "Device: " << devname.c_str() << "\n"; cl_command_queue queue = clCreateCommandQueue(context, devices[0], 0, 0); if(queue == 0) { std::cerr << "Can't create command queue\n"; clReleaseContext(context); return 0; } const int DATA_SIZE = 1048576; std::vector<float> a(DATA_SIZE), b(DATA_SIZE), res(DATA_SIZE); for(int i = 0; i < DATA_SIZE; i++) { a[i] = std::rand(); b[i] = std::rand(); } cl_mem cl_a = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(cl_float) * DATA_SIZE, &a[0], NULL); cl_mem cl_b = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(cl_float) * DATA_SIZE, &b[0], NULL); cl_mem cl_res = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(cl_float) * DATA_SIZE, NULL, NULL); if(cl_a == 0 || cl_b == 0 || cl_res == 0) { std::cerr << "Can't create OpenCL buffer\n"; clReleaseMemObject(cl_a); clReleaseMemObject(cl_b); clReleaseMemObject(cl_res); clReleaseCommandQueue(queue); clReleaseContext(context); return 0; } cl_program program = load_program(context, "shader.cl"); if(program == 0) { std::cerr << "Can't load or build program\n"; clReleaseMemObject(cl_a); clReleaseMemObject(cl_b); clReleaseMemObject(cl_res); clReleaseCommandQueue(queue); clReleaseContext(context); return 0; } cl_kernel adder = clCreateKernel(program, "adder", 0); if(adder == 0) { std::cerr << "Can't load kernel\n"; clReleaseProgram(program); clReleaseMemObject(cl_a); clReleaseMemObject(cl_b); clReleaseMemObject(cl_res); clReleaseCommandQueue(queue); clReleaseContext(context); return 0; } clSetKernelArg(adder, 0, sizeof(cl_mem), &cl_a); clSetKernelArg(adder, 1, sizeof(cl_mem), &cl_b); clSetKernelArg(adder, 2, sizeof(cl_mem), &cl_res); size_t work_size = DATA_SIZE; err = clEnqueueNDRangeKernel(queue, adder, 1, 0, &work_size, 0, 0, 0, 0); if(err == CL_SUCCESS) { err = clEnqueueReadBuffer(queue, cl_res, CL_TRUE, 0, sizeof(float) * DATA_SIZE, &res[0], 0, 0, 0); } if(err == CL_SUCCESS) { bool correct = true; for(int i = 0; i < DATA_SIZE; i++) { if(a[i] + b[i] != res[i]) { correct = false; break; } } if(correct) { std::cout << "Data is correct\n"; } else { std::cout << "Data is incorrect\n"; } } else { std::cerr << "Can't run kernel or read back data\n"; } clReleaseKernel(adder); clReleaseProgram(program); clReleaseMemObject(cl_a); clReleaseMemObject(cl_b); clReleaseMemObject(cl_res); clReleaseCommandQueue(queue); clReleaseContext(context); return 0; } 在附件中可以下载包括 Xcode project 以及 Visual Studio 2008 project 档的原始码。 std::vector<cl_platform_id> platforms(num); err = clGetPlatformIDs(num, &platforms[0], &num); if(err != CL_SUCCESS) { std::cerr << "Unable to get platform ID\n"; return 0; } cl_context_properties prop[] = { CL_CONTEXT_PLATFORM, reinterpret_cast<cl_context_properties>(platforms[0]), 0 };   cl_int err; cl_uint num; err = clGetPlatformIDs(0, 0, &num); if(err != CL_SUCCESS) { std::cerr << "Unable to get platforms\n"; return 0; } std::vector<cl_platform_id> platforms(num); err = clGetPlatformIDs(num, &platforms[0], &num); if(err != CL_SUCCESS) { std::cerr << "Unable to get platform ID\n"; return 0; } cl_context_properties prop[] = { CL_CONTEXT_PLATFORM, reinterpret_cast<cl_context_properties>(platforms[0]), 0 };   cl_int err; cl_uint num; err = clGetPlatformIDs(0, 0, &num); if(err != CL_SUCCESS) { std::cerr << "Unable to get platforms\n"; return 0; }