Skip to content

[release/2.10] [ROCm] derive the hipcub compatible cccl version from 'HIPCUB_CCCL_VERSION'#3439

Draft
jithunnair-amd wants to merge 1 commit into
ROCm:release/2.10from
jithunnair-amd:2.10-cherry-pick-pr-188072
Draft

[release/2.10] [ROCm] derive the hipcub compatible cccl version from 'HIPCUB_CCCL_VERSION'#3439
jithunnair-amd wants to merge 1 commit into
ROCm:release/2.10from
jithunnair-amd:2.10-cherry-pick-pr-188072

Conversation

@jithunnair-amd

Copy link
Copy Markdown
Collaborator

Cherry-pick of pytorch#188072 (commit c8e2473) onto release/2.10.

Summary

  • Introduce dependency on libhipcxx
  • Guard FpLimits for c10::BFloat16 backport to pre CCCL 3.x
  • Extend CUB_VERSION derivation for hipCUB using HIPCUB_CCCL_VERSION

Notes

This PR has merge conflicts that need to be resolved manually (conflict markers committed as requested).

Test plan

  • Build with ROCm
  • Verify hipCUB/CCCL compatibility glue compiles

Authored with assistance from Cursor

Made with Cursor

…RSION' (pytorch#188072)

Recreated pytorch#184552 from personal fork.

hipCUB is targeting to be compatible with CCCL 3.0 (ROCm/rocm-libraries#4079). This requires updating some of the CUB and hipCUB compatibility glue.

To help this transition hipCUB will be exposing `HIPCUB_CCCL_VERSION` in one of the following ROCm releases (>=7.14) (ROCm/rocm-libraries#7632), which will be bumped once hipCUB reaches parity with CCCL 3.0 and beyond.

This PR is created in anticipation for this future changes.

- Introduce dependency on [`libhipcxx`](https://github.com/ROCm/libhipcxx).
- Guard `FpLimits` for `c10::BFloat16` backport to pre CCCL 3.x. Numeric limits are handled by libhipcxx.
- Extend `CUB_VERSION` derivation for hipCUB. Keep old fallback for backwards compatibility. Only enable CUB V3 API once hipCUB reaches parity with CCCL 3.x.

Pull Request resolved: pytorch#188072
Approved by: https://github.com/jeffdaily

(cherry picked from commit c8e2473)
@rocm-repo-management-api

rocm-repo-management-api Bot commented Jul 14, 2026

Copy link
Copy Markdown

Jenkins build for a375156b72041e35467399ffaed4f3409e28b33c commit finished as FAILURE
Links: Pipeline Overview / Build artifacts / Test Results

Detected error during Pytorch building:

      |     ^
/var/lib/jenkins/pytorch/third_party/fbgemm/external//composable_kernel/include/ck/tensor_operation/gpu/device/impl/device_grouped_gemm_multiple_d_xdl_cshuffle_tile_loop.hpp:65:5: warning: failed to meet occupancy target given by 'amdgpu-waves-per-eu' in '_ZN2ck16tensor_operation6device34kernel_grouped_gemm_multiple_d_xdlINS_34GridwiseGemmMultiD_xdl_cshuffle_v3INS_13tensor_layout4gemm8RowMajorENS5_11ColumnMajorENS_5TupleIJS6_S7_EEES6_DB8_SA_ffNS8_IJffEEEtNS0_12element_wise11PassThroughESD_NSC_16MultiplyMultiplyELNS1_18GemmSpecializationE0ELi256ELi256ELi224ELi128ELi16ELi16ELi16ELi16ELi8ELi7ENS_8SequenceIJLi8ELi32ELi1EEEENSG_IJLi1ELi0ELi2EEEESI_Li2ELi16ELi16ELb0ELi0ESH_SI_SI_Li2ELi16ELi16ELb0ELi0ELi2ELi1ENSG_IJLi1ELi64ELi1ELi4EEEENSG_IJLi8ELi8ELi1EEEELNS_26BlockGemmPipelineSchedulerE0ELNS_24BlockGemmPipelineVersionE2ESA_SA_SA_SA_Lb0EEENS1_25GroupedGemmKernelArgumentILi2EEELSF_0ESA_SA_SB_tS6_S7_S9_S6_Li128ENS_25OffsettedBlockToCTileMap2INS_39BlockToCTileMap_Grouped_M00_N0_M01AdaptILi8ELi256ELi224EEEEESS_SD_SD_SE_LSL_0ELSM_2EEEvPU3AS4KviT13_T14_T15_': desired occupancy was 2, final occupancy is 1 [-Wpass-failed]
2 warnings generated when compiling for gfx942.
[7246/8132] Linking CXX static library lib/libfbgemm_genai.a
[7247/8132] Building HIPCC object caffe2/CMakeFiles/torch_hip.dir/__/aten/src/ATen/hip/torch_hip_generated_cub-RadixSortPairs-uint32-8.hip.o
FAILED: caffe2/CMakeFiles/torch_hip.dir/__/aten/src/ATen/hip/torch_hip_generated_cub-RadixSortPairs-uint32-8.hip.o /var/lib/jenkins/pytorch/build/caffe2/CMakeFiles/torch_hip.dir/__/aten/src/ATen/hip/torch_hip_generated_cub-RadixSortPairs-uint32-8.hip.o 
cd /var/lib/jenkins/pytorch/build/caffe2/CMakeFiles/torch_hip.dir/__/aten/src/ATen/hip && /opt/conda/envs/py_3.12/lib/python3.12/site-packages/cmake/data/bin/cmake -E make_directory /var/lib/jenkins/pytorch/build/caffe2/CMakeFiles/torch_hip.dir/__/aten/src/ATen/hip/. && /opt/conda/envs/py_3.12/lib/python3.12/site-packages/cmake/data/bin/cmake -D verbose:BOOL=OFF -D build_configuration:STRING=RELEASE -D generated_file:STRING=/var/lib/jenkins/pytorch/build/caffe2/CMakeFiles/torch_hip.dir/__/aten/src/ATen/hip/./torch_hip_generated_cub-RadixSortPairs-uint32-8.hip.o -P /var/lib/jenkins/pytorch/build/caffe2/CMakeFiles/torch_hip.dir/__/aten/src/ATen/hip/torch_hip_generated_cub-RadixSortPairs-uint32-8.hip.o.cmake
In file included from /var/lib/jenkins/pytorch/aten/src/ATen/hip/cub-RadixSortPairs-uint32-8.hip:2:
In file included from /var/lib/jenkins/pytorch/aten/src/ATen/hip/cub-RadixSortPairs.cuh:6:
In file included from /var/lib/jenkins/pytorch/aten/src/ATen/hip/cub.cuh:15:
/var/lib/jenkins/pytorch/aten/src/ATen/hip/cub_definitions.cuh:31:1: error: version control conflict marker in file

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants