CUDA 并行前缀和实战基于 Shuffle 指令__shfl_up_sync的 shfl_scan 示例深度解析【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples导读本文深入剖析 CUDA Samples 仓库中的 shfl_scan 示例位于 cpp/2_Concepts_and_Techniques/shfl_scan它展示了如何仅用 shuffle 内建指令__shfl_up_sync在不依赖任何并行扫描库的前提下完成 warp 内、block 内乃至跨 block 的并行前缀和Prefix Sum / Scan并进一步将其升级为两遍扫描的积分图像Integral Image即求和面积表算法。读完本文你将掌握 shuffle 扫描的金字塔式分解思想、__shfl_xor数据重组技巧以及这套方案在图像处理中的实际落地方式。示例概览一个只用 Shuffle 做 Scan的教学样本根据 README.md 的描述shfl_scan 的核心目的是演示如何使用 shuffle 内建指令__shfl_up_sync在线程块thread block内执行一次扫描scan操作。它属于Data-Parallel Algorithms数据并行算法与Performance Strategies性能策略两大关键概念类别。源码头注释shfl_scan.cu进一步明确示例分两个递进层次简单示例用 shuffle 执行一个纯前缀和scan操作进阶示例用 shuffle 计算积分图像integral image其中同时用到 shuffle 扫描shfl_up与 shuffle 异或shfl_xor两类指令。平台与硬件支持README 给出的官方支持范围如下维度支持情况操作系统Linux、WindowsCPU 架构x86_64、armv7l、aarch64依赖C11 CUDA详见 主 README 的 C11 CUDA 小节值得注意的两点补充README 中Supported SM Architectures一节为空但源码实际上限定了硬件要求。__shfl_*系列指令需要SM 3.0 及以上shfl_scan.cu 在main()中检测到设备计算能力低于 3.0 时会直接以EXIT_WAIVED退出并跳过测试仓库当前的 CMakeLists.txt 将CMAKE_CUDA_ARCHITECTURES设为75 80 86 87 89 90 100 110 120覆盖 Turing 至 Blackwell 系列并以cxx_std_17/cuda_std_17编译说明工程实际要求的语言标准高于 README 中标注的 C11。涉及的 CUDA Runtime APIREADME 列出了本示例用到的全部 CUDA Runtime APIcudaMemcpy、cudaFree、cudaMallocHost、cudaEventSynchronize、cudaEventRecord、cudaFreeHost、cudaGetDevice、cudaMemset、cudaMalloc、cudaEventElapsedTime、cudaGetDeviceProperties、cudaEventCreate。从源码可以印证它们的用途cudaEventCreate/Record/Synchronize/ElapsedTime用于对内核计时shfl_scan.cucudaMallocHost分配页锁定主机内存以加速拷贝shfl_scan.cucudaMemset负责清空部分和缓冲shfl_scan.cucudaGetDeviceProperties用于检查 SM 版本。构建与运行与仓库中其他示例一致shfl_scan 通过 CMake 构建并依赖仓库根目录的 Common 目录中的helper_cuda.h、helper_functions.hfindCudaDevice、checkCudaErrors、sdkCreateTimer等工具函数均来自这里。编译要求CMake 3.20 或更高见 CMakeLists.txt已安装匹配平台的 CUDA ToolkitREADME 前置条件该示例通过 cpp/2_Concepts_and_Techniques/CMakeLists.txt 的add_subdirectory(shfl_scan)挂入整个 2_Concepts_and_Techniques 的构建树。典型构建流程仓库根目录mkdir build cd build cmake ../cpp/2_Concepts_and_Techniques make shfl_scan ./shfl_scan运行时会依次执行两个子测试shuffle_simple_test简单前缀和与shuffle_integral_image_test1920×1080 合成图像的积分图像任一失败都会使程序以非零码退出。如果 GPU 计算能力低于 SM 3.0程序将输出提示并以EXIT_WAIVED结束。核心原理一warp 内扫描 —— 一次__shfl_up_sync的 log2(n) 步折叠shfl_scan 的灵魂在shfl_scan_test内核的前半段shfl_scan.cu__global__ void shfl_scan_test(int *data, int width, int *partial_sums NULL) { extern __shared__ int sums[]; int id ((blockIdx.x * blockDim.x) threadIdx.x); int lane_id id % warpSize; int warp_id threadIdx.x / warpSize; int value data[id]; #pragma unroll for (int i 1; i width; i * 2) { unsigned int mask 0xffffffff; int n __shfl_up_sync(mask, value, i, width); if (lane_id i) value n; } ... }这段代码实现了标准的Hillis-Steele 风格 warp 内扫描其要点是__shfl_up_sync(mask, value, i, width)让 lane 从lane_id - i处取得value其中mask是参与同步的线程掩码这里为全 1表示整个 warp 都参与width限定 warp 宽度此处为 32当偏移i超出 warp 边界时shuffle 返回的值无效因此用if (lane_id i)做守卫只有 lane 编号不小于i的线程才累加循环以i 1, 2, 4, ..., 16翻倍推进因此只需log2(warpSize) 5 步即可让每个 lane 都持有从 lane 0 到自身的前缀和。这正是 shuffle 相比共享内存扫描的最大优势——数据在寄存器间直接交换无共享内存读写、无__syncthreads()延迟更低、带宽压力更小。核心原理二从 warp 扩展到 block —— 共享内存金字塔前缀和不能止步于 warp因为每个 lane 的最终前缀和需要包含前面所有 warp的贡献。shfl_scan_test采用金字塔式分解shfl_scan.cu 的注释有明确描述warp 内扫描如上一节所示每个线程得到 warp 内的前缀和warp 和落共享内存每个 warp 的最后一条 lanethreadIdx.x % warpSize warpSize - 1把该 warp 的总和写入sums[warp_id]shfl_scan.cuwarp 和的扫描由 warp 0 对warp 和数组再做一次 shuffle 扫描得到每个 warp 的累计前缀和。这里用mask (1 (blockDim.x / warpSize)) - 1只激活前blockDim.x / warpSize条 lane避免越界参与shfl_scan.cuuniform add均匀加回每个 warp 读取前一个 warp 的累计和warp_id 0时取sums[warp_id - 1]加到本 warp 所有线程的value上shfl_scan.cu。整个过程中共享内存只存放每 warp 一个 int的小数组主要计算全部发生在寄存器与 warp 内最大限度地发挥了 shuffle 指令的威力。两次__syncthreads()仅用于保证共享内存写入与读取之间的顺序shfl_scan.cu 与 shfl_scan.cu。核心原理三跨 block 扫描 —— partial_sums 与 uniform_add单个 block 内部的金字塔扫描仍然只覆盖blockDim.x个元素。对于 65536 个元素的测试数据shuffle_simple_test采用两级内核 一次修正内核的流水shfl_scan.cushfl_scan_testgridSize, blockSize, shmem_sz(d_data, 32, d_partial_sums); shfl_scan_testp_gridSize, p_blockSize, shmem_sz(d_partial_sums, 32); uniform_addgridSize - 1, blockSize(d_data blockSize, d_partial_sums, n_elements);第一个内核按gridSize × blockSize256×256见 shfl_scan.cu对全部数据做 block 内扫描并把每个 block 的最终和写入d_partial_sums[blockIdx.x]shfl_scan.cu第二个内核仅对partial_sums数组256 个 block 和再次调用同一个shfl_scan_test得到 block 级的前缀和uniform_add内核把本 block 之前所有 block 的累计和读入共享内存广播变量buf加回给每个数据元素shfl_scan.cu。于是warp → block → grid三个层次全部由同一个 shuffle 扫描内核递归复用完成体现了同一算法在不同粒度上的可组合性。最终输出partial_sums[n_partialSums - 1]即为整个数组的总和Test Sum并通过CPUverify与 CPU 上的朴素前缀和做逐元素差分校验shfl_scan.cu同时用sdkCreateTimer记录 CPU 参考实现耗时以作对比。计时方面GPU 侧用cudaEventRecord包裹三次内核调用cudaEventElapsedTime得到毫秒耗时并按n_elements / (et / 1000.0f) / 1000000.0f换算成 MegaElements/s 吞吐量打印shfl_scan.cu。具体数值依赖硬件仓库未承诺任何基准结果读者可在自有 GPU 上实测。进阶实战用 Shuffle 生成积分图像Integral Image积分图像又称求和面积表是图像处理中的经典加速结构位置(x, y)存储从左上角到该点的矩形像素和之后任意矩形区域的求和只需三次加减。本示例在 shfl_integral_image.cuh 中给出了一个两遍扫描实现先水平行扫描再垂直列扫描测试数据为 1920×1080 的全 1 合成灰度图shfl_scan.cu。水平扫描每线程 16 个值 warp 内扫描shfl_intimage_rows内核让每个线程一次处理 16 个像素打包成 4 个uint4一个 block 处理一行。核心函数get_prefix_sumshfl_integral_image.cuh的流程是先把每个uint4的 4 个 8 位通道解包为 4 个独立的 32 位前缀和uint_to_uchar4 手工累加shfl_integral_image.cuh——之所以转成unsigned int是因为逐像素累加必然溢出 8 位对每线程内部的 16 个值做串行扫描得到 16 个局部前缀和shfl_integral_image.cuh跨线程借助 cooperative_groups 的cg::tiled_partition32(cta)得到 warp tile用tile.shfl_up(sum, i)把每线程第 16 个值即该线程 16 个像素的和在 warp 内做 shuffle 扫描并将扫描结果均匀加回全部 16 个前缀和shfl_integral_image.cuhwarp 之间仍走共享内存金字塔warp 最后一条 lane 写入sums[warp_id]warp 0 再做一次 warp 级 shuffle 扫描其余 warp 统一加上前序 warp 的累计和shfl_integral_image.cuh。__shfl_xor寄存器内的数据重组水平扫描后数据在寄存器中的分布是交错的每个线程持有 16 个连续像素的前缀和但写回全局内存时希望得到连续地址。shfl_intimage_rows用__shfl_xor在寄存器层面完成一次转置式重排shfl_integral_image.cuh其原理在代码注释中有非常清晰的图示GMEM[16] - x0 x1 x2 x3 y0 y1 y2 y3 z0 z1 z2 z3 w0 w1 w2 w3 但寄存器 r0..r3 在 4 个线程(0..3)中的布局为: threadId 0 1 2 3 r0 x0 y0 z0 w0 r1 x1 y1 z1 w1 r2 x2 y2 z2 w2 r3 x3 y3 z3 w3 执行 __shfl_xor 后: threadId 00 01 10 11 x0 y0 z0 w0 xor(01)-y1 x1 w1 z1 xor(10)-z2 w2 x2 y2 xor(11)-w3 z3 y3 x3__shfl_xor(mask, v, laneMask)让 lane 与lane_id ^ laneMask交换数据。经过xor(1)、xor(2)、xor(3)的两次组合每根列上的 x、y、z、w 分量被归拢到同一线程于是可以用uint4粒度做更大段的连续写入shfl_integral_image.cuh。源码注释还特意说明借助 lambda 形式的idxToElem访问辅助函数packed_result可以保持在寄存器中而不溢栈shfl_integral_image.cuh。垂直扫描把列求和也塞进 warpshfl_vertical_shfl内核完成列方向的累计shfl_integral_image.cuh。它的技巧是把当前 block 覆盖的 32×8 数据先写入共享内存sums[32][9]然后重排索引j threadIdx.x % 8、k threadIdx.x / 8 threadIdx.y * 4使原本位于不同列的 8 个元素被映射进同一条 warp从而再次用__shfl_up_sync(mask, partial_sum, i, 32)在 warp 内完成列的前缀和省去了共享内存逐级累加所需的多次__syncthreads()shfl_integral_image.cuh。外层循环以stepSum维护本 block 之前所有行的累计和逐块向下传播最终得到完整的积分图像。正确性验证行扫描后调用verifyDataRowSums因为输入是全 1 数据第 i 列的正确行前缀和就是i 1据此逐像素求差shfl_scan.cu两遍扫描全部完成后检查右下角元素finalSum h_image[w * h - 1]期望值恰为w * h 1920 × 1080相等即视为积分图像生成成功shfl_scan.cu。辅助设施util.h 与错误检查宏目录中的 util.h 提供了两个常用的错误检查宏CHECK_LAUNCH_ERROR()先查同步错误cudaGetLastError再查异步错误cudaDeviceSynchronize覆盖内核启动前与运行期两个阶段的失败CUDA_CHECK(call)包装任意 CUDA Runtime 调用失败时打印文件与行号后退出。虽然 shfl_scan.cu 主要使用 Common 目录中的checkCudaErrors这两个宏仍可视为本示例自带的、可复用的健壮性工具。小结从 README.md 的一句话主题出发shfl_scan 展示了三条可迁移到生产代码的实践warp 内扫描是性能基石__shfl_up_sync以 log2(n) 步完成扫描全程寄存器通信无共享内存与同步开销金字塔分解让算法可组合warp 和 → 共享内存 → 再次 shuffle 扫描 → uniform add同一套内核既能处理 block 内数据也能递归处理 partial sums 实现跨 block 扫描__shfl_xor不止用于扫描它还是寄存器级数据转置的利器能让内存访问更连续、写回更高效——这正是积分图像示例中水平扫描内核的精华所在。对于需要深入理解并行扫描原理、或想在图像/信号处理中实现高性能前缀和的开发者shfl_scan.cu 与 shfl_integral_image.cuh 提供了完整的、可直接对照阅读与改造的参考实现。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
CUDA 并行前缀和实战：基于 Shuffle 指令（__shfl_up_sync）的 shfl_scan 示例深度解析
CUDA 并行前缀和实战基于 Shuffle 指令__shfl_up_sync的 shfl_scan 示例深度解析【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples导读本文深入剖析 CUDA Samples 仓库中的 shfl_scan 示例位于 cpp/2_Concepts_and_Techniques/shfl_scan它展示了如何仅用 shuffle 内建指令__shfl_up_sync在不依赖任何并行扫描库的前提下完成 warp 内、block 内乃至跨 block 的并行前缀和Prefix Sum / Scan并进一步将其升级为两遍扫描的积分图像Integral Image即求和面积表算法。读完本文你将掌握 shuffle 扫描的金字塔式分解思想、__shfl_xor数据重组技巧以及这套方案在图像处理中的实际落地方式。示例概览一个只用 Shuffle 做 Scan的教学样本根据 README.md 的描述shfl_scan 的核心目的是演示如何使用 shuffle 内建指令__shfl_up_sync在线程块thread block内执行一次扫描scan操作。它属于Data-Parallel Algorithms数据并行算法与Performance Strategies性能策略两大关键概念类别。源码头注释shfl_scan.cu进一步明确示例分两个递进层次简单示例用 shuffle 执行一个纯前缀和scan操作进阶示例用 shuffle 计算积分图像integral image其中同时用到 shuffle 扫描shfl_up与 shuffle 异或shfl_xor两类指令。平台与硬件支持README 给出的官方支持范围如下维度支持情况操作系统Linux、WindowsCPU 架构x86_64、armv7l、aarch64依赖C11 CUDA详见 主 README 的 C11 CUDA 小节值得注意的两点补充README 中Supported SM Architectures一节为空但源码实际上限定了硬件要求。__shfl_*系列指令需要SM 3.0 及以上shfl_scan.cu 在main()中检测到设备计算能力低于 3.0 时会直接以EXIT_WAIVED退出并跳过测试仓库当前的 CMakeLists.txt 将CMAKE_CUDA_ARCHITECTURES设为75 80 86 87 89 90 100 110 120覆盖 Turing 至 Blackwell 系列并以cxx_std_17/cuda_std_17编译说明工程实际要求的语言标准高于 README 中标注的 C11。涉及的 CUDA Runtime APIREADME 列出了本示例用到的全部 CUDA Runtime APIcudaMemcpy、cudaFree、cudaMallocHost、cudaEventSynchronize、cudaEventRecord、cudaFreeHost、cudaGetDevice、cudaMemset、cudaMalloc、cudaEventElapsedTime、cudaGetDeviceProperties、cudaEventCreate。从源码可以印证它们的用途cudaEventCreate/Record/Synchronize/ElapsedTime用于对内核计时shfl_scan.cucudaMallocHost分配页锁定主机内存以加速拷贝shfl_scan.cucudaMemset负责清空部分和缓冲shfl_scan.cucudaGetDeviceProperties用于检查 SM 版本。构建与运行与仓库中其他示例一致shfl_scan 通过 CMake 构建并依赖仓库根目录的 Common 目录中的helper_cuda.h、helper_functions.hfindCudaDevice、checkCudaErrors、sdkCreateTimer等工具函数均来自这里。编译要求CMake 3.20 或更高见 CMakeLists.txt已安装匹配平台的 CUDA ToolkitREADME 前置条件该示例通过 cpp/2_Concepts_and_Techniques/CMakeLists.txt 的add_subdirectory(shfl_scan)挂入整个 2_Concepts_and_Techniques 的构建树。典型构建流程仓库根目录mkdir build cd build cmake ../cpp/2_Concepts_and_Techniques make shfl_scan ./shfl_scan运行时会依次执行两个子测试shuffle_simple_test简单前缀和与shuffle_integral_image_test1920×1080 合成图像的积分图像任一失败都会使程序以非零码退出。如果 GPU 计算能力低于 SM 3.0程序将输出提示并以EXIT_WAIVED结束。核心原理一warp 内扫描 —— 一次__shfl_up_sync的 log2(n) 步折叠shfl_scan 的灵魂在shfl_scan_test内核的前半段shfl_scan.cu__global__ void shfl_scan_test(int *data, int width, int *partial_sums NULL) { extern __shared__ int sums[]; int id ((blockIdx.x * blockDim.x) threadIdx.x); int lane_id id % warpSize; int warp_id threadIdx.x / warpSize; int value data[id]; #pragma unroll for (int i 1; i width; i * 2) { unsigned int mask 0xffffffff; int n __shfl_up_sync(mask, value, i, width); if (lane_id i) value n; } ... }这段代码实现了标准的Hillis-Steele 风格 warp 内扫描其要点是__shfl_up_sync(mask, value, i, width)让 lane 从lane_id - i处取得value其中mask是参与同步的线程掩码这里为全 1表示整个 warp 都参与width限定 warp 宽度此处为 32当偏移i超出 warp 边界时shuffle 返回的值无效因此用if (lane_id i)做守卫只有 lane 编号不小于i的线程才累加循环以i 1, 2, 4, ..., 16翻倍推进因此只需log2(warpSize) 5 步即可让每个 lane 都持有从 lane 0 到自身的前缀和。这正是 shuffle 相比共享内存扫描的最大优势——数据在寄存器间直接交换无共享内存读写、无__syncthreads()延迟更低、带宽压力更小。核心原理二从 warp 扩展到 block —— 共享内存金字塔前缀和不能止步于 warp因为每个 lane 的最终前缀和需要包含前面所有 warp的贡献。shfl_scan_test采用金字塔式分解shfl_scan.cu 的注释有明确描述warp 内扫描如上一节所示每个线程得到 warp 内的前缀和warp 和落共享内存每个 warp 的最后一条 lanethreadIdx.x % warpSize warpSize - 1把该 warp 的总和写入sums[warp_id]shfl_scan.cuwarp 和的扫描由 warp 0 对warp 和数组再做一次 shuffle 扫描得到每个 warp 的累计前缀和。这里用mask (1 (blockDim.x / warpSize)) - 1只激活前blockDim.x / warpSize条 lane避免越界参与shfl_scan.cuuniform add均匀加回每个 warp 读取前一个 warp 的累计和warp_id 0时取sums[warp_id - 1]加到本 warp 所有线程的value上shfl_scan.cu。整个过程中共享内存只存放每 warp 一个 int的小数组主要计算全部发生在寄存器与 warp 内最大限度地发挥了 shuffle 指令的威力。两次__syncthreads()仅用于保证共享内存写入与读取之间的顺序shfl_scan.cu 与 shfl_scan.cu。核心原理三跨 block 扫描 —— partial_sums 与 uniform_add单个 block 内部的金字塔扫描仍然只覆盖blockDim.x个元素。对于 65536 个元素的测试数据shuffle_simple_test采用两级内核 一次修正内核的流水shfl_scan.cushfl_scan_testgridSize, blockSize, shmem_sz(d_data, 32, d_partial_sums); shfl_scan_testp_gridSize, p_blockSize, shmem_sz(d_partial_sums, 32); uniform_addgridSize - 1, blockSize(d_data blockSize, d_partial_sums, n_elements);第一个内核按gridSize × blockSize256×256见 shfl_scan.cu对全部数据做 block 内扫描并把每个 block 的最终和写入d_partial_sums[blockIdx.x]shfl_scan.cu第二个内核仅对partial_sums数组256 个 block 和再次调用同一个shfl_scan_test得到 block 级的前缀和uniform_add内核把本 block 之前所有 block 的累计和读入共享内存广播变量buf加回给每个数据元素shfl_scan.cu。于是warp → block → grid三个层次全部由同一个 shuffle 扫描内核递归复用完成体现了同一算法在不同粒度上的可组合性。最终输出partial_sums[n_partialSums - 1]即为整个数组的总和Test Sum并通过CPUverify与 CPU 上的朴素前缀和做逐元素差分校验shfl_scan.cu同时用sdkCreateTimer记录 CPU 参考实现耗时以作对比。计时方面GPU 侧用cudaEventRecord包裹三次内核调用cudaEventElapsedTime得到毫秒耗时并按n_elements / (et / 1000.0f) / 1000000.0f换算成 MegaElements/s 吞吐量打印shfl_scan.cu。具体数值依赖硬件仓库未承诺任何基准结果读者可在自有 GPU 上实测。进阶实战用 Shuffle 生成积分图像Integral Image积分图像又称求和面积表是图像处理中的经典加速结构位置(x, y)存储从左上角到该点的矩形像素和之后任意矩形区域的求和只需三次加减。本示例在 shfl_integral_image.cuh 中给出了一个两遍扫描实现先水平行扫描再垂直列扫描测试数据为 1920×1080 的全 1 合成灰度图shfl_scan.cu。水平扫描每线程 16 个值 warp 内扫描shfl_intimage_rows内核让每个线程一次处理 16 个像素打包成 4 个uint4一个 block 处理一行。核心函数get_prefix_sumshfl_integral_image.cuh的流程是先把每个uint4的 4 个 8 位通道解包为 4 个独立的 32 位前缀和uint_to_uchar4 手工累加shfl_integral_image.cuh——之所以转成unsigned int是因为逐像素累加必然溢出 8 位对每线程内部的 16 个值做串行扫描得到 16 个局部前缀和shfl_integral_image.cuh跨线程借助 cooperative_groups 的cg::tiled_partition32(cta)得到 warp tile用tile.shfl_up(sum, i)把每线程第 16 个值即该线程 16 个像素的和在 warp 内做 shuffle 扫描并将扫描结果均匀加回全部 16 个前缀和shfl_integral_image.cuhwarp 之间仍走共享内存金字塔warp 最后一条 lane 写入sums[warp_id]warp 0 再做一次 warp 级 shuffle 扫描其余 warp 统一加上前序 warp 的累计和shfl_integral_image.cuh。__shfl_xor寄存器内的数据重组水平扫描后数据在寄存器中的分布是交错的每个线程持有 16 个连续像素的前缀和但写回全局内存时希望得到连续地址。shfl_intimage_rows用__shfl_xor在寄存器层面完成一次转置式重排shfl_integral_image.cuh其原理在代码注释中有非常清晰的图示GMEM[16] - x0 x1 x2 x3 y0 y1 y2 y3 z0 z1 z2 z3 w0 w1 w2 w3 但寄存器 r0..r3 在 4 个线程(0..3)中的布局为: threadId 0 1 2 3 r0 x0 y0 z0 w0 r1 x1 y1 z1 w1 r2 x2 y2 z2 w2 r3 x3 y3 z3 w3 执行 __shfl_xor 后: threadId 00 01 10 11 x0 y0 z0 w0 xor(01)-y1 x1 w1 z1 xor(10)-z2 w2 x2 y2 xor(11)-w3 z3 y3 x3__shfl_xor(mask, v, laneMask)让 lane 与lane_id ^ laneMask交换数据。经过xor(1)、xor(2)、xor(3)的两次组合每根列上的 x、y、z、w 分量被归拢到同一线程于是可以用uint4粒度做更大段的连续写入shfl_integral_image.cuh。源码注释还特意说明借助 lambda 形式的idxToElem访问辅助函数packed_result可以保持在寄存器中而不溢栈shfl_integral_image.cuh。垂直扫描把列求和也塞进 warpshfl_vertical_shfl内核完成列方向的累计shfl_integral_image.cuh。它的技巧是把当前 block 覆盖的 32×8 数据先写入共享内存sums[32][9]然后重排索引j threadIdx.x % 8、k threadIdx.x / 8 threadIdx.y * 4使原本位于不同列的 8 个元素被映射进同一条 warp从而再次用__shfl_up_sync(mask, partial_sum, i, 32)在 warp 内完成列的前缀和省去了共享内存逐级累加所需的多次__syncthreads()shfl_integral_image.cuh。外层循环以stepSum维护本 block 之前所有行的累计和逐块向下传播最终得到完整的积分图像。正确性验证行扫描后调用verifyDataRowSums因为输入是全 1 数据第 i 列的正确行前缀和就是i 1据此逐像素求差shfl_scan.cu两遍扫描全部完成后检查右下角元素finalSum h_image[w * h - 1]期望值恰为w * h 1920 × 1080相等即视为积分图像生成成功shfl_scan.cu。辅助设施util.h 与错误检查宏目录中的 util.h 提供了两个常用的错误检查宏CHECK_LAUNCH_ERROR()先查同步错误cudaGetLastError再查异步错误cudaDeviceSynchronize覆盖内核启动前与运行期两个阶段的失败CUDA_CHECK(call)包装任意 CUDA Runtime 调用失败时打印文件与行号后退出。虽然 shfl_scan.cu 主要使用 Common 目录中的checkCudaErrors这两个宏仍可视为本示例自带的、可复用的健壮性工具。小结从 README.md 的一句话主题出发shfl_scan 展示了三条可迁移到生产代码的实践warp 内扫描是性能基石__shfl_up_sync以 log2(n) 步完成扫描全程寄存器通信无共享内存与同步开销金字塔分解让算法可组合warp 和 → 共享内存 → 再次 shuffle 扫描 → uniform add同一套内核既能处理 block 内数据也能递归处理 partial sums 实现跨 block 扫描__shfl_xor不止用于扫描它还是寄存器级数据转置的利器能让内存访问更连续、写回更高效——这正是积分图像示例中水平扫描内核的精华所在。对于需要深入理解并行扫描原理、或想在图像/信号处理中实现高性能前缀和的开发者shfl_scan.cu 与 shfl_integral_image.cuh 提供了完整的、可直接对照阅读与改造的参考实现。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
