CUDA超大规模并行程序设计PPT课件

1、CUDA 超大规模并行程序设计超大规模并行程序设计赵开勇赵开勇.hk/kyzhaohttp:/ 从从GPGPU到到CUDAn 并行程序并行程序组织组织n 并行并行执行模型执行模型n CUDA基础基础n 存储器存储器n CUDA程序设计工具程序设计工具n 新一代新一代Fermi GPU33Graphic Processing Unit (GPU)n 用于个人计算机、工作站和游戏机的专用图像显示设备用于个人计算机、工作站和游戏机的专用图像显示设备l 显示卡显示卡 nVidia和和ATI (now AMD)是主要制造商是主要制造商 Intel准备通过准备通过Larra

2、bee进入这一市场进入这一市场l 主板集成主板集成 Intel443维图像流水线维图像流水线n一帧典型图像一帧典型图像l1M trianglesl3M verticesl25M fragmentsVertexProcessorFragmentProcessorRasterizerFramebufferTextureCPUGPUn30 frames/sl30M triangles/sl90M vertices/sl750M fragments/s55传统传统GPU架构架构Graphics programVertex processorsFragment processorsPixel opera

3、tionsOutput image66GPU的强大运算能力的强大运算能力数据级并行: 计算一致性专用存储器通道有效隐藏存储器延时77General Purpose Computing on GPU (GPGPU)88GPGPUn 核心思想核心思想l 用图形语言描述通用计算用图形语言描述通用计算问题问题l 把数据映射到把数据映射到vertex或者或者fragment处理器处理器n 但是但是l 硬件资源使用不充分硬件资源使用不充分l 存储器访问方式严重受限存储器访问方式严重受限l 难以调试和查错难以调试和查错l 高度图形处理和编程技巧高度图形处理和编程技巧99G80 GPUL2FBSPSPL1TF

4、Thread ProcessorVtx Thread IssueSetup / Rstr / ZCullGeom Thread IssuePixel Thread IssueInput AssemblerHostSPSPL1TFSPSPL1TFSPSPL1TFSPSPL1TFSPSPL1TFSPSPL1TFSPSPL1TFL2FBL2FBL2FBL2FBL2FBStreaming Multiprocessor (SM)Streaming Processor (SP)1010CUDA: Compute Unified Device Architecturen CUDA: 集成集成CPU + GP

5、U C应用程序应用程序n 通用并行计算模型通用并行计算模型l 单指令、多数据执行模式单指令、多数据执行模式 (SIMD) 所有线程执行同一段代码所有线程执行同一段代码(1000s threads on the fly) 大量并行计算资源处理不同数据大量并行计算资源处理不同数据l 隐藏存储器延时隐藏存储器延时 提升计算通信比例提升计算通信比例 合并相邻地址的内存访问合并相邻地址的内存访问 快速线程切换快速线程切换1 cycleGPU vs. 1000 cyclesCPU1111Evolution of CUDA-Enabled GPUsn Compute 1.0: basic CUDA comp

6、atibilityl G80n Compute 1.1: asynchronous memory copies and atomic global operationsl G84, G86, G92, G94, G96, and G98n Compute 1.2: dramatically improved memory coalescing rules, double the register count, intra-warp voting primitives, atomic shared memory operations l GT21Xn Compute 1.3: double pr

7、ecisionl GT2001212CUDA成功案例成功案例1313提纲提纲n 从从GPGPU到到CUDAn 并行程序组织并行程序组织n 并行并行执行模型执行模型n CUDA基础基础n 存储器存储器n CUDA程序设计工具程序设计工具n 新一代新一代Fermi GPU1414并行性的维度并行性的维度n 1维维ly = a + b /y, a, b vectorsn 2维维lP = M N /P, M, N matricesn 3维维lCT or MRI imaginga0a1anb0b1bny0y1yn+= =1515并行线程组织结构并行线程组织结构n Thread: 并行的基本单位并行的基本

8、单位n Thread block: 互相合作的线程组互相合作的线程组l Cooperative Thread Array (CTA)l 允许彼此同步允许彼此同步l 通过快速共享内存交换数据通过快速共享内存交换数据l 以以1维、维、2维或维或3维组织维组织l 最多包含最多包含512个线程个线程n Grid: 一组一组thread blockl 以以1维或维或2维组织维组织l 共享全局内存共享全局内存n Kernel: 在在GPU上执行的核心程序上执行的核心程序l One kernel one gridHostKernel 1Kernel 2DeviceGrid 1Block(0, 0)Block

9、(1, 0)Block(2, 0)Block(0, 1)Block(1, 1)Block(2, 1)Grid 2Block (1, 1)Thread(0, 1)Thread(1, 1)Thread(2, 1)Thread(3, 1)Thread(4, 1)Thread(0, 2)Thread(1, 2)Thread(2, 2)Thread(3, 2)Thread(4, 2)Thread(0, 0)Thread(1, 0)Thread(2, 0)Thread(3, 0)Thread(4, 0)1616Parallel Program Organization in CUDAThreadThrea

10、d blockGridSPSoftwareHardwareSMSMGPUTPCSMSMSMTPCSMSMSMTPCSMSMSM1717并行线程执行并行线程执行n 调用调用kernel function 需要指定执行配置需要指定执行配置n Threads和和blocks具有具有IDsl threadIdx: 1D, 2D, or 3Dl blockIdx: 1D, or 2Dl 由此决定相应处理数据由此决定相应处理数据_global_ void kernel(.);dim3 DimGrid(3, 2); / 6 thread blocks dim3 DimBlock(16, 16); / 256

11、 threads per block kernel (.);1818实例实例1: Element-Wise Addition/CPU program/sum of two vectors a and bvoid add_cpu(float *a, float *b, int N)for (int idx = 0; idxN; idx+) aidx += bidx;void main().fun_add(a, b, N);/CUDA program/sum of two vectors a and b _global_ void add_gpu(float *a, float *b, int N

12、)Int idx =blockIdx.x* blockDim.x+ threadIdx.x;if (idx N)aidx += bidx;void main().dim3 dimBlock (256);dim3 dimGrid( ceil( N / 256 );fun_add(a, b, N);1919提纲提纲n 从从GPGPU到到CUDAn 并行程序组织并行程序组织n 并行并行执行模型执行模型n CUDA基础基础n 存储器存储器n CUDA程序设计工具程序设计工具n 新一代新一代Fermi GPU2020CUDA Processing Flow2121并行线程执行并行线程执行n SM内以内以

13、(warp即即32 threads)为单为单位并行执行位并行执行l Warp内线程执行同一条指令内线程执行同一条指令l Half-warp是存储操作的基本单位是存储操作的基本单位WarpBlock 0Block 1Block 22222控制流控制流(Control Flow)n 同一同一warp内的分支语句可能执行内的分支语句可能执行不同的指令路径不同的指令路径l 不同指令路径的线程只能顺序执行不同指令路径的线程只能顺序执行 每次执行每次执行warp中一条可能的路径中一条可能的路径 N条指令路径条指令路径1/N throughputl 只需要考虑同一只需要考虑同一warp即可，不同即可，不同w

14、arp的不同的指令路径不具相关性的不同的指令路径不具相关性l G80上使用指令预测技术加速指令执上使用指令预测技术加速指令执行行2323控制流控制流(Control Flow)n 常见情况常见情况: 分支条件是分支条件是thread ID的函数时的函数时, 容易容易导致分支（导致分支（divergence）l Example with divergence: If (threadIdx.x 2) 在在thread block产生两条不同指令路径产生两条不同指令路径 Branch granularity 2) 也在也在thread block产生两条不同指令路径产生两条不同指令路径 Branch

15、 granularity is a whole multiple of warp size 同一同一warp的所有线程具备相同指令路径的所有线程具备相同指令路径2424线程同步线程同步nvoid _syncthreads();lBarrier synchronizationl同步同步thread block之内的所有线程之内的所有线程l避免访问共享内存时发生避免访问共享内存时发生RAW/WAR/WAW 冒险冒险(hazard)_shared_ float scratch256;scratchthreadID = beginthreadID;_syncthreads();int left = s

16、cratchthreadID -1;在此等待在此等待，直至所直至所有线程到达才开始有线程到达才开始执行下面的代码执行下面的代码2525Dead-Lock with _syncthreadsn Dead-lock iflSome threads have val larger than thresholdlAnd others not_global_ void compute(.)/ do some computation for valif( val threshold )return; _syncthreads(); / work with val & store it return

17、; 2626提纲提纲n 从从GPGPU到到CUDAn 并行程序组织并行程序组织n 并行并行执行模型执行模型n CUDA基础基础n 存储器存储器n CUDA程序设计工具程序设计工具n 新一代新一代Fermi GPU2727CUDA扩展语言结构扩展语言结构n Declspecsl global, device, l shared, local, constantn Keywordsl threadIdx, blockIdxl threadDim, blockDimn Intrinsicsl _syncthreadsn Runtime APIl Memory, symbol, execution m

18、anagementn Function launch_device_ float filterN; _global_ void convolve (float *image) _shared_ float regionM; . regionthreadIdx = imagei; _syncthreads() . imagej = result;/ Allocate GPU memoryvoid *myimage = cudaMalloc(bytes)/ 100 blocks, 10 threads per blockfoo (parameters);2828存储器空间存储器空间n R/W pe

19、r-thread registersl 1-cycle latencyn R/W per-thread local memoryl Slow register spilling to global memoryn R/W per-block shared memoryl 1-cycle latencyl “_shared_”l But bank conflicts may drag downn R/W per-grid global memoryl 500-cycle latencyl “_device_”l But coalescing accessing could hide latenc

20、yn Read only per-grid constant and texture memoriesl 500-cycle latency, but cached2929GPU Global Memory分配分配n cudaMalloc()l 分配显存中的分配显存中的l 两个参数两个参数 对象数组指针和数组尺寸对象数组指针和数组尺寸n cudaFree()l 释放显存中的释放显存中的 对象数组指针对象数组指针int blk_sz = 64;float* Md;int size = blk_sz * blk_sz * sizeof(float);cudaMalloc(void*)&Md

21、, size);cudaFree(Md);3030Host Device数据交换数据交换n cudaMemcpy()l Memory data transferl Requires four parameters Pointer to destination Pointer to source Number of bytes copied Type of transfer Host to Host, Host to Device, Device to Host, Device to DevicecudaMemcpy(Md, M.elements, size, cudaMemcpyHostToD

22、evice);cudaMemcpy(M.elements, Md, size, cudaMemcpyDeviceToHost);3131CUDA函数定义函数定义n _global_ 定义定义kernel函数函数l 必须返回必须返回voidn _device_ 函数函数l 不能用不能用&运算符取地址运算符取地址, 不支持递归调用不支持递归调用, 不支持静态变量不支持静态变量(static variable), 不支持可变长度参数函数调用不支持可变长度参数函数调用Executed on the:Only callable from the:_device_ float DeviceFunc

23、()devicedevice_global_ void KernelFunc()devicehost_host_ float HostFunc()hosthost3232CUDA数学函数数学函数npow, 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, etc.l只支持标量运算只支持标量运算l许多函数有一个快速

24、、较不精确的对应版许多函数有一个快速、较不精确的对应版本本 以以”_”为前缀，如为前缀，如_sin() 编译开关编译开关-use_fast_math强制生成该版本的目强制生成该版本的目标码标码3333实例实例2: 矩阵相乘矩阵相乘n 矩阵数据类型矩阵数据类型 不属于不属于CUDA!n 单精度浮点数单精度浮点数n width height个元素个元素n 矩阵元素在矩阵元素在elements中中l 1-D数组存放矩阵数据数组存放矩阵数据l Row-major storagetypedef struct int width; int height; float* elements; Matrix;A

25、BCWM.width = N.heightIM.heightM.widthN.width3434实例实例2: 矩阵相乘矩阵相乘n C = A B of size WIDTH x WIDTHn 一个线程处理一个矩阵元素一个线程处理一个矩阵元素l 简化简化: 假定假定 WIDTH x WIDTH 512 只需要只需要一个一个thread blockl 线程载入线程载入A的一行和的一行和B的一列的一列l A和和B的一对相应元素作一次乘法和的一对相应元素作一次乘法和一次加法一次加法ABCWIDTHWIDTHWIDTHWIDTH3535CUDA Implementation Host Side/ Mat

26、rix multiplication on the devicevoid Mul(const Matrix A, const Matrix B, Matrix C) int size = A.width A.width sizeof(float); / Load M and N to the device float *Ad, *Bd, *Cd; cudaMalloc(void*)&Ad, size);/matrix stored in linear order cudaMemcpy(Ad, A.elements, size, cudaMemcpyHostToDevice); cuda

27、Malloc(void*)&Bd, size); cudaMemcpy(Bd, B.elements, size, cudaMemcpyHostToDevice); / Allocate C on the device cudaMalloc(void*)&Cd, size);3636CUDA Implementation Host Side / Launch the device computation threads! dim3 dimGrid(1); dim3 dimBlock(M.width, M.width); Muld(Ad, Bd, Cd, M.width); /

28、Read P from the device copyFromDeviceMatrix(C.elements, Cd); cudaMemCopy(C, Cd, N * size, cudaMemcpyDeviceToHost); / Free device matrices cudaFree(Ad); cudaFree(Bd); cudaFree(Cd);3737CUDA Implementation Kernel/ Matrix multiplication kernel thread specification_global_ void Muld (float* Ad, float* Bd

29、, float* Cd, int width) / 2D Thread ID int tx = threadIdx.x; int ty = threadIdx.y; / cvalue is used to store the element of the matrix / that is computed by the thread float cvalue = 0;3838CUDA Implementation KernelABCWIDTHWIDTHWIDTHWIDTHtytx for (int k = 0; k width; +k) float ae = Adty * width + k;

30、 float be = Bd tx + k * width; cvalue += ae * be; / Write the matrix to device memory; / each thread writes one element Cdty * width + tx = cvalue;3939提纲提纲n 从从GPGPU到到CUDAn 并行程序组织并行程序组织n 并行并行执行模型执行模型n CUDA基础基础n 存储器存储器lShared memorylGlobal memoryn CUDA程序设计工具程序设计工具n 新一代新一代Fermi GPU4040共享存储器共享存储器(Shared

31、 Memory)n 设置于设置于streaming multiprocessor内部内部n 由一个线程块内部全部线程共享由一个线程块内部全部线程共享l 完全由软件控制完全由软件控制l 访问一个地址只需要访问一个地址只需要1个时钟周期个时钟周期4141共享存储器结构共享存储器结构n G80的共享存储器组织为的共享存储器组织为16 banksl Addressed in 4 bytesl Bank ID = 4-byte address % 16l 相邻相邻4-byte地址地址映射相邻映射相邻banksl 每一每一bank的带宽为的带宽为4 bytes per clock cyclen 对同一对同

32、一bank的同时访问导致的同时访问导致bank conflict l 只能顺序处理只能顺序处理l 仅限于同一线程块内的线程仅限于同一线程块内的线程Bank 15Bank 7Bank 6Bank 5Bank 4Bank 3Bank 2Bank 1Bank 000, 16, 32, 01, 17, 33, 02, 18, 34, 03, 19, 35, 15, 31, 47, 4242Bank Addressing实例实例n No Bank ConflictslLinear addressing stride = 1 (s=1)n No Bank ConflictslRandom 1:1 Perm

33、utation_shared_ float shared256;float foo = sharedthreadIdx.x;4343Bank Addressing实例实例n 2-way bank conflictslLinear addressing stride = 2 (s=2)n 8-way bank conflictslLinear addressing stride = 8 (s=8)_shared_ float shared256;float foo = shared2 * threadIdx.x;_shared_ float shared256;float foo = share

34、d8 * threadIdx.x;4444提纲提纲n 从从GPGPU到到CUDAn 并行程序组织并行程序组织n 并行并行执行模型执行模型n CUDA基础基础n 存储器存储器lShared memorylGlobal memoryn CUDA程序设计工具程序设计工具n 新一代新一代Fermi GPU4545全局内存全局内存(Global Memory)n 全局内存在全局内存在G80/G200上没有上没有缓存缓存l Constant memory和和texture memory有少量缓存有少量缓存n 存取延时存取延时l 400-600 clock cyclesn 非常容易成为性能瓶颈非常容易成为性

35、能瓶颈n 优化是提高性能的关键优化是提高性能的关键!4646Coalesced Global Memory Accesses4747Non-Coalesced Global Memory Accesses4848Non-Coalesced Global Memory Accesses4949Coalescing on 1.2 and Higher Devicesn Global memory access by threads in a half-warp can be coalescedl When the words accessed by all threads lie in the s

36、ame segment of size equal to: 32 bytes if all threads access 8-bit words 64 bytes if all threads access 16-bit words 128 bytes if all threads access 32-bit or 64-bit wordslAny pattern of addresses requested by the half-warp Including patterns where multiple threads access the same address5050Example

37、 of New Coalescing RulesAddress 0Thread 0Address 4Address Address 116Address 120Address 124Address 128Address Address 172Address 176Address 180Address 184Address 188Address 252Thread 1Thread 2Thread 3Thread Thread 14Thread 15Segment 0 (128B)Segment 1 (128B)Reduced to 32BSegment size is 32 bytes for

38、8-bit data, 64 bytes for 16-bit data, 128 bytes for 32-, 64- and 128-bit data.5151提纲提纲n 从从GPGPU到到CUDAn 并行程序组织并行程序组织n 并行并行执行模型执行模型n CUDA基础基础n 存储器存储器lShared memorylGlobal memoryn CUDA程序设计工具程序设计工具n 新一代新一代Fermi GPU5252下载下载CUDA软件软件n http:/ CUDA driverl硬件驱动硬件驱动n CUDA toolkitl工具包工具包n CUDA SDKl程序范例及动态链接库程序范

39、例及动态链接库n CUDA Visual Profilerl程序剖析工具程序剖析工具CPU(Host)CUDA Libraries (CUFFT& CUBLAS)CUDA Runtime LibrariesCUDA DriverApplicationGPU(Device)5353CUDA程序的编译程序的编译(compile)n CUDA源文件被源文件被nvcc处理处理l nvcc is a compiler drivern nvcc输出：输出：l PTX (Parallel Thread eXecution) Virtual ISA for multiple GPU hardware

40、Just-In-Time compilation by CUDA runtimel GPU binary Device-specific binary objectl Standard C code With explicit parallelismC/C+ CUDA ApplicationNVCCPTX CodeC/C+ CPU CodeGenericCUDA RuntimeSpecializedOther GPUsG80GT200CUDA Binary5454DEBUGn make dbg=1lCPU代码以代码以debug模式编译模式编译l可以用可以用debugger (e.g. gdb,

41、 visual studio)运行运行 但不能检查但不能检查GPU代码的中间结果代码的中间结果n make emu=1l在在CPU上以上以emulation方式顺序运行方式顺序运行l可以使用可以使用printf()打印中间结果打印中间结果 基本顺序执行基本顺序执行 但不能再现线程间的竞争但不能再现线程间的竞争(race)现象现象 浮点运算结果可能有微小的差别浮点运算结果可能有微小的差别5555检查资源使用检查资源使用n 使用使用-cubin flag编译开关编译开关n 检查检查.cubin文件的文件的”code”部分部分architecture sm_10abiversion 0modname

42、 cubincode name = BlackScholesGPUlmem = 0smem = 68reg = 20bar = 0bincode 0 xa0004205 0 x04200780 0 x40024c09 0 x00200780per thread local memoryper thread block shared memoryper thread registers5656CUDA Debugger: cuda-gdbn Released with CUDA 2.2l A ported version of GNU Debugger, gdbl Red Hat Enterpr

43、ise Linux 5.x 32-bit and 64-bitn Compiling with debug supportl nvcc g G foo.cu o foon Single-step individual warps (“next” or “step”)l Advances all threads in the same warpn Display device memory in the device kernell Data that resides in various GPU memory regions such as shared, local, and global

44、memoryn Switch to any CUDA block/threadl thread n Breaking into running applicationsl Ctrl+C to break into hanging programs5757“Nexus” GPU/CPU Development Suiten Major componentsl Nexus Debugger Source code debugger for GPU source code CUDA, DirectCompute, HLSL, l Nexus Analyzer System-wide event vi

45、ewer for both GPU & CPU eventsl Nexus Graphics Inspector For frame based, deep inspection of textures and geometryn Full integration with Visual Studiol Windows 7/Vistal Available on Oct. 29, 20095858提纲提纲n 从从GPGPU到到CUDAn 并行程序组织并行程序组织n 并行并行执行模型执行模型n CUDA基础基础n 存储器存储器lShared memorylGlobal memoryn C

46、UDA程序设计工具程序设计工具n 新一代新一代Fermi GPU59593 Major Generations of CUDA GPUsGPUG80GT200GT300 (Fermi)CUDA cores128240512Process (nm)904540Transistors681 Million1.4 Billion3.0 BillionDouble precision floating point capability None30 FMA ops/clock256 FMA ops/clockSingle precision floating point capability128 MAD ops/clock240 MAD ops/clock512 MAD ops/clockWarp scheduler112Special function units / SM224CUDA cores / SM8832Shared memory / SM16KB16KBConfigurable 48KB or 16KBL1 cache / SMNoneNoneConfigurable 16KB or 48KBL2 cach