(CUDA 编程4).CUDA硬件实现分析（二）------规行矩步

2009-08-19 01:19:13 作者：赵开勇 来源：本站 浏览次数：0 网友评论 0 条

（CUDA 编程4）CUDA硬件实现分析（二）------规行矩步

------GPU的革命

前面已经讲解了很多概念上的东西，其实CUDA的最重要的两个东西，就是线程和内存。只要掌握了这两个东西，CUDA的东西也就很简单了。它的编写语言是C扩展的，所以，就当C语言用就行了，只是主要它的特殊的几个标志就ok了。前面讲解了线程和内存的模型，大概，应该，似乎，可以在你的脑海里面有一个概念了吧。只要有这个概念，我的文章的目的就达到了。前面的《CUDA硬件实现分析（一）------安营扎寨-----GPU的革命》已经讲解了线程在CUDA的具体运行过程。下面我们一起来看看内存在CUDA的硬件实现中的一些规定。这也比较合理吧，大军安营扎寨了，就应该颁布规则制度，只有了解CUDA的规则制度，才能真正的把各个线程都管理好。才能在这个平台上让程序高效的运行。

这里我们先明确几个

一．            Threads，Warps, Blocks
1. 一个warp最多有32个threads。只有在总线程少于32的时候，才可能在一个warp里面少于32个线程。
2. 每一个block最多有16个warp。就是说一个一个block里面最多有512个thread。
3. 每一个Block在同一个SM上执行，也就是同一个block的warp都在同一个SM上运行。
4. G80有16个SM。
5. 所以最少16个blocks才能占全所有的SM。
6. 如果资源（看看前面讲解的线程都要从device哪里分什么资源）够线程分，一个SM上面可以跑多余一个block的线程。就是同时可以跑2个，3……个block的线程。

二．            访问速度
Register—HW 一个时间周期
Shared Memory------HW一个时钟周期
Local Memory --- DRAM，no cache， 慢
Global Memory --- DRAM， no cache， 慢
Constant Memory --- DRAM， cached， 1……10s……100s个周期，这个和cache的locality有关。
Texture Memory --- DRAM， cached， 1……10s……100s个周期，这个和cache的locality有关。
Instruction Memory（不可见）--- DRAM，cached

三．            CUDA程序架构
如图

四．            语言扩展
单从学习语言来说，我到觉得应该精学一种，然后其他的就触类旁通了。在学习新语言的时候，要想快速的入门，也有诀窍。1.变量的定义方式。2.函数的定义方式。3.逻辑控制方式（if，loop……）。只要把这3个东西弄明白了，管他啥新的语言，20分钟就可以入门……然后入门都可以了，那要慢慢的更入的研究，那就得看你对这门语言的了解了。其实万变不离其中。其实从计算机编程语言的角度出发，就是定义一些数据，然后对数据进行操作，so……学习语言就从这个角度入手，那就很简单了。像java或者C#等一些语言比C语言多的新的特性不外乎就是方便你开发而已。

所以我们这里再来说说CUDA的语言，不外乎就是扩展了C语言，为了方便在GPU显卡上运行，规定一个特定环境。就定义一些特定的变量，说明他们是在GPU上的。这里有说明内存，和函数，是在GPU上的……so，这样一来，CUDA就扩展了C语言的变量分配定义和函数定义。

上面这张图是来之Fall 2007 syllabus，上面已经说得很清楚各个变量定义的时候的位置和生存周期。其实就是在C语言的常规变量的时候，定义了变量的位置而已。

其中有一个约束限制，就是指针变量，在kernel里面的指针变量，只能指向从global上面分配的内存。

五．            内建的变量
所谓内建，就是CUDA自己在kernel里面定义的一些变量。就像我们以前计算线程id的时候，就利用了他的自己的变量，dim3 gridDim；dim3 blockDim；dim3 blockIdx；dim3 threadIdx；注意gridDim，他的gridDim.z在现在的CUDA1.1版本中没定义。

内建变量，[u]char[1..4], [u]short[1..4], [u]int[1..4], [u]long[1..4], float[1..4] 就是构建了有4个变量的struct；
             uint4 param; ---》等价位一个struct里面有4个int。
             int y = param.y;
dim3 就是unit3这样的struct。

六．            通用的数学函数

·          pow, sqrt, cbrt, hypot

·          exp, exp2, expm1

·          log, log2, log10, log1p

·          sin, cos, tan, asin, acos, atan, atan2

·          sinh, cosh, tanh, asinh, acosh, atanh

·          ceil, floor, trunc, round

但是这里要指出，有几个函数是不精确的，但是可以很快的运行：

–        __pow

–        __log, __log2, __log10

–        __exp

–        __sin, __cos, __tan

七．            在host部分的运行库（CUDA Runtime ）

1.       提供Device管理的api（有多个显卡的时候怎能来设置，这里我们还没讲当多显卡并行运行库）

2.       初始化调用的Runtime函数

3.       每一个host的线程（thread）只能调用一个device的函数在一个device上运行。就是同时不能有几个host（主机）线程在调用同一个device上的运行函数。

八．            内存管理函数
我们讲了那么多的内存，下面就来看看到底是什么样的函数：

–        cudaMalloc（），cudaFree（）；内存分配和释放（device上的）

–        内存copy：cudaMemcpy(), cudaMemcpy2D(), cudaMemcpyToSymbol(), cudaMemcpyFromSymbol()；

九．            线程同步函数

void __syncthreads();

同步同一个block里面的线程，让block里面的线程都允许到这一点的时候，就等待同一个block里面的其他线程，就像军训的时候，大家吃饭吃完了还不能一个个走，必须得一个桌子的人都吃完了才能走。还得列队一起走，呵呵，这就是同步。所以最好保证每一个kernel里面的处理都是很快的，这样才不会让其他thread等待太久，不然会挨骂的 - -！hoho

哥们来点实际的吧- - 肯定很多人都这么在吼了。嘿嘿，下面让我们看一段代码，example里面的transpose：

#define BLOCK_DIM 16

// This kernel is optimized to ensure all global reads and writes are coalesced,

// and to avoid bank conflicts in shared memory.  This kernel is up to 11x faster

// than the naive kernel below.  Note that the shared memory array is sized to

// (BLOCK_DIM+1)*BLOCK_DIM.  This pads each row of the 2D block in shared memory

// so that bank conflicts do not occur when threads address the array column-wise.

__global__ void transpose(float *odata, float *idata, int width, int height)

{

    __shared__ float block[BLOCK_DIM][BLOCK_DIM+1]; //(1)

    // read the matrix tile into shared memory

    unsigned int xIndex = blockIdx.x * BLOCK_DIM + threadIdx.x;

    unsigned int yIndex = blockIdx.y * BLOCK_DIM + threadIdx.y;

    if((xIndex < width) && (yIndex < height)) // (2)

    {

        unsigned int index_in = yIndex * width + xIndex;

        block[threadIdx.y][threadIdx.x] = idata[index_in];

    }

    __syncthreads(); //(3)

    // write the transposed matrix tile to global memory

    xIndex = blockIdx.y * BLOCK_DIM + threadIdx.x;

    yIndex = blockIdx.x * BLOCK_DIM + threadIdx.y;

    if((xIndex < height) && (yIndex < width))

    {

        unsigned int index_out = yIndex * height + xIndex;

        odata[index_out] = block[threadIdx.x][threadIdx.y];

    }

}

是不是在代码面前都感觉是那么的亲切啊～嘿嘿，有时候看算法的书看久了以后，看到代码实现，那才叫亲切啊>_<!

下面对加红的3处代码分别讲解：

（1）      shared__ float block[BLOCK_DIM][BLOCK_DIM+1];
为啥要用shared啊？知道为啥不，快啊～shared在block里面，是大家公用的，不用再从global哪里调换。还记得军训的时候食堂吃饭的时候那个内存讲解吗？桌子上的菜，肯定比你从中间几个桶里面去拿菜要快啊～

（2）      if((xIndex < width) && (yIndex < height))
由于CUDA在生成线程的时候，是一块一块的生成的，所以有的时候，x和y有可能越界。数据都是均匀的分配到每一个block里面去的，很可能就会出现有的blcok里面分配到的数据不一定完整，要是想不明白的（先画图想想），等后面的章节，举例子说明。

（3）      __syncthreads();
同步，就是等每一个thread都运行到这里的时候，再接下来运行。这里也避免了数据访问的冲突。

（4）      还有一个得讲讲，为啥下面的还重新计算一下x和y，因为线程被挂起来了以后，重新进入warp的时候，分配的id就不一定是一样的了。这里就有点像铁打的营盘流水的兵，每一波人都可以分配到同一张桌子吃饭，但是不一定是上一波人了～。碗筷还是那些碗筷，位置还是那个位置，坐的人不一样了。

上面的都是C语言扩展库的东西，就是说，这些是扩展C语言的东东，只要是基于C语言的规则都可以运行的，不管是host还是device。

好了，今天要讲的部分应该差不多了，要是再写多一些，或者就有会像手册第4章那样，看了一半很多人就不想看了～- -！太长，lz 帖子太长，顶lz- -！我可不想看到这样的留言，hoho～

从小到大，我们听过好多规则，但是好多时候听进去的有多少？很多玩游戏的人，游戏的规则都不清楚- -ps，俺现在还不清楚StarCraft里面的某些兵的战斗力是多少 - -！就玩- -！小小的bs一下自己。只是为了消遣罢了·

我始终觉得自己不适合玩游戏，我也知道自己玩游戏不会完成专业玩家，so～还是老老实实做自己的工作，做一些能养活自己的事情，或许能让自己更充实一些。在游戏里面充实以后，或许在现实中就会空虚了- -！到时候就不是游戏人生，那就是人生被游戏了。呵呵。想起高中的数学老师，虽然在很多journal上发表了文章，也自己出版了数学的几本书了。但是和别人打麻将的时候总是输多赢少。问他为啥数学这么厉害，玩麻将不行～他说，我玩麻将的时候只是消遣，根本就没用脑袋想麻将这回事。

或许我们做事情的时候，就应该专注做自己的一些方面吧，只要坚持下来做了，做好了，做精了，那就够了。小时候就从长辈哪里听到，半灌水叮当响这样的词，或许我们更应该低调做人，高调做事吧。