CUDA学习笔记（一）【转】

CUDA编程中，习惯称CPU为Host，GPU为Device。编程中最开始接触的东西恐怕是并行架构，诸如Grid、Block的区别会让人一头雾水，我所看的书上所讲述的内容比较抽象，对这些概念的内容没有细讲，因而在这里做一个整理。

Grid、Block和Thread的关系

Thread  ：并行运算的基本单位（轻量级的线程）
Block   ：由相互合做的一组线程组成。一个block中的thread能够彼此同步，快速交换数据，最多能够同时512个线程。
Grid     ：一组Block，有共享全局内存
Kernel ：在GPU上执行的程序，一个Kernel对应一个Grid。

其结构以下图所示：

?

1 2 3 4 5 6 7 8 9 10

/* 另外：Block和Thread都有各自的ID，记做blockIdx（1D，2D），threadIdx（1D，2D，3D） Block和Thread还有Dim，即blockDim与threadDim. 他们都有三个份量x，y，z 线程同步：void __syncthreads(); 能够同步一个Block内的全部线程 总结来讲，每一个 thread 都有本身的一份 register 和 local memory 的空间。 一组thread构成一个 block，这些 thread 则共享有一份shared memory。 此外，全部的 thread(包括不一样 block 的 thread)都共享一份 global memory、constant memory、和 texture memory。 不一样的 grid 则有各自的 global memory、constant memory 和 texture memory。 */

 

?

存储层次

1 2 3 4 5 6 7

per- thread register                             1 cycle per- thread local memory                     slow per-block shared memory                   1 cycle per-grid global memory                       500 cycle,not cached!! constant and texture memories            500 cycle, but cached and read-only 分配内存：cudaMalloc，cudaFree，它们分配的是global memory Hose-Device数据交换：cudaMemcpy

?

变量类型

1 2 3 4 5

__device__   // GPU的global memory空间，grid中全部线程可访问 __constant__ // GPU的constant memory空间，grid中全部线程可访问 __shared__   // GPU上的thread block空间，block中全部线程可访问 local        // 位于SM内，仅本thread可访问 // 在编程中，能够在变量名前面加上这些前缀以区分。

?

数据类型

1 2 3 4 5 6 7 8 9

// 内建矢量类型： int1，int2，int3，int4，float1，float2， float3，float4 ... // 纹理类型： texture<Type, Dim, ReadMode>texRef; // 内建dim3类型：定义grid和block的组织方法。例如： dim3 dimGrid(2, 2); dim3 dimBlock(4, 2, 2); // CUDA函数CPU端调用方法 kernelFoo<<<dimGrid, dimBlock>>>(argument);

?

函数定义

1 2 3 4 5 6 7 8 9 10

__device__ // 执行于Device，仅能从Device调用。限制，不能用&取地址；不支持递归；不支持static variable；不支持可变长度参数 __global__ // void： 执行于Device，仅能从Host调用。此类函数必须返回void __host__ // 执行于Host，仅能从Host调用，是函数的默认类型 // 在执行kernel函数时，必须提供execution configuration，即<<<....>>>的部分。 //   例如： __global__ void KernelFunc(...); dim3 DimGrid(100, 50); // 5000 thread blocks dim3 DimBlock(4, 8, 8); // 256 threads per block size_t SharedMemBytes = 64; // 64 bytes of shared memory KernelFunc<<< DimGrid, DimBlock, SharedMemBytes >>>(...);

?

数学函数

1 2	CUDA包含一些数学函数，如 sin ， pow 等。每个函数包含有两个版本， 例如正弦函数 sin ，一个普通版本 sin ，另外一个不精确但速度极快的__sin版本。

?

内置变量

1 2 3 4 5

/* gridDim, blockIdx, blockDim, threadIdx, wrapsize. 这些内置变量不容许赋值的 */

?

编写程序

1 2 3 4 5 6 7

/* 目前CUDA仅能良好的支持C，在编写含有CUDA代码的程序时， 首先要导入头文件cuda_runtime_api.h。文件名后缀为.cu，使用nvcc编译器编译。 目前最新的CUDA版本为5.0，能够在官方网站下载最新的工具包，网址为： https://developer.nvidia.com/cuda-downloads 该工具包内包含了ToolKit、样例等，安装起来比原先的版本也方便了不少。 */

?

相关扩展

1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91 92 93

1 GPU硬件 // i GPU一个最小单元称为Streaming Processor(SP)，全流水线单事件无序微处理器， 包含两个ALU和一个FPU，多组寄存器文件（ register file，不少寄存器的组合）， 这个SP没有cache。事实上，现代GPU就是一组SP的array，即SPA。 每个SP执行一个 thread   // ii 多个SP组成Streaming Multiprocessor(SM)。 每个SM执行一个block。每一个SM包含8个SP； 2个special function unit(SFU)： 这里面有4个FPU能够进行超越函数和插值计算 MultiThreading Issue Unit：分发线程指令 具备指令和常量缓存。 包含shared memory   // iii Texture Processor Cluster(TPC) ：包含某些其余单元的一组SM   2 Single-Program Multiple-Data （SPMD）模型    // i CPU以顺序结构执行代码， GPU以threads blocks组织并发执行的代码，即无数个threads同时执行   // ii 回顾一下CUDA的概念： 一个kernel程序执行在一个grid of threads blocks之中 一个threads block是一批相互合做的threads： 能够用过__syncthreads同步； 经过shared memory共享变量，不一样block的不能同步。   // iii Threads block声明： 能够包含有1到512个并发线程，具备惟一的blockID，能够是1,2,3D 同一个block中的线程执行同一个程序，不一样的操做数，能够同步，每一个线程具备惟一的ID   3 线程硬件原理   // i GPU经过Global block scheduler来调度block， 根据硬件架构分配block到某一个SM。 每一个SM最多分配8个block，每一个SM最多可接受768个 thread （能够是一个block包含512个 thread ， 也能够是3个block每一个包含256个 thread （3*256=768！））。 同一个SM上面的block的尺寸必须相同。每一个线程的调度与ID由该SM管理。   // ii SM满负载工做效率最高！考虑某个Block，其尺寸能够为8*8,16*16,32*32 8*8：每一个block有64个线程， 因为每一个SM最多处理768个线程，所以须要768/64=12个block。 可是因为SM最多8个block，所以一个SM实际执行的线程为8*64=512个线程。 16*16：每一个block有256个线程，SM能够同时接受三个block，3*256=768，满负载 32*32：每一个block有1024个线程，SM没法处理！   // iii Block是独立执行的，每一个Block内的threads是可协同的。   // iv 每一个线程由SM中的一个SP执行。 固然，因为SM中仅有8个SP，768个线程是以warp为单位执行的， 每一个warp包含32个线程，这是基于线程指令的流水线特性完成的。 Warp是SM基本调度单位，实际上，一个Warp是一个32路SIMD指令 。基本单位是half-warp。 如，SM满负载工做有768个线程，则共有768/32=24个warp ，每一瞬时，只有一组warp在SM中执行。 Warp所有线程是执行同一个指令， 每一个指令须要4个 clock cycle，经过复杂的机制执行。   // v 一个thread的一辈子： Grid在GPU上启动； block被分配到SM上； SM把线程组织为warp； SM调度执行warp； 执行结束后释放资源； block继续被分配....   4 线程存储模型   // i Register and local memory：线程私有，对程序员透明。 每一个SM中有8192个 register ，分配给某些block， block内部的 thread 只能使用分配的寄存器。 线程数多，每一个线程使用的寄存器就少了。   // ii shared memory：block内共享，动态分配。 如__shared__ float region[N]。 shared memory 存储器是被划分为16个小单元， 与half-warp长度相同，称为bank，每一个bank能够提供本身的地址服务。 连续的32位word映射到连续的bank。 对同一bank的同时访问称为bank conflict。 尽可能减小这种情形。   // iii Global memory：没有缓存！容易称为性能瓶颈，是优化的关键！ 一个half-warp里面的16个线程对global memory的访问能够被coalesce成整块内存的访问，若是： 数据长度为4,8或16bytes；地址连续；起始地址对齐；第N个线程访问第N个数据。 Coalesce能够大大提高性能。   // uncoalesced Coalesced方法：若是全部线程读取同一地址， 不妨使用constant memory； 若是为不规则读取可使用texture内存 若是使用了某种结构体，其大小不是4 8 16的倍数， 能够经过__align(X)强制对齐，X=4 8 16