第三部分：工作组和同步
0/5 (0投票)
在他的第三个教程中，GPGPU专家Rob Farber将介绍OpenCL™执行模型，并讨论如何协调工作组中工作项之间的计算。
这是关于使用OpenCL™进行可移植多线程编程系列文章的第三篇，将介绍OpenCL™执行模型，并讨论如何协调工作组中工作项之间的计算。上一篇文章（第二部分）介绍了OpenCL内存空间，并提供了一个通用的C++ API OpenCL示例，以方便在CPU和GPU设备上构建和试验您自己的OpenCL内核。
具体而言，本教程将探讨OpenCL执行模型，并将扩展第二部分的示例代码来
- 为OpenCL内核构建提供命令行参数，例如预处理器宏定义。
- 重构第二部分示例，以便更好地分离
OclTest类，并使其更易于创建独立的测试。仅进行了少量更改，但值得注意的是，已添加了支持使用1D、2D或3D工作组的测试的方法，通过在OclTest类中执行特定于测试的命令行预处理来本地化变量，此外还添加了逻辑，以便可以使用各种OpenCL测试内核而无需重新编译主机代码。 - 包含了一个OpenCL内核示例，该示例演示了如何使用2D工作组来分块全局内存，并利用同步的共享内存操作，以便工作组中的每个工作项都可以使用全局内存和多个共享内存区域执行简单的操作。该示例代码（包括注释）仅25行，重点在于OpenCL中并行2D代码中的索引和同步可以多么简单和直观。它应该可以轻松扩展以实现更复杂的内核。
OpenCL设计的内在美在于数据被本地化并与工作项或工作组中表达的计算相关联。这有效地将工作和数据分解为许多小的独立任务，这些任务可以排队到一个或多个设备上，以实现非常高的性能和与并行硬件的可伸缩性。
出于性能和便利性的考虑，开发人员可以将这些本地化数据映射到支持1D、2D和3D索引功能的工作组中。通过使用轻量级且高效的屏障，开发人员可以并发地利用高速本地内存中的数据来重用数据并支持复杂的内存访问模式，从而大大提高应用程序的性能。
重要的是要注意，OpenCL规范的3.4.3节非常明确地指出，同步只能发生在
- 单个工作组内的工作项之间。
- 在单个上下文中的命令队列中排队的命令之间。
本教程将重点关注单个工作组内工作项之间的同步。虽然不同工作组的工作项也可以通过使用原子全局内存事务来协调执行，但通常最好避免使用原子操作进行同步，除非作为最后的手段，因为它们会影响可伸缩性，需要使用较慢的全局内存，在代码中引入死锁，并限制可移植性，因为原子操作是OpenCL的一个扩展，只有部分OpenCL运行时支持。尽管如此，通过原子操作进行同步对于某些计算问题可能是一种有价值且必要的功能，将在未来的教程中进行讨论，同样也会讨论通过命令队列进行同步。
简而言之：不同工作组的工作项永远不应尝试同步或共享数据，因为运行时不保证所有工作项都在并发执行，并且这种同步很容易导致死锁。
OpenCL执行模型
OpenCL执行模型基于在1D、2D或3D网格，或NDRange（“N维范围”）上对计算内核进行并行执行。单个内核实例，或工作项，在局部网格的每个点上运行，而工作组则在全局网格上运行。
除了术语之外，OpenCL执行模型的一个基本方面是为多维空间中通过NDRanges定义的每个工作项定义一个唯一的全局ID和一组局部ID。通过这些唯一的标识符，OpenCL执行模型允许开发人员精确地识别内核的每个并行实例在索引空间中的位置，以便它可以执行正确实现应用程序所需的计算。（注意：程序员还可以指定工作组大小，或让运行时来决定。）
下图显示了一个3x3的网格，由单独着色的2维（2x2）工作组组成。每个工作项的局部坐标显示在对角线上。
在内核内部，通过调用get_global_id(index)来查找全局坐标，其中index为0、1或2，取决于网格的维度。通过get_local_id(index)查找工作组的局部坐标。通过get_work_dim()查找使用的维度数。其他内置工作组函数可以在Khronos文档的工作项内置函数中找到。
下图来自Tim Matteson的Supercomputing 2009教程，幻灯片15，展示了这些ID的各种示例。
工作组内的同步
区别在于，屏障要求所有线程在barrier()调用处停止，而内存围栏仅要求在mem_fence()调用之前发生的加载和/或存储被提交到内存。重要的是要理解，OpenCL编译器可以重新排序内存操作，以最佳地利用设备架构在本地和全局内存中。因此，程序员不能依赖源代码中内存访问的顺序，因为实际操作可能以不同的顺序发生。内存围栏操作使开发人员能够强制执行数据依赖关系。
仔细使用mem_fence()可以大大提高性能，因为这些操作为开发人员提供了尽可能长时间地保持工作项活跃的机会。规范甚至允许开发人员分别控制加载操作的顺序，使用read_mem_fence()，以及存储操作的顺序，使用write_mem_fence()。
相比之下，屏障字面上会造成执行阻塞，因为所有线程必须在任何线程继续执行之前到达屏障。Khronos规范在下面的引文中也指出，屏障在某些情况下可能导致死锁。（此外，规范还要求屏障执行内存围栏，对读写操作都进行围栏，以防止编译器重新排序内存操作。）
"在一个处理器上执行内核的工作组中的所有工作项在允许任何工作项继续执行到屏障之后之前，都必须执行此函数。所有工作组中的工作项在执行内核时都必须遇到此函数。如果屏障位于条件语句内，则如果任何工作项进入条件语句并执行屏障，则所有工作项都必须进入条件语句。如果屏障位于循环内，则在任何工作项被允许继续执行到屏障之后之前，所有工作项都必须为循环的每次迭代执行屏障。"
barrier函数还会排队一个内存围栏（读写），以确保对本地或全局内存的内存操作的正确排序。"
以下是一个简单的OpenCL内核fill_tile.cl，它演示了2D工作组网格的使用、工作组内的索引，以及在本地内存中的分配和同步。预处理器宏定义被明确使用，以显示多维数组中本地内存的静态分配。本地内存也可以通过setKernelArg()（使用C++绑定）或clSetKernelArg()（使用C）进行动态分配。
__kernel void fill_tiles(__global float* a,
__global float* b,
__global float* c)
{
// find our coordinates in the grid
int row = get_global_id(1);
int col = get_global_id(0);
// allocate local memory for the workgroup
__local float aTile[TILE_DIM_Y][TILE_DIM_X];
__local float bTile[TILE_DIM_Y][TILE_DIM_X];
// define the coordinates of this workitem thread
// in the 2D tile
int y = get_local_id(1);
int x = get_local_id(0);
aTile[y][x] = a[row*N + col];
bTile[y][x] = b[row*N + col];
barrier(CLK_LOCAL_MEM_FENCE);
//Note the change in tile location in bTile!
c[row*N + col] = aTile[x][y] * bTile[y][x];
}
请注意，局部内存aTile和bTile数组的分配是为整个工作组完成的，使用了以下几行：
// allocate local memory for the workgroup
__local float aTile[TILE_DIM_Y][TILE_DIM_X];
__local float bTile[TILE_DIM_Y][TILE_DIM_X];
然后，局部数组由工作组中的每个工作项填充。每个局部瓦片2D索引有一个工作项，因为工作组的创建大小为（TILE_DIM_X，TILE_DIM_Y）。因此，每个瓦片有TILE_DIM_X * TILE_DIM_Y个并发工作项。
如果没有屏障同步，aTile和bTile的内容可能未定义，因为使用bTile[y][x]进行计算的工作项可能比填充该位置的工作项（从全局内存填充）运行得早！请记住，局部网格中的每个坐标都有一个独立的工作项，它实际上与所有其他工作项并行运行。此外，屏障中的隐式内存围栏操作阻止了编译器重新排序内存加载和存储，从而使代码能够正确执行。请注意，工作组中不需要任何逻辑来处理其他工作组中的任何事情。
为了演示本地内存的速度，c的值是使用aTile和bTile数组之间的转置索引计算的。测试类确保TILE_DIM_Y和TILE_DIM_X相等，以便aTile和bTile是对称的。主机代码还确保预处理器宏定义值N正确描述了a、b和c中包含的M x N矩阵的行大小。
测试2D工作组的主机代码
以下是workGroup2D.cpp的列表。在代码中搜索NEW，或者查看本文档中黄色的高亮注释，以了解与第二部分中使用的通用代码的区别。几乎所有的更改都发生在OclTest类中，而不是在main()中。最重要的更改是添加了两个新方法，getGlobalWorkItems()和getWorkItemsInWorkGroup()，用于指定创建使用工作组的多维测试所需的NDRanges。有关此源代码的更多信息，请参阅第二部分。
#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>
#include <string>
#include <cmath>
using namespace std;
// Note: Search for NEW to find the changes from part2.
class OclTest {
private:
// NEW Experiment with the size of the tile to explore
// performance on various devices (CPU, GPU, etc)
static const int TILE_DIM_X=16;
static const int TILE_DIM_Y=TILE_DIM_X;
cl::Kernel kernel;
cl_int mItems, nItems, nTiles;
cl_int vLen, vSize;
float *h_vecA, *h_vecB, *h_vecC;
cl::Buffer d_vecA, d_vecB, d_vecC;
public:
OclTest( cl::Context& context, cl::vector<cl::Device>& devices,
const char* kernelFile, int argc, char *argv[])
{
// NEW: parse commandline arguments so all test variables
// are now contained within the class.
if(argc < 2) {
cerr << "Use: cpu|gpu kernel sizeM sizeN" << endl;
exit(EXIT_FAILURE);
}
mItems = atoi(argv[0]);
nItems = atoi(argv[1]);
nTiles = mItems * nItems;
vLen = (mItems*TILE_DIM_Y)*(nItems*TILE_DIM_X);
vSize = vLen * sizeof(float);
// NEW: Demonstrate using defines in the ocl build
string buildOptions;
{ // create preprocessor defines for the kernel
char buf[256];
sprintf(buf,"-D TILE_DIM_X=%d -D TILE_DIM_Y=%d -D N=%d",
TILE_DIM_X, TILE_DIM_Y,nItems*TILE_DIM_Y);
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;
program.build(devices, 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>(devices[0])
<< endl;
exit(-1);
}
//NEW get the name of the kernel from the filename
string kernelName = string(kernelFile)
.substr(0,string(kernelFile).find(".cl"));
cerr << "specified kernel: " << kernelName << endl;
kernel = cl::Kernel(program, kernelName.c_str());
// set up the kernel inputs
h_vecA = new float[vLen];
h_vecB = new float[vLen];
h_vecC = new float[vLen];
d_vecA = cl::Buffer(context, CL_MEM_READ_ONLY, vSize);
d_vecB = cl::Buffer(context, CL_MEM_READ_WRITE, vSize);
d_vecC = cl::Buffer(context, CL_MEM_READ_WRITE, vSize);
kernel.setArg(0, d_vecA);
kernel.setArg(1, d_vecB);
kernel.setArg(2, d_vecC);
}
inline void initData(cl::CommandQueue& queue, cl::Event& event,
int seed)
{
srand(seed);
//NEW initialize data between 0 - 1
for(int i=0; i < vLen; i++) h_vecA[i] = rand()/(float)RAND_MAX;
for(int i=0; i < vLen; i++) h_vecB[i] = rand()/(float)RAND_MAX;
queue.enqueueWriteBuffer(d_vecA, CL_TRUE, 0, vSize, h_vecA);
queue.enqueueWriteBuffer(d_vecB, CL_TRUE, 0, vSize, h_vecB);
}
inline cl::Kernel& getKernel() { return(kernel); }
//NEW methods to return information for queuing work-groups
cl::NDRange getGlobalWorkItems() {
return( cl::NDRange(nItems*TILE_DIM_X, mItems*TILE_DIM_Y) );
}
cl::NDRange getWorkItemsInWorkGroup() {
return( cl::NDRange(TILE_DIM_X, TILE_DIM_Y) );
}
//NEW test for results from the fill_tile.cl kernel
inline int goldenTest(cl::CommandQueue& queue, cl::Event& event)
{
event.wait();
#ifdef PROFILING
cl_ulong start=
event.getProfilingInfo<CL_PROFILING_COMMAND_START>();
cl_ulong end=
event.getProfilingInfo<CL_PROFILING_COMMAND_END>();
double time = 1.e-9 * (end-start);
cout << "Time for kernel to execute " << time << endl;
#endif
// bring data back to the host via a blocking read
queue.enqueueReadBuffer(d_vecC, CL_TRUE, 0, vSize, h_vecC);
for(int row=0; row < mItems; row++)
for(int col=0; col < nItems; col++) {
float a[TILE_DIM_Y][TILE_DIM_X];
float b[TILE_DIM_Y][TILE_DIM_X];
float c[TILE_DIM_Y][TILE_DIM_X];
// fill a and b arrays
for(int y=0; y< TILE_DIM_Y; y++) {
int rindex = (row*TILE_DIM_Y+y)*nItems*TILE_DIM_Y;
for(int x=0; x < TILE_DIM_X; x++) {
a[y][x] = h_vecA[rindex + (col*TILE_DIM_X + x)];
b[y][x] = h_vecB[rindex + (col*TILE_DIM_X + x)];
c[y][x] = h_vecC[rindex + (col*TILE_DIM_X + x)];
}
}
// double check
for(int y=0; y< TILE_DIM_Y; y++)
for(int x=0; x< TILE_DIM_X; x++) {
if( c[y][x] != (a[x][y]*b[y][x]) ) {
cerr << "Error on c[" << y << "][" << x << "]";
cerr << " " << c[y][x] << " " << (a[x][y]*b[y][x]) << endl;
return(1);
}
}
}
return(0);
}
};
void displayPlatformInfo(cl::vector< cl::Platform > platformList,
int deviceType)
{
// print out some device specific information
cout << "Platform number is: " << platformList.size() << endl;
string platformVendor;
platformList[0].getInfo((cl_platform_info)CL_PLATFORM_VENDOR,
&platformVendor);
cout << "device Type "
<< ((deviceType==CL_DEVICE_TYPE_GPU)?"GPU":"CPU") << endl;
cout << "Platform is by: " << platformVendor << "\n";
}
int main(int argc, char* argv[])
{
int seed=4;
if( argc < 2) {
cerr
<< "Use: {cpu|gpu} kernelFile"
<< endl;
exit(EXIT_FAILURE);
}
// handle command-line arguments
const string platformName(argv[1]);
int deviceType = platformName.compare("cpu")?
CL_DEVICE_TYPE_GPU:CL_DEVICE_TYPE_CPU;
const char* kernelFile = argv[2];
try {
cl::vector< cl::Platform > platformList;
cl::Platform::get(&platformList);
displayPlatformInfo(platformList, deviceType);
cl_context_properties cprops[3] =
{CL_CONTEXT_PLATFORM,
(cl_context_properties)(platformList[0])(), 0};
cl::Context context(deviceType, cprops);
cl::vector<cl::Device> devices =
context.getInfo<CL_CONTEXT_DEVICES>();
#ifdef PROFILING
cl::CommandQueue queue(context, devices[0],
CL_QUEUE_PROFILING_ENABLE);
#else
cl::CommandQueue queue(context, devices[0], 0);
#endif
OclTest test(context, devices, kernelFile, argc-3, argv+3);
cl::Event event;
test.initData(queue, event, seed);
queue.enqueueNDRangeKernel(test.getKernel(),
cl::NullRange, // offset starts at 0,0
test.getGlobalWorkItems(), // number of work groups
test.getWorkItemsInWorkGroup(), // workgroup size
NULL, &event);
if(test.goldenTest(queue, event) == 0) {
cout << "test passed" << endl;
} else {
cout << "TEST FAILED!" << endl;
}
} catch (cl::Error error) {
cerr << "caught exception: " << error.what()
<< '(' << error.err() << ')' << endl;
}
return EXIT_SUCCESS;
}
构建和性能
要在Linux下构建主机程序，请将workGroup2D.cpp复制到一个文件中，然后使用以下命令
export ATISTREAMSDKROOT=$HOME/AMD/ati-stream-sdk-v2.2-lnx64 export LD_LIBRARY_PATH=$ATISTREAMSDKROOT/lib/x86:\ $ATISTREAMSDKROOT/lib/x86_64:$LD_LIBRARY_PATH g++ -I $ATISTREAMSDKROOT/include workGroup2D.cpp -L \ $ATISTREAMSDKROOT/lib/x86_64 -lOpenCL -o workGroup2D
将fill_tiles.cl的源代码复制到同一目录。
以下是使用ATI Radeon HD 5870 GPU和AMD Phenom™ II X6 1055T处理器产生的测试结果。
export ATISTREAMSDKROOT=$HOME/AMD/ati-stream-sdk-v2.2-lnx64 export LD_LIBRARY_PATH=$ATISTREAMSDKROOT/lib/x86:\ $ATISTREAMSDKROOT/lib/x86_64:$LD_LIBRARY_PATH ./workGroup2D gpu fill_tiles.cl 300 400 Platform number is: 1 device Type GPU Platform is by: Advanced Micro Devices, Inc. buildOptions -D TILE_DIM_X=16 -D TILE_DIM_Y=16 -D N=6400 specified kernel: fill_tiles Time for kernel to execute 0.00356587 test passed ./workGroup2D cpu fill_tiles.cl 300 400 Platform number is: 1 device Type CPU Platform is by: Advanced Micro Devices, Inc. buildOptions -D TILE_DIM_X=16 -D TILE_DIM_Y=16 -D N=6400 specified kernel: fill_tiles Time for kernel to execute 1.0187 test passed
虽然此内核不是一个好的性能测试，但它仍然能提供GPU和CPU性能差异的感觉。本示例代码和内核旨在方便用户试验不同的瓦片大小和计算。不妨试试，看看工作组和共享内存如何使您的应用程序受益。是否存在导致银行冲突的大小？
摘要
OpenCL推进了“可移植并行”的概念，因为它不仅仅是一种创建可在CPU、GPU、DSP和其他设备上运行的内核的语言。它还定义了使用工作组和共享数据协调并发并行计算的能力。
关键的OpenCL协调概念包括：
- NDRange（“N维范围”）可以定义1D、2D或3D工作组的大小，这很方便、高效，并且可以使代码更具可读性。
- 内核被实例化为工作项，工作项被分组到工作组中。开发人员可以指定工作组大小，或将其留给运行时决定。
- 编译器可以重新排序加载和存储。使用mem_fence()来强制执行数据依赖。
- 屏障也提供了出色的轻量级同步机制，但它们要求所有线程到达屏障才能继续执行。
- 在某些情况下，围栏操作可能更有效。
- 工作组是独立的。屏障和内存围栏不会跨工作组同步。可以使用原子操作，但由于死锁和可伸缩性问题，只能谨慎使用。
额外资源
- Tim Mattson（Intel）、Ian Buck（Nvidia）、Michael Houston（AMD）和Ben Gaster（AMD）的《OpenCL：异构并行计算机编程的标准平台》。
- CMSoft的《OpenCL 99高级》教程。
- AMD OpenCL论坛.