CUDA编程

CUDA编程(1)周斌@NVIDIA&USTC2014年10月致谢某些幻灯片来自DavidKirk和Wen-meiHwu’sUIUC课件大部分幻灯片来自PatrickCozziUniversityofPennsylvaniaCIS565GPU架构概览GPU特别适用于密集计算，高度可并行计算图形学晶体管主要被用于:执行计算而不是:缓存数据控制指令流GPU架构概览晶体管用途Imagefrom:让我们开始用CUDA来编程!GPU计算的历史2001/2002–研究人员把GPU当做数据并行协处理器GPGPU这个新领域从此诞生2007–NVIDIA发布CUDACUDA–全称ComputeUniformDeviceArchitecture统一计算设备架构GPGPU发展成GPUComputing2008–Khronos发布OpenCL规范CUDA的一些信息层次化线程集合Ahierarchyofthreadgroups共享存储Sharedmemories同步BarriersynchronizationCUDA术语Host–即主机端通常指CPU采用ANSI标准C语言编程Device–即设备端通常指GPU(数据可并行)采用ANSI标准C的扩展语言编程Host和Device拥有各自的存储器CUDA编程包括主机端和设备端两部分代码CUDA术语Kernel–数据并行处理函数通过调用kernel函数在设备端创建轻量级线程线程由硬件负责创建并调度类似于OpenGL的shader?CUDA核函数（Kernels）在N个不同的CUDA线程上并行执行ThreadIDExecutionConfigurationDeclarationSpecifierCUDA程序的执行Imagefrom:线程层次ThreadHierarchiesGrid–一维或多维线程块(block)一维或二维Block–一组线程一维，二维或三维一个Grid里面的每个Block的线程数是一样的block内部的每个线程可以：同步synchronize访问共享存储器sharedmemory线程层次ThreadHierarchiesImagefrom:线程层次ThreadHierarchiesBlock–一维，二维或三维例如:索引数组，矩阵，体线程层次ThreadHierarchiesThreadID:Scalarthreadidentifier线程索引:threadIdx一维Block:ThreadID==ThreadIndex二维Block(Dx,Dy)ThreadIDofindex(x,y)==x+yDy三维Block(Dx,Dy,Dz)ThreadIDofindex(x,y,z)==x+yDy+zDxDy线程层次ThreadHierarchies1ThreadBlock2DBlock2DIndex线程层次ThreadHierarchiesThreadBlock线程块线程的集合G80andGT200:多达512个线程Fermi:多达1024个线程位于相同的处理器核共享所在核的存储器线程层次ThreadHierarchiesThreadBlock线程块线程的集合G80andGT200:多达512个线程Fermi:多达1024个线程位于相同的处理器核心（SM）共享所在核心的存储器Imagefrom:线程层次ThreadHierarchies块索引:blockIdx维度:blockDim一维或二维或三维2DThreadBlock16x16Threadsperblock线程层次ThreadHierarchies例如:N=32每个块有16x16个线程(跟N无关)threadIdx([0,15],[0,15])Grid里面有2x2个线程块BlockblockIdx([0,1],[0,1])blockDim=16i=[0,1]*16+[0,15]线程层次ThreadHierarchies线程块之间彼此独立执行任意顺序:并行或串行被任意数量的处理器以任意顺序调度处理器的数量具有可扩展性线程层次ThreadHierarchiesImagefrom:线程层次ThreadHierarchies一个块内部的线程共享容量有限的低延迟存储器同步执行合并访存__syncThreads()Barrier–块内线程一起等待所有线程执行都某处语句轻量级Imagefrom:线程层次ThreadHierarchiesCUDA内存传输Imagefrom:Host可以从device往返传输数据Globalmemory全局存储器Constantmemory常量存储器Imagefrom:内存传输cudaMalloc()在设备端分配globalmemorycudaFree()释放存储空间Imagefrom:内存传输Codefrom:内存传输Codefrom:内存传输Codefrom:内存传输cudaMemcpy()内存传输HosttohostHosttodeviceDevicetohostDevicetodeviceHostDeviceGlobalMemoryCUDA内存传输Codefrom:内存传输Codefrom:(host)Destination(device)CUDA内存传输Codefrom:内存传输MatrixMultiply矩阵相乘算法提示向量点乘行优先或列优先?每次点乘结果输出一个元素MatrixMultiply矩阵相乘样例Imagefrom:P=M*N假定MandN是方阵1,000x1,000矩阵1,000,000点乘Each1,000multiplesand1,000addsMatrixMultiply矩阵相乘样例MatrixMultiply:CPU实现Codefrom:(float*M,float*N,float*P,intwidth)‏{for(inti=0;iwidth;++i)‏for(intj=0;jwidth;++j){floatsum=0;for(intk=0;kwidth;++k){floata=M[i*width+k];floatb=N[k*width+j];sum+=a*b;}P[i*width+j]=sum;}}MatrixMultiply:CUDA算法框架Codefrom:::CUDA算法框架Codefrom::CUDA算法框架第1步在算法框架中添加CUDAmemorytransfersMatrixMultiply:CUDA算法框架MatrixMultiply:数据传输Codefrom:分配输入MatrixMultiply:数据传输Codefrom:分配输出MatrixMultiply:数据传输Codefrom: