文章目录1 软硬件架构知识2 kernel执行的性能部分3 并行性能表现3.1 线程束分支效率1 - 分化分支 / 分支数3.2 SM的实际占用率活跃的warp数量 / SM支持最大并发warp数量3.3 全局内存读取速度全局内存读取字节数 / 内核执行时间s3.4 全局内存读取效率实际使用的全局内存加载字节数 / 总全局内存加载字节数)3.5 线程束上指令数平均值指标越小越好越大可能是指令分化。3.6 线程束被阻塞的比例同步导致的线程停顿和__syncthreads()有关 )3.7 全局内存的写入吞吐量存储的全局内存数据量 / 执行时间s3.8 DRAM读取数据的吞吐量3.9 全局内存读取事务3.11 全局内存每次加载或储存所需的内存事务数3.10 全局内存写事务3.11 共享内存读写事务数量3.12 暂时没有用到的指标4 同步4.1 __syncthreads:4.2 cudaDeviceSynchronize():4.3 *warp内32个线程存在隐士同步*5 使用nsight compute测试性能博主公司重组求推荐大模型部署ai infra base上海的工作。my phone 156012371031 软硬件架构知识cuda 内存中最重要的是cuda的共享内存和全局内存。全局内存大概几个GB访问速度慢声明周期是整个内核的执行期间。shared memory是SM流式处理器上的高速内存只有十几KB一个block中所有的thread都可以访问。执行速度块声明周期短仅在当前block执行期间存在。shared memory的大小一般是固定的比如32KB// 这里是一个使用共享内存的例子__global__ voidkernel(float*A,float*B,float*C){int idxthreadIdx.xblockIdx.x*blockDim.x;// 将数据从全局内存加载到共享内存__shared__ float shared_A[THREAD_NUM];shared_A[threadIdx.x]A[idx];__syncthreads();// 在共享内存上进行计算shared_A[threadIdx.x]*2;__syncthreads();// 将结果写回全局内存C[idx]shared_A[threadIdx.x];}同一个block中的线程可以使用共享内存通讯不同block中的内存只能使用全局内存通讯。cudaMemcpy是隐式同步操作host会被阻塞到内存操作结束。cudaDeviceSynchronize 会阻塞设备的所有请求操作。nsight compute用于分析launch的kernel nsight systems用于分析程序整体流程。MFLOPS代表每秒百万次浮点运算 GFLOPS代表每秒十亿次浮点运算TFLOPS代表每秒万亿次浮点运算OPS 是 Operations Per Second 的缩写TOPS 代表每秒进行一万亿次运算。我的板卡的算力是 1.13TFLOPS, Pascal架构 显存带宽是48.1GB/s SM数量3。华为昇腾910算力 320TFLOPfp16, 640TOPSint8HBM带宽为400GB/s鲲云caisa4.0 算里 16TOPSint8关于cuda中的索引threadIdx.x// 线程在块中的索引第几个线程blockIdx.x// block在grid中的索引第几个blockblockDim.x// block在x轴上的维度x轴方向上有多少个线程gridDim.x// grid在x轴上有多少个blockblockDim.x*gridDim.x// x轴上的thread总数dim3的结构的定义dim3block(3);dim3grid((nElemblock.x-1)/block.x);printf(grid.x %d,grid.x);// grid是dim3实例,grid.x表示x轴上的block的数量printf(block.x %d,block.x);// block是dim3,block.x表示x轴上thread的数量关于二维抽象矩阵的索引这里的索引是为了把线程和块的索引映射到矩阵上只是一种抽象概念。idxthreadIdx.xblockIdx.x*blockDim.x idythreadIdx.yblockIdx.y*blockDim.y内存一维索引这里是为了将矩阵坐标映射到全局内存上硬件中的内存是线性存储的。idxidy*nxidx这里有一张blockthread索引-坐标(ix, iy)-全局内存索引(ifx)的对应关系图。一个不错的实用索引图SM中的重要组件cuda核心寄存器文件存放数组常量索引kernel中定义的变量local memory(较大的结构体和本地数组寄存器无法存放的变量)const memory(SM中有专门的const mem用__const__来声明)共享内存__shared__修饰 shared memory在kernel函数中声明生命周期和现成块一致一级缓存warp调度器指令分发单元加载存储单元不理解todo存储器作用域声明期RegisterThreadKernelLocal MemoryThreadKernelShared MemoryBlockKernelGlobal MemoryGridApplicationConstantGridApplication2 kernel执行的性能部分不同的grid和block配置速度也会不同sumMatrixOnGPU-2D.cu的速度测试使用sys/time.h中的函数gettimeofday测试非nvprof工具。二维grid 二维block的速度 单位s sumMatrixOnGPU2D add matrix at device 0.003266 sumMatrixOnGPU2D (512,1024), (32,16) 一维grid 一维block的速度 sumMatrixOnGPU1D add matrix at device 0.000003 sumMatrixOnGPU1D (128,1), (128,1) 二维grid 一维block的速度 sumMatrixOnGPU2dGrid1dBlock add matrix at device 0.000001 sumMatrixOnGPU2dGrid1dBlock (64,16384), (256,1)item含义block.x, block.yblock是dim3的定义表示一个线程块在xy轴上的维度即xy轴上有多少个线程。blockDim.x表示每个线程块在x维度上的线程数。blockIdx.x表示当前线程块在网格中的x维位置。threadIdx.x表示当前线程在其线程块内的x维索引。warp线程束是一个抽象的概念数量是32.SIMD, SIMT一个不同点SIMD要求所有的线程要求所有元素在一个统一的同步组中执行SIMT允许线程束中多个线程独立执行。SIMT中每个线程可以有自己的寄存器状态。线程块之内所有的线程有同步的方法保证线程块所有的线程都能同步完成但是没有线程块之间的同步方法。多个线程块可以分配到同一个SM中。Fermi架构中一个SM含有32个CUDA核心SM有两个线程束调度器和两个指令调度单元。当一个线程block被指定给一个SM时block中所有的线程被分成线程束两个调度器选择两个线程束再把一个指令从线程束中发到一个组上组里有16个 CUDA核心16个加载/ 存储单元。 Fermi架构每个SM可以同时处理48个线程束即可以在一个SM上同时常驻1536个线程。3种常见限制内核性能的因素存储带宽计算资源指令和内存延迟防止线程束分化的一个小技巧 if((tid/warpSize) % 2 0) 如果可能以warp为单位安排if…else分支可以避免线程束中任务分化。NVCC编译的用-g -G停止编译器的优化用工具nvprof --metrics branch_efficiency ./app 可以测试得到具体的branch效率 另外发现-O0无法停止分支预测。SM中有多个线程束从一个线程束切换到另一个线程束是没有开销的线程束的整个生命周期都是保存在芯片内的。SM中线程束的数量是由线程消耗的共享内存和寄存器决定的。高效内核的原则 1- 保持每个块中的线程数量是32的倍数。 2- 避免块太小块至少有128或256个线程。 3- 块的数量要多于SM的数量SM需要多个块保持活跃。4- 实验得到最佳的执行配置。3 并行性能表现3.1 线程束分支效率1 - 分化分支 / 分支数nvprof--metrics branch_efficiency./simpleDivergence Invocations Metric Name Metric Description Min Max Avg DeviceNVIDIA GeForce MX150 (0)Kernel:mathKernel1(float*)1branch_efficiency Branch Efficiency83.33%83.33%83.33%Kernel:mathKernel2(float*)1branch_efficiency Branch Efficiency100.00%100.00%100.00%Kernel:mathKernel3(float*)1branch_efficiency Branch Efficiency100.00%100.00%100.00%3.2 SM的实际占用率活跃的warp数量 / SM支持最大并发warp数量nvprof--metrics achieved_occupancy./sumMatrix Invocations Metric Name Metric Description Min Max Avg DeviceNVIDIA GeForce MX150 (0)Kernel:sumMatrixOnGPU2D(float*,float*,float*,int,int)1achieved_occupancy Achieved Occupancy0.8130790.8130790.813079提高SM实际占用率的思路增加block并发的数量优化block大小通常128,256,512, block太大无法同时并行多个block。block太小SM利用率不足。减少block中寄存器的使用变量可以考虑放在全局内存中。减少非必要的共享内存的使用。增加计算密度比如做循环展开增加计算操作和内存访问比率。如果kernel是内存访问密集的可以考虑使用共享内存或者常量内存来缓存数据。3.3 全局内存读取速度全局内存读取字节数 / 内核执行时间snvprof--metrics gld_throughput./sumMatrix Invocations Metric Name Metric Description Min Max Avg DeviceNVIDIA GeForce MX150 (0)Kernel:sumMatrixOnGPU2D(float*,float*,float*,int,int)1gld_throughput Global Load Throughput2.3684GB/s2.3684GB/s2.3684GB/s3.4 全局内存读取效率实际使用的全局内存加载字节数 / 总全局内存加载字节数)nvprof--metrics gld_efficiency./sumMatrix Invocations Metric Name Metric Description Min Max Avg DeviceNVIDIA GeForce MX150 (0)Kernel:sumMatrixOnGPU2D(float*,float*,float*,int,int)1gld_efficiency Global Memory Load Efficiency100.00%100.00%100.00%3.5 线程束上指令数平均值指标越小越好越大可能是指令分化。下面例子中reduceNeighbored的指令数是reduceNeighboredLess的两倍多797.25 vs 355.56 ).nvprof--metrics inst_per_warp./reduceInteger Invocations Metric Name Metric Description Min Max Avg DeviceNVIDIA GeForce MX150 (0)Kernel:reduceNeighbored(int*,int*,unsignedint)1inst_per_warp Instructions per warp797.250000797.250000797.250000Kernel:reduceNeighboredLess(int*,int*,unsignedint)1inst_per_warp Instructions per warp355.562500355.562500355.5625003.6 线程束被阻塞的比例同步导致的线程停顿和__syncthreads()有关 )可以看到最后5次迭代线程束展开后阻塞比例降低最后6个stride展开后使用的__syncthreads减少nvprof--metrics stall_sync./reduceInteger Kernel:reduceUnrolling8(int*,int*,unsignedint)1stall_sync Issue StallReasons(Synchronization)39.87%39.87%39.87%Kernel:reduceUnrollWarps8(int*,int*,unsignedint)1stall_sync Issue StallReasons(Synchronization)24.10%24.10%24.10%3.7 全局内存的写入吞吐量存储的全局内存数据量 / 执行时间snvprof--metricsgst_throughput ./app3.8 DRAM读取数据的吞吐量dram_read_throughput 更广泛它包括所有与 DRAM 相关的读取操作不仅限于全局内存读取gld_throughput还可能包括其他类型的内存读取例如共享内存或常量内存。据说dram_read_transactions包含gld_ghroughput.DRAM是out-chip, DRAM包含全局内存常量内存纹理内存本地内存onchip的内存寄存器共享内存。3.9 全局内存读取事务单位事务次数通常表示为“次数”transactionsnvprof--metrics gld_transactions./appgld_transactions 反映了对内存进行访问的次数而gld_throughput 关注的是数据传输的速率。例子Kernel:reduceSmem(int*,int*,unsignedint)gld_transactions Global Load Transactions838861083886108388610gst_transactions Global Store Transactions1310721310721310723.11 全局内存每次加载或储存所需的内存事务数nvprof--metrics gld_transactions_per_request,gst_transactions_per_request./appKernel:copyGmem(float*,float*,int,int)gld_transactions_per_request Global Load Transactions Per Request16.00000416.00000416.000004gst_transactions_per_request Global Store Transactions Per Request4.0000004.0000004.0000003.10 全局内存写事务nvprof--metrics gst_transactions./app3.11 共享内存读写事务数量共享内存容易出现bank冲突但是这里的事务数大意味着数据分布比较差和bank冲突没有直接关系。比如做transpose时候使用共享内存按照行写共享内存按照列读取共享内存warp内就会出现bank串行访问的冲突需要多个内存事务完成呢个。// shared memory 读操作nvprof--metrics shared_load_transactions_per_request./app// 写操作nvprof--metrics shared_store_transactions_per_request./app输出结果如Kernel:setRowReadRow(int*)shared_load_transactions_per_request Shared Memory Load Transactions Per Request1.0000001.0000001.000000shared_store_transactions_per_request Shared Memory Store Transactions Per Request1.0000001.0000001.000000Kernel:setColReadCol(int*)shared_load_transactions_per_request Shared Memory Load Transactions Per Request32.00000032.00000032.000000shared_store_transactions_per_request Shared Memory Store Transactions Per Request32.00000032.00000032.000000最理想情况一个请求只发出一个内存事务。setColReadCol的内存事务是32因为出现储存体bank冲突。3.12 暂时没有用到的指标ipc:每个周期执行的指令 issued_ipc:每个周期发出的指令 flop_hp_efficiency:实现的半精度浮点运算与理论峰值的比值 l2_utilization:L2 缓存利用率相对于理论峰值利用率的级别范围为0到10inst_fp_32:非谓词线程算术、比较等执行的单精度浮点指令数 stall_texture:由于纹理子系统被充分利用或有太多未完成的请求而发生的停顿百分比 local_load_throughput:本地内存加载吞吐量 local_store_throughput:本地内存存储吞吐量 shared_load_throughput:共享内存负载吞吐量 shared_store_throughput:共享内存存储吞吐量 shared_store_transactions:共享内存存储事务数 shared_load_transactions_per_request:每次共享内存加载时执行的平均共享内存加载事务数 inst_replay_overhead:每条指令执行的平均重放次数 warp_execution_efficiency:每个 warp 的平均活动线程数与 SM 支持的每个 warp 的最大线程数之比 inst_per_warp:每个 warp 执行的平均指令数nvprof metrics指标官方介绍4 同步4.1 __syncthreads:用来保证一个线程块中所有的线程操作都完成。4.2 cudaDeviceSynchronize():用来保证设备上所有的cuda操作完成。线程块之间没有同步指令只能串行启动内核。4.3warp内32个线程存在隐士同步warp内的线程必须要保证同步运行。5 使用nsight compute测试性能todo
)
