Skip to content

Build against CUDA 13 / CCCL 3 and fix gradcheck on torch >= 2.0 - #637

Open
kevintsq wants to merge 1 commit into
NVIDIA:masterfrom
kevintsq:cuda13-blackwell
Open

Build against CUDA 13 / CCCL 3 and fix gradcheck on torch >= 2.0#637
kevintsq wants to merge 1 commit into
NVIDIA:masterfrom
kevintsq:cuda13-blackwell

Conversation

@kevintsq

@kevintsq kevintsq commented Aug 6, 2026

Copy link
Copy Markdown

CUDA 13 ships CCCL 3, which relocates the bundled Thrust/CUB headers and removes a number of APIs the 0.5.4 sources relied on. Five changes are needed to compile:

  • setup.py: CUDA >= 13 moved Thrust/CUB from /include to /include/cccl, so neither nvcc nor the host compiler finds them implicitly. Detect and add that directory.
  • coordinate_map_functors.cuh: thrust::unary_function is gone in CCCL 3. It only supplied argument_type/result_type typedefs, so declare them directly instead of inheriting.
  • coordinate_map_gpu.cuh: assigning the unique_ptr returned by map_type::create() to a shared_ptr makes libstdc++ call an unqualified __to_address(). map_type mentions cuda::std::pair under CCCL 3, so ADL also finds cuda::std::__to_address and the call is ambiguous. Build the shared_ptr from the released pointer to bypass that overload.
  • cudf/detail/nvtx/ranges.hpp: the vendored nvtx3.hpp predates the NVTX v3 headers CUDA >= 12 ships and PyTorch pulls in. Both declare an inline versioning namespace inside nvtx3, making every name in it ambiguous. These ranges are profiling annotations only, so build without them.
  • concurrent_unordered_map.cuh, coordinate_map_gpu.cu, spmm.cu: CCCL 3 dropped transitive includes, and CUDA 13 replaced the device-ordinal cudaMemPrefetchAsync overload with a cudaMemLocation one.

Note that spmm.cu keeps upstream's single sort_by_key over a zip iterator. Splitting it into two sort_by_key calls on the same key array is not equivalent: the first sorts the keys in place, so the second computes an identity permutation and the values stop matching their (row, col) pairs.

Separately, utils/gradcheck.py unconditionally forwarded check_sparse_nnz= to torch.autograd.gradcheck, which dropped that parameter in torch 2.0, so every gradient check raised TypeError. Forward it only when accepted, and map it to the replacement masked= argument otherwise.

Verified on CUDA 13.0 / PyTorch 2.9.1+cu130 / Python 3.13 / gcc 13.3 / sm_120 (RTX 5090 and RTX PRO 6000 Blackwell). Convolution, pooling and coo_spmm agree bit-exactly between the CPU and CUDA backends forward and backward, and coo_spmm matches a reference COO product on unsorted input. The test suite's remaining failures are pre-existing test-code rot against torch 2.x and the pre-0.5 SparseTensor API, not regressions.

CUDA 13 ships CCCL 3, which relocates the bundled Thrust/CUB headers and
removes a number of APIs the 0.5.4 sources relied on. Five changes are
needed to compile:

- setup.py: CUDA >= 13 moved Thrust/CUB from <cuda>/include to
  <cuda>/include/cccl, so neither nvcc nor the host compiler finds them
  implicitly. Detect and add that directory.
- coordinate_map_functors.cuh: thrust::unary_function is gone in CCCL 3.
  It only supplied argument_type/result_type typedefs, so declare them
  directly instead of inheriting.
- coordinate_map_gpu.cuh: assigning the unique_ptr returned by
  map_type::create() to a shared_ptr makes libstdc++ call an unqualified
  __to_address(). map_type mentions cuda::std::pair under CCCL 3, so ADL
  also finds cuda::std::__to_address and the call is ambiguous. Build the
  shared_ptr from the released pointer to bypass that overload.
- cudf/detail/nvtx/ranges.hpp: the vendored nvtx3.hpp predates the NVTX v3
  headers CUDA >= 12 ships and PyTorch pulls in. Both declare an inline
  versioning namespace inside nvtx3, making every name in it ambiguous.
  These ranges are profiling annotations only, so build without them.
- concurrent_unordered_map.cuh, coordinate_map_gpu.cu, spmm.cu: CCCL 3
  dropped transitive includes, and CUDA 13 replaced the device-ordinal
  cudaMemPrefetchAsync overload with a cudaMemLocation one.

Note that spmm.cu keeps upstream's single sort_by_key over a zip iterator.
Splitting it into two sort_by_key calls on the same key array is not
equivalent: the first sorts the keys in place, so the second computes an
identity permutation and the values stop matching their (row, col) pairs.

Separately, utils/gradcheck.py unconditionally forwarded check_sparse_nnz=
to torch.autograd.gradcheck, which dropped that parameter in torch 2.0, so
every gradient check raised TypeError. Forward it only when accepted, and
map it to the replacement masked= argument otherwise.

Verified on CUDA 13.0 / PyTorch 2.9.1+cu130 / Python 3.13 / gcc 13.3 /
sm_120 (RTX 5090 and RTX PRO 6000 Blackwell). Convolution, pooling and
coo_spmm agree bit-exactly between the CPU and CUDA backends forward and
backward, and coo_spmm matches a reference COO product on unsorted input.
The test suite's remaining failures are pre-existing test-code rot against
torch 2.x and the pre-0.5 SparseTensor API, not regressions.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
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.

1 participant