CUDA 矩阵转置性能优化实战：cuda-samples transpose 示例的合并访问、Bank Conflict 与 Partition Camping 深入剖析-北京尧图网络

CUDA 矩阵转置性能优化实战cuda-samples transpose 示例的合并访问、Bank Conflict 与 Partition Camping 深入剖析【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples矩阵转置是线性代数与高性能计算中的经典算子也是衡量 GPU 访存优化技巧的试金石。本文以 NVIDIA cuda-samples 仓库中 transpose 示例为核心系统讲解它从朴素实现到高度优化实现的完整演进路径全局内存合并访问coalescing、共享内存 bank conflict 消除、以及 partition camping 规避等关键技术。读完本文你将掌握这套以拷贝为性能上界、逐级优化的转置性能研究方法论并能直接复现、扩展该示例的基准测试流程。示例概览它在仓库中的位置与目标该示例位于 cpp/6_Performance/transpose/归属于仓库的6_Performance性能专题目录。目录下共包含三个文件文件作用transpose.cu全部设备端device与主机端host代码内含 8 个 kernel 与完整的计时、验证主程序CMakeLists.txt该示例的 CMake 构建脚本doc/MatrixTranspose.pdfNVIDIA 官方性能研究白皮书对该性能研究有详细描述从源码文件头注释可以确认其研究目标见 transpose.cu该文件同时包含矩阵转置的 device 与 host 代码。它执行多个转置 kernel通过合并访问、消除共享内存 bank conflict、以及规避 partition camping 逐步提升性能。其中若干 kernel 执行的是拷贝操作用于代表转置所能达到的最佳性能上界。示例默认对 1024×1024 的单精度浮点矩阵进行转置通过统一的计时框架在同一 GPU 上横向对比多种转置策略的有效带宽GB/s并与参考解做正确性校验。仓库 6_Performance 目录 README 也将其定位为展示不同性能实现以达到高性能的性能研究示例。为什么矩阵转置是性能难题访存模式分析一个常规的矩阵转置B[j][i] A[i][j]意味着如果按行主序row-major存储那么读取 A 时按行连续访问天然合并但写入 B 时却是按列访问——同一 warp 内 32 个线程写出的地址在全局内存中彼此相隔width个元素形成典型的未合并访问uncoalesced access。未合并访问会被硬件拆分成多次内存事务严重浪费显存带宽。这正是转置性能优化要解决的核心矛盾。CUDA 的解法是经典的tile分块共享内存shared memory转置先把数据块合并读入共享内存完成块内转置再合并写回全局内存。本示例在此基础上进一步把共享内存自身的 bank conflict 和全局内存分区的 partition camping 也逐一消除。八个 kernel 的性能演进路线主程序通过一个for (int k 0; k 8; k)循环依次调度 8 个 kernel见 transpose.cu。它们可划分为三组1. 拷贝基线性能上界转置必须读一次、写一次全部数据因此能跑多快的上界就是同规模纯拷贝的带宽。示例用两个拷贝 kernel 作为参考基准copysimple copy完全不使用共享内存直接按行连续读取、连续写入是全局内存带宽的原始表现。copySharedMemshared memory copy数据经共享内存中转后再写回用于量化引入共享内存本身的代价。转置 kernel 若能逼近这两个拷贝 kernel 的带宽即认为已达到接近硬件极限的水平。2. 三个完整转置 kernel逐步优化三个 kernel 使用相同的分块策略每块处理TILE_DIM × TILE_DIM的 tile用TILE_DIM × BLOCK_ROWS个线程每个线程处理TILE_DIM/BLOCK_ROWS个元素。代码中定义为见 transpose.cu#define TILE_DIM 32 #define BLOCK_ROWS 16即每块 32×32 tile、32×16 512 线程、每线程处理 2 个元素。transposeNaivenaive朴素的逐元素转置。读入与写出都未做任何合并优化是性能最差的下限参照。transposeCoalescedcoalesced第一步引入共享内存。线程按合并方式读入数据块存入__shared__ float tile[TILE_DIM][TILE_DIM]同步后用转置后的下标交换threadIdx.x与threadIdx.y合并写回全局内存见 transpose.cu。这一版解决了全局内存的未合并访问问题但写共享内存和读共享内存两个阶段都存在bank conflict尚未达到最优。transposeNoBankConflictsoptimized把共享内存数组从[TILE_DIM][TILE_DIM]改为[TILE_DIM][TILE_DIM 1]即每行多出一个元素的padding填充从而错开行首地址消除共享内存 bank conflict见 transpose.cu。这是转置优化中最经典也最有效的一招后文会详述其原理。transposeDiagonaldiagonal在前者基础上再叠加对角块重排diagonal reordering用于规避全局内存的 partition camping 现象见 transpose.cu。它把blockIdx.x重新解释为沿对角线方向的距离blockIdx.y对应不同对角线通过求模运算把对角坐标映射回笛卡尔坐标。该 kernel 既合并访问、又无 bank conflict、还能缓解 partition camping是完整转置的最终形态。3. 两个部分转置 kernel性能剖析工具transposeCoarseGrainedcoarse-grained与transposeFineGrainedfine-grained并不执行完整转置而是分别只完成合并读入共享内存或共享内存中转置后写回中的一半工作用于单独剖析转置两个阶段的性能特征。正因如此源码注释明确说明它们会无法通过参考解校验见 transpose.cu主程序在调度它们时也会跳过正确性比较见 transpose.cu。共享内存 Bank Conflict 与 Padding 技巧的底层原理GPU 共享内存被划分为 32 个 bank每个 bank 的带宽为每周期 4 字节。当一个 warp 内的多个线程同时访问同一 bank时访问会被硬件串行化这就是 bank conflict。在transposeCoalesced中线程写共享内存时threadIdx.y i相同的行内线程连续写tile[threadIdx.yi][threadIdx.x]32 个线程刚好覆盖一行 32 个float正好映射到 32 个 bank无冲突但在读转置后数据时线程访问tile[threadIdx.x][threadIdx.yi]此时threadIdx.x相同的线程每隔 16 个线程因为BLOCK_ROWS 16会访问到同一行即命中同一 bank产生 2 路 bank conflict每 16 个线程冲突一次共 2 组。反之写阶段在另一维度也存在类似冲突。transposeNoBankConflicts的解法是将数组声明为__shared__ float tile[TILE_DIM][TILE_DIM 1];每行多出的 1 个元素使相邻行的起始地址错开一个 bank于是本应落在同一 bank 的不同行元素被分散到相邻 bank冲突被彻底消除。代价仅仅是每块多出TILE_DIM个元素的共享内存性价比极高。这也是transposeDiagonal继续沿用[TILE_DIM][TILE_DIM 1]声明的原因。全局内存 Partition Camping 与对角重排在较新的 GPU 架构上全局内存的访问还会受partition camping影响同一时间片内若多个 block 的访问恰好集中命中 L2 分区的少数几个分区会造成局部拥塞拉低有效带宽。transposeDiagonal的思路是改变 block 的执行调度顺序让沿矩阵对角线分布的 block 尽量在同一时间运行从而把对全局内存分区的压力均匀打散。代码中的核心映射逻辑为见 transpose.cuif (width height) { blockIdx_y blockIdx.x; blockIdx_x (blockIdx.x blockIdx.y) % gridDim.x; } else { int bid blockIdx.x gridDim.x * blockIdx.y; blockIdx_y bid % gridDim.y; blockIdx_x ((bid / gridDim.y) blockIdx_y) % gridDim.x; }对于方形矩阵直接交换逻辑坐标并取模对于非方形矩阵则先把二维块 ID 展平为一维bid再做对角映射。映射之后kernel 其余部分与transposeNoBankConflicts完全一致仅需把blockIdx.x/blockIdx.y替换为blockIdx_x/blockIdx_y改动极小、收益显著。主程序执行流程计时、缩放与校验main()的执行逻辑见 transpose.cu可归纳为解析参数读取-device、-dimX、-dimY、-help等命令行选项。查询设备通过cudaGetDevice/cudaGetDeviceProperties获取 SM 版本与 SM 数量。计算缩放因子以每 SM 192 核为基线按 GPU 实际总核心数计算scale_factor max(192.0f / (cores_per_sm * sm_count), 1.0f)见 transpose.cu用于在小规模 GPU 上自动缩小测试矩阵保证 tile 数量可比。每 SM 核心数由 helper_cuda.h 中的_ConvertSMVer2Cores按架构查表得出。计算测试矩阵规模默认 1024×1024若指定-dimX/-dimY则以命令行参数为准非方形矩阵或尺寸不是TILE_DIM整数倍时会报错退出见 transpose.cu。分配与初始化数据host 侧用h_idata[i] (float)i填充device 侧cudaMalloc后cudaMemcpy上传同时用纯 CPU 的computeTransposeGold生成参考解见 transpose.cu。循环调度 8 个 kernel每个 kernel 先 warmup 一次避免启动开销计入计时再用 CUDA event 计时 100 次重复启动NUM_REPS 100最后回传结果与参考解比对。逐轮清零d_odata在进入下一轮 kernel 前将输出缓冲区显式清零防止上一个 kernel 的残留数据造成compareData假阳性见 transpose.cu。值得注意的细节是该示例使用 CUDAcooperative groups的cg::this_thread_block()与cg::sync(cta)替代传统的__syncthreads()做块内同步这也是较新 CUDA 编程模型的推荐写法见 transpose.cu。有效带宽的计算方法计时使用 CUDA event 对start/stop打点测得NUM_REPS次 kernel 启动的总耗时kernelTime毫秒单次耗时取kernelTime / NUM_REPS。有效带宽公式见 transpose.cufloat kernelBandwidth 2.0f * 1000.0f * mem_size / (1024 * 1024 * 1024) / (kernelTime / NUM_REPS);其中mem_size sizeof(float) * size_x * size_y见 transpose.cu。系数2.0f表示转置需读一次、写一次共两倍数据量1000.0f将毫秒换算为秒除以1024³将字节换算为 GB。最终输出格式为transpose simple copy , Throughput 870.1234 GB/s, Time 0.00921 ms, Size 1048576 fp32 elements, NumDevsUsed 1, Workgroup 512输出中Workgroup 512即TILE_DIM × BLOCK_ROWS 32 × 16的线程块规模。通过对比不同 kernel 的 GB/s 数值可以直观看到从 naive 到 diagonal 的逐级提速幅度。正确性验证机制每个完整转置 kernel 的结果都会通过cudaMemcpy回传 host与computeTransposeGold生成的参考解比对。比对函数是 helper_image.h 中的compareData采用绝对误差阈值epsilon 0.01逐元素校验copy 类 kernel 的参考解就是输入数据本身coarse/fine-grained 因非完整转置而跳过校验见 transpose.cu。任一 kernel 校验失败都会打印*** xxx kernel FAILED ***并将success置为 false最终决定进程以Test passed或Test failed退出。命令行参数运行transpose可执行文件时支持以下选项见 transpose.cu参数含义默认值-devicen指定使用的 GPU 设备编号n 0, 1, 2…由findCudaDevice自动选择-dimXrow_dim_size矩阵行维度x 方向大小由 GPU 规模自动推导上限约 1024-dimYcol_dim_size矩阵列维度y 方向大小与dimX相同-help打印上述帮助信息后退出—约束条件矩阵必须是方形size_x size_y且尺寸必须为TILE_DIM 32的整数倍否则程序直接报错退出见 transpose.cu。此外若2 * mem_size超过设备显存也会被拒绝见 transpose.cu。构建与运行该示例通过 CMake 构建。其 CMakeLists.txt 要求 CMake 3.20 及以上默认编译架构覆盖75 80 86 87 89 90 100 110 120对应从 Volta 到 Blackwell 的 SM 架构默认开启-lineinfo便于调试工具定位行号ENABLE_CUDA_DEBUG时可切换为-G以支持 cuda-gdb 调试并启用了CUDA_SEPARABLE_COMPILATION与 C17/CUDA 17 标准。在 Linux 上可按仓库主 README 的构建说明从任意子目录构建该示例mkdir build cd build cmake .. make -j$(nproc)编译产物位于build/bin/${TARGET_ARCH}/${TARGET_OS}/${BUILD_TYPE}目录Linux x86_64 Release 对应build/bin/x64/linux/release见 README.md。随后运行./transpose # 默认 1024x1024 矩阵 ./transpose -dimX512 -dimY512 # 指定矩阵规模 ./transpose -device0 -dimX2048 -dimY2048Windows 下则使用 Visual Studio 的x64 Native Tools Command Prompt for VS以cmake .. -G Visual Studio 16 2019 -A x64配置并生成解决方案后编译运行见 README.md。支持环境一览根据示例 README.md支持的 SM 架构SM 5.0 / 5.2 / 5.3 / 6.0 / 6.1 / 7.0 / 7.2 / 7.5 / 8.0 / 8.6 / 8.7 / 8.9 / 9.0覆盖 Maxwell 至 Hopper/Blackwell 之前的各代架构支持的操作系统Linux、Windows支持的 CPU 架构x86_64、armv7l前置条件安装对应平台的 CUDA Toolkit示例 README 中的官方指引链接涉及的 CUDA Runtime APIcudaMemcpy、cudaMalloc、cudaFree、cudaGetLastError、cudaEventSynchronize、cudaEventRecord、cudaGetDevice、cudaEventDestroy、cudaEventElapsedTime、cudaGetDeviceProperties、cudaEventCreate全部与 transpose.cu 中的调用一一对应。进一步阅读该示例配套的官方白皮书 doc/MatrixTranspose.pdf 对上述性能研究合并访问、bank conflict、partition camping 及对角重排有更详尽的数值分析与硬件原理阐述是深入理解本示例的最佳延伸资料。若想横向了解仓库中其他性能优化主题可继续阅读 cpp/6_Performance/README.md 中列出的 alignedTypes、UnifiedMemoryPerf、cudaGraphsPerfScaling 等姊妹示例。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考

CUDA 矩阵转置性能优化实战：cuda-samples transpose 示例的合并访问、Bank Conflict 与 Partition Camping 深入剖析

相关推荐

Ubuntu切换默认内核版本：Grub配置全流程实操指南

Monty 沙箱文件系统挂载深度解析：monty-fs 的 MountTable 与结构化沙箱边界

Python文件操作三剑客：os、pathlib与shutil实战指南

更多资讯

360°视频编码测试条件与参考配置全解析

Transformer电价预测实战：从注意力机制到超长序列建模

VSCode搭建Arduino UNO开发环境：代码补全与编译烧录全攻略

攻击流量样本分析实战：从PCAP捕获到威胁狩猎与检测规则落地

ESP32音频开发实战：I2S+WAV+MicroPython零基础入门

高通Camx相机调试：UMD/KMD日志开关、图像Dump与离线合成实战

今日更新

基于MATLAB的CRI显色指数计算：从SPD光谱到Ra的完整流程

JSP+Servlet+MySQL博客系统源码部署与优化全攻略

蜂窝调度算法 MATLAB 实现：比例公平、轮询与最大 C/I 对比

本周精选

AI驱动的自动化专利分析平台技术解析

LeetCode-Go 题解：1252. Cells with Odd Values in a Matrix（奇数值单元格计数）双解法剖析

在线答疑系统Java毕设实战：Spring Boot前后端分离与状态流转

本月热门

基于YOLO的焊缝缺陷检测：从数据标注到工业部署全流程实践

分布式服务的决策记录

创业产品的首版范围

读完文章，想聊聊您的网站？