(CUDA 编程4).CUDA硬件实现分析（二）------规行矩步
作者：赵开勇 来源:http://www.hpctech.com/2009/0818/202.html
前面已经讲解了很多概念上的东西，其实CUDA的最重要的两个东西，就是线程和内存。只要掌握了这两个东西，CUDA的东西也就很简单了。它的编写语言是C扩展的，所以，就当C语言用就行了，只是主要它的特殊的几个标志就ok了。
（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];
}
}
扫一扫
3335
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
