第五部分：OpenCL 缓冲区和内存亲和性
0/5 (0投票)
Rob Farber 在这一系列关于使用 OpenCL™ 进行可移植多线程编程的第五篇文章中，讨论了 OpenCL™ 缓冲区，并演示了如何在多设备、多 GPU 环境中将计算与数据关联起来。
上一篇文章，本系列的第 4 部分，介绍了 OpenCL™ 运行时，并演示了如何在异构设备的命令队列之间执行并发计算。
从第 4 部分改编的示例源代码，用于将用户指定的任务数量分发到多个 CPU 和 GPU 命令队列中。阻塞 I/O 和映射 I/O 都将用于将数据关联到每个设备上运行的内核。提供的 OpenCL 内核，它只将数字加到自身，故意保持简单，以便将注意力集中在队列和数据管理问题上。读者可以通过命令行参数轻松提供其他更复杂的内核，以便在利用 (1) 仅主机处理器、(2) 一个或多个 GPU 设备或 (3) 所有设备混合使用时，在多种环境中评估它们。与第 4 部分中的示例一样，并发执行通过 OpenMP pragma 简洁地启动。
本文中的源代码继续使用简单而有用的预处理器功能，将 C++ 模板类型传递给 OpenCL 内核。虽然本文中的示例并未修改数据类型来评估性能，但如果读者选择在混合多设备环境中评估替代内核的性能，则保留了此功能作为一种便利。
内存亲和性
OpenCL 使用宽松的内存一致性模型，正如本系列第 2 部分中所讨论的那样。OpenCL 内存设计的内在之美在于，数据是本地化的，并由程序员与工作项或工作组关联。然后，可以将这些工作项排队到许多设备上，以实现非常高的性能和可伸缩性。该模型确实要求程序员负责确保所有任务都能看到一致的内存视图。
OpenCL 在上下文中使用缓冲区对象来在设备之间共享数据，这与那些编写传统共享内存机器的程序员的预期编程模型不同。 缓冲区对象提供了程序员和 OpenCL 运行时协同工作的基础，以创建一个可以在多种机器和设备配置上无需编译即可运行的单一程序。例如，本文中的可执行文件在单处理器、单 GPU 系统以及包含 CPU 和多 GPU 的大型系统上无需重新编译即可运行。
值得注意的是，许多 OpenCL 程序员在遇到缓冲区最大尺寸限制时会感到困惑，这个限制通常远小于 GPU 上的内存大小。为了避免这个限制，有必要考虑将计算分区到跨越一个或多个设备上运行。然后，缓冲区就成为表达工作分布的自然方式——而不是将内存映像移动到单个大规模并行设备上的机制。为了保持效率，OpenCL 提供了映射缓冲区和异步缓冲区，以便在传输额外数据时计算可以继续进行。通过这种方式，OpenCL 程序员可以创建应用程序，在各种用户硬件配置上，在无需重新编译的情况下，提供高并发性和可移植的并行性。
通过在命令队列上排队一个或多个传输来执行显式的、由程序员发起的传输。示例包括
C API
clEnqueueReadBuffer(),clEnqueueReadImage()clEnqueueWriteBuffer(),clEnqueueWriteImage()clEnqueueCopyBuffer(),clEnqueueCopyImage()
C++ API
cl::enqueueReadBuffer(),cl::enqueueWriteBuffer()
数据传输可以是阻塞的，在这种情况下，队列会等待传输完成；也可以是异步的，需要使用事件在传输完成后进行通知。使用异步数据传输通过允许计算与数据移动重叠来提高应用程序性能——从而减少解决问题所需的时间。由于 PCIe 总线是全双工的，这意味着它可以同时在两个方向传输数据，因此有潜力实现 2 倍的数据传输带宽。
或者，可以通过将缓冲区映射到主机地址空间来隐式传输对象数据的区域。这些传输可以异步发生，也可以按需发生，这意味着只有计算所需的数据部分才会被移动并缓存到设备上。API 示例包括
C API
clEnqueueMapBuffer(),clEnqueueMapImage()clEnqueueUnmapMemObject()
C++ API
cl::Buffer()（通过下面讨论的各种标志）cl::enqueueMapBuffer(),cl::enqueueMapImage()cl::enqueueUnmapMemObject();
确定需要哪些标志或标志组合才能正确创建缓冲区并可能将其映射到设备可能会令人困惑。（以下内容摘录自 Khronos.org 消息板用户 bwatt 的精彩解释。）
在映射内存或使用 C++ OpenCL 包装器创建缓冲区时，可以使用三个标志
CL_MEM_ALLOC_HOST_PTRCL_MEM_COPY_HOST_PTRCL_MEM_USE_HOST_PTR
使用这三个标志可以创建五种有效组合
- 非映射，需要手动数据传输
- 未指定标志
CL_MEM_COPY_HOST_PTR- 映射缓冲区
CL_MEM_ALLOC_HOST_PTR- (
CL_MEM_ALLOC_HOST_PTR|CL_MEM_COPY_HOST_PTR) CL_MEM_USE_HOST_PTR
让 OpenCL 分配内存（选项 1a 和 2a）最有可能以可移植的方式提供良好的性能，因为缓冲区可以内部分配以最好地符合对齐、固定内存和其他特定于设备的性能标准。如果移植现有应用程序，则可能需要使用已分配的内存区域，这意味着CL_MEM_USE_HOST_PTR可能是唯一选项。从性能角度来看，这可能不是最佳选项。
选择队列模型
Derek Gerstmann 在 SIGGRAPH ASIA 2009 上提供了一份精彩的演示文稿，“高级 OpenCL 事件模型使用”，其中讨论了几个用例，包括
- 1 个顺序队列、1 个上下文、1 个设备。
- 1 个乱序队列、1 个上下文、1 个设备。
- 2 个顺序队列、2 个独立上下文、2 个设备。
- 2 个顺序队列、1 个合并上下文、2 个设备。
本文中的示例使用单个上下文，每个设备一个命令队列，即“协作多设备使用模型”。这意味着在单个上下文中创建的所有对象都由所有命令队列共享。OpenCL 规范的附录 A A.1 节“共享 OpenCL 对象”表明
- 使用上下文创建的 OpenCL 内存对象、程序对象和内核对象，可以通过使用相同上下文创建的多个命令队列进行共享。
- 命令队列可以在与该命令队列关联的设备上缓存内存对象状态的更改。
- 应用程序需要实现主机处理器上的适当线程同步，以确保共享对象的（应用程序认为正确的）状态更改按照正确的顺序发生（当多个命令队列在多个线程中更改共享对象的某个状态时）。
简而言之，要小心多个设备如何使用缓冲区，因为运行时可能会引入副本和其他行为，这些行为会影响性能甚至程序正确性。在 Khronos.org 消息板上这个帖子中，就这个主题进行了良好的近期讨论。基于这次讨论，本文中的示例代码为数据传输和内存映射到设备分配了每个设备的缓冲区。
示例
testSum.cpp 的源代码根据用户在命令行上指定的尺寸，在主机内存中创建一个二维整数数组。该数组填充了随机数，并映射到一个或多个设备。OpenCL 内核simpleAdd.cl，只是将每个数组元素加到自身。
正如在第 4 部分中所讨论的，C++ 模板头文件testSum.hpp 使用模板特化来获取多个原子数据类型（float、double、int 等）的模板参数的类型名。使用此方法可以保持代码的简单性，并且不会因过多的 C++ 技巧而使其复杂化。由于内核是在实例化类内部构建的，因此可以通过预处理器定义将此类型名传递给 OpenCL 内核构建。
testSum.hpp 的完整代码如下。请注意，它已进行调整以支持多个内核调用，基本上删除了initData方法中的所有代码。除此之外，这段代码与前两篇文章中使用的 C++ 模板文件非常相似。
#include <cmath>
#include <algorithm>
// The following defines specialized templates to provide a string
// containing the typename
template<class T>
struct TypeName {
string getName();
private:
T *t;
};
template<> string TypeName<double>::getName() {return(string("double")); }
template<> string TypeName<float>::getName() {return(string("float")); }
template<> string TypeName<unsigned long>::getName() {return(string("ulong"));}
template<> string TypeName<long>::getName() { return(string("long")); }
template<> string TypeName<unsigned int>::getName() {return(string("uint"));}
template<> string TypeName<int>::getName() {return(string("int")); }
template<> string TypeName<unsigned char>::getName() {return(string("uchar"));}
template<> string TypeName<char>::getName() {return(string("char")); }
// specification of the OclTest template
template <typename TYPE1>
class OclTest {
private:
cl::Kernel kernel;
string myType;
cl::Event event;
// variables for the test
int vecsize;
public:
cl::Event *getEventPtr() { return &event;}
OclTest() {}
OclTest( cl::CommandQueue& queue, const char* kernelFile,
int argc, char *argv[])
{
cl::Device device = queue.getInfo<CL_QUEUE_DEVICE>();
cl::Context context = queue.getInfo<CL_QUEUE_CONTEXT>();
cout << "---------- building OpenCL kernel ("
<< kernelFile << ") -----" << endl;
myType= TypeName<TYPE1>().getName();
cout << " My type is " << myType.c_str() << endl;
// Demonstrate using defines in the ocl build
string buildOptions;
{ // create preprocessor defines for the kernel
char buf[256];
sprintf(buf,"-D TYPE1=%s ", myType.c_str());
buildOptions += string(buf);
}
// build the program from the source in the file
ifstream file(kernelFile);
string prog(istreambuf_iterator<char>(file),
(istreambuf_iterator<char>()));
cl::Program::Sources source( 1, make_pair(prog.c_str(),
prog.length()+1));
cl::Program program(context, source);
file.close();
try {
cerr << " buildOptions " << buildOptions << endl;
cl::vector<cl::Device> foo;
foo.push_back(device);
program.build(foo, buildOptions.c_str() );
} catch(cl::Error& err) {
// Get the build log
cerr << "Build failed! " << err.what()
<< '(' << err.err() << ')' << endl;
cerr << "retrieving log ... " << endl;
cerr
<< program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(device)
<< endl;
exit(-1);
}
// Get the name of the kernel from the filename
string kernelName = string(kernelFile)
.substr(0,string(kernelFile).find(".cl"));
kernel = cl::Kernel(program, kernelName.c_str());
}
inline void initData(int _vecsize)
{
vecsize = _vecsize;
}
inline cl::Kernel& getKernel() { return(kernel); }
// Methods to return information for queuing work-groups
cl::NDRange getGlobalWorkItems() {
return( cl::NDRange( vecsize ) );
}
cl::NDRange getWorkItemsInWorkGroup() {
// Only one work item per workgroup
return( cl::NDRange(1, 1) );
}
};
TestSum.cpp
testSum.cpp 的重点是演示在单个上下文中创建多个命令队列，同时为多个设备上的多个内核调用维护数据亲和性。使用 C++ OpenCL 包装器可以使代码相当紧凑地找到所有设备并将它们放入一个上下文
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);
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);
cout << "Using the following device(s) in one context" << endl;
for(int i=0; i < devices.size(); i++) {
cout << " " << devices[i].getInfo<CL_DEVICE_NAME>() << endl;
}
这段代码已经在包含两个 ATI Radeon HD 5870 GPU 和 AMD Phenom™ II X6 1055T 处理器（运行最新的 AMD 加速并行处理 (APP) SDK（以前称为 ATI Stream））的系统上进行了测试。（注意：第 3 部分中的代码在多 GPU 系统上也正常工作。）
为每个设备创建单独的命令队列的指定方式也很紧凑
// Create the separate command queues to perform work
cl::vector< cl::CommandQueue > contextQueues;
for(int i=0; i < devices.size(); i++) {
#ifdef PROFILING
cl::CommandQueue queue(context, devices[i],CL_QUEUE_PROFILING_ENABLE);
#else
cl::CommandQueue queue(context, devices[i],0);
#endif
contextQueues.push_back( queue );
}
C++ 预处理器条件编译用于在映射缓冲区和未映射缓冲区之间进行选择。设置预处理器变量USE_MAP将使用隐式将数据传输到 OpenCL 设备的映射缓冲区来编译代码。默认使用显式阻塞传输。
int nDevices = contextQueues.size();
unsigned int* vec = new uint[nvec*vecsize];
int vecBytes=vecsize*sizeof(uint);
// Fill the host memory with random data for the sums
srand(0);
for(int i=0; i < (nvec*vecsize); i++) vec[i] = (rand()&0xffffff);
// Create a separate buffer for each device in the context
#ifdef USE_MAP
// This maps all of the host data into memory so it does not need
// to be manuually copied.
cl::vector< cl::Buffer > d_vec;
for(int i=0; i < contextQueues.size(); i++) {
d_vec.push_back(cl::Buffer(context, CL_MEM_COPY_HOST_PTR,
nvec* vecBytes, vec) );
}
int vecOffset=vecBytes; // the buffer is of size vec, so use row offset
#else
cl::vector< cl::Buffer > d_vec;
for(int i=0; i < contextQueues.size(); i++) {
d_vec.push_back(cl::Buffer(context, CL_MEM_READ_WRITE, vecBytes) );
}
int vecOffset=0; // the buffer is the size of one vector so no offset
#endif
与第 4 部分中的示例一样，OpenMP 已被用于在所有设备上并行地命令队列。这提供了一种简单简洁的方式来通过 OpenMP pragma 引入并行性，如下面的代码片段所示。有关 OpenMP 的更多信息可以在互联网上找到。一个很好的起点是维基百科文章OpenMP。请注意，内核在每个队列上运行多次。
#pragma omp parallel for
for(int i=0; i < contextQueues.size(); i++) {
test[i].initData(vecsize);
test[i].getKernel().setArg(0,vecsize);
test[i].getKernel().setArg(1,d_vec[i]);
for(int j=i; j < nvec; j += nDevices) {
#ifdef USE_MAP
test[i].getKernel().setArg(2,j); // set the offset for the kernel
#else
test[i].getKernel().setArg(2,0);
// manually transfer the data to the device
contextQueues[i].enqueueWriteBuffer(d_vec[i], CL_TRUE,0, vecBytes,
&vec[j*vecsize]);
#endif
contextQueues[i].enqueueNDRangeKernel(
test[i].getKernel(),
cl::NullRange, // offset starts at 0,0
test[i].getGlobalWorkItems(), // number of work groups
test[i].getWorkItemsInWorkGroup(), // workgroup size
NULL, test[i].getEventPtr());
// manually transfer the data from the device
contextQueues[i].enqueueReadBuffer(d_vec[i], CL_TRUE,
j * vecOffset,
vecBytes,
&vec[j*vecsize]);
}
contextQueues[i].finish(); // wait for everything to finish
}
在所有设备完成后，主机将双重检查结果并打印成功或失败
// perform the golden test
{
int i;
srand(0);
for(i=0; i < (nvec*vecsize); i++) {
unsigned int r = (rand()&0xffffff);
r += r;
if(r != vec[i]) break;
}
if(i == (nvec*vecsize)) {
cout << "test passed" << endl;
} else {
cout << "TEST FAILED!" << endl;
}
}
下面是 testSum.cpp 的完整代码。
#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
#include <CL/cl.hpp>
#include <fstream>
#include <iostream>
using namespace std;
#include "testSum.hpp"
int main(int argc, char* argv[])
{
if( argc < 5) {
cerr << "Use: {cpu|gpu|both} kernelFile nvec vecsize" << endl;
exit(EXIT_FAILURE);
}
// handle command-line arguments
const string platformName(argv[1]);
const char* kernelFile = argv[2];
int nvec = atoi(argv[3]);
int vecsize = atoi(argv[4]);
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 if(platformName.compare("both")==0) {
deviceType.push_back(CL_DEVICE_TYPE_GPU);
deviceType.push_back(CL_DEVICE_TYPE_CPU);
} else { cerr << "Invalid device type!" << endl; return(1); }
// create the contexts 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);
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);
cout << "Using the following device(s) in one context" << endl;
for(int i=0; i < devices.size(); i++) {
cout << " " << devices[i].getInfo<CL_DEVICE_NAME>() << endl;
}
// Create the separate command queues to perform work
cl::vector< cl::CommandQueue > contextQueues;
for(int i=0; i < devices.size(); i++) {
#ifdef PROFILING
cl::CommandQueue queue(context, devices[i],CL_QUEUE_PROFILING_ENABLE);
#else
cl::CommandQueue queue(context, devices[i],0);
#endif
contextQueues.push_back( queue );
}
// Create tests for all the queues
cl::vector< OclTest<uint> > test;
for(int i=0; i < contextQueues.size(); i++) {
test.push_back(OclTest<uint>(contextQueues[i],
kernelFile, argc-3, argv+3));
}
int nDevices = contextQueues.size();
unsigned int* vec = new uint[nvec*vecsize];
int vecBytes=vecsize*sizeof(uint);
// Fill the host memory with random data for the sums
srand(0);
for(int i=0; i < (nvec*vecsize); i++) vec[i] = (rand()&0xffffff);
// Create a separate buffer for each device in the context
#ifdef USE_MAP
// This maps all of the host data into memory so it does not need
// to be manuually copied.
cl::vector< cl::Buffer > d_vec;
for(int i=0; i < contextQueues.size(); i++) {
d_vec.push_back(cl::Buffer(context, CL_MEM_COPY_HOST_PTR,
nvec* vecBytes, vec) );
}
int vecOffset=vecBytes; // the buffer is of size vec, so use row offset
#else
cl::vector< cl::Buffer > d_vec;
for(int i=0; i < contextQueues.size(); i++) {
d_vec.push_back(cl::Buffer(context, CL_MEM_READ_WRITE, vecBytes) );
}
int vecOffset=0; // the buffer is the size of one vector so no offset
#endif
// run the tests
#pragma omp parallel for
for(int i=0; i < contextQueues.size(); i++) {
test[i].initData(vecsize);
test[i].getKernel().setArg(0,vecsize);
test[i].getKernel().setArg(1,d_vec[i]);
for(int j=i; j < nvec; j += nDevices) {
#ifdef USE_MAP
test[i].getKernel().setArg(2,j); // set the offset for the kernel
#else
test[i].getKernel().setArg(2,0);
// manually transfer the data to the device
contextQueues[i].enqueueWriteBuffer(d_vec[i], CL_TRUE,0, vecBytes,
&vec[j*vecsize]);
#endif
contextQueues[i].enqueueNDRangeKernel(
test[i].getKernel(),
cl::NullRange, // offset starts at 0,0
test[i].getGlobalWorkItems(), // number of work groups
test[i].getWorkItemsInWorkGroup(), // workgroup size
NULL, test[i].getEventPtr());
// manually transfer the data from the device
contextQueues[i].enqueueReadBuffer(d_vec[i], CL_TRUE,
j * vecOffset,
vecBytes,
&vec[j*vecsize]);
}
contextQueues[i].finish(); // wait for everything to finish
}
// perform the golden test
{
int i;
srand(0);
for(i=0; i < (nvec*vecsize); i++) {
unsigned int r = (rand()&0xffffff);
r += r;
if(r != vec[i]) break;
}
if(i == (nvec*vecsize)) {
cout << "test passed" << endl;
} else {
cout << "TEST FAILED!" << endl;
}
}
delete [] vec;
} catch (cl::Error error) {
cerr << "caught exception: " << error.what()
<< '(' << error.err() << ')' << endl;
}
return EXIT_SUCCESS;
}
以下命令使用最近发布的 AMD 2.5 版本 SDK 来构建和运行代码。请注意 C++ 预处理器定义USE_MAP的指定。
echo "---------------" g++ -I $ATISTREAMSDKROOT/include -fopenmp testSum.cpp -L $ATISTREAMSDKROOT/lib/x86_64 -lOpenCL -o testSum g++ -D USE_MAP -I $ATISTREAMSDKROOT/include -fopenmp testSum.cpp -L $ATISTREAMSDKROOT/lib/x86_64 -lOpenCL -o testSum
simpleAdd.cl
simpleAdd.cl 的完整源代码列表如下
inline __kernel void simpleAdd(int veclen, __global TYPE1* c, int offset)
{
// get the index of the test we are performing
int index = get_global_id(0);
c[index + offset*veclen] += c[index + offset*veclen];
}
命令行选项允许设置以下内容
- cpu，gpu，both
- cpu：仅在处理器核心上运行。
- gpu：在所有 GPU 设备上运行。
- both：在 CPU 和 GPU 设备上运行。
- OpenCL 内核的文件名。
- 数组的行数。这还定义了将有多少个内核调用跨设备发生（每行一个内核调用）。
- 每行的列数。这定义了将有多少数据跨 PCIe 总线传输。
以下显示了在所有设备上运行时的一个示例输出
$ ./testSum both simpleAdd.cl 300 10240 Using the following device(s) in one context Cypress Cypress AMD Phenom(tm) II X6 1055T Processor ---------- building OpenCL kernel (simpleAdd.cl) ----- My type is uint buildOptions -D TYPE1=uint ---------- building OpenCL kernel (simpleAdd.cl) ----- My type is uint buildOptions -D TYPE1=uint ---------- building OpenCL kernel (simpleAdd.cl) ----- My type is uint buildOptions -D TYPE1=uint test passed
摘要
本文中的示例表明，可以在不重新编译的情况下简洁地指定一个可以在多种设备配置上运行的应用程序。这是 OpenCL 可移植并行性的本质。
将数据与计算关联起来是一项关键能力，它允许将工作排队到许多设备上，以实现非常高的性能和可伸缩性。OpenCL 程序员需要重新定义他们的思维方式，而不是从内存映像的角度出发，而是从需要时可以在设备之间移动的缓冲区的角度出发。通过这种方式，可以利用 OpenCL 的映射和数据传输功能。这与大多数 SMP（共享多处理器）传统架构所呈现的单个大型内存的整体视图不同。通过缓冲区对象，OpenCL 程序员和运行时可以协同工作，使单个程序能够在单处理器、混合 CPU/GPU 系统或包含 CPU 和多 GPU 的系统上运行，而无需重新编译。