刘超俊-cuda编程.pptx

CUDA编程介绍——刘超俊主要内容GPU和CPU的比较Cuda编程模型核函数的定义以及调用线程结构线程通信（共享内存）及同步 --向量点积运算的实例ControlALUALUALUALUDRAMCacheDRAMGPU与CPU的比较CPU的内部结构：控制单元、逻辑单元和存储单元运算器：算术运算(如加减乘除)和逻辑运算( 如逻辑加、逻辑乘和非运算) 控制器：读取各种指令,并对指令进行分析,作出相应的控制。若干个寄存器：可直接参与运算并存放运算的中间结果。 GPU采用了大量的执行单元，这些执行单元可以轻松的加载并行处理，而不像CPU那样的单线程处理。另外，现代的GPU也可以在每个指令周期执行更多的单一指令。 CPUGPU为什么使用GPU1：强大的处理能力2：高显存带宽. . .. . .GPU编程模型并发执行（主机）串行指令（设备）并行指令（主机）串行指令（设备）并行指令在CUDA编程中，并行指令通过定义运行在GPU中的核函数实现。核函数的定义及调用核函数定义核函数调用与标准c语言的两点区别：1、定义一个核函数时，在函数返回类型前面加_global_；2、在调用核函数时，在核函数函数名后面有两个参数设置，并用 括起。线程结构线程格构成线程块构成线程线程结构目前一个核函数中只能有一个线程格，而线程块和线程可以是一维、二维或者三维的组织形式，CUDA中使用了一种数据类型dim3来定义这种组织形式，dim3是基于uint3定义的矢量类型，相当于由3个unsigned int型组成的结构体，其有内建变量threadIdx和blockIdx。这样就可以使用一维、二维或者三维的索引来标识线程。threadID:对于一维的线程块：threadID=threadIdx.x;对于大小为（Dx,Dy）的二维线程块： threadID=threadIdx.x+threadIdx.y*Dx;对于大小为（Dx,Dy,Dz）的三维线程块： threadID=threadIdx.x+threadIdx.y*Dx+threadIdx.z*Dx*Dy;线程结构// Kernel definitionglobal void MatAdd(float A[N][N], float B[N][N],float C[N][N]){ int i = threadIdx.x; int j = threadIdx.y; C[i ][ j ] = A[i][ j ] + B[i ][ j ];}int main(){...// Kernel invocation with one block of N * N * 1 threadsint numBlocks = 1;dim3 threadsPerBlock(N, N);MatAddnumBlocks, threadsPerBlock(A, B, C);...}由于硬件限制，每个线程块中的线程数量是有限制的，在目前的GPU上，一个线程块中的线程数一般不超过1024个线程。线程结构当一个线程块的线程数无法满足现实应用的需求时，我们可以启用多个线程块。此时总的线程数等于每个块内的线程数乘以线程块数。注意：整型数据相除的特点线程结构由于硬件的限制，线程块的数量同样是有限制的，通常不多于65536个块。当硬件最大线程数无法满足现实需求时，采用的策略是启用最大线程数，然后在某些线程中进行多次计算。现假设硬件最大线程块数目为65536个，每个块最大线程数为1024。若要对两个矩阵元素个数大于65536*1024矩阵求和，有如下程序：// Kernel definition_global_ void MatAdd(float A[N][N], float B[N][N],float C[N][N]){int i = blockIdx.x * blockDim.x + threadIdx.x;int j = blockIdx.y *blockDim.y + threadIdx.y;while(iNjN){C[i ][ j ] = A[i][ j ] + B[i ][ j ];i+=blockDim.x*gridDim.x;j+=blockDim.y*gridDim.y;}}int main(){...// Kernel invocationdim3 threadsPerBlock(32, 32); //32*32=1024dim3 numBlocks(256, 256); //256*256=65536MatAddnumBlocks, threadsPerBlock(A, B, C);...}共享内存及同步CUDA C中存在一种共享内存的机制，同一个块内的线程可以共享内存，实现块内多个线程的通信与协作。在访问共享内存时