CUDA的设备,流处理器核线程块threadblock线程网格‌gridDim块block和多gpu设备同步数据概念CUDA的设备,流处理器核线程块threadblock线程网格‌gridDim块block和多gpu设备同步数据概念CUDA的设备,流处理器核线程块threadblock线程网格‌gridDim块block和多gpu设备同步数据概念前言一、介绍CUDA编程的并行原理了解线程、线程块、网格等概念了解CUDA的同步机制blockSize线程块尺寸‌gridSize网格尺寸‌二、在CUDA编程中核函数(kernel function)内部预定义了以下关键变量用于标识线程的组织结构和执行环境1、 threadIdx‌线程‌2、blockIdx‌线程块3、blockDim‌线程块维度4、gridDim‌ 网格5、 warpSize‌三、CUDA流的应用CUDASTREAMCUDA流的使用、同步用CUDA流完成矩阵运算1 流Streams在GPU上按顺序的操作2 多GPU 编码总结前言SIMT和SIMDCUDA执行的是SIMT架构单指令多线程架构SIMT和SIMDSingle Instruction, Multiple Data类似SIMT应该算是SIMD的升级版更灵活但效率略低SIMT是NVIDIA提出的GPU新概念。二者都通过将同样的指令广播给多个执行官单元来实现并行。一个主要的不同就是SIMD要求所有的vector element在一个统一的同步组里同步的执行而SIMT允许线程们在一个warp中独立的执行。一、介绍CUDA编程的并行原理了解线程、线程块、网格等概念了解CUDA的同步机制基于硬件的支持通过cuda来实现对底层GPU的调用关于这部分内容首先需要熟悉一些关键名词。thread一个CUDA的并行程序会被以许多个thread来执行每个thread都有自己的register 和 local memory 的空间。block数个thread会组成一个block同一个block中的thread可以同步运行他们通过shared memory来进行通信。grid多个block则会再构成grid每个grid会有自己的global memory、constant memory 和 texture memory。warpwarp是SM的基本执行单元一个block里面的线程通过warp进行调用使用SIMT模式。如A100机器每个warp可以执行32个thread。....// 2Dint thread16;int grid(numRows()*numCols()thread-1)/(thread*thread);constdim3blockSize(thread,thread);constdim3gridSize(grid);rgba_to_greyscalegridSize,blockSize(d_rgbaImage,d_greyImage,numRows(),numCols());.....blockSize线程块尺寸‌类型为dim3表示每个线程块包含的线程数量13在示例中dim3 blockSize(thread, thread)创建了二维线程块每块包含thread×thread个线程5每个线程块最大线程数限制为102457同一线程块内的线程可通过共享内存通信gridSize网格尺寸‌类型为dim3表示网格中包含的线程块数量在示例中dim3 gridSize(grid)创建了一维网格包含grid个线程块3最大网格维度为65535x/y/z方向‌执行配置gridSize, blockSize‌该语法指定核函数启动时的并行执行结构35总线程数 gridSize.x * gridSize.y * gridSize.z * blockSize.x * blockSize.y * blockSize.z13二、在CUDA编程中核函数(kernel function)内部预定义了以下关键变量用于标识线程的组织结构和执行环境uint3 __device_builtin__ __STORAGE__ threadIdx;// 线程uint3 __device_builtin__ __STORAGE__ blockIdx;// 线程块dim3 __device_builtin__ __STORAGE__ blockDim;// 线程块维度dim3 __device_builtin__ __STORAGE__ gridDim;// 网格int __device_builtin__ __STORAGE__ warpSize;// 计算x开始坐标位置//unsigned int xblockIdx.x*blockDim.xthreadIdx.x;unsigned int yblockIdx.y*blockDim.ythreadIdx.y;1、 threadIdx‌线程类型为uint3表示当前线程在其所属线程块(block)中的三维索引。通过threadIdx.x、threadIdx.y、threadIdx.z可访问各维度坐标‌2、blockIdx‌线程块类型为uint3表示当前线程块在网格(grid)中的三维索引。通过blockIdx.x、blockIdx.y、blockIdx.z可访问各维度坐标3、blockDim‌线程块维度类型为dim3表示线程块的维度即每个线程块包含的线程数。通过blockDim.x、blockDim.y、blockDim.z可获取各维度大小4、gridDim‌ 网格类型为dim3表示网格的维度即网格包含的线程块数。通过gridDim.x、gridDim.y、gridDim.z可获取各维度大小5、 warpSize‌类型为int表示GPU的线程束(warp)大小通常为32。线程束是GPU调度的基本单位同一线程束内的线程执行相同指令这些变量共同定义了CUDA的线程层次结构网格(grid)由多个线程块(block)组成线程块由多个线程(thread)组成总线程数计算公式gridDim.x * gridDim.y * gridDim.z * blockDim.x * blockDim.y * blockDim.z三、CUDA流的应用CUDASTREAMCUDA流的使用、同步用CUDA流完成矩阵运算1 流Streams在GPU上按顺序的操作定义流创建流使用流销毁流//1. 定义流cudaStream_t s1;//2. 创建流cudaStreamCreate(s1);//3. 使用流cudaMemcpyAsync(....,s1);//4. 销毁流cudaStreamDestroy(s1);2 多GPU 编码两个GPU之间数据访问是通过PCIeint gpu10;int gpu21;void*d_ANULL;//1. 设置当前使用gpu索引idcudaSetDevice(gpu1);cudaMalloc(d_A,1024);int accessible0;cudaDeviceCanAccessPeer(accessible,gpu1,gpu2);if(accessible){cudaSetDevice(gpu2);//设置gpu2 可以访问gpu1的内存地址cudaDeviceEnablePeerAccess(gpu1,0);// kernelx, y, z核心函数}两个GPU设备之间的数据拷贝函数 字节cudaMemcpyPeerAsync(void* dst, int dstDevice, const void* src, int srcDevice, size_t count, cudaStream_t stream );如果两个设备允许字节在最短的PCIe路径PCIe路径上传输如果两个设备不允许CUDA驱动通过CPU驱动通过CPU memory传输CUDADeviceQuery(RuntimeAPI)version(CUDARTstaticlinking)Detected1CUDACapabledevice(s)Device0:NVIDIA GeForce MX150CUDADriver Version/Runtime Version12.9/12.9CUDACapability Major/Minor version number:6.1Total amountofglobal memory:2048MBytes(2147352576bytes)MapSMtoCoresforSM6.1isundefined.Default to use192Cores/SMMapSMtoCoresforSM6.1isundefined.Default to use192Cores/SM(3)Multiprocessorsx(192)CUDACores/MP:576CUDACoresGPUClock rate:1532MHz(1.53GHz)Memory Clock rate:3004Mhz Memory Bus Width:64-bitL2Cache Size:524288bytes Max Texture DimensionSize(x,y,z)1D(131072),2D(131072,65536),3D(16384,16384,16384)Max Layered TextureSize(dim)x layers 1D(32768)x2048,2D(32768,32768)x2048Total amountofconstant memory:65536bytes Total amountofshared memory per block:49152bytes Total numberofregisters available per block:65536Warp size:32Maximum numberofthreads per multiprocessor:2048Maximum numberofthreads per block:1024Maximum sizesofeach dimensionofa block:1024x1024x64Maximum sizesofeach dimensionofa grid:2147483647x65535x65535Maximum memory pitch:2147483647bytes Texture alignment:512bytes Concurrent copy and kernel execution:Yeswith5copyengine(s)Run time limit on kernels:Yes IntegratedGPUsharing Host Memory:No Support host page-locked memory mapping:Yes Alignment requirementforSurfaces:Yes Device hasECCsupport:Disabled Device supports UnifiedAddressing(UVA):Yes DevicePCIBusID/PCIlocationID:2/0Compute Mode:Default(multiple host threads can use::cudaSetDevice()withdevice simultaneously)deviceQuery,CUDADriverCUDART,CUDADriver Version12.9,CUDARuntime Version12.9,NumDevs1,Device0NVIDIAGeForceMX150总结CUDA项目源码地址https://github.com/chensongpoixs/ccuda_sample