第一部分：OpenCL™ – 可移植的并行性
0/5 (0投票)
对 GPGPU 编程感兴趣？阅读 Rob Farber 的《大规模并行编程》系列。了解如何充分利用您的 CPU、GPU、APU、DSP 等。
本系列的第一篇文章是关于使用 OpenCL™ 进行可移植多线程编程的，简要讨论了该标准的理念，并演示了如何下载和使用 ATI Stream 软件开发工具包 (SDK) 来构建和运行 OpenCL 程序。
OpenCL 是一项跨厂商的标准，它拥有巨大的潜力来利用现代处理器、嵌入式设备和图形处理器 (GPU) 的大规模并行性。由于其广泛的行业支持，OpenCL 有可能成为可移植多核和多线程应用程序的事实上的软件。编写一个可以运行在从嵌入式系统和掌上设备到工作站和超级计算机等各种平台上的单一应用程序的吸引力是不可否认的。
需要注意的一个关键点是，在 OpenCL 中，编译器内置于运行时，这提供了卓越的灵活性和可移植性，因为 OpenCL 应用程序可以在运行时选择和使用系统中的不同 OpenCL 设备。甚至有可能创建如今可以在尚未发明的设备上（无需修改）使用的 OpenCL 应用程序可执行文件！
OpenCL 的挑战主要在于如何同时设计并行程序并编写它们，使其在各种异构硬件平台上都健壮且性能良好。熟练的开发人员是利用多核处理器和多线程 OpenCL 应用程序及硬件投资的关键——尤其是随着标准的快速发展。这些个体能够跟上标准的最新进展，并能够考虑不同厂商 SDK 和实现的差异。
本教程系列专注于帮助熟练的 C 和 C++ 程序员快速上手，以便他们能够使用 OpenCL 编写高效的可移植并行程序。出于实际考虑，本系列将使用运行在多核处理器和 GPU 以及两者的混合异构系统上的示例。随着系列的继续（取决于读者的兴趣），OpenCL 的可移植性将通过运行在各种平台（嵌入式、笔记本电脑、台式机和集群）上使用主要厂商的操作系统的工作示例得到强调。
这些教程旨在教授人们如何用 OpenCL 来思考和编程。单个教程将侧重于传达基本概念、语法和开发工具知识，以便读者能够快速了解如何利用特定的技术功能或发布特性。将提供完整的示例，以最大程度地减少挫败感，并方便复制、粘贴、构建并立即开始使用代码。将至少为一种平台提供明确的构建说明。在可能的情况下，将包含来自不同平台的性能数据，以便开发人员了解可移植性和不同平台性能的能力。
请注意，OpenCL 仍然非常新。在其 OpenCL 概述中，管理 OpenCL 标准的 Khronos Group 指出，Apple Computer 最初于 2008 年 6 月提出了应该有一个 OpenCL 标准。第一个 v1.0 厂商实现于 2009 年下半年开始出货。v1.1 标准于 2010 年 6 月发布，同时发布了兼容的厂商实现。
OpenCL 背后的理念
OpenCL 的核心理念是一个可移植的执行模型，它允许一个内核在问题域的每个点上执行。内核是在程序中声明并在 OpenCL 设备上执行的函数。它由应用于程序中定义的任何函数的 __kernel 限定符标识。内核可以以数据并行或任务并行的方式运行。
本质上，OpenCL 开发人员可以想象以数据并行的方式使用 GPU SIMD（单指令多数据）或处理器 SSE（流式 SIMD 扩展）指令处理向量或其他数据结构，或者以任务并行的方式允许多个不同的任务驻留在处理器核心或 MIMD（多指令多数据）架构上。可以通过 函数限定符（如 vec_type_hint() 或 work_group_size_hint()）提供 OpenCL 编译器提示。
- 当独立线程可以处理不同的函数时，选择任务并行模型。任务级并发要求将封装在函数中的独立工作映射到单个线程，这些线程异步执行。
- 对于计算密集型循环，其中重复执行相同、独立的运算，请选择数据并行线程模型。数据并行意味着对不同数据重复应用相同的独立运算。
OpenCL 应用程序运行在主机上，主机通过队列将工作提交给计算设备。此模型隐含着主机和其一个或多个 OpenCL 设备之间的某种形式的数据传输。应用程序按顺序将内核执行实例排队，每个设备一个队列。然而，设备上既可以进行按顺序执行，也可以进行乱序执行。
以下是 OpenCL 的一些核心术语
- 工作项：OpenCL 设备上的基本工作单元。
- 内核：工作项的代码，基本上是一个 C 函数。
- 程序：内核和其他函数的集合。
- 上下文：工作项执行的环境，包括设备及其内存和命令队列。
虽然 OpenCL 应用程序本身可以用 C 或 C++ 编写，但应用程序内核的源代码是用 ISO C99 C 语言规范的变体编写的。这些内核通过内置的运行时编译器进行编译，或者如果需要，可以保存以便稍后加载。OpenCL C 语言用于内核的
- ISO C99 标准的一个子集，它消除了某些特性，如头文件、函数指针、递归、可变长度数组和位字段。
- ISO C99 标准的一个超集，增加了
- 工作项和工作组
- 向量类型
- 同步。
- 地址空间限定符。
- 它还包括一组丰富的内置函数，以方便 OpenCL 的功能，例如
- 图像处理。
- 工作项处理。
- 专用数学例程和其他操作。
其他信息来源
- ATI Stream SDK v2：下载、文档、OpenCL Zone 和 论坛。
- The Khronos Group：开发者资源、手册页、规范和 论坛。
- Apple Computer
- IBM
- NVIDIA
- Intel 将在其 OpenCL SDK 在年底前提供。
安装 ATI Stream SDK v2
ATI Stream SDK v2 的完整下载和安装说明可以在 此处找到。以下是简要概述
- 下载适用于您的操作系统的 ATI Stream SDK（Linux 或 Windows）。在本文中，我们使用适用于 64 位 Linux 的最新版本 ATI Stream SDK v2.2。Linux 用户还必须安装 ICD 信息，这允许多个供应商之间的跨平台支持正常工作。
- 解压目录和相关文件。在 Linux 上使用 tar。Microsoft 用户将运行安装可执行文件。在 Linux 上
- mkdir AMD
- cd AMD
- tar -xzf ati-stream-sdk-v2.2-lnx64.tgz
- 下载并安装 ICD 信息。对于 Linux，这是文件 icd-registration.tgz。
- 以 root 用户身份，切换到根目录并解压
- (cd /; tar –xzf icd-registration.tgz)
- 设置适当的环境变量并构建示例。在 Linux 上，
- export ATISTREAMSDKROOT=$HOME/AMD/ati-stream-sdk-v2.2-lnx64
- export ATISTREAMSDKSAMPLESROOT=/$HOME /AMD/ati-stream-sdk-v2.2-lnx64
- export LD_LIBRARY_PATH=$ATISTREAMSDKROOT/lib/x86:$ATISTREAMSDKROOT/lib/x86_64:$LD_LIBRARY_PATH
- cd ati-stream-sdk-v2.2-lnx64/
- make
示例构建完成后，可以通过在示例目录中运行 CLIinfo 应用程序来查看可用的设备：./samples/opencl/bin/x86_64/CLInfo。
将出现类似以下内容的输出，表明该系统同时拥有 CPU 和 GPU 设备
Number
of platforms: 1
Platform Profile: FULL_PROFILE
Platform Version: OpenCL 1.1 ATI-Stream-v2.2 (302)
Platform Name: ATI Stream
Platform Vendor: Advanced Micro Devices, Inc.
Platform Extensions: cl_khr_icd cl_amd_event_callback
Platform Name: ATI Stream
Number
of devices: 2
Device Type: CL_DEVICE_TYPE_CPU
Device ID: 4098
Max compute units: 6
Max work items dimensions: 3
Max work items[0]: 1024
Max work items[1]: 1024
Max work items[2]: 1024
Max work group size: 1024
Preferred vector width char: 16
Preferred vector width short: 8
Preferred vector width int: 4
Preferred vector width long: 2
Preferred vector width float: 4
Preferred vector width double: 0
Max clock frequency: 800Mhz
Address bits: 64
Max memory allocation: 1073741824
Image support: No
Max size of kernel argument: 4096
Alignment (bits) of base address: 1024
Minimum alignment (bytes) for any datatype: 128
Single precision floating point capability
Denorms: Yes
Quiet NaNs: Yes
Round to nearest even: Yes
Round to zero: Yes
Round to +ve and infinity: Yes
IEEE754-2008 fused multiply-add: No
Cache type: Read/Write
Cache line size: 64
Cache size: 65536
Global memory size: 3221225472
Constant buffer size: 65536
Max number of constant args: 8
Local memory type: Global
Local memory size: 32768
Profiling timer resolution: 1
Device endianess: Little
Available: Yes
Compiler available: Yes
Execution capabilities:
Execute OpenCL kernels: Yes
Execute native function: Yes
Queue properties:
Out-of-Order: No
Profiling : Yes
Platform ID: 0x7f47e30e2b20
Name: AMD Phenom(tm) II X6 1055T Processor
Vendor: AuthenticAMD
Driver version: 2.0
Profile: FULL_PROFILE
Version: OpenCL 1.1 ATI-Stream-v2.2 (302)
Extensions: cl_amd_fp64
cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics
cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics
cl_khr_int64_base_atomics cl_khr_int64_extended_atomics
cl_khr_byte_addressable_store cl_khr_gl_sharing cl_ext_device_fission
cl_amd_device_attribute_query cl_amd_printf
Device Type: CL_DEVICE_TYPE_GPU
Device ID: 4098
Max compute units: 20
Max work items dimensions: 3
Max work items[0]: 256
Max work items[1]: 256
Max work items[2]: 256
Max work group size: 256
Preferred vector width char: 16
Preferred vector width short: 8
Preferred vector width int: 4
Preferred vector width long: 2
Preferred vector width float: 4
Preferred vector width double: 0
Max clock frequency: 850Mhz
Address bits: 32
Max memory allocation: 134217728
Image support: Yes
Max number of images read arguments: 128
Max number of images write arguments: 8
Max image 2D width: 8192
Max image 2D height: 8192
Max image 3D width: 2048
Max image 3D height: 2048
Max image 3D depth: 2048
Max samplers within kernel: 16
Max size of kernel argument: 1024
Alignment (bits) of base address: 32768
Minimum alignment (bytes) for any datatype: 128
Single precision floating point capability
Denorms: No
Quiet NaNs: Yes
Round to nearest even: Yes
Round to zero: Yes
Round to +ve and infinity: Yes
IEEE754-2008 fused multiply-add: Yes
Cache type: None
Cache line size: 0
Cache size: 0
Global memory size: 536870912
Constant buffer size: 65536
Max number of constant args: 8
Local memory type: Scratchpad
Local memory size: 32768
Profiling timer resolution: 1
Device endianess: Little
Available: Yes
Compiler available: Yes
Execution capabilities:
Execute OpenCL kernels: Yes
Execute native function: No
Queue properties:
Out-of-Order: No
Profiling : Yes
Platform ID: 0x7f47e30e2b20
Name: Cypress
Vendor: Advanced Micro Devices, Inc.
Driver version: CAL 1.4.736
Profile: FULL_PROFILE
Version: OpenCL 1.1 ATI-Stream-v2.2 (302)
Extensions: cl_amd_fp64
cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics
cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics
cl_khr_3d_image_writes cl_khr_byte_addressable_store cl_khr_gl_sharing
cl_amd_device_attribute_query cl_amd_printf cl_amd_media_ops
Passed!
第一个应用程序和 OpenCL 内核
以下应用程序 first.cpp 是用 C++ 编写的。但是，它的编写方式非常直接的过程式，以说明创建可以在 GPU 或 CPU 上工作的 OpenCL 应用程序所需的步骤。该应用程序只是创建一个随机值向量，然后将其传输到 OpenCL 内核，内核将这些值平方。然后主机检索这些值并仔细检查结果。
要使用 OpenCL，开发人员必须
- 定义平台。
- 在平台上执行代码。
- 在内存中移动数据。
- 编写（并构建）程序。在此应用程序中，使用了 clCreateProgramWithSource() 来构建编译 OpenCL 内核。
以下是应用程序中包含的实际 OpenCL 内核。请注意
- 内核源代码包含在一个常量字符字符串中。商业开发人员应考虑使用 strings 等命令，常量字符串很可能在可执行映像中被其他人看到。AMD 提供了一篇知识库文章，解释如何使用 二进制内核生成来解决此问题，
- 这是一个并行内核。每个线程通过调用 get_global_id() 来获取其 ID，该 ID 用作向量的索引。
const char *KernelSource = "\n" \
"__kernel void square( \n" \
" __global float* input, \n" \
" __global float* output, \n" \
" const unsigned int count) \n" \
"{ \n" \
" int i = get_global_id(0); \n" \
" if(i < count) \n" \
" output[i] = input[i] * input[i]; \n" \
"} \n" \
"\n";
相比之下，串行函数看起来会像下面这样
void SerialSource(int n, float* input, float* output)
{
for (int i=0; i<n; i++)
output[i] = input[i] * input[i];
}
以下是 first.cpp 的完整源代码
#include <iostream>
using namespace std;
#define __NO_STD_VECTOR // Use cl::vector and cl::string and
#define __NO_STD_STRING // not STL versions, more on this later
#include <CL/cl.h>
#define DATA_SIZE (1024*1240)
const char *KernelSource = "\n" \
"__kernel void square( \n" \
" __global float* input, \n" \
" __global float* output, \n" \
" const unsigned int count) \n" \
"{ \n" \
" int i = get_global_id(0); \n" \
" if(i < count) \n" \
" output[i] = input[i] * input[i]; \n" \
"} \n" \
"\n";
int main(int argc, char* argv[])
{
int devType=CL_DEVICE_TYPE_GPU;
if(argc > 1) {
devType = CL_DEVICE_TYPE_CPU;
cout << "Using: CL_DEVICE_TYPE_CPU" << endl;
} else {
cout << "Using: CL_DEVICE_TYPE_GPU" << endl;
}
cl_int err; // error code returned from api calls
size_t global; // global domain size for our calculation
size_t local; // local domain size for our calculation
cl_platform_id cpPlatform; // OpenCL platform
cl_device_id device_id; // compute device id
cl_context context; // compute context
cl_command_queue commands; // compute command queue
cl_program program; // compute program
cl_kernel kernel; // compute kernel
// Connect to a compute device
err = clGetPlatformIDs(1, &cpPlatform, NULL);
if (err != CL_SUCCESS) {
cerr << "Error: Failed to find a platform!" << endl;
return EXIT_FAILURE;
}
// Get a device of the appropriate type
err = clGetDeviceIDs(cpPlatform, devType, 1, &device_id, NULL);
if (err != CL_SUCCESS) {
cerr << "Error: Failed to create a device group!" << endl;
return EXIT_FAILURE;
}
// Create a compute context
context = clCreateContext(0, 1, &device_id, NULL, NULL, &err);
if (!context) {
cerr << "Error: Failed to create a compute context!" << endl;
return EXIT_FAILURE;
}
// Create a command commands
commands = clCreateCommandQueue(context, device_id, 0, &err);
if (!commands) {
cerr << "Error: Failed to create a command commands!" << endl;
return EXIT_FAILURE;
}
// Create the compute program from the source buffer
program = clCreateProgramWithSource(context, 1,
(const char **) &KernelSource,
NULL, &err);
if (!program) {
cerr << "Error: Failed to create compute program!" << endl;
return EXIT_FAILURE;
}
// Build the program executable
err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
if (err != CL_SUCCESS) {
size_t len;
char buffer[2048];
cerr << "Error: Failed to build program executable!" << endl;
clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG,
sizeof(buffer), buffer, &len);
cerr << buffer << endl;
exit(1);
}
// Create the compute kernel in the program
kernel = clCreateKernel(program, "square", &err);
if (!kernel || err != CL_SUCCESS) {
cerr << "Error: Failed to create compute kernel!" << endl;
exit(1);
}
// create data for the run
float* data = new float[DATA_SIZE]; // original data set given to device
float* results = new float[DATA_SIZE]; // results returned from device
unsigned int correct; // number of correct results returned
cl_mem input; // device memory used for the input array
cl_mem output; // device memory used for the output array
// Fill the vector with random float values
unsigned int count = DATA_SIZE;
for(int i = 0; i < count; i++)
data[i] = rand() / (float)RAND_MAX;
// Create the device memory vectors
//
input = clCreateBuffer(context, CL_MEM_READ_ONLY,
sizeof(float) * count, NULL, NULL);
output = clCreateBuffer(context, CL_MEM_WRITE_ONLY,
sizeof(float) * count, NULL, NULL);
if (!input || !output) {
cerr << "Error: Failed to allocate device memory!" << endl;
exit(1);
}
// Transfer the input vector into device memory
err = clEnqueueWriteBuffer(commands, input,
CL_TRUE, 0, sizeof(float) * count,
data, 0, NULL, NULL);
if (err != CL_SUCCESS) {
cerr << "Error: Failed to write to source array!" << endl;
exit(1);
}
// Set the arguments to the compute kernel
err = 0;
err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &input);
err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &output);
err |= clSetKernelArg(kernel, 2, sizeof(unsigned int), &count);
if (err != CL_SUCCESS) {
cerr << "Error: Failed to set kernel arguments! " << err << endl;
exit(1);
}
// Get the maximum work group size for executing the kernel on the device
err = clGetKernelWorkGroupInfo(kernel, device_id,
CL_KERNEL_WORK_GROUP_SIZE,
sizeof(local), &local, NULL);
if (err != CL_SUCCESS) {
cerr << "Error: Failed to retrieve kernel work group info! "
<< err << endl;
exit(1);
}
// Execute the kernel over the vector using the
// maximum number of work group items for this device
global = count;
err = clEnqueueNDRangeKernel(commands, kernel,
1, NULL, &global, &local,
0, NULL, NULL);
if (err) {
cerr << "Error: Failed to execute kernel!" << endl;
return EXIT_FAILURE;
}
// Wait for all commands to complete
clFinish(commands);
// Read back the results from the device to verify the output
//
err = clEnqueueReadBuffer( commands, output,
CL_TRUE, 0, sizeof(float) * count,
results, 0, NULL, NULL );
if (err != CL_SUCCESS) {
cerr << "Error: Failed to read output array! " << err << endl;
exit(1);
}
// Validate our results
//
correct = 0;
for(int i = 0; i < count; i++) {
if(results[i] == data[i] * data[i])
correct++;
}
// Print a brief summary detailing the results
cout << "Computed " << correct << "/" << count << " correct values" << endl;
cout << "Computed " << 100.f * (float)correct/(float)count
<< "% correct values" << endl;
// Shutdown and cleanup
delete [] data; delete [] results;
clReleaseMemObject(input);
clReleaseMemObject(output);
clReleaseProgram(program);
clReleaseKernel(kernel);
clReleaseCommandQueue(commands);
clReleaseContext(context);
return 0;
}
first.cpp 如何工作的细节超出了本文的初始介绍范围。乍一看，OpenCL 代码对于仅执行平方向量中值的简单任务而言显得相当冗长。然而，这是一个 CPU 和 GPU 应用程序，可以在许多不同的设备类型上运行，如下所示。它不受厂商限制，并有可能在未来的设备上保持不变地运行。这种可移植性以及利用大规模并行硬件架构的能力是 OpenCL 的优势。
在 Linux 上构建可执行文件很简单
- 复制源代码并将其放入名为 first.cpp 的文件中。
- 设置 OpenCL 主目录的环境变量
OCL_HOME=../ati-stream-sdk-v2.2-lnx64
- g++ -I $OCL_HOME/include -L $OCL_HOME/lib/x86_64 first.cpp -l OpenCL
以下显示了这段代码在 CPU 和 GPU 设备上的运行情况
bda:~/AMD/test$ time ./a.out Using: CL_DEVICE_TYPE_GPU Computed 1269760/1269760 correct values Computed 100% correct values real 0m0.354s user 0m0.260s sys 0m0.150s bda:~/AMD/test$ time ./a.out CPU Using: CL_DEVICE_TYPE_CPU Computed 1269760/1269760 correct values Computed 100% correct values real 0m0.261s user 0m0.200s sys 0m0.040s
请注意，GPU 花费的时间比 CPU 多。原因是将数据从主机传输到 GPU 设备花费了时间，而 CPU 能够立即开始执行计算。经验表明，高性能 GPU 编程需要三个步骤
- 获取数据并将其保留在 GPU 上，以消除 PCI 总线数据传输瓶颈。
- 给 GPU 足够的工作量。启动内核确实需要少量开销。然而，现代 GPU 速度如此之快，以至于它们可以在内核启动期间执行大量工作。因此，内核启动被排队到设备上。
- 优化计算以最大程度地减少访问 GPU 内存的瓶颈。同样，GPU 硬件速度如此之快，以至于在计算硬件本地的内存（例如，寄存器等）中重用数据非常重要，以防止计算受到 GPU 内存系统的瓶颈。
OpenCL 示例程序 PCIeBandwidth 和 GlobalMemoryBandwidth 的输出显示了在 Linux Ubuntu 10.04 下，PCIe 总线与 ATI Radeon™ HD 5800 系列图形处理器全局内存带宽之间的相对速度差异。
- ./samples/opencl/bin/x86/PCIeBandwidth
Host to device : 2.44032 GB/s Device to host : 1.26776 GB/s
- ./samples/opencl/bin/x86/GlobalMemoryBandwidth
Global Memory Read AccessType : single VectorElements : 4 Bandwidth : 169.918 GB/s Global Memory Read AccessType : linear VectorElements : 4 Bandwidth : 154.875 GB/s Global Memory Read AccessType : linear(uncached) VectorElements : 4 Bandwidth : 118.425 GB/s Global Memory Write AccessType : linear VectorElements : 4 Bandwidth : 177.312 GB/s
C++ 程序员会注意到示例代码采用了非常过程化的方法。这是故意的，目的是为了增强为 C 和 C++ 开发人员创建 OpenCL 应用程序所需的步骤。ATI Stream SDK 示例包含一个简单的 Template.cpp 示例，该示例使用了更接近 C++ 的方法。此示例位于
ati-stream-sdk-v2.2-lnx64/samples/opencl/cl/app/Template
在此目录中键入 make 将构建示例，然后可以通过键入
build/debug/x86_64/Template 来启动它。Template_Kernels.cl 文件可以修改以执行您自己的计算。该代码的文档位于 docs 目录中。虽然是一个好的开始，但这些 OpenCL 示例在拥有通用 C++ 数据并行项目的简洁性和强大功能之前，还有一段路要走。现有项目表明，可以使用类似 STL 的向量类来创建简洁易读的数据并行应用程序。尝试使用此示例代码进行实验，以创建易于使用但又通用的模板。
摘要
架构和平衡比是理解 OpenCL 设备性能的关键概念。特别是，主机和 OpenCL 设备之间的链路带宽，以及设备的内存带宽，都可以成为关键性能指标。在某些情况下，这些硬件特性中的任何一个都可能使将计算从主机迁移到 OpenCL 设备变得过于昂贵。然而，许多问题要求每个传输的数据项都有足够的计算量才能极大地加速 OpenCL 应用程序。有关机器特性或平衡比如何定义应用程序性能的更多信息，我建议阅读我在Scientific Computing 上的入门文章 HPC Balance and Common Sense 以及广泛的 Atkins Report。
最后，本教程提供了 OpenCL 的快速基本介绍以及一个可以在 CPU 和 GPU 设备类型上运行的示例代码。试试看它如何工作。
OpenCL 和 OpenCL 徽标是 Apple Inc. 的商标，经 Khronos 许可使用。