第九部分：OpenCL 扩展和设备分裂
4.75/5 (4投票s)
本文讨论了 OpenCL 扩展，
本系列关于OpenCL™便携式并行性的上一篇文章（第8部分）演示了如何通过通用“即插即用工具”框架将OpenCL™整合到异构工作流中。该框架可以在单个工作站、机器网络或云计算框架内流式传输任意消息（向量、数组以及任意复杂嵌套结构）。创建可扩展工作流的能力至关重要，因为数据处理和转换可能与用于生成所需结果的计算问题一样复杂和耗时。
本文讨论了OpenCL扩展，这些扩展为程序员提供了额外的功能，例如双精度算术和_设备分割_。（设备分割提供了一个接口，可以将单个OpenCL设备细分为多个设备——每个设备都有一个单独的异步命令队列。）
OpenCL扩展可以由供应商、OpenCL工作组的子集或整个OpenCL工作组定义。最便携的扩展是那些由整个OpenCL工作组正式批准的KHR扩展，而供应商扩展的便携性最差，并且可能与特定设备或产品线绑定。无论由谁提供定义，都不能保证任何平台上都提供某个扩展。
以下是OpenCL扩展的三种类型和命名约定
- KHR扩展：KHR扩展由OpenCL工作组正式批准，并附带一套符合性测试，以帮助确保一致的行为。提供KHR扩展是为了支持某些OpenCL设备上可用但并非所有设备都可用的功能。Microsoft DirectX扩展是仅在支持Microsoft Windows的设备上可用的重要功能的一个示例。KHR扩展具有形式为_cl_khr_<name>_的唯一名称。
- EXT扩展：EXT扩展由一个或多个OpenCL工作组成员开发。不需要符合性测试。可以合理地将这些视为“进行中的工作”扩展，用于在正式批准为KHR扩展之前评估可用性、价值和便携性。EXT扩展具有形式为_cl_ext_<name>_的唯一名称。
- 供应商扩展：这些扩展由供应商提供，用于公开特定于供应商设备或产品线的功能。供应商扩展应被视为高度不可移植。AMD设备属性查询是提供有关AMD设备的额外信息的一个示例。供应商扩展被分配形式为_cl_<vendor>_<name>_的唯一名称。因此，AMD扩展将具有名称字符串_cl_amd_<name>_。
指令_#pragma OPENCL EXTENSION_控制OpenCL编译器的行为，以允许或不允许扩展。例如，本系列的第4部分通过以下行在AMD设备上启用了双精度计算。（注意，在AMD SDK 2.6版本中，_cl_amd_fp64_可以更新为_cl_khr_fp64_。）
#pragma OPENCL EXTENSION cl_amd_fp64 : enable
扩展pragma的语法是
#pragma OPENCL EXTENSION <extention_name> : <behavior>
<behavior>令牌可以是以下之一
- 启用：如果支持，则启用扩展；如果不支持指定的扩展或使用令牌“all”，则报告错误。
- 禁用：OpenCL实现/编译器的行为就像指定的扩展不存在一样。
- 全部：仅使用和支持OpenCL的核心功能，所有扩展都被忽略。如果不支持指定的扩展，则编译器会发出警告。
默认情况下，编译器要求所有扩展都必须显式启用，就好像它已经提供了以下pragma一样
#pragma OPENCL EXTENSION all : disable
2011年12月版本的“AMD加速并行处理OpenCL”指南列出了以下KHR扩展的可用性
- cl_khr_global_int32_base_atomics：全局内存中32位整数的基本原子操作。
- cl_khr_global_int32_extended_atomics：全局内存中32位整数的扩展原子操作。
- cl_khr_local_int32_base_atomics：本地内存中32位整数的基本原子操作。
- cl_khr_local_int32_extended_atomics：本地内存中32位整数的扩展原子操作。
- cl_khr_int64_base_atomics：全局内存和本地内存中64位整数的基本原子操作。
- cl_khr_int64_extended_atomics：全局内存和本地内存中64位整数的扩展原子操作。
- cl_khr_3d_image_writes：支持内核写入3D图像。
- cl_khr_byte_addressable_store：这消除了在内核程序中不允许写入宽度小于32位的指针（或数组元素）的限制。
- cl_khr_gl_sharing：允许将OpenGL上下文或共享组与CL上下文关联以实现互操作性。
- cl_khr_icd：OpenCL可安装客户端驱动程序（ICD），允许开发人员从系统中可能安装的多个OpenCL运行时中进行选择。（自SDK v2 for AMD Accelerated Parallel Processing起，此扩展自动启用。）
- cl_khr_d3d10_sharing：允许将D3D10上下文或共享组与CL上下文关联以实现互操作性。
Khronos OpenCL API注册表的1.2版本列出了以下扩展的可用性。可以单击超链接或访问Khronos OpenCL工作组文档“OpenCL扩展规范”以查找有关各个1.2版本扩展的更详细信息。
- cl_nv_d3d9_sharing
- cl_nv_d3d10_sharing
- cl_nv_d3d11_sharing
- cl_khr_icd
- cl_khr_d3d10_sharing
- cl_amd_device_attribute_query
- cl_amd_fp64：（注意：cl_khr_fp64在当前的AMD版本中有效。）
- cl_amd_media_ops
- cl_ext_migrate_memobject
- cl_ext_device_fission
- cl_ext_atomic_counters_32
- cl_ext_atomic_counters_64
-
cl_intel_dx9_media_sharing
AMD还提供了以下其他供应商特定扩展，如2011年12月版本的“AMD加速并行处理OpenCL”指南附录A.8中所述。
- cl_amd_event_callback
- cl_amd_media_ops
- cl_amd_printf
设备分割
默认情况下，每个OpenCL内核都尝试根据_数据并行_计算模型使用设备上的所有计算资源。换句话说，相同的内核用于处理设备上所有计算资源上的数据。相比之下，_任务并行_模型使用可用的计算资源在同一设备上运行一个或多个独立的内核。任务并行和数据并行都是构建代码以加速应用程序性能的有效方法，因为某些问题更适合用任务并行解决，而另一些则更适合用数据并行解决。一般来说，从软件和硬件角度来看，有效地实现任务并行更为复杂。OpenCL默认行为是根据数据并行模型使用所有可用的计算资源，这是一个很好的选择，因为它将为单个内核提供最大的加速。
在AMD平台上，有两种方法可以限制在多核处理器上运行内核时使用的核心数量。
- AMD OpenCL运行时检查环境变量**CPU_MAX_COMPUTE_UNITS**。如果定义了该变量，AMD运行时会将OpenCL应用程序使用的处理器核心数量限制为该变量指定的值。只需将此环境变量设置为一个介于1到系统中多核处理器总数之间的数字。注意：此变量不会影响GPU等其他设备，也不能保证与所有供应商运行时一起使用。
- EXT设备分割扩展_cl_ext_device_fission_提供了一个OpenCL内部接口，用于将设备细分为多个子设备。程序员随后可以在每个子设备上创建一个命令队列，并排队内核，这些内核仅在子设备内的资源（例如处理器核心）上运行。每个子设备与其他子设备异步运行。目前，设备分割仅适用于多核处理器（AMD和Intel）和Cell宽带引擎。不支持GPU。
（注意：可以限制工作组和工作项的数量，因此OpenCL内核只使用多核处理器的几个核心，然后依靠操作系统有效地调度多个应用程序运行。由于多种原因，不推荐使用此方法，其中包括它实际上硬编码了对资源的刻意浪费使用。此外，此技巧依赖于外部因素（如操作系统）来实现高效操作。此外，此技巧不适用于GPU。）
Ben Gaster的题为“OpenCL的设备分割扩展”的网络研讨会幻灯片讨论了容器并行管道中的设备分割。他指出，用户希望细分设备时有三种通用用例
- 保留设备的一部分用于高优先级/延迟敏感任务。
- 更直接地控制工作到单个计算单元的分配。
- 沿着共享硬件特征（如缓存）细分计算设备。
通常，这些用例需要一定程度的额外控制才能获得超越标准 OpenCL 1.1 API 所提供的最佳性能。正确使用此接口需要对设备有一些详细了解。
AMD SDK 示例提供了一个在多核处理器上使用设备分割的示例。在标准安装中，此示例位于 /opt/AMDAPP/samples/cl/app/DeviceFission。Ben Gaster 还在其 2011 年 3 月向 Khronos Group 提交的演示文稿“OpenCL 设备分割”中提供了一些关于利用设备分割所需基础知识的精彩幻灯片。
OpenCL 即插即用工具框架中的设备分割
正如本教程第8部分所述，数据预处理可能与生成所需结果的实际计算一样复杂和耗时。一个“即插即用”框架（如下图所示，并在第8部分中更详细地讨论）自然地利用了多核处理器的并行性，因为管道中的每个元素都是一个独立的应用程序。操作系统调度程序确保任何具有执行工作所需数据的应用程序都将运行——通常在单独的处理器核心上。在某些情况下，希望对工作流进行分区，以便多个即插即用OpenCL应用程序可以在单独的核心上运行而互不干扰。也许任务对延迟敏感，或者开发人员希望在UNIX下使用numactl等命令将应用程序绑定到特定的处理核心，以更好地利用缓存。
第8部分中的_dynOCL.cc_的以下源代码已修改为使用设备分割。更改已用颜色突出显示。简而言之，更改包括
- 使用C++定义来启用C++设备分割绑定。
- 检查设备是否支持cl_ext_device_fission扩展。
- 细分设备。
//Rob Farber
#include <cstdlib>
#include <sys/types.h>
#include <dlfcn.h>
#include <string>
#include <iostream>
#include "packetheader.h"
#define PROFILING // Define to see the time the kernel takes
#define __NO_STD_VECTOR // Use cl::vector instead of STL version
#define __CL_ENABLE_EXCEPTIONS // needed for exceptions
// 1. enable the C++ bindings for Device Fission
#define USE_CL_DEVICE_FISSION 1
#include <CL/cl.hpp>
#include <fstream>
using namespace std;
void *lib_handle;
typedef char* (*initFini_t)(const char*, const char*, uint32_t*, uint32_t*);
typedef char* (*func_t)(const char*, const char*, uint32_t*, uint32_t*, char*);
typedef void (*dynFree_t)(char*);
typedef void (*oclSetup_t)(const char*, cl::CommandQueue*);
int main(int argc, char **argv)
{
if(argc < 3) {
cerr << "Use: sourcefilename cpu|gpu oclSource" << endl;
return -1;
}
string base_filename(argv[1]);
base_filename = base_filename.substr(0,base_filename.find_last_of("."));
// build the shared object or dll
string buildCommand("gcc -fPIC -shared -I $ATISTREAMSDKROOT/include ");
buildCommand += string(argv[1])
+ string(" -o ") + base_filename + string(".so ");
cerr << "Compiling with \"" << buildCommand << "\"" << endl;
if(system(buildCommand.c_str())) {
cerr << "compile command failed!" << endl;
cerr << "Build command " << buildCommand << endl;
return -1;
}
// load the library -------------------------------------------------
string nameOfLibToLoad("./");
nameOfLibToLoad += base_filename;
nameOfLibToLoad += ".so";
lib_handle = dlopen(nameOfLibToLoad.c_str(), RTLD_LAZY);
if (!lib_handle) {
cerr << "Cannot load library: " << dlerror() << endl;
return -1;
}
// load the symbols -------------------------------------------------
initFini_t dynamicInit= NULL;
func_t dynamicFunc= NULL;
initFini_t dynamicFini= NULL;
dynFree_t dynamicFree= NULL;
// reset errors
dlerror();
// load the function pointers
dynamicFunc= (func_t) dlsym(lib_handle, "func");
const char* dlsym_error = dlerror();
if (dlsym_error) { cerr << "sym load: " << dlsym_error << endl; return -1;}
dynamicInit= (initFini_t) dlsym(lib_handle, "init");
dlsym_error = dlerror();
if (dlsym_error) { cerr << "sym load: " << dlsym_error << endl; return -1;}
dynamicFini= (initFini_t) dlsym(lib_handle, "fini");
dlsym_error = dlerror();
if (dlsym_error) { cerr << "sym load: " << dlsym_error << endl; return -1;}
dynamicFree= (dynFree_t) dlsym(lib_handle, "dynFree");
dlsym_error = dlerror();
if (dlsym_error) { cerr << "sym load: " << dlsym_error << endl; return -1;}
// add a function to specify the ocl context and kernel file
oclSetup_t oclSetupFunc;
oclSetupFunc = (oclSetup_t) dlsym(lib_handle, "oclSetup");
dlsym_error = dlerror();
if (dlsym_error) { cerr << "sym load: " << dlsym_error << endl; return -1;}
// --------------------------------------------------------------
// Setup OCL context
//
const string platformName(argv[2]);
const char* oclKernelFile = argv[3];
int ret= -1;
cl::vector<int> deviceType;
cl::vector< cl::CommandQueue > contextQueues;
// crudely parse the command line arguments.
if(platformName.compare("cpu")==0)
deviceType.push_back(CL_DEVICE_TYPE_CPU);
else if(platformName.compare("gpu")==0)
deviceType.push_back(CL_DEVICE_TYPE_GPU);
else { cerr << "Invalid device type!" << endl; return(1); }
// create the context and queues
try {
cl::vector< cl::Platform > platformList;
cl::Platform::get(&platformList);
// Get all the appropriate devices for the platform the
// implementation thinks we should be using.
// find the user-specified devices
cl::vector<cl::Device> devices;
for(int i=0; i < deviceType.size(); i++) {
cl::vector<cl::Device> dev;
platformList[0].getDevices(deviceType[i], &dev);
// 2. check if the device supports Device Fission
for(int j=0; j < dev.size(); j++) {
if(dev[j].getInfo<CL_DEVICE_EXTENSIONS>().
find("cl_ext_device_fission") == std::string::npos) {
cerr << "Device Fission NOT on device" << endl;
return(-1);
} else
cerr << "Have DEVICE_FISSION" << endl;
}
for(int j=0; j < dev.size(); j++) devices.push_back(dev[j]);
}
// set a single context
cl_context_properties cprops[] = {CL_CONTEXT_PLATFORM, NULL, 0};
cl::Context context(devices, cprops);
cerr << "Using the following device(s) in one context" << endl;
for(int i=0; i < devices.size(); i++) {
cerr << " " << devices[i].getInfo<CL_DEVICE_NAME>() << endl;
}
// Create the separate command queues to perform work
// 3. Device Fission : equally sub-divide the device
cl_device_partition_property_ext props[] = {
CL_DEVICE_PARTITION_EQUALLY_EXT,
1,
CL_PROPERTIES_LIST_END_EXT,
0
};
cl::vector<cl::Device> sdevices;
devices[0].createSubDevices(props, &sdevices);
cerr << "Sub-divided into " << sdevices.size() << " devices" << endl;
for(int i=0; i < sdevices.size(); i++) {
#ifdef PROFILING
cl::CommandQueue queue(context, sdevices[i],CL_QUEUE_PROFILING_ENABLE);
#else
cl::CommandQueue queue(context, sdevices[i],0);
#endif
contextQueues.push_back( queue );
}
} catch (cl::Error error) {
cerr << "caught exception: " << error.what()
<< '(' << error.err() << ')' << endl;
return(-1);
}
oclSetupFunc(oclKernelFile, &contextQueues[0]);
// --------------------------------------------------------------
// work with protobufs
//
//enable C++ binary cin and cout
if (!setPacket_binaryIO()) {
cerr << "Cannot set binary mode for cin and cout!" << endl;
return -1;
}
uint32_t size, type;
char *retBlob;
// handle initialization and put information on output stream when told
if( (retBlob=(*dynamicInit)(argv[0], base_filename.c_str(),&size, &type)) ) {
writePacketHdr(size, type, &std::cout);
cout.write(retBlob, size);
(dynamicFree)(retBlob);
}
// read stream from cin and put information on output stream when told
while(readPacketHdr(&size, &type, &std::cin)) {
char *blob = new char[size];
cin.read(blob, size);
retBlob =(*dynamicFunc)(argv[0], base_filename.c_str(), &size, &type, blob);
if(retBlob) {
writePacketHdr(size, type, &std::cout);
cout.write(retBlob, size);
// optimization: if retBlob == blob then allocated was by this program
if(retBlob != blob) (dynamicFree)(retBlob);
}
delete [] blob;
}
// handle finalization (fini) and put information on output stream when told
if( retBlob = (*dynamicFini)(argv[0], base_filename.c_str(),&size, &type) ) {
writePacketHdr(size, type, &std::cout);
cout.write(retBlob, size);
(dynamicFree)(retBlob);
}
// unload the library -----------------------------------------------
dlclose(lib_handle);
return 0;
}
要构建和使用此示例，只需将此源代码替换为第8部分中的_dynOCL.cc_。
出于测试目的，以下OpenCL内核_longAdd.cl_可以替换第8部分命令中的_simpleAdd.cl_，以演示设备分割正在工作。
inline __kernel void init(int veclen, __global TYPE1* c, int offset)
{
}
inline __kernel void func(int veclen, __global TYPE1* c, int offset)
{
// get the index of the test we are performing
int index = get_global_id(0);
// loop performing busywork to show processor activity
int n=100000;
for(int j=0; j < n; j++)
for(int i=0; i < n; i++) {
TYPE1 tmp = c[index + offset*veclen];
c[index + offset*veclen] += c[index + offset*veclen];
c[index + offset*veclen] -= tmp;
}
}
inline __kernel void fini(int veclen, __global TYPE1* c, int offset)
{
}
以下是运行Ubuntu 10.10的6核AMD Phenom™ II X6 1055T处理器上的系统监视器图形输出，该输出演示了OpenCL在所有核心上运行的默认行为。如前所述，在第8部分脚本中，_longAdd.cl_源代码替换了_simpleAdd.cl_。请注意，当应用程序开始运行时，所有六个处理器的处理器利用率都会跳升。
利用本教程中_dynOCL.cc_的设备分割版本，我们看到只有一个处理核心（在这种情况下是橙色线）实现了高利用率。
摘要
OpenCL 扩展为程序员提供了额外的功能，例如双精度算术和设备分割。供应商扩展的便携性最差，但它们确实提供了一个重要的途径来公开 API 以利用设备功能。KHR 扩展是最通用的，因为它们需要正式批准和一套测试来定义标准行为。EXT 扩展可以看作是“正在进行中”的 API，最终可能会获得 KHR 扩展的正式地位。
借助设备分割扩展，程序员可以通过API细分多核处理器，以更好地利用系统功能。第8部分介绍的Google protobuf流式传输框架可以轻松扩展以利用设备分割。通过诸如numactl等操作系统命令，程序员甚至可以将此流式传输框架中的OpenCL应用程序绑定到特定的处理核心。通过扩展，OpenCL应用程序程序员可以使用设备分割来进一步优化本教程系列第7和第8部分中讨论的OpenCL插件和通用工作流。