第四部分：使用 OpenCL 队列协调计算

这是关于使用 OpenCL™ 进行可移植多线程编程系列的第四篇文章，将讨论 OpenCL™ 运行时，并演示如何在异构设备的工作队列之间执行并发计算。前一篇文章（第 3 部分）介绍了 OpenCL 执行模型，并讨论了如何在工作组中的工作项之间共享数据和协调计算。

我们将改编第 3 部分的示例源代码，以利用多个命令队列，从而在 CPU 和 GPU 设备上进行并发内核执行的实验。本文将提供一个完整的示例，演示如何使用多个队列，并为读者提供一个框架，以便在海量并行环境中试验随机数生成器 (RNG)。并发执行通过一个 OpenMP pragma 简洁地启动。

本文的源代码还演示了一种简单而有用的预处理器功能，可将 C++ 模板类型传递给 OpenCL 内核。这使得为了性能评估而更改原子类型（例如 float、double、int、long 等）变得像编辑模板类型参数一样简单。例如，将 OclTestClass<float> 更改为 OclTestClass<double> 将在 C++ OclTestClass 和该类构建的 OpenCL 内核中同时使用双精度功能！

OpenCL 开发人员必须理解 OpenCL 运行时模型背后的思想，该模型允许开发人员创建无需重新编译即可自动调整以在各种硬件配置上高效运行的应用程序二进制文件。简而言之，程序员可以使用 OpenCL 命令队列执行和事件来指定任意排队命令之间的运行时依赖关系——实际上是定义了依赖图中的链接，运行时可以优化这些链接以最好地利用可用的硬件配置和资源。

正如前两篇教程所讨论的，OpenCL 使得开发人员能够指定计算内核，这些内核将在运行时重新编译以利用各种设备（例如 CPU、GPU、DSP、Cell 等）。这一强大功能体现了 OpenCL 可移植并行性的核心方面，因为它允许应用程序利用不同的硬件，甚至是尚未设计的硬件，而无需程序员干预，甚至无需重建应用程序二进制文件。

然而，除非开发人员也能协调工作——同样无需重新编译或程序员干预——以最好地利用应用程序在运行时碰巧可用的任何硬件配置，否则在多种异构设备上运行的能力用途有限。在任务之间定义依赖图体现了 OpenCL 设计促进“可移植并行性”的第二个核心方面，因为开发人员可以定义关键依赖关系，同时让运行时能够选择独立任务的最佳执行顺序，以最好地利用给定的硬件配置。通过这种方式，OpenCL 应用程序能够在单个 CPU/GPU 系统、具有多个 GPU 的系统，甚至使用 MPI（消息传递接口）或其他分布式通信框架的分布式集群上运行而无需更改。

在混合 CPU/GPU 环境中定义依赖关系、分配数据和计算、协调任务以及负载均衡是当前研究中非常活跃和令人兴奋的领域。假设程序员将工作划分为若干个独立任务，人们对创建易于使用的方法和软件运行时包有浓厚兴趣，这些包能够：

1. 将数据与计算相关联并绑定，以及在多个设备内存地址空间之间定义高效的数据布局。
2. 定义并强制执行任务之间的数据依赖关系。
3. 在多个异构计算单元之间以透明的方式进行负载均衡。
4. 将上述功能扩展到线程化和分布式消息传递环境中。
5. 在此基础设施之上构建类似于 BLAS 和 LAPACK 库的核心库。

正如论文《StarPU：异构多核架构任务调度的统一平台》所指出的，OpenCL 提供了一个重要的基础：

“OpenCL 倡议显然是一次有价值的尝试，旨在为 CPU、GPGPU 以及可能的其他加速器提供通用的编程接口。然而，OpenCL API 是一个非常低级别的接口，它基本上提供了显式卸载任务或在协处理器之间移动数据的原语。它不支持任务调度或全局数据一致性，因此不能被视为一个真正的‘运行时系统’，而更像一个虚拟设备驱动程序。”

目前，大量的研究和 API 规范制定工作正在通过一些激动人心的项目进行，例如 MAGMA（GPU 和多核架构上的矩阵代数），它有潜力成为像 BLAS 和 LAPACK 一样的新标准数学库 API，还有像 Vancouver 这样的项目，旨在基于 OpenCL 的可移植并行性将计算带入百亿亿次（exascale）时代。以下所有项目都允许用户下载并免费使用其软件和/或阅读其论文。

• StarPU 是一个用于异构多核架构上任务调度的统一平台，可免费下载。该平台似乎在下面讨论的 MAGMA 等知名项目中获得了越来越多的关注。
• MAGMA 项目提供了论文和一个可免费下载的用于混合计算的软件包。以下 LAPACK 工作组说明，“更快、更廉价、更好——为 GPU 开发线性代数软件的混合方法论”，讨论了 MAGMA 如何利用 StarPU 来创建新一代线性代数库。
• 论文《Maestro：面向 OpenCL 设备的数据编排与调优》讨论了一个使用 OpenCL 进行数据编排的开源库。编排数据移动以高效地与计算相结合是 Vancouver 项目的必要组成部分，该项目旨在为百亿亿次计算创建一个软件栈。Jeffrey Vetters 在 SC10 演讲的幻灯片第 3 页指出了 Maestro 与 OpenCL 相关的目标和方法。

• Dague（有向无环图环境）提供了关键功能，包括分布式多级动态调度器、异步通信引擎和数据依赖引擎。该库可免费下载。
• MOSIX 提供了一个 OpenCL 分布式功能的示例，在论文《一个用于在拥有众多 GPU 设备的集群上进行基于 OpenCL 的异构计算的软件包》中有所描述。
• 虽然不严格与 OpenCL 相关，但值得注意的是，诸如 MPC（多处理器计算）项目正在努力将 MPI 从面向进程的运行时过渡到面向线程的运行时，以减少开销并提高性能，这有望将面向线程的 OpenCL 代码迁移到分布式计算集群。

统一运行时系统的潜力和好处在 Benedict Gaster 2010 年 10 月为 OpenCL 编程网络研讨会系列所做演讲的以下幻灯片（幻灯片 3）中得到了总结：

OpenCL 运行时

主机程序通常实现将计算内核绑定在一起以创建应用程序的粘合代码。如第 3 部分所述，不同的工作项可以组合成工作组，这也允许工作组内的工作项之间进行共享和同步。假设大部分计算工作将在 OpenCL 内核中进行。因此，主机应用程序的逻辑通常侧重于构建内核、设置 OpenCL 上下文和设备、通过一个或多个命令队列向设备分派工作，以及协调所有内核和命令队列之间的应用程序工作负载任务同步，以便满足所有数据依赖关系——无论执行顺序如何——同时还给予运行时尽可能多的调度工作的灵活性。

内核和命令队列都在一个上下文中运行，该上下文包含内核执行环境的状态。正如在第 2 部分中讨论的，OpenCL 使用一种宽松的内存一致性模型，要求程序员确保内存在并发和异步的 OpenCL 操作中得到一致的使用。这使得程序员需要负责工作组内的操作同步、命令队列中任务之间的同步，以及跨多个命令队列的同步。

命令队列是协调内核执行和强制数据依赖关系的两种机制之一。例如，如果一个对某组数据进行操作的内核在初始化设备上数据的传输操作之前运行，那将是非常糟糕的。程序崩溃、不正确和不确定的结果是使用未初始化内存的明显后果。（第二种机制是工作组中工作项之间的原子同步操作。）

本系列教程之前的示例使用了最简单的顺序执行、单设备和单队列实现。这简化了代码，因为每个排队的任务都将按照其放入队列的顺序在设备上顺序执行。由于数据传输首先入队，然后是内核调用，因此无需额外的代码或复杂性即可保证满足数据依赖关系。

当存在多个队列和/或指定了乱序命令执行时，就需要同步。在这些情况下，必须使用事件来强制执行一个上下文内和队列之间的同步。（当任务在多线程主机应用程序中异步添加到队列时，也会出现一种特殊情况，因为插入顺序无法保证。）

队列行为的类型（顺序 vs. 乱序）在创建队列时指定。默认情况下，使用 clCreateCommandQueue() 创建的队列是按顺序发出命令的。但是，在

队列命令属性中指定 CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE 将创建一个乱序队列。此外，可以使用 clSetCommandQueueProperty() 在队列创建后更改其排队行为。

虽然 OpenCL 规范允许为每个设备创建多个队列，但它并未定义运行时如何选择下一个要运行的任务，或者在设备上是否可以并发执行内核——假设有足够的资源。经验丰富的并行程序员知道，异步和并发的命令调度是一个非常好的主意，因为它给了运行时最大的机会，通过在一个或多个设备上并发运行多个任务来最大化性能。实际上，这利用了运行时的调度能力来优化性能，就像线程级并行允许 GPGPU 上的线程级调度器实现高性能一样。

以下主机 API 调用用于指定等待列表、返回事件以及查询事件信息：

• clWaitForEvents(num_events, *event_list)：阻塞直到事件完成。
• clEnqueueMarker(queue, *event)：为一个在队列中移动的标记返回一个事件。
• clEnqueueWaitForEvents(queue, num_events, *event_list)：向队列中插入一个“WaitForEvents”操作。
• clGetEventInfo()：获取命令类型和状态，由 CL_QUEUED、CL_SUBMITTED、CL_RUNNING、CL_COMPLETE 或一个错误码表示。
• clGetEventProfilingInfo()：返回来自命令队列的性能分析信息，如提交、开始和结束时间。

通过这些 API 调用，可以指定一个上下文中非常复杂的依赖图。如引言中所述，这为 OpenCL 开发人员提供了极其强大、通用的功能！以下是来自AMD 教程幻灯片第 28 页的一张图，展示了如何使用等待事件（由箭头表示）来支持两个上下文之间的依赖关系——每个上下文包含两个设备：

1. 来自 Q1,1 的 C3 依赖于来自 Q1,2 的 C1 和 C2
2. 来自 Q1,4 的 C1 依赖于来自 Q1,2 的 C2
3. 在 Q1,4 中，C3 依赖于 C2

非常重要的一点是，OpenCL 规范的 3.4.3 节明确指出，同步只能发生在：

• 单个工作组内的多个工作项之间。
• 入队到单个上下文中的命令队列的命令之间。（添加斜体是为了强调同步只能在上下文内部发生。）

对于队列之间和跨上下文的更粗粒度的同步，clFlush() 和 clFinish() 提供了一种强制刷新并等待所有排队命令完成的简单粗暴的方法。

Derek Gerstmann 为 SIGGRAPH ASIA 2009 做了一场精彩的演讲，“高级 OpenCL 事件模型用法”，强调了在单个上下文中命令队列之间同步的必要性。我强烈建议您在阅读本文后回顾他的幻灯片，因为他详细介绍了一系列使用模型，并提供了易于理解的代码片段，展示了如何在各种用例中使用事件，包括：

• 1 个顺序队列，1 个上下文，1 个设备。
• 1 个乱序队列，1 个上下文，1 个设备。
• 2 个顺序队列，2 个独立上下文，2 个设备。
• 2 个顺序队列，1 个组合上下文，2 个设备。

注意事项和有用提示

• 建议不要在单个上下文中混合使用多个异构设备，正如此 AMD 教程的幻灯片 6 所指出的。性能下降可能由多种因素造成。例如，某些 OpenCL 实现使用单个线程来支持一个上下文。混合多种设备类型可能会导致服务设备请求时出现延迟，因为线程可能变得太慢而无法快速响应设备。根据应用程序的不同，这可能会导致性能下降。
• 目前，一些 OpenCL 实现在为单个设备定义多个队列时表现出较差的性能。例如，Markus Stürmer 在他的博客文章《如果你想使用多个主机线程来供给多个 OpenCL 设备，会发生一些有趣的事情》中指出，为每个设备运行两个命令队列可能会对性能产生负面影响。
• 并行内存传输对于获得良好的流式传输性能是必需的，并且对实时应用程序至关重要。这个帖子，“尝试从 SDK 2.3 获得异步传输”，讨论了许多注意事项。
• 帖子“SDK 2.3 和多 GPU”指出，现在在 Linux 下可以使用两个甚至三个 GPU。在 SDK 2.3 之前，多 GPU 在 Linux 下无法工作。但是，所有 GPU 都需要连接到显示器或虚拟 VGA 连接器。
• 截至 SDK 2.3 版本，ATI 驱动程序不支持乱序执行。

示例

关于并行计算机上的随机数生成器（RNG）存在一些悬而未决的问题。最简单的实现之一是为并行计算机上的每个线程提供一个单独的 RNG，并为每个线程使用不同的种子。正如以下文章指出的，这样的实现可能会也可能不会在随机数序列中引入人为因素和不希望的相关性（链接）。

为了帮助读者开始研究并行 RNG 这个有趣的领域，本文提供了一个使用 OpenCL 的简单框架，该框架可以在多种平台上运行。George Marsaglia 的 MWC（带进位乘法）算法，如 Code Project 上所述，被用作一个好的 RNG 的简单示例。一个顺序数生成器被用作简单的反例，因为顺序数肯定不是随机的！来自 ENT：一个伪随机数序列测试程序 的单比特熵测试被用来评估 RNG 的性能。一个更完整的测试将是 Marsaglia 的 DIEHARD 测试集。

从代码的角度来看，这个示例重构了第 3 部分的源代码，为 CPU 和 GPU 设备分别定义了一个队列。新增了一个命令行参数“both”，用于在 CPU 和 GPU 资源上同时运行。为了灵活性和性能测试，特别值得关注的是，创建 OclTest 类时使用的 C++ 模板类型参数，在类内部构建 OpenCL 内核时作为预处理器定义传递。只需更改 OclTest 类的模板类型，如下所示，就可以使用 float、double 或其他变量类型来测试 OclTest。无需对 OpenCL 内核进行任何更改。

OclTest<unsigned long>  test1(…) // a 64-bit int test class and kernel
OclTest<float>  test1(…) // a single-precision test class and kernel
OclTest<double> test1(…) // a double-precision test class and kernel

此外，还利用 OpenMP 在两个设备上并行地并发排队命令。这提供了一种通过 OpenMP pragma 引入并行性的简单而简洁的方法，如下面的代码片段所示：

    // initialize the data and run the tests
#pragma omp parallel for
    for(int i=0; i < contextQueues.size(); i++) {
      test[i].initData( contextQueues[i], (i*1000000));
      
      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());
      contextQueues[i].finish();
    }

然而，正如这个帖子《在 OpenCL 程序中使用多个 GPU》所指出的，并行排队命令可能不足以启动并发执行。在 AMD SDK 的 2.3 版本中，需要调用队列的 finish 函数来强制并发执行。

contextQueues[i].finish();

更多关于 OpenMP 的信息可以在互联网上找到。一个很好的起点是维基百科的文章，OpenMP。

为了方便起见，第 3 部分的代码被分成了主文件 testRNG.cpp，它处理所有指定队列并向其分配工作的粘合代码。使用了 cl::vector 模板使代码简单易读。OpenCL 测试类模板 OclTest.hpp 通过以下行包含进来：

#include "OclTest1.hpp"

以下是 testRNG.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 "OclTest1.hpp"
 
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 << endl;
}
 
int main(int argc, char* argv[])
{
  int seed=4;
 
  if( argc < 2) {
    cerr << "Use: {cpu|gpu|both} kernelFile" << endl;
    exit(EXIT_FAILURE);
  }
 
  // handle command-line arguments
  const string platformName(argv[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 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); }
 
  const char* kernelFile = argv[2];
 
  // create the contexts and queues
  try {
    cl::vector< cl::Platform > platformList;
    cl::Platform::get(&platformList);
    
    for(int i=0; i < deviceType.size(); i++) {
      displayPlatformInfo(platformList, deviceType[i]);
      
      cl_context_properties cprops[3] = 
      {CL_CONTEXT_PLATFORM, 
       (cl_context_properties)(platformList[0])(), 0};
      
      cl::Context context(deviceType[i], cprops);
      cl::vector<cl::Device> devices = 
      context.getInfo<CL_CONTEXT_DEVICES>();
      for(int j=0; j < devices.size(); j++ ) {
#ifdef PROFILING
      cl::CommandQueue queue(context, devices[j],CL_QUEUE_PROFILING_ENABLE);
#else
      cl::CommandQueue queue(context, devices[j],0);
#endif
      contextQueues.push_back( queue );
      }
    }
    
    // Create tests for all the queues
    cl::vector< OclTest<ulong> > test;
    for(int i=0; i < contextQueues.size(); i++) {
      test.push_back(OclTest<ulong>(contextQueues[i], kernelFile, argc-3, argv+3));
    }
 
    // initialize the data and run the tests
#pragma omp parallel for
    for(int i=0; i < contextQueues.size(); i++) {
      test[i].initData( contextQueues[i], (i*1000000));
      
      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());
      contextQueues[i].finish(); // needed for concurrency
    }
    // perform the golden tests
    for(int i=0; i < contextQueues.size(); i++) {
      if(test[i].goldenTest( contextQueues[i] ) == 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;
}

C++ 模板头文件 OclTest.hpp 提供了一个简单的模板特化用法，用于获取几种原子类型的参数类型名。使用此方法可以保持代码简单，而不会因为过多的 C++ 魔法而使代码复杂化。

// 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")); }

由于内核是在实例化的类内部构建的，我们可以通过预处理器定义将类型传递给 OpenCL 内核构建过程。

    // 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);
    }

然后内核只需使用预处理器定义中的类型。注意：双精度是 OpenCL 的一个扩展，需要在 OpenCL 源代码中启用。以下 pragma 在使用 AMD SDK 时启用双精度。

#pragma OPENCL EXTENSION cl_amd_fp64 : enable

OclTest 类现在是一个类模板，包含在 testRNG.cpp 中。

为简单起见，此代码为每个线程执行一次测试。包含“1”的位数计数将返回给主机，用于计算随机数生成器的单比特熵。修改代码以计算单字节熵也很简单。

另请注意，现在对构造函数的调用仅传递命令队列。这使示例代码进一步脱离了第 1 部分的原始 C 代码，因为所有设备和上下文信息都可以如下面的代码片段所示确定：

cl::Device queueDevice = queue.getInfo<CL_QUEUE_DEVICE>();
std::string deviceName = queueDevice.getInfo<CL_DEVICE_NAME>();

以下是 OclTest.hpp 的完整源代码

#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:
  // Experiment with RNG and performance
  // performance on various devices (CPU, GPU, etc)
  cl::Kernel kernel;
  cl_int nTests;
  cl_int vLen, vSize;
  TYPE1  *h_oneCnt;
  cl::Buffer d_oneCnt;
  string myType;
  cl::Event event;
  long seed1,seed2;
  int seedOffset;
  unsigned long  nIter;
  
  // used to calculate entropy
  const static double log2of10=3.32192809488736234787;
  inline double rt_log2(double x) {
    return log2of10 * log10(x);
  }
 
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>();
    myType= TypeName<TYPE1>().getName();
    cout << "My type is " << myType.c_str() << endl;
 
    if(argc < 4) {
      cerr << "Ocl kernel use: nTests seed1 seed2 nIterPerTest" << endl;
      exit(EXIT_FAILURE);
    }
    nTests = atol(argv[0]);
    seed1 = atol(argv[1])+seedOffset;
    seed2 = atol(argv[2])+seedOffset;
    nIter = atol(argv[3]);
    seedOffset = 0;
 
    vLen = nTests;
    vSize = vLen * sizeof(TYPE1);
 
    // 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"));
    cerr << "specified kernel: " <<  kernelName << endl;
    kernel = cl::Kernel(program, kernelName.c_str());
    
    // set up the kernel inputs
    h_oneCnt = new TYPE1[vLen];
    d_oneCnt = cl::Buffer(context, CL_MEM_READ_WRITE, vSize);
    
    kernel.setArg(0, nIter);
    kernel.setArg(1, seed1);
    kernel.setArg(2, seed2);
    kernel.setArg(3, d_oneCnt);
  }
  
  inline void initData(cl::CommandQueue& queue, int seed) 
  {
    seedOffset = seed;
  }
 
  inline cl::Kernel& getKernel() { return(kernel); }
  //NEW methods to return information for queuing work-groups
  cl::NDRange getGlobalWorkItems() {
    return( cl::NDRange( vLen ) ); 
  }
  cl::NDRange getWorkItemsInWorkGroup() {
      // Only one work item per workgroup
    return( cl::NDRange(1, 1) ); 
  }
 
  inline int goldenTest(cl::CommandQueue& queue) 
  {
    event.wait();
#ifdef PROFILING
    cl::Device queueDevice = queue.getInfo<CL_QUEUE_DEVICE>();
    std::string deviceName = queueDevice.getInfo<CL_DEVICE_NAME>();
    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);
    double startTime = start * 1.e-9;
    double endTime = end * 1.e-9;
    cout << "Kernel (start,end) " << startTime << "," << endTime
       << " Time for kernel to execute " << time 
       << " device \"" << deviceName << "\"" << endl;
#endif
 
    // bring data back to the host via a blocking read
    queue.enqueueReadBuffer(d_oneCnt, CL_TRUE, 0, vSize, h_oneCnt);
 
    // Find the minimum entropy and display some info
    double minBitEntropy=10;
    double minBitMean=10;
    for(int i=0; i < vLen; i++) {
      // calculate entropy and arithmetic mean
      double totalBitCount = 32*nIter;
      double bitEntropy=0.;
      double bitProb[2];
      double bitMean = h_oneCnt[i]/totalBitCount;
      bitProb[1] = ((double)h_oneCnt[i])/totalBitCount;
      bitProb[0] = 1. - bitProb[1];
      if( (bitProb[0] > 0.f) && (bitProb[1] > 0.f) ) {
      bitEntropy = bitProb[0] * rt_log2(1./bitProb[0])
        + bitProb[1] * rt_log2(1./bitProb[1]);
      }
      minBitEntropy = min(minBitEntropy,bitEntropy);
      minBitMean = min(minBitMean,bitMean);
    }
    cout << "Minimum bitEntropy " << minBitEntropy 
       << " Minimum bitMean " << minBitMean << endl;
    cout << "Optimum compression would reduce " 
       << "the size of this random bit sequence by " 
       << (1.-minBitEntropy) << '%' << endl;
 
    if(minBitEntropy >= 0.9999) return(0);
    else return(1);
  }
};

以下命令使用最近发布的 AMD 2.3 版 SDK 构建并运行代码：

echo "---------------"
export ATISTREAMSDKROOT=$HOME/AMD/ati-stream-sdk-v2.3-lnx64
export ATISTREAMSDKSAMPLESROOT=$HOME/AMD/ati-stream-sdk-v2.3-lnx64
g++ -I $ATISTREAMSDKROOT/include -fopenmp testRNG.cpp -L $ATISTREAMSDKROOT/lib/x86_64 -lOpenCL -o testRNG

以下是一个用 OpenCL 编写的简单随机数生成器，simpleRNG.cl。

inline void SetSeed(unsigned int m_z, unsigned int u, unsigned int m_w,  unsigned int v)
{
  m_z +=u;
  m_w +=v;
}
 
inline unsigned int GetUint(unsigned int m_z, unsigned int m_w)
{
 m_z = 36969 * (m_z & 65535) + (m_z >> 16);
 m_w = 18000 * (m_w & 65535) + (m_w >> 16);
 return (m_z << 16) + (m_w & 65535);
}
 
inline 
__kernel void simpleRNG(unsigned long n, long seed1,
                  long seed2,
                        __global TYPE1* c)
{
  unsigned int m_z=521288629;
  unsigned int m_w=362436069;
 
  // get the index of the test we are performing
  int index = get_global_id(0);
 
  // Change TYPE1 via the template parameter to OclTest
  TYPE1 oneCnt=0;
 
  // set the seed for the random generator
  SetSeed(m_z, seed1+index, m_w, seed2+index);
 
  // Generate the random numbers and count the bits
  for(unsigned long iter=0; iter < n; iter++) {
    unsigned int rnd = GetUint(m_z, m_w);
    // uncomment to make test fail by setting high 
    //rnd = rnd % 0xffff;
    for(int i=0; i< 32; i++) {
      if( (rnd&0x01) ) oneCnt++;
      rnd = rnd >> 1;
    }
  }
  c[index] = oneCnt; 
}

请注意，如下表所示，两个测试在 CPU 和 GPU 上并发运行。在任一设备完成其任务之前，两个设备上的运行都已经开始。

	开始时间	结束时间
GPU	42549.7	42557.0
CPU	42548.4	42563.1

为了进行比较，我们提供了一个顺序数生成器作为“随机”数生成器的示例。以下是 badRNG.cl 的代码：

inline 
__kernel void badRNG(unsigned long n, long seed1,
                 long seed2,
                 __global TYPE1* c)
{
  // get the index of the test we are performing
  int index = get_global_id(0);
 
  // Change TYPE1 via the template parameter to OclTest
  TYPE1 oneCnt=0;
 
  // Generate sequential numbers and count the bits
  for(unsigned long iter=0; iter < n; iter++) {
    unsigned int rnd = iter;
    for(int i=0; i< 32; i++) {
      if( (rnd&0x01) ) oneCnt++;
      rnd = rnd >> 1;
    }
  }
  c[index] = oneCnt;

}

正如预期的那样，这产生了很差的结果。检查开始时间表明，这些运行也是在两个设备上同时进行的。

bda$ ./testRNG both badRNG.cl 1024 1 2 1000000
Platform number is: 1
device Type GPU
Platform is by: Advanced Micro Devices, Inc.
Platform number is: 1
device Type CPU
Platform is by: Advanced Micro Devices, Inc.
My type is ulong
buildOptions -D TYPE1=ulong 
specified kernel: badRNG
My type is ulong
buildOptions -D TYPE1=ulong 
specified kernel: badRNG
Kernel (start,end) 42775.5,42783.1 Time for kernel to execute 7.58799 device "Cypress"
Minimum bitEntropy 0.891907 Minimum bitMean 0.308906
Optimum compresion would reduce the size of this random bit sequence by 0.108093%
TEST FAILED!
Kernel (start,end) 42774.1,42788 Time for kernel to execute 13.8672 device "AMD Phenom(tm) II X6 1055T Processor"
Minimum bitEntropy 0.891907 Minimum bitMean 0.308906
Optimum compresion would reduce the size of this random bit sequence by 0.108093%
TEST FAILED!

摘要

OpenCL 推进了“可移植并行性”的概念，因为它不仅仅是一种创建可在 CPU、GPU、DSP 和其他设备上运行的内核的语言。它还定义了一个平台 API 来协调异构并行计算。虽然技术文献中充斥着并行协调语言和 API，但 OpenCL 在促进协调多个异构设备上并发运行的内核方面的能力是独一无二的。这两种核心能力（可移植内核和以灵活方式定义和强制数据依赖关系的能力）相结合，使开发人员能够创建可在各种硬件平台和配置上高效运行的应用程序。

OpenCL 协调的关键概念包括：

• 每个设备都有自己的异步工作队列。
• 通过来自不同（或相同）设备的事件句柄在 OpenCL 计算之间进行同步。
• 支持使用所有可用计算资源的算法和系统。
• 将“原生函数”入队以与 C/C++ 代码集成。（原生函数是指由库或操作系统提供的函数。）

附加资源

• Derek Gerstmann 在 SIGGRAPH ASIA 2009 上的演讲“高级 OpenCL 事件模型用法”。
• 将 BigDFT 移植到 OpenCL
• Dominik Behr 为 PPRAM 2009 编写的 AMD OpenCL 教程，发布于 GPGPU.org。
• “OpenCL™：CPU 和 GPU 的并行计算” 作者：Lee Howes (AMD)。
• “OpenCL：编程异构并行计算机的标准平台” 作者：Tim Mattson (Intel)、Ian Buck (Nvidia)、Michael Houston (AMD) 和 Ben Gaster (AMD)。