尧图网站建设 
cuda-samples 实战基于 CCCL 3.3 的 cub::DeviceTransform N 输入 / M 输出融合变换示例解析【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samplescubDeviceTransform 是 cuda-samples 仓库中演示 CCCL 3.3 扩展后的cub::DeviceTransformN-to-M 形态的官方示例一次设备端调用读取 N 条输入序列通过用户自定义算子返回cuda::std::tuple同时写出 M 条输出序列。本文以该示例为骨架结合仓库源码逐行剖析两种典型用法N3→1 的融合运算与 N2→2 的和差一次成型并讲解其 CMake 构建方式、CCCL 依赖获取机制与运行验证逻辑帮助读者直接上手并迁移到自己的 CUDA 工程中。示例概述与核心知识点示例位于仓库的cpp/4_CUDA_Libraries/cubDeviceTransform/目录由三个文件组成README.md、CMakeLists.txt 与核心实现 cubDeviceTransform.cu。它演示了cub::DeviceTransform在 CCCL 3.3 中引入的N-input / M-output扩展形态覆盖以下关键概念CCCL 3.3NVIDIA 的 CUDA C Core Libraries包含 CUB、libcu、Thrust 的统一版本管理CUB Device Algorithms设备级算法库cub::DeviceTransform::Transform负责以设备级并行方式对序列逐元素应用变换Fused Elementwise Transforms融合逐元素变换将多个输入的读取与多个输出的写入合并为一次设备端调用减少内核启动开销与中间存储Counting Iterators计数迭代器cuda::counting_iterator可在不显式分配数据的情况下提供递增整数序列。示例演示了两个具体场景场景输入数 N输出数 M变换语义场景一31result[i] (a[i] b[i]) * c[i]其中c为从 100 开始的计数迭代器场景二22(sum[i], diff[i]) (a[i] b[i], a[i] - b[i])一趟完成和与差两个场景的结果都会在主机端用参考实现逐一比对并打印OK/FAIL程序返回码随之决定EXIT_SUCCESS/EXIT_FAILURE因此该示例天然自带可回归验证的测试属性。场景一N3 输入 → 1 输出融合乘法与计数迭代器第一个用例run_n_to_one_transform()位于 cubDeviceTransform.cu核心逻辑如下thrust::device_vectorint a {0, -2, 5, 3}; thrust::device_vectorfloat b {5.2f, 3.1f, -1.1f, 3.0f}; auto counting cuda::counting_iteratorint{100}; thrust::device_vectorint result(a.size()); auto op [] __host__ __device__(int x, float y, int z) - int { return static_castint((x y) * z); }; checkCudaErrors(cub::DeviceTransform::Transform( cuda::std::tuple{a.begin(), b.begin(), counting}, result.begin(), a.size(), op)); checkCudaErrors(cudaDeviceSynchronize());这段代码揭示了 N-to-M 形态的三个关键设计点输入侧使用cuda::std::tuple打包迭代器N3 个输入分别来自thrust::device_vectorint a、thrust::device_vectorfloat b以及一个cuda::counting_iteratorint{100}。三个输入的数据类型可以不同int与float混合由算子op的参数列表统一接收计数迭代器代替显式数据cuda::counting_iteratorint{100}会在迭代时依次产生100, 101, 102, ...完全不需要在设备上分配一块保存这些整数的内存。它来自 libcu 的头文件 cuda/iterator是只读合成序列的经典用法算子必须是__host__ __device__可调用对象示例使用--extended-lambda编译选项支持在设备端执行扩展 lambda见 CMakeLists.txtop同时被用于设备端变换与主机端参考计算的思路保持一致。调用完成后代码通过thrust::host_vector把设备数据拷回主机并在主机端按同一公式计算期望值expected[i] static_castint((ha[i] hb[i]) * static_castint(100 i));对照示例数据可推算输出a {0, -2, 5, 3}、b {5.2f, 3.1f, -1.1f, 3.0f}、计数序列{100, 101, 102, 103}得到result {520, 111, 397, 618}注意float乘法结果被static_castint截断取整。程序会把实际值与期望值一并打印并标注OK。场景二N2 输入 → M2 输出一趟同时产出和与差第二个用例run_n_to_m_transform()位于 cubDeviceTransform.cu是 N-to-M 形态最具代表性的一趟多出写法thrust::device_vectorint a {1, 5, 10, 7, 3}; thrust::device_vectorint b {4, 2, 8, 1, 9}; thrust::device_vectorint sum(a.size()); thrust::device_vectorint diff(a.size()); auto op [] __host__ __device__(int x, int y) - cuda::std::tupleint, int { return {x y, x - y}; }; checkCudaErrors(cub::DeviceTransform::Transform(cuda::std::tuple{a.begin(), b.begin()}, cuda::std::tuple{sum.begin(), diff.begin()}, a.size(), op)); checkCudaErrors(cudaDeviceSynchronize());与场景一相比这里出现了两个对称的 tuple输入 tuplecuda::std::tuple{a.begin(), b.begin()}对应 N2 条输入序列输出 tuplecuda::std::tuple{sum.begin(), diff.begin()}对应 M2 条输出序列算子返回值cuda::std::tupleint, inttuple 中的第 k 个元素会按位写入第 k 条输出序列。这就是融合变换的收益所在原本需要启动两次内核一次求和、一次求差并遍历两次输入的工作现在只需一次设备端调用、一次内存遍历即可同时得到sum与diff对于派生多个统计量、多通道处理等场景可以显著降低启动开销与访存次数。根据示例数据可得sum {5, 7, 18, 8, 12}、diff {-3, 3, 2, 6, -6}主机端逐元素比对后打印结果与OK。从源码结构看cubDeviceTransform.cu 的注释N 与 M 并不要求相等M 可以是 1场景一也可以是任意多个场景二为 2输入与输出序列数完全由调用时传入的 tuple 长度决定。设备端调用链与错误处理main()函数cubDeviceTransform.cu展示了本示例涉及的 CUDA Runtime APIint devID findCudaDevice(argc, (const char **)argv); cudaDeviceProp props; checkCudaErrors(cudaGetDeviceProperties(props, devID)); printf(Device: %s (Compute Capability %d.%d)\n\n, props.name, props.major, props.minor);findCudaDevice来自仓库公共头文件 Common/helper_cuda.h若命令行传入--deviceN则使用指定设备否则自动挑选 GFlops 最高的设备并执行cudaSetDevicecudaGetDeviceProperties用于读取设备名称与计算能力major.minor并打印到终端checkCudaErrors定义于 Common/helper_cuda.h包装每次 CUDA 调用出错时打印文件名、行号、错误码与错误字符串并退出保证任何 Runtime 错误包括cub::DeviceTransform::Transform与cudaDeviceSynchronize都能被立即捕获。cudaDeviceSynchronize用于确保设备端变换在主机端读取结果thrust::host_vector拷贝之前完成这也是在 CUDA 默认流之外使用设备算法时的标准同步手段。构建方式与 CCCL 依赖管理该示例的构建由 CMakeLists.txt 驱动并已通过 cpp/4_CUDA_Libraries/CMakeLists.txt 中的add_subdirectory(cubDeviceTransform)挂入上层构建树因此既可以在仓库根目录统一构建也可以单独构建该目标。CCCL 依赖的自动获取是构建配置中最关键的一环set(CCCL_SAMPLES_CCCL_TAG v3.3.3 CACHE STRING Tag/branch of NVIDIA/cccl to fetch for the CCCL samples) if(NOT TARGET CCCL::CCCL) include(${CMAKE_CURRENT_SOURCE_DIR}/../../../cmake/CPM.cmake) if(DEFINED CCCL_SOURCE_DIR AND NOT CCCL_SOURCE_DIR STREQUAL ) CPMAddPackage(NAME CCCL SOURCE_DIR ${CCCL_SOURCE_DIR}) else() CPMAddPackage( NAME CCCL GIT_REPOSITORY https://github.com/NVIDIA/cccl GIT_TAG ${CCCL_SAMPLES_CCCL_TAG} ) endif() endif()依赖 CCCL3.3 及以上版本示例默认通过 cmake/CPM.cmakeCPM 包管理器在配置阶段拉取并固定pinned到v3.3.3标签若需使用本地已检出的 CCCL 源码可用-DCCCL_SOURCE_DIR/path/to/cccl覆盖此时 CPM 直接引用本地目录而不再联网拉取顶层 CMakeLists.txt 中 MSVC 编译器会自动追加-Xcompiler/Zc:preprocessor以满足 CCCL 对标准符合预处理器的要求。其余构建要点包括配置项值 / 说明cmake_minimum_required3.20find_package(CUDAToolkit REQUIRED)依赖已安装的 CUDA ToolkitCMAKE_CUDA_ARCHITECTURES默认75 80 86 87 89 90 100 110 120顶层与示例 CMake 均设置编译选项-Wno-deprecated-gpu-targetsENABLE_CUDA_DEBUG开启时加-G支持 cuda-gdb否则加-lineinfoLambda 支持--extended-lambdaCUDA 侧标准cxx_std_17与cuda_std_17链接库CUDA::cudart、CCCL::CCCL公共头文件include_directories(../../../Common)即仓库 Common 目录可分离编译CUDA_SEPARABLE_COMPILATION ON典型构建流程在仓库根目录执行mkdir build cd build cmake .. -DCMAKE_BUILD_TYPERelease make -j$(nproc) cubDeviceTransform运行可执行文件支持--deviceN指定 GPU后程序依次输出设备信息、两个场景的变换结果与期望值最终以Done或FAILED收尾退出码为EXIT_SUCCESS0或EXIT_FAILURE非 0。支持平台与前置条件根据 README.md 的声明支持的 SM 架构SM 7.0 / 7.5 / 8.0 / 8.6 / 8.9 / 9.0 / 10.0 / 11.0 / 12.0即 Volta 及之后各代架构仓库默认编译架构列表见上文构建表格实际以 CMake 配置为准支持的操作系统Linux、Windows支持的 CPU 架构x86_64、aarch64前置条件安装与平台匹配的 CUDA Toolkit并确保 CCCL 3.3 依赖可用默认由 CPM 在配置阶段自动拉取。示例涉及的 API 集中在三类CCCL CUB 的cub::DeviceTransform::Transform、CCCL libcu 的cuda::counting_iterator与cuda::std::tuple以及 CUDA Runtime API 的cudaDeviceSynchronize/cudaGetDeviceProperties三者通过cub/device/device_transform.cuh、cuda/iterator等头文件组织见 cubDeviceTransform.cu 的 include 区。小结从示例到自有工程的迁移要点cubDeviceTransform 为在自有 CUDA 工程中使用 N-to-M 形态的cub::DeviceTransform提供了可直接套用的范式输入输出均以cuda::std::tuple传递迭代器输入输出数量任意组合如 N3→1、N2→2且允许输入序列类型异构算子返回cuda::std::tuple...返回的每个元素对应一条输出序列一次设备调用完成多路读-算-写适合在单次遍历中派生多个结果合成迭代器如cuda::counting_iterator可以替代真实数据序列减少显存占用与拷贝主机端参考实现比对 checkCudaErrors包装是该仓库所有示例通用的自验证模式可直接复用到生产代码的单元测试中。无论是多通道图像处理、多统计量归并计算还是需要同时输出若干派生字段的数据流水线这种单次内核、多路输出的融合写法都是降低启动开销、提升访存效率的实用手段。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 cuda-samples 实战：基于 CCCL 3.3 的 cub::DeviceTransform N 输入 / M 输出融合变换示例解析
cuda-samples 实战基于 CCCL 3.3 的 cub::DeviceTransform N 输入 / M 输出融合变换示例解析【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samplescubDeviceTransform 是 cuda-samples 仓库中演示 CCCL 3.3 扩展后的cub::DeviceTransformN-to-M 形态的官方示例一次设备端调用读取 N 条输入序列通过用户自定义算子返回cuda::std::tuple同时写出 M 条输出序列。本文以该示例为骨架结合仓库源码逐行剖析两种典型用法N3→1 的融合运算与 N2→2 的和差一次成型并讲解其 CMake 构建方式、CCCL 依赖获取机制与运行验证逻辑帮助读者直接上手并迁移到自己的 CUDA 工程中。示例概述与核心知识点示例位于仓库的cpp/4_CUDA_Libraries/cubDeviceTransform/目录由三个文件组成README.md、CMakeLists.txt 与核心实现 cubDeviceTransform.cu。它演示了cub::DeviceTransform在 CCCL 3.3 中引入的N-input / M-output扩展形态覆盖以下关键概念CCCL 3.3NVIDIA 的 CUDA C Core Libraries包含 CUB、libcu、Thrust 的统一版本管理CUB Device Algorithms设备级算法库cub::DeviceTransform::Transform负责以设备级并行方式对序列逐元素应用变换Fused Elementwise Transforms融合逐元素变换将多个输入的读取与多个输出的写入合并为一次设备端调用减少内核启动开销与中间存储Counting Iterators计数迭代器cuda::counting_iterator可在不显式分配数据的情况下提供递增整数序列。示例演示了两个具体场景场景输入数 N输出数 M变换语义场景一31result[i] (a[i] b[i]) * c[i]其中c为从 100 开始的计数迭代器场景二22(sum[i], diff[i]) (a[i] b[i], a[i] - b[i])一趟完成和与差两个场景的结果都会在主机端用参考实现逐一比对并打印OK/FAIL程序返回码随之决定EXIT_SUCCESS/EXIT_FAILURE因此该示例天然自带可回归验证的测试属性。场景一N3 输入 → 1 输出融合乘法与计数迭代器第一个用例run_n_to_one_transform()位于 cubDeviceTransform.cu核心逻辑如下thrust::device_vectorint a {0, -2, 5, 3}; thrust::device_vectorfloat b {5.2f, 3.1f, -1.1f, 3.0f}; auto counting cuda::counting_iteratorint{100}; thrust::device_vectorint result(a.size()); auto op [] __host__ __device__(int x, float y, int z) - int { return static_castint((x y) * z); }; checkCudaErrors(cub::DeviceTransform::Transform( cuda::std::tuple{a.begin(), b.begin(), counting}, result.begin(), a.size(), op)); checkCudaErrors(cudaDeviceSynchronize());这段代码揭示了 N-to-M 形态的三个关键设计点输入侧使用cuda::std::tuple打包迭代器N3 个输入分别来自thrust::device_vectorint a、thrust::device_vectorfloat b以及一个cuda::counting_iteratorint{100}。三个输入的数据类型可以不同int与float混合由算子op的参数列表统一接收计数迭代器代替显式数据cuda::counting_iteratorint{100}会在迭代时依次产生100, 101, 102, ...完全不需要在设备上分配一块保存这些整数的内存。它来自 libcu 的头文件 cuda/iterator是只读合成序列的经典用法算子必须是__host__ __device__可调用对象示例使用--extended-lambda编译选项支持在设备端执行扩展 lambda见 CMakeLists.txtop同时被用于设备端变换与主机端参考计算的思路保持一致。调用完成后代码通过thrust::host_vector把设备数据拷回主机并在主机端按同一公式计算期望值expected[i] static_castint((ha[i] hb[i]) * static_castint(100 i));对照示例数据可推算输出a {0, -2, 5, 3}、b {5.2f, 3.1f, -1.1f, 3.0f}、计数序列{100, 101, 102, 103}得到result {520, 111, 397, 618}注意float乘法结果被static_castint截断取整。程序会把实际值与期望值一并打印并标注OK。场景二N2 输入 → M2 输出一趟同时产出和与差第二个用例run_n_to_m_transform()位于 cubDeviceTransform.cu是 N-to-M 形态最具代表性的一趟多出写法thrust::device_vectorint a {1, 5, 10, 7, 3}; thrust::device_vectorint b {4, 2, 8, 1, 9}; thrust::device_vectorint sum(a.size()); thrust::device_vectorint diff(a.size()); auto op [] __host__ __device__(int x, int y) - cuda::std::tupleint, int { return {x y, x - y}; }; checkCudaErrors(cub::DeviceTransform::Transform(cuda::std::tuple{a.begin(), b.begin()}, cuda::std::tuple{sum.begin(), diff.begin()}, a.size(), op)); checkCudaErrors(cudaDeviceSynchronize());与场景一相比这里出现了两个对称的 tuple输入 tuplecuda::std::tuple{a.begin(), b.begin()}对应 N2 条输入序列输出 tuplecuda::std::tuple{sum.begin(), diff.begin()}对应 M2 条输出序列算子返回值cuda::std::tupleint, inttuple 中的第 k 个元素会按位写入第 k 条输出序列。这就是融合变换的收益所在原本需要启动两次内核一次求和、一次求差并遍历两次输入的工作现在只需一次设备端调用、一次内存遍历即可同时得到sum与diff对于派生多个统计量、多通道处理等场景可以显著降低启动开销与访存次数。根据示例数据可得sum {5, 7, 18, 8, 12}、diff {-3, 3, 2, 6, -6}主机端逐元素比对后打印结果与OK。从源码结构看cubDeviceTransform.cu 的注释N 与 M 并不要求相等M 可以是 1场景一也可以是任意多个场景二为 2输入与输出序列数完全由调用时传入的 tuple 长度决定。设备端调用链与错误处理main()函数cubDeviceTransform.cu展示了本示例涉及的 CUDA Runtime APIint devID findCudaDevice(argc, (const char **)argv); cudaDeviceProp props; checkCudaErrors(cudaGetDeviceProperties(props, devID)); printf(Device: %s (Compute Capability %d.%d)\n\n, props.name, props.major, props.minor);findCudaDevice来自仓库公共头文件 Common/helper_cuda.h若命令行传入--deviceN则使用指定设备否则自动挑选 GFlops 最高的设备并执行cudaSetDevicecudaGetDeviceProperties用于读取设备名称与计算能力major.minor并打印到终端checkCudaErrors定义于 Common/helper_cuda.h包装每次 CUDA 调用出错时打印文件名、行号、错误码与错误字符串并退出保证任何 Runtime 错误包括cub::DeviceTransform::Transform与cudaDeviceSynchronize都能被立即捕获。cudaDeviceSynchronize用于确保设备端变换在主机端读取结果thrust::host_vector拷贝之前完成这也是在 CUDA 默认流之外使用设备算法时的标准同步手段。构建方式与 CCCL 依赖管理该示例的构建由 CMakeLists.txt 驱动并已通过 cpp/4_CUDA_Libraries/CMakeLists.txt 中的add_subdirectory(cubDeviceTransform)挂入上层构建树因此既可以在仓库根目录统一构建也可以单独构建该目标。CCCL 依赖的自动获取是构建配置中最关键的一环set(CCCL_SAMPLES_CCCL_TAG v3.3.3 CACHE STRING Tag/branch of NVIDIA/cccl to fetch for the CCCL samples) if(NOT TARGET CCCL::CCCL) include(${CMAKE_CURRENT_SOURCE_DIR}/../../../cmake/CPM.cmake) if(DEFINED CCCL_SOURCE_DIR AND NOT CCCL_SOURCE_DIR STREQUAL ) CPMAddPackage(NAME CCCL SOURCE_DIR ${CCCL_SOURCE_DIR}) else() CPMAddPackage( NAME CCCL GIT_REPOSITORY https://github.com/NVIDIA/cccl GIT_TAG ${CCCL_SAMPLES_CCCL_TAG} ) endif() endif()依赖 CCCL3.3 及以上版本示例默认通过 cmake/CPM.cmakeCPM 包管理器在配置阶段拉取并固定pinned到v3.3.3标签若需使用本地已检出的 CCCL 源码可用-DCCCL_SOURCE_DIR/path/to/cccl覆盖此时 CPM 直接引用本地目录而不再联网拉取顶层 CMakeLists.txt 中 MSVC 编译器会自动追加-Xcompiler/Zc:preprocessor以满足 CCCL 对标准符合预处理器的要求。其余构建要点包括配置项值 / 说明cmake_minimum_required3.20find_package(CUDAToolkit REQUIRED)依赖已安装的 CUDA ToolkitCMAKE_CUDA_ARCHITECTURES默认75 80 86 87 89 90 100 110 120顶层与示例 CMake 均设置编译选项-Wno-deprecated-gpu-targetsENABLE_CUDA_DEBUG开启时加-G支持 cuda-gdb否则加-lineinfoLambda 支持--extended-lambdaCUDA 侧标准cxx_std_17与cuda_std_17链接库CUDA::cudart、CCCL::CCCL公共头文件include_directories(../../../Common)即仓库 Common 目录可分离编译CUDA_SEPARABLE_COMPILATION ON典型构建流程在仓库根目录执行mkdir build cd build cmake .. -DCMAKE_BUILD_TYPERelease make -j$(nproc) cubDeviceTransform运行可执行文件支持--deviceN指定 GPU后程序依次输出设备信息、两个场景的变换结果与期望值最终以Done或FAILED收尾退出码为EXIT_SUCCESS0或EXIT_FAILURE非 0。支持平台与前置条件根据 README.md 的声明支持的 SM 架构SM 7.0 / 7.5 / 8.0 / 8.6 / 8.9 / 9.0 / 10.0 / 11.0 / 12.0即 Volta 及之后各代架构仓库默认编译架构列表见上文构建表格实际以 CMake 配置为准支持的操作系统Linux、Windows支持的 CPU 架构x86_64、aarch64前置条件安装与平台匹配的 CUDA Toolkit并确保 CCCL 3.3 依赖可用默认由 CPM 在配置阶段自动拉取。示例涉及的 API 集中在三类CCCL CUB 的cub::DeviceTransform::Transform、CCCL libcu 的cuda::counting_iterator与cuda::std::tuple以及 CUDA Runtime API 的cudaDeviceSynchronize/cudaGetDeviceProperties三者通过cub/device/device_transform.cuh、cuda/iterator等头文件组织见 cubDeviceTransform.cu 的 include 区。小结从示例到自有工程的迁移要点cubDeviceTransform 为在自有 CUDA 工程中使用 N-to-M 形态的cub::DeviceTransform提供了可直接套用的范式输入输出均以cuda::std::tuple传递迭代器输入输出数量任意组合如 N3→1、N2→2且允许输入序列类型异构算子返回cuda::std::tuple...返回的每个元素对应一条输出序列一次设备调用完成多路读-算-写适合在单次遍历中派生多个结果合成迭代器如cuda::counting_iterator可以替代真实数据序列减少显存占用与拷贝主机端参考实现比对 checkCudaErrors包装是该仓库所有示例通用的自验证模式可直接复用到生产代码的单元测试中。无论是多通道图像处理、多统计量归并计算还是需要同时输出若干派生字段的数据流水线这种单次内核、多路输出的融合写法都是降低启动开销、提升访存效率的实用手段。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考