OpenCL™ 入门教程
本文旨在帮助您更轻松地理解和实现 OpenCL™。
Benedict R. Gaster，AMD 架构师，OpenCL™
OpenCL™ 是一项新兴技术，虽然规范已发布（www.khronos.org/registry/cl/），但目前很少有文档提供带有示例的基础介绍。本文旨在帮助您更轻松地理解和实现 OpenCL™。
请注意
- 我在 AMD 工作，因此我将在 Windows® 和 Linux® 上测试我们实现的所有示例代码；然而，我的目的是说明 OpenCL™ 的使用，而不考虑平台。所有示例均使用纯 OpenCL™ 编写，应能在任何实现上良好运行。
-
我已经尽力提供可在非 AMD OpenCL™ 实现上“开箱即用”的示例，但我不会在非 AMD 实现上进行测试；因此，示例在这些系统上可能无法按预期工作。如果发生这种情况，请通过我们的 Stream Computing 论坛告知我，我将尽最大努力纠正代码并发布更新。
以下“Hello World”教程提供了 OpenCL™ 的简单介绍。我希望在本次教程之后，继续推出更多涵盖以下主题的教程：
- 使用平台和设备层构建健壮的 OpenCL™
- 程序编译和内核对象
- 管理缓冲区
- 内核执行
- 内核编程 - 基础
- 内核编程 - 同步
- 矩阵乘法 - 一个案例研究
- 内核编程 - 内置函数
OpenCL™ 中的“Hello World”程序
以下是关于 OpenCL™ 示例如何编写的一些注意事项
-
OpenCL™ 规定了一个主机 API，该 API 定义为与 C89 兼容，并且不提及 C++ 或其他编程语言绑定。目前，有几项工作正在进行中，以开发其他语言的绑定（参见本文末尾的链接），并且特别地，一直以来都在大力推动开发 C++ 绑定。在本教程及后续教程中，我将仅使用 C++ 绑定，并以此类术语来描述 OpenCL™。有关相应的 C API，请参阅 OpenCL™ 1.0 规范。或者，您可以查看 C++ 绑定的源代码，以了解特定的 C++ 绑定使用了哪个底层 OpenCL™ 函数以及带有何种参数。
- OpenCL™ 定义了一种类 C 语言，用于编程计算设备程序。这些程序通过 API 调用传递给 OpenCL™ 运行时，这些调用期望
char *类型的值。通常，将这些程序保留在单独的源文件中更为方便。对于本教程及后续教程，我假设设备程序存储在文件名为name_kernels.cl的文件中，其中name会根据上下文而变化，但后缀_kernels.cl不会。相应的设备程序在运行时加载并传递给 OpenCL™ API。有许多其他方法可以实现这一点；这里选择这种方法是为了提高可读性。
对于这个第一个 OpenCL™ 程序，我们从主机应用程序的源代码开始。
头文件
与 C++ 中使用的任何其他外部 API 一样，在使用 OpenCL™ API 时必须包含一个头文件。通常，它位于主包含目录内的 CL 目录中。对于我们拥有的 C++ 绑定（将纯 C API 替换为 cl.h）
#include <utility>
#define __NO_STD_VECTOR // Use cl::vector instead of STL version
#include <CL/cl.hpp>
对于我们的程序，我们使用少量额外的 C++ 头文件，这些头文件与 OpenCL™ 无关。
#include <cstdio>
#include <cstdlib>
#include <fstream>
#include <iostream>
#include <string>
#include <iterator>
由于我们将动态请求 OpenCL™ 设备返回“Hello World\n”字符串，因此我们将其定义为一个常量以用于计算。
const std::string hw("Hello World\n");
错误
大多数 OpenCL™ API 调用的一个共同属性是，它们要么将错误代码（cl_int 类型）作为函数本身的结果返回，要么将错误代码存储在用户作为参数传递给调用的位置。与任何可能失败的 API 调用一样，对于应用程序来说，正确检查其在错误情况下的行为非常重要。在大多数情况下，我们不会关心从错误中恢复；为了简单起见，我们定义一个函数 checkErr 来查看某个调用是否已成功完成。在这种情况下，OpenCL™ 返回 CL_SUCCESS 值。如果不是，它会输出一条用户消息并退出；否则，它将简单地返回。
inline void
checkErr(cl_int err, const char * name)
{
if (err != CL_SUCCESS) {
std::cerr << "ERROR: " << name
<< " (" << err << ")" << std::endl;
exit(EXIT_FAILURE);
}
}
C++ 中错误处理的一个常见范例是通过使用异常，OpenCL™ C++ 绑定提供了这样的接口。稍后的教程将介绍异常和其他 C++ 绑定提供的可选功能。现在，让我们看看我们第一个 OpenCL™ 应用程序所需剩余的一个函数，“main”。
OpenCL™ 上下文
初始化和使用 OpenCL™ 的第一步是创建一个上下文。其余的 OpenCL™ 工作（创建设备和内存、编译和运行程序）在此上下文内执行。一个上下文可以有多个关联的设备（例如，CPU 或 GPU 设备），并且在上下文中，OpenCL™ 保证设备之间存在宽松的内存一致性。我们将在稍后的教程中详细介绍这一点；现在，我们使用单个设备 CL_DEVICE_TYPE_CPU，用于 CPU 设备。如果我们假设 OpenCL™ 实现支持该设备，我们也可以使用 CL_DEVICE_TYPE_GPU 或其他支持的设备类型。但在创建上下文之前，我们必须首先将 OpenCL™ 运行时排队，以确定存在哪些平台，即不同供应商的 OpenCL™ 实现。cl::Platform 类提供了静态方法 cl::Platform::get 来实现此目的，并返回一个平台列表。现在我们选择第一个平台并使用它来创建上下文。构造函数 cl::Context 应该成功，在这种情况下，err 的值为 CL_SUCCESS。
int
main(void)
{
cl_int err;
cl::vector< cl::Platform > platformList;
cl::Platform::get(&platformList);
checkErr(platformList.size()!=0 ? CL_SUCCESS : -1, "cl::Platform::get");
std::cerr << "Platform number is: " << platformList.size() << std::endl;
std::string platformVendor;
platformList[0].getInfo((cl_platform_info)CL_PLATFORM_VENDOR, &platformVendor);
std::cerr << "Platform is by: " << platformVendor << "\n";
cl_context_properties cprops[3] =
{CL_CONTEXT_PLATFORM, (cl_context_properties)(platformList[0])(), 0};
cl::Context context(
CL_DEVICE_TYPE_CPU,
cprops,
NULL,
NULL,
&err);
checkErr(err, "Conext::Context()");
在深入研究“真正”工作发生的计算设备之前，我们首先分配一个 OpenCL™ 缓冲区来保存将在设备上运行的内核的结果，即字符串“Hello World\n”。目前，我们仅在主机上分配了一些内存，并要求 OpenCL™ 直接使用此内存，在创建缓冲区时传递标志 CL_MEM_USE_HOST_PTR。
char * outH = new char[hw.length()+1];
cl::Buffer outCL(
context,
CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR,
hw.length()+1,
outH,
&err);
checkErr(err, "Buffer::Buffer()");
OpenCL™ 设备
在 OpenCL™ 中，许多操作是针对给定上下文执行的。例如，缓冲区（内存的一维区域）和图像（内存的二维和三维区域）的分配都是上下文操作。但也有设备特定的操作。例如，程序编译和内核执行是按设备进行的，为此需要特定的设备句柄。那么我们如何获得设备的句柄呢？我们只需从上下文中查询即可。OpenCL™ 提供了将特定对象的有关信息排队的能力，使用 C++ API 时，它以 object.getInfo<CL_OBJECT_QUERY>() 的形式出现。在从上下文中获取设备的特定情况下
cl::vector<cl::Device> devices;
devices = context.getInfo<CL_CONTEXT_DEVICES>();
checkErr(
devices.size() > 0 ? CL_SUCCESS : -1, "devices.size() > 0");
现在我们有了上下文关联的设备列表，在本例中是一个 CPU 设备，我们需要加载并构建计算程序（我们打算在设备或更通用的设备上运行的程序）。以下代码的前几行只是从磁盘加载 OpenCL™ 设备程序，将其转换为字符串，并使用辅助构造函数创建一个 cl::Program::Sources 对象。给定一个 cl::Program::Sources 类型的对象，创建一个 cl::Program 对象并将其与上下文关联，然后为一组特定的 *设备* 构建它。
std::ifstream file("lesson1_kernels.cl");
checkErr(file.is_open() ? CL_SUCCESS:-1, "lesson1_kernel.cl");
std::string prog(
std::istreambuf_iterator<char>(file),
(std::istreambuf_iterator<char>()));
cl::Program::Sources source(
1,
std::make_pair(prog.c_str(), prog.length()+1));
cl::Program program(context, source);
err = program.build(devices,"");
checkErr(file.is_open() ? CL_SUCCESS : -1, "Program::build()");
给定的*程序*可以有许多入口点，称为内核，要调用一个内核，我们必须构建一个内核对象。假定内核名称（表示为字符串）与计算程序中用 __kernel 属性定义的函数之间存在直接映射。在这种情况下，我们可以构建一个 cl::kernel 对象，即 kernel。内核参数使用 C++ API 通过 kernel.setArg() 设置，该函数接受特定参数的索引和值。
cl::Kernel kernel(program, "hello", &err);
checkErr(err, "Kernel::Kernel()");
err = kernel.setArg(0, outCL);
checkErr(err, "Kernel::setArg()");
现在样板代码已完成，是时候计算结果了（包含字符串“Hello World\n”的输出缓冲区）。所有设备计算都使用命令队列完成，命令队列是给定设备的虚拟接口。每个命令队列与给定设备都有一对一的映射；它是使用对 cl::CommandQueue 类构造函数的调用与关联的上下文创建的。给定一个 cl::CommandQueue *queue*，可以使用 queue.enqueuNDRangeKernel 排队*内核*。这会将一个*内核*排队以便在关联的设备上执行。内核可以在 1D、2D 或 3D 索引域上执行，如果资源足够，这些索引会并行执行。启动域中的元素（索引）总数称为 global 工作大小；单个元素称为 work-item。当需要 work-item 之间的通信时，可以将 Work-item 分组到 work-group 中。Work-groups 由子索引函数（称为 local 工作大小）定义，该函数描述了与全局启动域指定的维度相对应的每个维度的大小。内核启动有许多需要考虑的因素，我们将在未来的教程中更详细地介绍。目前，足以注意到对于 Hello World，每个 work-item 都计算结果字符串中的一个字母；并且足以启动 hw.length()+1，其中 hw 是我们在程序开头定义的 const std::string。我们需要额外的 work-item 来处理 NULL 终止符。
cl::CommandQueue queue(context, devices[0], 0, &err);
checkErr(err, "CommandQueue::CommandQueue()");
cl::Event event;
err = queue.enqueueNDRangeKernel(
kernel,
cl::NullRange,
cl::NDRange(hw.length()+1),
cl::NDRange(1, 1),
NULL,
&event);
checkErr(err, "ComamndQueue::enqueueNDRangeKernel()");
上面 enqueueNDRangeKernel 调用的最后一个参数是 cl::Event 对象，可用于查询与之关联的命令的状态（例如，已完成）。它支持 wait() 方法，该方法会阻塞直到命令完成。这对于确保在将结果读回主机内存（使用 queue.enqueueReadBuffer()）之前内核已完成执行是必需的。将计算结果读回主机内存后，只需将结果输出到 stdout 并退出程序。
event.wait();
err = queue.enqueueReadBuffer(
outCL,
CL_TRUE,
0,
hw.length()+1,
outH);
checkErr(err, "ComamndQueue::enqueueReadBuffer()");
std::cout << outH;
return EXIT_SUCCESS;
}
最后，为了使程序完整，lesson1_kernels.cl 设备程序实现需要定义外部入口点 hello。内核实现很简单：它使用 get_global_id() 根据启动域计算一个唯一的索引，将其用作 hw 字符串的索引，然后将其值写入输出数组 out。
#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable
__constant char hw[] = "Hello World\n";
__kernel void hello(__global char * out)
{
size_t tid = get_global_id(0);
out[tid] = hw[tid];
}
为了健壮性，最好检查线程 ID (tid) 是否在 hw 的范围内；目前，我们假设对 queue.enqueueNDRangeKernel() 的相应调用是正确的。
构建和运行
在 Linux 上，使用一个命令构建 OpenCL™ 程序就足够了；例如
gcc –o hello_world –Ipath-OpenCL-include –Lpath-OpenCL-libdir lesson1.cpp –lOpenCL
运行
LD_LIBRARY_PATH=path-OpenCL-libdir ./hello_world
在 Windows 上，使用 Visual Studio 命令窗口，示例如下
cl /Fehello_world.exe /Ipath-OpenCL-include lesson.cpp path-OpenCL-libdir/OpenCL.lib
假设 OpenCL.dll 在路径中，然后运行
.\hello_world
会在 stdout 上输出以下字符串
Hello World
这完成了我们的 OpenCL™ 入门教程。我们请求您的反馈、评论和问题。请访问我们的 Stream 论坛。
有用链接
以下列表提供了除 C 之外的 OpenCL™ 特定编程绑定的链接。我尚未测试过这些绑定，也无法保证其正确性，但希望它们会有所帮助
- OpenCL™ 规范和头文件
http://www.khronos.org/registry/cl/ - OpenCL™ 技术论坛
http://www.khronos.org/message_boards/viewforum.php?f=28 - 本教程中使用的 C++ 绑定可以在 Khronos 的 OpenCL™ 网页上找到，以及完整的文档
http://www.khronos.org/registry/cl/ - Python 绑定可以在这里找到
http://pyopencl.next-touch.com/ - C# 绑定可以在这里找到
http://www.khronos.org/message_boards/viewtopic.php?f=28&t=1932 - OpenCL™ 简介
http://ati.amd.com/technology/streamcomputing/intro_opencl.html