Add compatibility for CUB/CCCL 3.0+ API

[omoYang]

PR Category

Environment Adaptation

PR Types

Others

Description

1. Migrate APIs removed in CCCL 3.0 to their replacements so the same source compiles against both old CCCL 2.x and CCCL 3.x:

-- cub::Sum/Equality/Max -> cuda::std::plus/equal_to/cuda::maximum (aliased in cub.h under CUB_VERSION >= 300000)

• cub::CTA_SYNC() -> __syncthreads()
• cub::TransformInputIterator/CountingInputIterator -> thrust::transform_iterator/counting_iterator
• cub::LaneId()/LaneMaskLt() -> cuda::ptx::get_sreg_laneid/lanemask_lt
• cub::BaseTraits 5-param -> 4-param (guarded by CUB_VERSION)
• thrust_exec_check_disable -> _CCCL_EXEC_CHECK_DISABLE

2. Add cccl.cmake CUDA 13 branch that checks out CCCL v3.1.0 and applies util_device.cuh.v3.1.0.patch (inline -> forceinline).

是否引起精度变化

否

[Paddle-Bot]

这些 Thrust iterator 头文件只在 __NVCC__ 分支中引入，但本次替换后的 thrust::counting_iterator / thrust::transform_iterator 也会在 __HIPCC__ 路径编译（例如 fused_token_prune_kernel.cu、reduce_function.h）。hipcub/hipcub.hpp 只提供 hipCUB/rocPRIM 迭代器，不会声明这些 thrust 类型，因此 HIP 编译会在这些翻译单元报类型未定义。请把两个 Thrust 头移到 CUDA/HIP 共用分支，或为 HIP 使用对应的 rocThrust 头/命名空间。

[Paddle-Bot]

这里会把所有 CUDA 12.x（包括仓库覆盖的 CUDA 12.0.1 镜像和 Windows CUDA 12.0）切换到 CCCL v3.1.0。CCCL 3.x 的兼容矩阵只承诺每个 12.x 系列的最新补丁版本，12.0/12.1 并不在该支持范围；因此这些仍受支持的配置会被强制使用未验证的 CCCL 版本并可能在配置或编译阶段失败。请按实际验证过的 CUDA/CCCL 组合收窄条件，或为旧的 12.x 保留兼容的 CCCL 回退。

[Paddle-Bot]

Paddle-Bot Review Board （review完成)

序号	位置	优先级	规则来源	状态
1	HIP Thrust 头文件		仓库规则：功能正确性与兼容性	🚧
2	CUDA 12.x CCCL 分支		仓库规则：功能正确性与兼容性	🚧

^{Powered by Nyanpasu with gpt-5.6-sol xhigh, please check the suggestions carefully.}

[Paddle-Bot]

Failed CI看板

流水线名称	问题标签	修复建议	日志片段
CI #34965026844 / Distribute-stable-build / Build	PR代码/CCCL-NVTX头文件冲突	【与当前改动有关】避免同一编译单元同时包含旧版 nvToolsExt.h 和 CCCL 3.x 的 nvtx3/nvToolsExt.h，统一 NVTX 兼容入口后重跑 CUDA 构建	Job
CI #34965026844 / Linux-DCU / Build	PR代码/ROCm transform_iterator 不兼容	【与当前改动有关】为 HIP/ROCm 保留 hipCUB/原有 iterator 实现，仅在 NVCC + CCCL 3.x 路径使用 thrust::transform_iterator，并补充 ROCm 编译验证	Job
Coverage #34965026915 / Coverage build	PR代码/CCCL-NVTX头文件冲突	【与当前改动有关】修复 paddle/phi/backends/dynload/nvtx.h 与 CCCL 3.x NVTX3 头文件的互斥包含关系	Job
CI-Build #34965026977 / PR-CI-Inference / Build	PR代码/CCCL-NVTX头文件冲突	【与当前改动有关】确保 CCCL 的 thrust/detail/config.h 引入 NVTX3 后不再继续引入旧版 nvToolsExt.h	Job
CI-Build #34965026977 / Linux-build / Build	PR代码/CCCL-NVTX头文件冲突	【与当前改动有关】在所有 CUDA 构建配置中统一 NVTX 头文件版本或增加兼容保护，避免 nvtx* 类型重复定义	Job
CI-H #34965026941 / Coverage build	构建被取消	【无法判断】当前快照只有 The operation was canceled，无法判断是流水线级联取消、资源问题还是代码问题；应重跑该 job 并获取完整取消上下文	Job
Bot Approval Required #34965026473 / Require review-bot approval	当前 head 缺少审批	【与当前改动无关】为 SHA 52b72343c4d5c... 获取 risemeup1111、Paddle-Bot 或允许的人工审批账号的有效 APPROVED 状态	Job
Approval #34965026494 / Check approval	审批检查失败，详情缺失	【无法判断】补充日志中两条 approved errors 对应的文件和规则后再判断是否需要修改代码或审批配置	Job

日志分析报告

失败的测试 case:

1. CI #34965026844 / Distribute-stable-build / Build
   步骤：CUDA 编译。
   关键错误：/usr/local/cuda/.../nvToolsExt.h 与 nvtx3/nvToolsExt.h 重复定义 nvtxPayloadType_t、nvtxEventAttributes_v2、nvtxResourceGenericType_t 等类型；ninja exit 7。

2. CI #34965026844 / Linux-DCU / Build
   步骤：HIP/ROCm 编译 lamb_kernel.cu。
   关键错误：rocprim transform_iterator 找不到 SquareFunctor 的 value_type；SquaredL2Norm 调用无匹配函数；gfx906/gfx928 编译失败。

3. Coverage #34965026915 / Coverage build
   步骤：编译 paddle/phi/api/profiler/profiler.cc。
   关键错误：旧版 nvToolsExt.h 与 CCCL NVTX3 头文件产生多项重复声明；make exit 7。

4. CI-Build #34965026977 / PR-CI-Inference / Build
   步骤：编译 profiler.cc。
   关键错误：CCCL libcudacxx 的 nvtx3.h 与 paddle/phi/backends/dynload/nvtx.h 引入的 nvToolsExt.h 冲突。

5. CI-Build #34965026977 / Linux-build / Build
   步骤：Linux CUDA 编译。
   关键错误：同样出现 nvToolsExt.h 与 nvtx3/nvToolsExt.h 的重复定义；make exit 7。

6. CI-H #34965026941 / Coverage build
   步骤：cmake 后执行 make -j20，进度约 68%。
   关键错误：The operation was canceled；快照未提供实际编译错误。

7. Bot Approval Required #34965026473 / Require review-bot approval
   步骤：检查当前 head 的机器人和人工审批状态。
   关键错误：当前 head 未获得要求的 APPROVED 状态，脚本 exit 1。

8. Approval #34965026494 / Check approval
   步骤：执行 approval 检查。
   关键错误：There are 2 approved errors；快照未提供两条错误的具体内容。

根本原因分析:

1. Jobs 1、3、4、5 属于同一条根因链。PR 在 CUDA 12.0 以上切换到 CCCL v3.1.0，并引入 CCCL 3.x 的 Thrust/CUDA 头文件。日志显示 CCCL 的 cuda/__nvtx/nvtx3.h 与 Paddle 现有 paddle/phi/backends/dynload/nvtx.h 引入的旧版 nvToolsExt.h 同时进入 profiler.cc 等编译单元，最终造成 NVTX 类型和函数重复声明。

2. Job 2 是另一条与 PR 直接相关的根因链。PR 将 cub::TransformInputIterator 改为 thrust::transform_iterator，但相关修改没有完整隔离 HIP/ROCm 路径。ROCm 的 rocPRIM iterator 约束与 CUDA Thrust iterator 不同，导致 SquareFunctor 的 value_type 推导失败，并进一步造成 SquaredL2Norm 模板调用失败。

3. Job 6 只有取消信息，没有可验证的失败编译或测试错误，不能据此归因到 PR，也不能确认是前述构建失败导致的级联取消。

4. Jobs 7、8 是审批流程检查，不是源码构建测试。Job 7 明确表明当前 head 缺少有效审批；Job 8 只给出两条错误的计数，缺少具体规则和文件，因此无法判断其与 PR diff 的关系。

修复建议:

1. 【与当前改动有关】Distribute-stable-build / Build、Coverage build、PR-CI-Inference / Build、Linux-build / Build：统一 NVTX 头文件兼容策略，保证 CCCL 3.x 的 NVTX3 和旧版 nvToolsExt.h 不会在同一编译单元同时声明；随后重新验证所有 CUDA 构建矩阵。

2. 【与当前改动有关】Linux-DCU / Build：对 reduce_function.h、squared_l2_norm.h 及其他 iterator 替换点增加 CUDA/HIP 分支或后端别名，ROCm 路径继续使用兼容的 hipCUB/rocPRIM iterator，并执行 Linux-DCU 编译验证。

3. 【无法判断】CI-H / Coverage build：针对当前 SHA 重跑被取消的 job，并取得完整取消前后的日志后再判断是否需要代码修复。

4. 【与当前改动无关】Require review-bot approval：补齐当前 head 的机器人或允许人工审批。

5. 【无法判断】Check approval：获取两条 approved errors 的完整内容、对应文件和审批规则，再决定是修改 PR 还是补充审批。

Powered by Nyanpasu with gpt-5.6-luna 默认推理级别, please check the suggestions carefully.