AMD OpenCL 大学课程

<span style="font-family: Consolas; color: rgb(0, 0, 0);">AMD OpenCL大学课程是非常好的入门级OpenCL教程，通过看教程中的PPT，我们能够很快的了解OpenCL机制以及编程方法。下载地址：<a target=_blank target="_blank" title="http://developer.amd.com/zones/OpenCLZone/universities/Pages/default.aspx" href="http://developer.amd.com/zones/OpenCLZone/universities/Pages/default.aspx" style="color: rgb(51, 102, 153); text-decoration: none;">http://developer.amd.com/zones/OpenCLZone/universities/Pages/default.aspx</a></span>

<span style="font-family: Consolas; color: rgb(0, 0, 0);">     教程中的英文很简单，我相信学OpenCL的人都能看得懂，而且看原汁原味的英文表述，更有利于我们了解各种术语的来龙去脉。</span>

<span style="font-family: Consolas; color: rgb(0, 0, 0);">     我把这些教程翻译成自己的中文表述，主要是强化理解需要，其实我的英文很烂。</span>

<span style="font-family: Consolas; color: rgb(0, 0, 0);">一、并行计算概述</span>

<span style="font-family: Consolas; color: rgb(0, 0, 0);">    在计算机术语中，并行性是指：<span style="color: rgb(155, 0, 211);">把一个复杂问题,分解成多个能同时处理的子问题的能力</span>。要实现并行计算，首先我们要有物理上能够实现并行计算的硬件设备，比如多核CPU，每个核能同时实现算术或逻辑运算。</span>

<span style="font-family: Consolas; color: rgb(0, 0, 0);">    通常，我们通过GPU实现两类并行计算：</span>

<span style="font-family: Consolas; color: rgb(0, 0, 0);">      任务并行：<span style="color: rgb(0, 0, 255);">把一个问题分解为能够同时执行的多个任务</span>。</span>

<span style="font-family: Consolas; color: rgb(0, 0, 0);">      数据并行：<span style="color: rgb(0, 0, 255);">同一个任务内，它的各个部分同时执行</span>。</span>

<span style="font-family: Consolas; color: rgb(0, 0, 0);">   下面我们通过一个农场主雇佣工人摘苹果的例子来描述不同种类的并行计算。</span>

<a target=_blank target="_blank" href="http://images.cnblogs.com/cnblogs_com/mikewolf2002/201201/201201301920165978.png" style="color: rgb(51, 102, 153); text-decoration: none;"><img title="image" alt="image" src="http://images.cnblogs.com/cnblogs_com/mikewolf2002/201201/201201301920189174.png" border="0" height="176" width="167" style="border: 0px none; max-width: 100%; padding-left: 0px; padding-right: 0px; display: inline; padding-top: 0px;" /></a>

• 摘苹果的工人就是硬件上的并行处理单元（process elements)。

• 树就是要执行的任务。

• 苹果就是要处理的数据。

  串行的任务处理就如下图所示，一个工人背着梯子摘完所有树上的苹果（<span style="color: rgb(0, 0, 255);">一个处理单元处理完所有任务的数据</span>）。

    数据并行就好比农场主雇佣了好多工人来摘完一个树上的苹果（多个处理单元并行完成一个任务中的数据），这样就能很快摘完一颗树上的苹果。

   农场主也可以为每棵树安排一个工人，这就好比任务并行。在每个任务内，由于只有一个工人，所以是串行执行的，但任务之间是并行的。

对一个复杂问题，影响并行计算的因素很多。通常，我们都是通过分解问题的方式来实施并算法行。

这又包括两方面内容：

• 任务分解：把算法分解成很多的小任务，就像前面的例子中，把果园按苹果树进行划分，这时我们并不关注数据，也就是说不关注每个树上到底有多少个苹果。
• 数据分解：就是把很多数据，分成不同的、离散的小块，这些数据块能够被并行执行，就好比前面例子中的苹果。

   通常我们按照算法之间的依赖关系来分解任务，这样就形成了一个任务关系图。一个任务只有没有依赖任务的时候，才能够被执行。

    这有点类似于数据结构中的有向无环图，两个没有连通路径的任务之间可以并行执行。下面再给一个烤面包的例子，如果所示，预热烤箱和购买面粉糖两个任务之间可以并行执行。

  

     对大多数科学计算和工程应用来说，数据分解一般都是基于输出数据,例如：

• 在一副图像中，对一个滑动窗口（例如：3*3像素)内的像素实施滤波操作，可以得到一个输出像素的卷积。
• 第一个输入矩阵的第i行乘以第二个输入矩阵的第j列，得到的向量和即为输出矩阵第i行，第j列的元素。

这种方法对于输入和输出数据是一对一，或者多对一的对应关系比较有效。

    也有的数据分解算法是基于输入数据的，这时，输入数据和输出数据一般是一对多的关系，比如求图像的直方图，我们要把每个像素放到对应的槽中(bins，对于灰度图，bin数量通常是256）。一个搜索函数，输入可能是多个数据，输出却只有一个值。对于这类应用，我们一般用每个线程计算输出的一部分，然后通过同步以及原子操作得到最终的值，OpenCL中求最小值的kernel函数就是典型代表[可以看下ATI Stream Computing OpenCL programming guide第二章中求最小值的kernel例子]。

     通常来说，怎样分解问题和具体算法有关，而且还要考虑自己使用的硬件和软件，比如AMD GPU平台和Nvdia GPU平台的优化就有很多不同。

二、常用基于硬件和软件的并行

    在上个实际90年代，并行计算主要研究如何在cpu上实施指自动的指令级并行。

• 同时发射多条指令（之间没有依赖关系），并行执行这些指令。
• 在本教程中，我么不讲述自动的硬件级并行，感兴趣的话，可以看看计算机体系结构的教程。

    高层的并行，比如线程级别的并行，一般很难自动化，需要程序员告诉计算机，该做什么，不该做什么。这时，程序员还要考虑硬件的具体指标，通常特定硬件都是适应于某一类并行编程，比如多核cpu就适合基于任务的并行编程，而GPU更适应于数据并行编程。

Hardware type	Examples	Parallelism
Multi-core superscalar processors	Phenom II CPU	Task
Vector or SIMD processors	SSE units (x86 CPUs)	Data
Multi-core SIMD processors	Radeon 5870 GPU	Data

 

    现代的GPU有很多独立的运算核（processor）组成，在AMD GPU上就是stream core，这些core能够执行SIMD操作（单指令，多数据），所以特别适合数据并行操作。通常GPU上执行一个任务，都是把任务中的数据分配到各个独立的core中执行。

     在GPU上，我们一般通过循环展开，Loop strip mining 技术，来把串行代码改成并行执行的。比如在CPU上，如果我们实现一个向量加法，代码通常如下：

 1: for(i = 0; i < n; i++)

 2: {

 3: C[i] = A[i] + B[i];

 4: }

在GPU上，我们可以设置n个线程，每个线程执行一个加法，这样大大提高了向量加法的并行性。

 1: __kernel void VectorAdd(__global const float* a, __global const float* b, __global float* c, int n)

 2: {

 3: int i = get_global_id(0);

 4: c[i] = a[i] + b[i];

 5: }

    上面这个图展示了向量加法的SPMD（单指令多线程）实现，从图中可以看出如何实施Loop strip mining 操作的。

    GPU的程序一般称作Kernel程序，它是一种SPMD的编程模型（the Single Program Multiple Data ）。SPMD执行同一段代码的多个实例，每个实例对数据的不同部分进行操作。

     在数据并行应用中，用loop strip mining来实现SPMD是最常用的方法：

• • 在分布式系统中，我们用Message Passing Interface (MPI)来实现SPMD。
  • 在共享内存并行系统中，我们用POSIX线程来实现SPMD。
  • 在GPU中，我们就是用Kernel来显现SPMD。

    在现代的CPU上，创建一个线程的开销还是很大的，如果要在CPU上实现SPMD，每个线程处理的数据块就要尽量大点，做更多的事情，以便减少平均线程开销。但在GPU上，都是轻量级的线程，创建、调度线程的开销比较小，所以我们可以做到把循环完全展开，一个线程处理一个数据。

GPU上并行编程的硬件一般称作SIMD。通常，发射一条指令后，它要在多个ALU单元中执行（ALU的数量即使simd的宽度），这种设计减少了控制流单元以级ALU相关的其他硬件数量。

SIMD的硬件如下图所示：

 

 

    在向量加法中，宽度为4的SIMD单元，可以把整个循环分为四个部分同时执行。在工人摘苹果的例子中，工人的双手类似于SIMD的宽度为2。另外，我们要知道，现在的GPU硬件上都是基于SIMD设计，GPU硬件隐式的把SPMD线程映射到SIMD core上。对开发有人员来说，我们并不需要关注硬件执行结果是否正确，我们只需要关注它的性能就OK了。

    CPU一般都支持并行级的原子操作，这些操作保证不同的线程读写数据，相互之间不会干扰。有些GPU支持系统范围的并行操作，但会有很大开销，比如Global memory的同步。

1、OpenCL架构

   OpenCL可以实现混合设备的并行计算，这些设备包括CPU，GPU，以及其它处理器，比如Cell处理器，DSP等。使用OpenCL编程，可以实现可移植的并行加速代码。[但由于各个OpenCL device不同的硬件性能，可能对于程序的优化还要考虑具体的硬件特性]。

   通常OpenCL架构包括四个部分：

• 平台模型(Platform Model)
• 执行模型(Execution Model)
• 内存模型(Memory Model)
• 编程模型(Programming Model)

2、OpenCL平台模型

   不同厂商的OpenCL实施定义了不同的OpenCL平台，通过OpenCL平台，主机能够和OpenCL设备之间进行交互操作。现在主要的OpenCL平台有AMD、Nvida，Intel等。OpenCL使用了一种Installable Client Driver模型，这样不同厂商的平台就能够在系统中共存。在我的计算机上就安装有AMD和Intel两个OpenCL Platform[现在的OpenCL driver模型不允许不同厂商的GPU同时运行]。

    OpenCL平台通常包括一个主机(Host)和多个OpenCL设备（device），每个OpenCL设备包括一个或多个CU(compute units)，每个CU包括又一个或多个PE（process element）。 每个PE都有自己的程序计数器(PC)。主机就是OpenCL运行库宿主设备，在AMD和Nvida的OpenCL平台中，主机一般都指x86 CPU。

   对AMD平台来说，所有的CPU是一个设备，CPU的每一个core就是一个CU，而每个GPU都是独立的设备。

  

3、OpenCL编程的一般步骤

  下面我们通过一个实例来了解OpenCL编程的步骤，假设我们用的是AMD OpenCL平台（因为本人的GPU是HD5730），安装了AMD Stream SDK 2.6,并在VS2008中设置好了include，lib目录等。

    首先我们建立一个控制台程序，最初的代码如下：

 1: #include "stdafx.h"

 2: #include <CL/cl.h>

 3: #include <stdio.h>

 4: #include <stdlib.h>

 5: 

 6: #pragma comment (lib,"OpenCL.lib")

 7: 

 8: int main(int argc, char* argv[])

 9: {

 10: return 0;

 11: }

 

第一步，我们要选择一个OpenCL平台，所用的函数就是

    通常，这个函数要调用2次，第一次得到系统中可使用的平台数目，然后为（Platform）平台对象分配空间，第二次调用就是查询所有的平台，选择自己需要的OpenCL平台。代码比较长，具体可以看下AMD Stream SDK 2.6中的TemplateC例子，里面描述如何构建一个robust的最小OpenCL程序。为了简化代码，使程序看起来不那么繁琐，我直接调用该函数，选取系统中的第一个OpenCL平台，我的系统中安装AMD和Intel两家的平台，第一个平台是AMD的。另外，我也没有增加错误检测之类的代码，但是增加了一个status的变量，通常如果函数执行正确，返回的值是0。

 1: #include "stdafx.h"

 2: #include <CL/cl.h>

 3: #include <stdio.h>

 4: #include <stdlib.h>

 5: 

 6: #pragma comment (lib,"OpenCL.lib")

 7: 

 8: int main(int argc, char* argv[])

 9: {

 10: cl_uint status;

 11: cl_platform_id platform;

 12: 

 13: status = clGetPlatformIDs( 1, &platform, NULL );

 14: 

 15: return 0;

 16: }

第二步是得到OpenCL设备，

     这个函数通常也是调用2次，第一次查询设备数量，第二次检索得到我们想要的设备。为了简化代码，我们直接指定GPU设备。

 

 1: #include "stdafx.h"

 2: #include <CL/cl.h>

 3: #include <stdio.h>

 4: #include <stdlib.h>

 5: 

 6: #pragma comment (lib,"OpenCL.lib")

 7: 

 8: int main(int argc, char* argv[])

 9: {

 10: cl_uint status;

 11: cl_platform_id platform;

 12: 

 13: status = clGetPlatformIDs( 1, &platform, NULL );

 14: 

 15: cl_device_id device;

 16: 

 17: clGetDeviceIDs( platform, CL_DEVICE_TYPE_GPU,

 18: 1,

 19: &device,

 20: NULL);

 21: 

 22: return 0;

 23: }

下面我们来看下OpenCL中Context的概念：

通常，Context是指管理OpenCL对象和资源的上下文环境。为了管理OpenCL程序，下面的一些对象都要和Context关联起来：

 

—设备（Devices）:执行Kernel程序对象。

—程序对象（Program objects）: kernel程序源代码

—Kernels:运行在OpenCL设备上的函数。

—内存对象（Memory objects）: device处理的数据对象。

—命令队列（Command queues）: 设备之间的交互机制。

•  

注意：创建一个Context的时候，我们必须把一个或多个设备和它关联起来。对于其它的OpenCL资源，它们创建时候，也要和Context关联起来，一般创建这些资源的OpenCL函数的输入参数中，都会有context。

这个函数中指定了和context关联的一个或多个设备对象,properties参数指定了使用的平台，如果为NULL,厂商选择的缺省值被使用，这个函数也提供了一个回调机制给用户提供错误报告。

现在的代码如下：

 1: #include "stdafx.h"

 2: #include <CL/cl.h>

 3: #include <stdio.h>

 4: #include <stdlib.h>

 5: 

 6: #pragma comment (lib,"OpenCL.lib")

 7: 

 8: int main(int argc, char* argv[])

 9: {

 10: cl_uint status;

 11: cl_platform_id platform;

 12: 

 13: status = clGetPlatformIDs( 1, &platform, NULL );

 14: 

 15: cl_device_id device;

 16: 

 17: clGetDeviceIDs( platform, CL_DEVICE_TYPE_GPU,

 18: 1,

 19: &device,

 20: NULL);

 21: cl_context context = clCreateContext( NULL,

 22: 1,

 23: &device,

 24:                

 25: 

 26: return 0;

 27: }

接下来，我们要看下命令队列。在OpenCL中，命令队列就是主机的请求，在设备上执行的一种机制。

• 在Kernel执行前，我们一般要进行一些内存拷贝的工作，比如把主机内存中的数据传输到设备内存中。

另外要注意的几点就是：对于不同的设备，它们都有自己的独立的命令队列；命令队列中的命令(kernel函数）可能是同步的，也可能是异步的，它们的执行顺序可以是有序的，也可以是乱序的。

命令队列在device和context之间建立了一个连接。

命令队列properties指定以下内容:

• 是否乱序执行（在AMD GPU中，好像现在还不支持乱序执行）
• 是否启动profiling。Profiling通过事件机制来得到kernel执行时间等有用的信息，但它本身也会有一些开销。

 

如下图所示，命令队列把设备和context联系起来，尽管它们之间不是物理连接。

添加命令队列后的代码如下：

 1: #include "stdafx.h"

 2: #include <CL/cl.h>

 3: #include <stdio.h>

 4: #include <stdlib.h>

 5: 

 6: #pragma comment (lib,"OpenCL.lib")

 7: 

 8: int main(int argc, char* argv[])

 9: {

 10: cl_uint status;

 11: cl_platform_id platform;

 12: 

 13: status = clGetPlatformIDs( 1, &platform, NULL );

 14: 

 15: cl_device_id device;

 16: 

 17: clGetDeviceIDs( platform, CL_DEVICE_TYPE_GPU,

 18: 1,

 19: &device,

 20: NULL);

 21: cl_context context = clCreateContext( NULL,

 22: 1,

 23: &device,

 24: NULL, NULL, NULL);

 25: 

 26: cl_command_queue queue = clCreateCommandQueue( context,

 27: device,

 28: CL_QUEUE_PROFILING_ENABLE, NULL );

 29: 

 30: return 0;

 31: }

 

OpenCL内存对象：

    OpenCL内存对象就是一些OpenCL数据，这些数据一般在设备内存中，能够被拷入也能够被拷出。OpenCL内存对象包括buffer对象和image对象。

buffer对象：连续的内存块----顺序存储，能够通过指针、行列式等直接访问。

image对象：是2维或3维的内存对象，只能通过read_image() 或 write_image()来读取。image对象可以是可读或可写的，但不能同时既可读又可写。

    该函数会在指定的context上创建一个buffer对象，image对象相对比较复杂，留在后面再讲。

flags参数指定buffer对象的读写属性，host_ptr可以是NULL，如果不为NULL，一般是一个有效的host buffer对象，这时，函数创建OpenCL buffer对象后，会把对应host buffer的内容拷贝到OpenCL buffer中。

     在Kernel执行之前，host中原始输入数据必须显式的传到device中，Kernel执行完后，结果也要从device内存中传回到host内存中。我们主要通过函数clEnqueue{Read|Write}{Buffer|Image}来实现这两种操作。从host到device，我们用clEnqueueWrite，从device到host，我们用clEnqueueRead。clEnqueueWrite命令包括初始化内存对象以及把host 数据传到device内存这两种操作。当然，像前面一段说的那样，也可以把host buffer指针直接用在CreateBuffer函数中来实现隐式的数据写操作。

     这个函数初始化OpenCL内存对象，并把相应的数据写到OpenCL内存关联的设备内存中。其中，blocking_write参数指定是数拷贝完成后函数才返回还是数据开始拷贝后就立即返回（阻塞模式于非阻塞模式）。Events参数指定这个函数执行之前，必须要完成的Event（比如先要创建OpenCL内存对象的Event）。

 

OpenCL程序对象：

   程序对象就是通过读入Kernel函数源代码或二进制文件，然后在指定的设备上进行编译而产生的OpenCL对象。

     这个函数通过源代码(strings)，创建一个程序对象，其中counts指定源代码串的数量，lengths指定源代码串的长度（为NULL结束的串时，可以省略）。当然，我们还必须自己编写一个从文件中读取源代码串的函数。

     对context中的每个设备，这个函数编译、连接源代码对象，产生device可以执行的文件，对GPU而言就是设备对应shader汇编。如果device_list参数被提供，则只对这些设备进行编译连接。options参数主要提供一些附加的编译选项，比如宏定义、优化开关标志等等。

     如果程序编译失败，我们能够根据返回的状态，通过调用clGetProgramBuildInfo来得到错误信息。

加上创建内存对象以及程序对象的代码如下:

 1: 

 2: #include "stdafx.h"

 3: #include <CL/cl.h>

 4: #include <stdio.h>

 5: #include <stdlib.h>

 6: #include <time.h>

 7: #include <iostream>

 8: #include <fstream>

 9: 

 10: using namespace std;

 11: #define NWITEMS 262144

 12: 

 13: #pragma comment (lib,"OpenCL.lib")

 14: 

Kernel对象：

    Kernel就是在程序代码中的一个函数，这个函数能在OpenCL设备上执行。一个Kernel对象就是kernel函数以及其相关的输入参数。

 

Kernel对象通过程序对象以及指定的函数名字创建。注意：函数必须是程序源代码中存在的函数。

运行时编译：

    在运行时，编译程序和创建kernel对象是有时间开销的，但这样比较灵活，能够适应不同的OpenCL硬件平台。程序动态编译一般只需一次，而Kernel对象在创建后，可以反复调用。

 

创建Kernel后，运行Kernel之前，我们还要为Kernel对象设置参数。我们可以在Kernel运行后，重新设置参数再次运行。

arg_index指定该参数为Kernel函数中的第几个参数(比如第一个参数为0，第二个为1,…)。内存对象和单个的值都可以作为Kernel参数。下面是2个设置Kernel参数的例子：

clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&d_iImage);

clSetKernelArg(kernel, 1, sizeof(int), (void*)&a);

在Kernel运行之前，我们先看看OpenCL中的线程结构：

大规模并行程序中，通常每个线程处理一个问题的一部分，比如向量加法，我们会把两个向量中对应的元素加起来，这样，每个线程可以处理一个加法。

下面我看一个16个元素的向量加法：两个输入缓冲A、B，一个输出缓冲C

在这种情况下，我们可以创建一维的线程结构去匹配这个问题。

每个线程把自己的线程id作为索引，把相应元素加起来。

     OpenCL中的线程结构是可缩放的，Kernel的每个运行实例称作WorkItem(也就是线程），WorkItem组织在一起称作WorkGroup，OpenCL中，每个Workgroup之间都是相互独立的。

通过一个global id(在索引空间，它是唯一的）或者一个workgroup id和一个work group内的local id，我就能标定一个workitem。

在kernel函数中，我们能够通过API调用得到global id以及其他信息：

get_global_id(dim)

get_global_size(dim)

这两个函数能得到每个维度上的global id。

get_group_id(dim)

get_num_groups(dim)

get_local_id(dim)

get_local_size(dim）

这几个函数用来计算group id以及在group内的local id。

get_global_id(0) = column, get_global_id(1) = row

get_num_groups(0) * get_local_size(0) == get_global_size(0)

 

OpenCL内存模型

    OpenCL的内存模型定义了各种各样内存类型，各种内存模型之间有层级关系。各种内存之间的数据传输必须是显式进行的，比如从host memory到device memory，从global memory到local memory等等。

    WorkGroup被映射到硬件的CU上执行（在AMD 5xxx系列显卡上，CU就是simd，一个simd中有16个pe，或者说是stream core），OpenCL并不提供各个workgroup之间的一致性，如果我们需要在各个workgroup之间共享数据或者通信之类的，要自己通过软件实现。

Kernel函数的写法

每个线程（workitem）都有一个kenerl函数的实例。下面我们看下kernel的写法：

 1: __kernel void vecadd(__global const float* A, __global const float* B, __global float* C)

 2: {

 3: int id = get_global_id(0);

 4: C[id] = A[id] + B[id];

 5: }

每个Kernel函数都必须以__kernel开始，而且必须返回void。每个输入参数都必须声明使用的内存类型。通过一些API，比如get_global_id之类的得到线程id。

内存对象地址空间标识符有以下几种：

__global – memory allocated from global address space

__constant – a special type of read-only memory

__local – memory shared by a work-group

__private – private per work-item memory

__read_only/__write_only – used for images

Kernel函数参数如果是内存对象，那么一定是__global,__local或者constant。

 

运行Kernel

   首先要设置线程索引空间的维数以及workgroup大小等。

   我们通过函数clEnqueueNDRangeKerne把Kernel放在一个队列里，但不保证它马上执行，OpenCL driver会管理队列，调度Kernel的执行。注意：每个线程执行的代码都是相同的，但是它们执行数据却是不同的。

 

   该函数把要执行的Kernel函数放在指定的命令队列中，globald大小（线程索引空间）必须指定，local大小（work group）可以指定，也可以为空。如果为空，则系统会自动根据硬件选择合适的大小。event_wait_list用来选定一些events，只有这些events执行完后，该kernel才可能被执行，也就是通过事件机制来实现不同kernel函数之间的同步。

   当Kernel函数执行完毕后，我们要把数据从device memory中拷贝到host memory中去。

释放资源：

    大多数的OpenCL资源都是指针，不使用的时候需要释放掉。当然，程序关闭的时候这些对象也会被自动释放掉。

    释放资源的函数是：clRelase{Resource} ，比如: clReleaseProgram(), clReleaseMemObject()等。

 

错误捕捉：

    如果OpenCL函数执行失败，会返回一个错误码，一般是个负值，返回0则表示执行成功。我们可以根据该错误码知道什么地方出错了，需要修改。错误码在cl.h中定义，下面是几个错误码的例子.

CL_DEVICE_NOT_FOUND -1

CL_DEVICE_NOT_AVAILABLE -2

CL_COMPILER_NOT_AVAILABLE -3

CL_MEM_OBJECT_ALLOCATION_FAILURE -4

…

下面是一个OpenCL机制的示意图

程序模型

    数据并行：work item和内存对象元素之间是一一映射关系；workgroup可以显示指定，也可以隐式指定。

    任务并行：kernel的执行独立于线程索引空间；用其他方法表示并行，比如把不同的任务放入队列，用设备指定的特殊的向量类型等等。

    同步：workgroup内work item之间的同步；命令队列中不同命令之间的同步。

完整代码如下：

 1: #include "stdafx.h"

 2: #include <CL/cl.h>

 3: #include <stdio.h>

 4: #include <stdlib.h>

 5: #include <time.h>

 6: #include <iostream>

 7: #include <fstream>

 8: 

 9: using namespace std;

 10: #define NWITEMS 262144

 11: 

 12: #pragma comment (lib,"OpenCL.lib")

 13: 

 14: //把文本文件读入一个string中

 15: int convertToString(const char *filename, std::string& s)

 16: {

 17: size_t size;

 18: char* str;

 19: 

 20: std::fstream f(filename, (std::fstream::in | std::fstream::binary));

 21: 

 22: if(f.is_open())

 23: {

 24: size_t fileSize;

 25: f.seekg(0, std::fstream::end);

 26: size = fileSize = (size_t)f.tellg();

 27: f.seekg(0, std::fstream::beg);

 28: 

 29: str = new char[size+1];

 30: if(!str)

 31: {

 32: f.close();

 33: return NULL;

 34: }

 35: 

 36: f.read(str, fileSize);

 37: f.close();

 38: str[size] = '\0';

 39: 

 40: s = str;

 41: delete[] str;

 42: return 0;

 43: }

 44: printf("Error: Failed to open file %s\n", filename);

 45: return 1;

 46: }

 47: 

 48: int main(int argc, char* argv[])

 49: {

 50: //在host内存中创建三个缓冲区

 51: float *buf1 = 0;

 52: float *buf2 = 0;

 53: float *buf = 0;

 54: 

 55: buf1 =(float *)malloc(NWITEMS * sizeof(float));

 56: buf2 =(float *)malloc(NWITEMS * sizeof(float));

 57: buf =(float *)malloc(NWITEMS * sizeof(float));

 58: 

 59: //初始化buf1和buf2的内容

 60: int i;

 61: srand( (unsigned)time( NULL ) );

 62: for(i = 0; i < NWITEMS; i++)

 63: buf1[i] = rand()%65535;

 64: 

 65: srand( (unsigned)time( NULL ) +1000);

 66: for(i = 0; i < NWITEMS; i++)

 67: buf2[i] = rand()%65535;

 68: 

 69: for(i = 0; i < NWITEMS; i++)

 70: buf[i] = buf1[i] + buf2[i];

 71: 

 72: cl_uint status;

 73: cl_platform_id platform;

 74: 

 75: //创建平台对象

 76: status = clGetPlatformIDs( 1, &platform, NULL );

 77: 

 78: cl_device_id device;

 79: 

 80: //创建GPU设备

 81: clGetDeviceIDs( platform, CL_DEVICE_TYPE_GPU,

 82: 1,

 83: &device,

 84: NULL);

 85: //创建context

 86: cl_context context = clCreateContext( NULL,

 87: 1,

 88: &device,

 89: NULL, NULL, NULL);

 90: //创建命令队列

 91: cl_command_queue queue = clCreateCommandQueue( context,

 92: device,

 93: CL_QUEUE_PROFILING_ENABLE, NULL );

 94: //创建三个OpenCL内存对象，并把buf1的内容通过隐式拷贝的方式

 95: //拷贝到clbuf1,buf2的内容通过显示拷贝的方式拷贝到clbuf2

 96: cl_mem clbuf1 = clCreateBuffer(context,

 97: CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,

 98: NWITEMS*sizeof(cl_float),buf1,

 99: NULL );

 100: 

 101: cl_mem clbuf2 = clCreateBuffer(context,

 102: CL_MEM_READ_ONLY ,

 103: NWITEMS*sizeof(cl_float),NULL,

 104: NULL );

 105: 

 106: status = clEnqueueWriteBuffer(queue, clbuf2, 1,

 107: 0, NWITEMS*sizeof(cl_float), buf2, 0, 0, 0);

 108: 

 109: cl_mem buffer = clCreateBuffer( context,

 110: CL_MEM_WRITE_ONLY,

 111: NWITEMS * sizeof(cl_float),

 112: NULL, NULL );

 113: 

 114: const char * filename = "add.cl";

 115: std::string sourceStr;

 116: status = convertToString(filename, sourceStr);

 117: const char * source = sourceStr.c_str();

 118: size_t sourceSize[] = { strlen(source) };

 119: 

 120: //创建程序对象

 121: cl_program program = clCreateProgramWithSource(

 122: context,

 123: 1,

 124: &source,

 125: sourceSize,

 126: NULL);

 127: //编译程序对象

 128: status = clBuildProgram( program, 1, &device, NULL, NULL, NULL );

 129: if(status != 0)

 130: {

 131: printf("clBuild failed:%d\n", status);

 132: char tbuf[0x10000];

 133: clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0x10000, tbuf, NULL);

 134: printf("\n%s\n", tbuf);

 135: return -1;

 136: }

 137: 

 138: //创建Kernel对象

 139: cl_kernel kernel = clCreateKernel( program, "vecadd", NULL );

 140: //设置Kernel参数

 141: cl_int clnum = NWITEMS;

 142: clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*) &clbuf1);

 143: clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*) &clbuf2);

 144: clSetKernelArg(kernel, 2, sizeof(cl_mem), (void*) &buffer);

 145: 

 146: //执行kernel

 147: cl_event ev;

 148: size_t global_work_size = NWITEMS;

 149: clEnqueueNDRangeKernel( queue,

 150: kernel,

 151: 1,

 152: NULL,

 153: &global_work_size,

 154: NULL, 0, NULL, &ev);

 155: clFinish( queue );

 156: 

 157: //数据拷回host内存

 158: cl_float *ptr;

 159: ptr = (cl_float *) clEnqueueMapBuffer( queue,

 160: buffer,

 161: CL_TRUE,

 162: CL_MAP_READ,

 163: 0,

 164: NWITEMS * sizeof(cl_float),

 165: 0, NULL, NULL, NULL );

 166: //结果验证，和cpu计算的结果比较

 167: if(!memcmp(buf, ptr, NWITEMS))

 168: printf("Verify passed\n");

 169: else printf("verify failed");

 170: 

 171: if(buf)

 172: free(buf);

 173: if(buf1)

 174: free(buf1);

 175: if(buf2)

 176: free(buf2);

 177: 

 178: //删除OpenCL资源对象

 179: clReleaseMemObject(clbuf1);

 180: clReleaseMemObject(clbuf2);

 181: clReleaseMemObject(buffer);

 182: clReleaseProgram(program);

 183: clReleaseCommandQueue(queue);

 184: clReleaseContext(context);

 185: return 0;

 186: }

 187: 

也可以在http://code.google.com/p/imagefilter-opencl/downloads/detail?name=amdunicourseCode1.zip&can=2&q=#makechanges上下载完整版本。

GPU架构

内容包括：

1.OpenCLspec和多核硬件的对应关系

• AMD GPU架构
• Nvdia GPU架构
• Cell Broadband Engine

2.一些关于OpenCL的特殊主题

• OpenCL编译系统
• Installable client driver

 

首先我们可能有疑问，既然OpenCL具有平台无关性，我们为什么还要去研究不同厂商的特殊硬件设备呢?

• 了解程序中的循环和数据怎样映射到OpenCL Kernel中，便于我们提高代码质量，获得更高的性能。
• 了解AMD和Nvdia显卡的区别。
• 了解各种硬件的区别，可以帮助我们使用基于这些硬件的一些特殊的OpenCL扩展，这些扩展在后面课程中会讲到。

3、传统的CPU架构

•     对单个线程来说，CPU优化能获得最小时延，而且CPU也适合处理控制流密集的工作，比如if、else或者跳转指令比较多的任务。
• 控制逻辑单元在芯片中占用的面积要比ALU单元多。
• 多层次的cache设计被用来隐藏时延（可以很好的利用空间和时间局部性原理）
• 有限的寄存器数量使得同时active的线程不能太多。
• 控制逻辑单元记录程序的执行、提供指令集并行（ILP）以及最小化CPU管线的空置周期（stalls，在该时钟周期，ALU没做什么事）。

4、现代的GPGPU架构

 

• 对于现代的GPU，通常的它的控制逻辑单元比较简单（和cpu相比），cache也比较小
• 线程切换开销比较小，都是轻量级的线程。
• GPU的每个“核”有大量的ALU以及很小的用户可管理的cache。[这儿的核应该是指整个GPU]。
• 内存总线都是基于带宽优化的。150GB/s的带宽可以使得大量ALU同时进行内存操作。

5、AMD GPU硬件架构

现在我们简单看下AMD 5870显卡(cypress)的架构

• 20个simd引擎，每个simd引擎包含16个simd。
• 每个simd包含16个stream core
• 每个stream core都是5路的乘法-加法运算单元（VLIW processing）。
• 单精度运算可以达到 Teraflops。
• 双精度运算可以达到544Gb/s

上图为一个simd引擎的示意图，每个simd引擎由一系列的stream core组成。

• 每个stream core是一个5路的VLIW处理器，在一个VLIW指令中，可以最多发射5个标量操作。标量操作在每个pe上执行。
• CU（8xx系列cu对应硬件的simd）内的stream core执行相同的VLIW指令。
• 在CU（或者说simd）内同时执行的work item放在一起称作一个wave，它是cu中同时执行的线程数目。在5870中wave大小是64，也就是说一个cu内，最多有64个work item在同时执行。

注：5路的运算对应(x,y,z,w),以及T（超越函数），在cayman中，已经取消了T，改成四路了。

 

我们现在看下AMD GPU硬件在OpenCL中的对应关系：

• 一个workitme对应一个pe，pe就是单个的VLIW core
• 一个cu对应多个pe，cu就是simd引擎。

上图是AMD GPU的内存架构（原课件中的图有点小错误，把Global memory写成了LDS)

• 对每个cu来说，它使用的内存包括onchip的LDS以及相关寄存器。在5870中，每个LDS是32K，共32个bank，每个bank 1k，读写单位4 byte。
• 对没给cu来说，有8K的L1 cache。（for 5870）
• 各个cu之间共享的L2 cache，在5870中是512K。
• fast Path只能执行32位或32位倍数的内存操作。
• complete path能够执行原子操作以及小于32位的内存操作。

AMD GPU的内存架构和OpenCL内存模型之间的对应关系：

• LDS对应local memeory，主要用来在一个work group内的work times之间共享数据。steam core访问LDS的速度要比Global memory快一个数量级。
• private memory对应每个pe的寄存器。
• constant memory主要是利用了L1 cache

注意：对AMD CPU，constant memory的访问包括三种方式：Direct-Addressing Patterns，这种模式要求不包括行列式，它的值都是在kernel函数初始化的时候就决定了，比如传入一个固定的参数。Same Index Patterns，所有的work item都访问相同的索引地址。Globally scoped constant arrays，行列式会被初始化，如果小于16K，会使用L1 cache，从而加快访问速度。

当所有的work item访问不同的索引地址时候，不能被cache，这时要在global memory中读取。

 

 

6、Nvdia GPU Femi架构

 

GTX480-Compute 2.0 capability：

• 有15个core或者说SM（Streaming Multiprocessors ）。
• 每个SM,一般有32 cuda处理器。
• 共480个cuda处理器。
• 带ECC的global memory
• 每个SM内的线程按32个单位调度执行，称作warp。每个SM内有2个warp发射单元。
• 一个cuda核由一个ALU和一个FPU组成，FPU是浮点处理单元。

SIMT和SIMD

SIMT是指单指令、多线程。

• 硬件决定了多个ALU之间要共享指令。
• 通过预测来处理多个线程间的Diverage(是指同一个warp中的指令执行路径产生不同）。
• NV把一个warp中执行的指令当作一个SIMT。SIMT指令指定了一个线程的执行以及分支行为。

SIMD指令可以得到向量的宽度，这点和X86 SSE向量指令比较类似。

SIMD的执行和管线相关

• 所有的ALU执行相同的指令。
• 根据指令可以管线分为不同的阶段。当第一条指令完成的时候（4个周期），下条指令开始执行。

Nvida GPU内存机制：

• 每个SM都有L1 cache，通过配置，它可以支持shared memory，也可以支持global memory。
• 48 KB Shared / 16 KB of L1 cache，16 KB Shared / 48 KB of L1 cache
• work item之间数据共享通过shared memory
• 每个SM有32K的register bank
• L2(768K)支持所有的操作，比如load,store等等
• Unified path to global for loads and stores 

和AMD GPU类似，Nv的GPU 内存模型和OpenCL内存模型的对应关系是：

• shared memory对应local memory
• 寄存器对应private memory

7、Cell Broadband Engine

 

由索尼，东芝，IBM等联合开发，可用于嵌入式平台，也可用于高性能计算（SP3次世代游戏主机就用了cell处理器）。

• Bladecenter servers提供OpenCL driver支持
• 如图所示，cell处理器由一个Power Processing Element (PPE) 和多个Synergistic Processing Elements (SPE)组成。
• Uses the IBM XL C for OpenCL compiler 11
• Cell Power/VMX CPU 的设备类型是CL_DEVICE_TYPE_CPU，Cell SPU 的设备类型是CL_DEVICE_TYPE_ACCELERATOR。
• OpenCL Accelerator设备和CPU共享内存总线。
• 提供一些扩展，比如Device Fission、Migrate Objects来指定一个OpenCL对象驻留在什么位置。
• 不支持OpenCL image对象，原子操作，sampler对象以及字节内存地址。

8、OpenCL编译系统

• LLVM-底层的虚拟机
• Kernel首先在front-end被编译成LLVM IR
• LLVM是一个开源的编译器，具有平台独立性，可以支持不同厂商的back_end编译,网址：http://llvm.org

9、Installable Client Driver

• ICD支持不同厂商的OpenCL实施在系统中共存。
• 代码紧被链接接到libOpenCL.so
• 应用程序可在运行时选择不同的OpenCL实施（就是选择不同platform）
• 现在的GPU驱动还不支持跨厂商的多个GPU设备同时工作。
• 通过clGetPlatformIDs() 和clGetPlatformInfo() 来检测不同厂商的OpenCL平台。

本节主要讲述GPU的memory架构。优化基于GPU device的kernel程序时，我们需要了解很多GPU的memory知识，比如内存合并，bank conflit（冲突）等等，这样才能针对具体算法做一些优化工作。

1、GPU总线寻址介绍

 

   假定X是一个指向整数（32位整数）数组的指针，数组的首地址为0x00001232。一个线程要访问元素X[0],

   int tmp = X[0];

   

    假定memory总线宽度为256位(HD5870就是如此,即为32字节），因为基于字节地址的总线要访问memeory，必须和总线宽度对齐，也就是说按必须32字节对齐来访问memory，比如访问0x00000000,0x00000020,0x00000040,…等，所以我们要得到地址0x00001232中的数据，比如访问地址0x00001220,这时，它会同时得到0x00001220到 0x0000123F 的所有数据。因为我们只是取的一个32位整数，所以有用的数据是4个字节，其它28的字节的数据都被浪费了，白白消耗了带宽。

   

 

2、合并内存访问

    为了利用总线带宽，GPU通常把多个线程的内存访问尽量合并到较少的内存请求命令中去。

    假定下面的OpenCL kernel代码：int tmp = X[get_global_id(0)];

数组X的首地址和前面例子一样，也是0x00001232，则前16个线程将访问地址：0x00001232 到 0x00001272。假设每个memory访问请求都单独发送的话，则有16个request，有用的数据只有64字节，浪费掉了448字节（16*28）。

    假定多个线程访问32个字节以内的地址，它们的访问可以通过一个memory request完成，这样可以大大提高带宽利用率，在专业术语描述中这样的合并访问称作coalescing。

   例如上面16个线程访问地址0x00001232 到 0x00001272，我们只需要3次memory requst。

   在HD5870显卡中，一个wave中16个连续线程的内存访问会被合并，称作quarter-wavefront，是重要的硬件调度单位。

   下面的图是HD5870中，使用memory访问合并以及没有使用合并的bandwidth比较：

   下图是GTX285中的比较：

3、Global memory的bank以及channel访问冲突

   我们知道内存由bank，channel组成，bank是实际存储数据的单元，一个mc可以连接多个channel，形成单mc，多channel的连接方式。在物理上，不同bank的数据可以同时访问，相同的bank的数据则必须串行访问，channel也是同样的道理。但由于合并访问的缘故，对于global memory来说，bank conflit影响要小很多，除非是非合并问，不同线程访问同一个bank。理想情况下，我们应该做到不同的workgroup访问的不同的bank，同一个group内，最好用合并操作。

   下面我简单的画一个图，不知道是否准确，仅供参考：

 

    

     在HD5870中，memory地址的低8位表示一个bank中的数据，接下来的3位表示channel（共8个channel），bank位的多少依赖于显存中bank的多少。

4、local memory的bank conflit

   bank访问冲突对local memory操作有更大的影响（相比于global memory），连续的local memory访问地址，应该映射到不同的bank上，

     在AMD显卡中，一个产生bank访问冲突wave将会等待所有的local memory访问完成，硬件不能通过切换到另一个wave来隐藏local memory访问时延。所以对local memory访问的优化就很重要。HD5870显卡中，每个cu（simd）有32bank，每个bank 1k，按4字节对齐访问。如果没有bank conflit，每个bank能够没有延时的返回一个数据，下面的图就是这种情况。

   如果多个memory访问对应到一个bank上，则conflits的数量决定时延的大小。下面的访问方式将会有3倍的时延。

  但是，如果所有访问都映射到一个bank上，则系统会广播数据访问，不会产生额外时延。

GPU线程及调度

     本节主要讲述OpenCL中的Workgroup如何在硬件设备中被调度执行。同时也会讲一下同一个workgroup中的workitem，如果它们执行的指令发生diverage（就是执行指令不一致）对性能的影响。学习OpenCL并行编程，不仅仅是对OpenCL Spec本身了解，更重要的是了解OpenCL硬件设备的特性，现阶段来说，主要是了解GPU的的架构特性，这样才能针对硬件特性优化算法。

     现在OpenCL的Spec是1.1，随着硬件的发展，相信OpenCL会支持更多的并行计算特性。基于OpenCL的并行计算才刚刚起步，…

1、workgroup到硬件线程

     在OpenCL中，Kernel函数被workgroup中的workitem（线程，我可能混用这两个概念）执行。在硬件层次，workgroup被映射到硬件的cu（compute unit）单元来执行具体计算，而cu一般由更多的SIMT（单指令，线程）pe（processing elements）组成。这些pe执行具体的workitem计算，它们执行同样的指令，但操作的数据不一样，用simd的方式完成最终的计算。

    由于硬件的限制，比如cu中pe数量的限制，实际上workgroup中线程并不是同时执行的，而是有一个调度单位，同一个workgroup中的线程，按照调度单位分组，然后一组一组调度硬件上去执行。这个调度单位在nv的硬件上称作warp,在AMD的硬件上称作wavefront，或者简称为wave。

  上图显示了workgroup中，线程被划分为不同wave的分组情况。wave中的线程同步执行相同的指令，但每个线程都有自己的register状态，可以执行不同的控制分支。比如一个控制语句

if(A)

{

… //分支A

}

else

{

  … //分支B

}

    假设wave中的64个线程中，奇数线程执行分支A，偶数线程执行分支B，由于wave中的线程必须执行相同的指令，所以这条控制语句被拆分为两次执行[编译阶段进行了分支预测]，第一次分支A的奇数线程执行，偶数线程进行空操作，第二次偶数线程执行，奇数线程空操作。硬件系统有一个64位mask寄存器，第一次是它为01…0101,第二次会进行反转操作10…1010，根据mask寄存器的置位情况，来选择执行不同的线程。可见对于分支多的kernel函数，如果不同线程的执行发生diverage的情况太多，会影响程序的性能。

2、AMD wave调度

    AMD GPU的线程调度单位是wave，每个wave的大小是64。指令发射单元发射5路的VLIW指令，每个stream core(SC)执行一条VLIW指令，16个stream core在一个时钟周期执行16条VLIW指令。每个时钟周期，1/4wave被完成，整个wave完成需要四个连续的时钟周期。

    另外还有以下几点值得我们了解：

• 发生RAW hazard情况下，整个wave必须stall 4个时钟周期，这时，如果其它的wave可以利用，ALU会执行其它的wave以便隐藏时延，8个时钟周期后，如果先前等待wave已经准备好了，ALU会继续执行这个wave。
• 两个wave能够完全隐藏RAW时延。第一个wave执行时候，第二个wave在调度等待数据，第一个wave执行完时，第二个wave可以立即开始执行。

3、nv warp调度

     work group以32个线程为单位，分成不同warp，这些warp被SM调度执行。每次warp中一半的线程被发射执行，而且这些线程能够交错执行。可以用的warp数量依赖于每个block的资源情况。除了大小不一样外，wave和warp在硬件特性上很相似。

4、Occupancy开销

    在每个cu中，同时激活的wave数量是受限制的，这和每个线程使用register和local memory大小有关，因为对于每个cu，register和local memory总量是一定的。

    我们用术语Occupancy来衡量一个cu中active wave的数量。如果同时激活的wave越多，能更好的隐藏时延，在后面性能优化的章节中，我们还会更具体讨论Occupancy。

5、控制流和分支预测(prediction)

   前面我说了if else的分支执行情况，当一个wave中不同线程出现diverage的时候，会通过mask来控制线程的执行路径。这种预测（prediction）的方式基于下面的考虑：

• 分支的代码都比较短
• 这种prediction的方式比条件指令更高效。
• 在编译阶段，编译器能够用predition替换switch或者if else。

  prediction 可以定义为：根据判断条件，条件码被设置为true或者false。

__kernel 
void test() {

 int tid= get_local_id(0) ;
 if( tid %2 == 0)
Do_Some_Work() ;
 else
Do_Other_Work() ; 
}

例如上面的代码就是可预测的，

Predicate = True for threads 0,2,4….

Predicate = False for threads 1,3,5….

下面在看一个控制流diverage的例子

• 在case1中，所有奇数线程执行DoSomeWork2(),所有偶数线程执行DoSomeWorks,但是在每个wave中，if和else代码指令都要被发射。
• 在case2中，第一个wave执行if，其它的wave执行else，这种情况下，每个wave中，if和else代码只被发射一个。

   在prediction下，指令执行时间是if，else两个代码快执行时间之和。

6、Warp voting

   warp voting是一个warp内的线程之间隐式同步的机制。

    比如一个warp内线程同时写Local meory某个地址，在线程并发执行时候，warp voting机制可以保证它们的前后顺序正确。更详细的warp voting大家可以参考cuda的资料。

  

    在OpenCL编程中，由于各种硬件设备不同，导致我们必须针对不同的硬件进行优化，这也是OpenCL编程的一个挑战，比如warp和wave数量的不同，使得我们在设计workgroup大小时候，必须针对自己的平台进行优化，如果选择32，对于AMD GPU，可能一个wave中32线程是空操作，而如果选择64，对nv GPU来说，可能会出现资源竞争的情况加剧，比如register以及local meomory的分配等等。这儿还不说混合CPU device的情况,OpenCL并行编程的道路还很漫长，期待新的OpenCL架构的出现。

性能优化

1、线程映射

   所谓线程映射是指某个线程访问哪一部分数据，其实就是线程id和访问数据之间的对应关系。

合适的线程映射可以充分利用硬件特性，从而提高程序的性能，反之，则会降低performance。

   请参考Static Memory Access Pattern Analysis on a Massively Parallel GPU这篇paper，文中讲述线程如何在算法中充分利用线程映射。这是我在google中搜索到的下载地址：http://www.ece.neu.edu/~bjang/patternAnalysis.pdf

   使用不同的线程映射，同一个线程可能访问不同位置的数据。下面是几个线程映射的例子：

      我们考虑一个简单的串行矩阵乘法：这个算法比较适合输出数据降维操作，通过创建N*M个线程，我们移去两层外循环，这样每个线程执行P个加法乘法操作。现在需要我们考虑的问题是，线程索引空间究竟应该是M*N还是N*M？

    当我们使用M*N线程索引空间时候，Kernel如下图所示：

   而使用N*M线程索引空间时候，Kernel如下图所示：

    使用两种映射关系，程序执行结果是一样的。下面是在nv的卡GeForce 285 and 8800 GPUs上的执行结果。可以看到映射2（及N*M线程索引空间），程序的performance更高。

    performance差异主要是因为在两种映射方式下，对global memory访问的方式有所不同。在行主序的buffer中，数据都是按行逐个存储，为了保证合并访问，我们应该把一个wave中连续的线程映射到矩阵的列（第二维），这样在A*B=C的情况下，会把矩阵B和C的内存读写实现合并访问，而两种映射方式对A没有影响（A又i3决定顺序）。

   完整的源代码请从：http://code.google.com/p/imagefilter-opencl/downloads/detail?name=amduniCourseCode4.zip&can=2&q=#makechanges下载，程序中我实现了两种方式的比较。结果确实第二种方式要快一些。

   下面我们再看一个矩阵转置的例子，在例子中，通过改变映射方式，提高了global memory访问的效率。

   矩阵转置的公式是:Out(x,y) = In(y,x)

   从上图可以看出，无论才去那种映射方式，总有一个buffer是非合并访问方式（注：在矩阵转置时，必须要把输入矩阵的某个元素拷贝到临时位置，比如寄存器，然后才能拷贝到输出矩阵）。我们可以改变线程映射方式，用local memory作为中间元素，从而实现输入，输出矩阵都是global memory合并访问。

  下面是AMD 5870显卡上，两种线程映射方式实现的矩阵转置性能比较：

    完整代码：http://code.google.com/p/imagefilter-opencl/downloads/detail?name=amduniCourseCode5.zip&can=2&q=#makechanges

2、Occupancy

    前面的教程中，我们提到过Occupancy的概念，它主要用来描述CU中资源的利用率。

    OpenCL中workgroup被映射到硬件的CU中执行，在一个workgroup中的所有线程执行完之后，这个workgroup才算执行结束。对一个特定的cu来说，它的资源（比如寄存器数量，local memory大小，最大线程数量等）是固定的，这些资源都会限制cu中同时处于调度状态的workgroup数量。如果cu中的资源数量足够的的话，映射到同一个cu的多个workgroup能同时处于调度状态，其中一个workgroup的wave处于执行状态，当处于执行状态的workgroup所有wave因为等待资源而切换到等待状态的话，不同workgroup能够从就绪状态切换到ALU执行，这样隐藏memory访问时延。这有点类似操作系统中进程之间的调度状态。我简单画个图，以供参考：

• 对于一个比较长的kernel，寄存器是主要的资源瓶颈。假设kernel需要的最大寄存器数目为35，则workgroup中的所有线程都会使用35个寄存器，而一个CU（假设为5870）的最大寄存器数目为16384，则cu中最多可有16384/35=468线程，此时，一个workgroup中的线程数目（workitem）不可能超过468，
• 考虑另一个问题，一个cu共16384个寄存器，而workgroup固定为256个线程，则使用的寄存器数量可达到64个。

    每个CU的local memory也是有限的，对于AMD HD 5XXX显卡，local memory是32K，NV的显卡local memory是32-48K（具体看型号）。和使用寄存器的情况相似，如果kernel使用过多的local memory，则workgroup中的线程数目也会有限制。

   GPU硬件还有一个CU内的最大线程数目限制：AMD显卡256，nv显卡512。

   NV的显卡对于每个CU内的激活线程有数量限制，每个cu 8个或16个warp，768或者1024个线程。

   AMD显卡对每个CU内的wave数量有限制，对于5870，最多496个wave。

   这些限制都是因为有限的资源竞争引起的，在nv cuda中，可以通过可视化的方式查看资源的限制情况。

3、向量化

   向量化允许一个线程同时执行多个操作。我们可以在kernel代码中，使用向量数据类型，比如float4来获得加速。向量化在AMD的GPU上效果更为明显，这是因为AMD的显卡的stream core是(x,y,z,w)这样的向量运算单元。

   下图是在简单的向量赋值运算中，使用float和float4的性能比较。

    kernel代码为：