Skip to content

Fix CuTe tuple protocol compatibility across CCCL versions - #3386

Open
xiufanl wants to merge 3 commits into
NVIDIA:mainfrom
xiufanl:codex/cute-cccl-tuple-compat
Open

Fix CuTe tuple protocol compatibility across CCCL versions#3386
xiufanl wants to merge 3 commits into
NVIDIA:mainfrom
xiufanl:codex/cute-cccl-tuple-compat

Conversation

@xiufanl

@xiufanl xiufanl commented Jul 14, 2026

Copy link
Copy Markdown
Contributor

Summary

Recent CCCL versions no longer provide the private cuda/std/__tuple_dir/structured_bindings.h aggregation header, while older CCCL layouts do not always expose the tuple protocol primary declarations through the newer split headers. CuTe's NVRTC path needs to support both layouts.

Newer CCCL versions also provide cuda::std::get for integer_sequence. When CuTe aliases that type directly, its associated namespace brings the newer overload into unqualified lookup and conflicts with CuTe's tuple protocol.

Changes

  • Use the public cuda/std/tuple entry point for NVRTC compilation.
  • Detect the available CCCL header layout and retain fallback tuple primary declarations only when neither the legacy nor split headers are available.
  • Remove direct dependencies on the private structured-bindings header from CuTe tuple-like containers.
  • Define a CuTe-owned integer_sequence and adapt the configured STL make_integer_sequence result into it.
  • Keep tuple algorithms on their original unqualified get path; ordinary tuple-like ADL remains unchanged, while CuTe integer sequences no longer associate with the CCCL namespace.
  • Add focused regression coverage for two- and three-input transforms and fold operations on integer sequences, including a type-ownership assertion.

Compatibility

  • Preserves older CUDA/CCCL support by continuing to use the configured STL namespace to generate integer packs.
  • Uses CCCL-provided tuple declarations when supported and retains the existing fallback only for older layouts.
  • Avoids requiring newer library APIs or compiler-specific forwarding helpers.

Testing

  • Added focused CuTe tuple regression cases in test/unit/cute/core/tuple.cpp.
  • git diff --check passes; full CUDA validation is pending on the updated commit.

return f(cute::get<I>(static_cast<T&&>(t))...);
} else {
return f(get<I>(static_cast<T&&>(t))...);
}

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Would much rather implement CuTe's own integer_sequence and/or fix this overload properly than WAR it and require this repetition everywhere.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks - agreed. I replaced the per-call branching with a CuTe-owned integer_sequence and adapted the configured STL make_integer_sequence result into it. Since the associated namespace is now cute, the original unqualified get calls work without colliding with newer cuda::std::get, so tuple_algorithms.hpp is back to its direct ADL path. I also kept the focused transform/fold coverage and added a type-ownership assertion.

@xiufanl
xiufanl force-pushed the codex/cute-cccl-tuple-compat branch 2 times, most recently from dfbbe14 to 8c97b17 Compare July 15, 2026 10:38
Use CCCL's public tuple entry point during NVRTC compilation while preserving fallback declarations for older header layouts.

Give CuTe its own integer_sequence type and adapt the configured STL generator into it. This prevents newer CCCL integer_sequence get overloads from entering ADL, keeps the original tuple algorithm path, and adds transform and fold regression coverage.
@xiufanl
xiufanl force-pushed the codex/cute-cccl-tuple-compat branch from 8c97b17 to 1061020 Compare July 15, 2026 19:36
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