65.9K
CodeProject 正在变化。 阅读更多。
Home

第二部分：OpenCL™ – 内存空间

manythreads
emptyStarIconemptyStarIconemptyStarIconemptyStarIconemptyStarIcon

0/5 (0投票)

2010年10月27日

CPOL

14分钟阅读

viewsIcon

136271

在其第二篇教程中，GPGPU专家Rob Farber讨论了OpenCL™内存空间和OpenCL内存层次结构，以及如何开始思考工作项和工作组。本教程还提供了一个通用示例，以方便实验各种OpenCL内核。

这一系列关于使用OpenCL™进行便携式多线程编程的第二篇文章将讨论内存空间。前一篇文章介绍了OpenCL标准背后的思想，并演示了如何下载和使用ATI Stream软件开发套件（SDK）来构建和运行OpenCL程序。

本教程还提供了一个通用示例，以方便实验各种OpenCL内核。

OpenCL宽松内存一致性模型

OpenCL内存模型隐含的观点是，内核（在设备上运行的工作项的代码）驻留在单独的内存空间中。只有在通过某种形式的数据传输将数据移入设备内存后，才能访问数据。每个工作项都可以使用私有内存、本地内存、常量内存和全局内存。

本质上，OpenCL使用所谓的宽松内存一致性模型（Khronos OpenCL工作组，2008a，第25页），该模型

  • 允许工作项访问私有内存中的数据。
  • 允许在工作组执行期间，工作项共享本地内存。然而，只有在各种同步点（如内核中的屏障（只能用于同步工作组元素之间对本地内存的视图）和队列事件）之后，内存才保证一致。
  • 不同的工作组无法相互通信或同步，因为跨越工作组项集合的变量的一致性无法保证。
  • 正如将在后续文章中讨论的，数据依赖性也可以通过工作队列和原子操作来定义和满足。

为了很好地解释内存一致性模型，我建议阅读

OpenCL内存设计的内在之美在于，数据由程序员本地化并与工作项或工作组相关联。然后，这些工作项可以被排队到一个或多个设备上，以实现非常高的性能和可扩展性。每个设备内部都可以潜在地支持大量并发执行线程。（实际上，我们将术语线程和工作项互换使用。）

对于大多数开发人员来说，一个巨大的挑战在于他们不习惯以数据分区和将计算任务拆分成单独工作项的方式来思考。编写单线程应用程序或多核SMP系统的应用程序，为程序员提供了直观简单的缓存一致内存模型。换句话说，所有驻留在本地处理器缓存中的数据都保证在处理器之间是一致的。因此，程序员可以自由地忽略数据分区问题，因为他们可以假定任何核心都可以访问通用地址空间中的任何字节。不幸的是，保持缓存一致性会引入显著的通信开销，并限制程序员和硬件设计者优化性能的能力。因此，SMP系统的成本更高，并且由于扩展瓶颈，随着核心数量的增加，性能会很快受到限制。我个人在包含四颗四核芯片（共16核）的小型SMP系统上，看到了由于保持缓存一致性而导致的性能显著下降。

简而言之，OpenCL宽松一致性模型意味着不同的工作项在计算进行过程中可能看到全局内存的不同视图。需要同步来确保工作组内工作项之间的数据一致性，而对所有内存空间的读写在工作项内是始终按顺序进行的。

OpenCL为应用程序开发人员提供的“胡萝卜”是能够开发出极具可扩展性、在SMP系统上实现极高性能、并在廉价GPU硬件上可能实现数量级性能提升的应用程序。混合CPU/GPU系统现在非常普遍。为这些平台开发应用程序的需求为OpenCL的便携式并行性提供了动力，因为GPU支持数百到数千个并发执行线程。同样，大众市场多核SMP系统的数量也在变化，拥有数十到数百个核心的产品已经上市或将在不久的将来上市。为了利用这种硬件能力，OpenCL程序员必须学会以数据分区和本地化计算的方式来思考计算问题——换句话说，使用工作单元和工作组。

实现高性能的规则

正如本系列第一篇文章中所讨论的，高性能应用程序遵循三个通用规则。可以看出，适当使用设备内存空间和层次结构至关重要

  1. 将数据获取并保留在GPU上，以消除PCI总线数据传输瓶颈。SDK示例显示的性能差异说明了在ATI Radeon™ HD 5870 GPU上的情况
    • PCIe带宽：2.4 GB/s
    • 全局内存带宽：118 – 169 GB/s
  2. 给GPU足够的工作量。

    启动内核确实需要少量开销。然而，现代GPU速度非常快，可以在内核启动期间执行大量工作。

  3. 优化计算以最小化访问GPU内存的瓶颈。

例如，ATI Radeon HD 5870 GPU所能提供的最佳读取内存带宽为169 GB/s。一个执行一次浮点运算的内核，每次读取向量中的一个元素，最多只能提供42 GFLOP/s的性能，这比该GPU声称的2,720 GFLOP/s峰值性能要慢得多。

第一点强调了对数据进行适当分区以将其移动到一个或多个设备内存空间并保留在那里的必要性。正如第3点所示，数据重用是性能的关键。我们将在以后的文章中讨论第二点。

OpenCL内存层次结构

OpenCL定义了一个内存类型层次结构，如下图所示，摘自AMD OpenCL入门。正确使用可以带来显著的性能提升。然而，OpenCL标准仅规定了不同类型内存的访问级别。许多重要的性能细节可能因供应商而异。好消息是，这种灵活性允许每个供应商交付最高性能的产品，但对于希望利用OpenCL的可移植性并同时在不同设备上获得最高性能的开发人员来说，这带来了一个挑战。然而，供应商的实现似乎受到现有硬件架构（尤其是图形处理器）的强烈影响。因此，要创建高性能的可移植OpenCL应用程序，最好以GPU架构为中心进行思考。

image001.jpg

全局内存：全局内存通常是计算设备上容量最大的内存子系统。对于大多数设备，全局内存的容量将以千兆字节为单位。全局内存虽然容量大且对GPU或SMP系统上的所有线程可见，但应被视为最慢的内存子系统，并且在使用上存在一些限制。这些限制因设备而异，这会使代码设计复杂化。

  • 要确定设备上的全局内存量，请使用clGetDeviceInfo，将CL_DEVICE_GLOBAL_MEM_SIZE作为param_name传入，或检查AMD SDK示例中的CLIinfo的输出。
  • 全局内存应被视为“流式内存”。这意味着当流式传输连续内存地址或能够利用内存子系统全部带宽的内存访问模式时，将获得最佳性能。此类内存操作称为合并（coalesced）。
    • AMD OpenCL编程指南指出，“对于最优的合并内存模式，GPU优化算法的一个常见访问模式是，同一波前（wavefront）中的工作项从同一缓存行访问内存位置。”（波前描述了N个硬件线程，它们并行执行相同的指令。这与工作组不同，工作组是用于对线程进行分组的软件抽象。）
    • 合并内存访问的规则因设备而异，并且可能非常复杂，如此处讨论的那样。
  • 某些内存访问模式可能导致银行冲突，这会极大地降低应用程序性能。（出于性能考虑，内存子系统被组织成银行，以将流式带宽提高与银行数量相关的倍数。当多个线程尝试同时访问同一内存银行时，会发生银行冲突。在这种情况下，每个银行一次只能服务单个请求，这会导致顺序减速，从而严重影响性能。）请注意，银行冲突非常依赖于设备。
    • 某些硬件会极大地加速特殊的访问情况，但在其他设备上可能会导致性能不佳。
      • 从全局内存进行的广播读取在不同设备上的表现可能好或差。例如，此演示文稿的幻灯片69指出，ATI Radeon HD 5870 GPU具有特殊的硬件来执行广播，当许多线程尝试从全局内存的公共位置读取时。这种内存访问在Fermi之前的NVIDIA GPU上表现不佳。
        • 如果可能，请使用常量内存来实现跨几乎所有GPU类型的广播操作的良好性能。
      • 归约操作在许多应用程序中都很常见。因此，高性能归约是人们高度关注的主题。以下是一项研究，讨论了AMD如何设计其硬件以在ATI Radeon HD 5870 GPU上支持从全局内存的快速归约。

私有内存：这是工作项内部使用的内存，类似于GPU多处理器或CPU核心中的寄存器。

  • 私有内存速度很快，无需同步原语即可使用。它在编译时由JIT编译器为给定的内核和卡分配和分区。
  • 正如此线程中所讨论的，私有内存的位置和大小未在OpenCL规范中定义。这种模糊性使得难以决定使用多少私有内存，除了“尽可能少”之外，还可以使用基准测试来定义设备的适当用量。
  • 在某些设备（如GPU）上使用过多的私有内存会导致应用程序性能急剧下降，因为它会溢出到较慢的内存。根据设备的不同，私有内存可以溢出到缓存内存。没有缓存内存的GPU会将私有内存溢出到全局内存，导致性能大幅下降。

本地内存：OpenCL本地内存比全局内存快得多——通常在芯片上。

  • 要确定设备上的本地内存量，请使用clGetDeviceInfo，并将CL_DEVICE_LOCAL_MEM_SIZE作为param_name传入，或检查AMD SDK示例中的CLIinfo的输出。
  • 本地内存用于实现合并访问、在工作组中的工作项之间共享数据以及减少对低带宽全局内存的访问。

常量内存：常量内存顾名思义，是一个只读内存区域。

  • 在NVIDIA的GPU设备上，常量内存是一个专门的内存区域，适用于广播操作。在AMD设备上，这是一个利用硬件优化来广播数据的全局内存区域。
  • 要确定常量内存缓冲区的大小，请使用clGetDeviceInfo并传入CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE，或检查AMD SDK示例中的CLIinfo的输出。
  • 可移植代码必须保守地假设，在函数内部或程序范围内用__constant限定符声明的每个变量都算作一个单独的常量参数。OpenCL 1.1的在线Khronos文档指出，实现不要求将这些声明聚合到最少数量的常量参数中。

额外资源

以下是OpenCL地址限定符（这与访问限定符不同）

  • __global：从全局地址空间分配的内存，图像默认是全局的。
  • __constant：一个只读内存区域。
  • __local：由工作组共享的内存。
  • __private：每个工作项私有的内存。

注意：内核参数必须是__global、__constant或__local。使用不同地址限定符强制转换的指针是未定义的。

示例

本文将使用cl.hpp——OpenCL 1.1 C++绑定头文件，它实现了C++包装器API，以创建更短、更简洁的程序。相比之下，本系列中的第一个示例使用了cl.h，它定义了OpenCL 1.1核心API。

以下代码second.cpp从文件中读取OpenCL内核源代码，构建它，运行它，然后对结果进行双重检查。默认行为是显示内核运行所需的时间，这使得比较成为可能。由于OpenCL内核源代码在一个文件中，因此很容易尝试不同的内核来理解各种内存类型在CPU和GPU硬件上的性能影响。然而，这个示例提供了一个可用且易于泛化的框架来试验其他数据集和计算内核。

进入main()后，解析命令行参数。

int main(int argc, char* argv[])
{
  int seed=4;
 
  if( argc < 4) {
    cerr 
      << "Use: {cpu|gpu} kernelFile n_iter vectorSize (k)" 
      << 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];
  unsigned int n_iter = atoi(argv[3]);
  unsigned int vecLen = 1000 * atoi(argv[4]);

main()中的其余代码，在一个try{...}catch块的范围内执行构建和运行内核的所有剩余工作，以处理任何异常。请注意

  • PROFILING默认已定义。
  • 生成数据、加载和构建内核以及执行黄金检查的工作由类OclTest执行。
  • C++包装器API用于以比本系列上一教程中讨论的first.cpp（使用了核心OpenCL API）更少的源代码行完成更多工作。同样，源代码更易读。
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, n_iter, vecLen);
 
    cl::Event event;
    test.initData(queue, event, seed);
 
    queue.enqueueNDRangeKernel(test.getKernel(), 
			       cl::NullRange, cl::NDRange(vecLen),
			       cl::NDRange(1, 1), 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;
}

类OclTest将执行测试所需的工作分解开。可以看出

  1. 构造函数。OclTest(…)
    • 从源文件构建内核。
    • 分配三个主机向量，创建设备向量，并将设备向量分配给内核。
  2. 方法initData(…)
    • 用随机数据填充主机向量。
    • 将h_vecA和h_vecB从主机传输到设备的传输排入队列。
  3. 方法getKernel()返回用于在CommandQueue上排队的内核。
  4. 在main(…)将内核排队后，调用goldenTest(…)方法。此方法等待命令队列完成，然后双重检查设备上的计算是否正确。
class OclTest {
private:
  int nIter, vLen,vSize;
  cl::Kernel kernel;
  int  *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 n_iter, int vecLen) {
    nIter = n_iter;
    vLen = vecLen;
    vSize = vLen * sizeof(int);
    
    // 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 {
      program.build(devices);
    } 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);
    }
    kernel = cl::Kernel(program, "vec_iii_1d");
    
    // set up the kernel inputs
    h_vecA = new int[vLen];
    h_vecB = new int[vLen];
    h_vecC = new int[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);
    for(int i=0; i < vLen; i++) h_vecA[i] = rand();
    for(int i=0; i < vLen; i++) h_vecB[i] = abs(rand()) % nIter;
    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); }
 
  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 i=0; i < vLen; i++) 
      if(h_vecC[i] != h_vecA[i] + h_vecB[i]) 
	return(1);
    return(0);
  }
};

second.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>

#include <string>

using namespace std;
 
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";
}
 
class OclTest {
private:
  int nIter, vLen,vSize;
  cl::Kernel kernel;
  int  *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 n_iter, int vecLen) {
    nIter = n_iter;
    vLen = vecLen;
    vSize = vLen * sizeof(int);
    
    // 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 {
      program.build(devices);
    } 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);
    }
    kernel = cl::Kernel(program, "vec_iii_1d");
    
    // set up the kernel inputs
    h_vecA = new int[vLen];
    h_vecB = new int[vLen];
    h_vecC = new int[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);
    for(int i=0; i < vLen; i++) h_vecA[i] = rand();
    for(int i=0; i < vLen; i++) h_vecB[i] = abs(rand()) % nIter;
    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); }
 
  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 i=0; i < vLen; i++) 
      if(h_vecC[i] != h_vecA[i] + h_vecB[i]) 
	return(1);
    return(0);
  }
};
 
int main(int argc, char* argv[])
{
  int seed=4;
 
  if( argc < 4) {
    cerr 
      << "Use: {cpu|gpu} kernelFile n_iter vectorSize (k)" 
      << 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];
  unsigned int n_iter = atoi(argv[3]);
  unsigned int vecLen = 1000 * atoi(argv[4]);
 
  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, n_iter, vecLen);
 
    cl::Event event;
    test.initData(queue, event, seed);
 
    queue.enqueueNDRangeKernel(test.getKernel(), 
			       cl::NullRange, cl::NDRange(vecLen),
			       cl::NDRange(1, 1), 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下构建可执行文件非常直接

  1. 复制源代码并将其放在一个名为second.cpp的文件中。
  2. 设置OpenCL主目录的环境变量。以下是使用bash shell时的一个可能示例
    export OCL_HOME=../ati-stream-sdk-v2.2-lnx64
  3. 编译second.cpp:
    g++ -I $OCL_HOME/include -L
$OCL_HOME/lib/x86_64 second.cpp -l OpenCL –o second
  4. 指定共享库的搜索路径
export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:$OCL_HOME/lib/x86_64

本教程中将要进行的整数加法测试的OpenCL内核相当直接。基本上，将d_vecA加到d_vecB，并将结果放入d_vecC。然而，测试将通过使用for循环将值加1而不是执行加法来演示私有内存和全局内存之间的性能差异。使用了整数运算以获得精确结果。

以下是try_slow_iii.cl内核。（“iii”表示该内核使用三个整数向量）。此内核速度很慢，因为处理器必须在全局内存中原地递减b数组。

__kernel void vec_iii_1d(__global int *a, __global int *b, 
                         __global int *c)
{
  size_t tid = get_global_id(0);
  c[tid] = a[tid];
  while(b[tid] > 0) {
    b[tid]--;
    c[tid]++;
  }
}

相比之下，try_fast_iii.cl内核使用私有内存来加速增量和减量操作。具体来说，变量tmp包含每个tid值的a向量的值。类似地，私有变量i加载了b的值。

__kernel void vec_iii_1d(__global int *a, __global int *b,
                         __global int *c)
{
  size_t tid = get_global_id(0);
  int tmp = a[tid];
  for(int i=b[tid]; i > 0; i--) tmp++;
  c[tid] = tmp;
}

以下显示了如何使用可执行文件以及使用ATI Radeon HD 5870 GPU和AMD Phenom™ II X6 1055T处理器生成的一些计时结果

./second gpu try_slow_iii.cl 100 10000
Platform number is: 1
device Type GPU
Platform is by: Advanced Micro Devices, Inc.
Time for kernel to execute 11.1207
test passed
 
./second gpu try_fast_iii.cl 100 10000
Platform number is: 1
device Type GPU
Platform is by: Advanced Micro Devices, Inc.
Time for kernel to execute 0.0445558
test passed
 
./second cpu try_fast_iii.cl 100 10000
Platform number is: 1
device Type CPU
Platform is by: Advanced Micro Devices, Inc.
Time for kernel to execute 23.7571
test passed
 
./second cpu try_slow_iii.cl 100 10000
Platform number is: 1
device Type CPU
Platform is by: Advanced Micro Devices, Inc.
Time for kernel to execute 23.5225
test passed

计时结果显示了在GPU上使用私有内存与全局内存之间的显着性能差异。CPU运行表明，AMD Phenom II X6 1055T处理器在此测试中不受全局内存与私有内存之间差异的显着影响。总的来说，GPU快得多。

  try_fast_iii.cl try_slow_iii.cl
GPU 0.0445558 11.1207
CPU（所有核心） 23.7571 23.5225

摘要

令人兴奋的是，OpenCL在技术社区中越来越普及。例如，最新的Ubuntu Linux发行版Maverick Meerkat（10.10版）增加了OpenCL支持。（AMD硬件用户可能需要应用一个热修复来纠正最近的Linux安全更新。请注意，驱动程序下载近100 MB。）

专家级地利用OpenCL内存空间来分区数据和本地化计算工作，是区分优秀OpenCL开发者与其他开发者的三大特性之一，可能是*最*具决定性的特性。这些人是能够创建高性能应用程序的程序员，这些应用程序可以扩展到大量线程和许多设备。他们也是同样的人，通过理解给定应用程序域的OpenCL设备性能，可以创建*可移植*的高性能应用程序。正如本系列文章第一篇中所指出的，平衡比有助于量化特定应用程序或作业组合所需的设备特性。正如非常清楚的，主机与OpenCL设备之间的带宽，以及设备的内存带宽，是大多数应用程序的关键性能指标。

最后，本教程提供了OpenCL C++包装器API的快速基本介绍，以及几个可以在CPU和GPU设备类型上运行的示例内核。这个框架非常通用，允许测试各种其他OpenCL内核。

OpenCL 和 OpenCL 徽标是 Apple Inc. 的商标，经 Khronos 许可使用。

© . All rights reserved.