Skip to content

[Backport 3.5] Update cuda::ptx for CUDA 13.4 - #11054

Open
pciolkosz wants to merge 1 commit into
NVIDIA:branch/3.5.xfrom
pciolkosz:backport-10909-branch-3.5.x
Open

[Backport 3.5] Update cuda::ptx for CUDA 13.4#11054
pciolkosz wants to merge 1 commit into
NVIDIA:branch/3.5.xfrom
pciolkosz:backport-10909-branch-3.5.x

Conversation

@pciolkosz

@pciolkosz pciolkosz commented Aug 28, 2026

Copy link
Copy Markdown
Contributor

Backports #10909 onto branch/3.5.x

* Add cuda::ptx for Rubin CTK 13.4

* Add NV TARGET SM107

* Disable `tensormap.replace.assumelarge` ctk13.4

* Update ptx ctk 13.4

* add new dot variants

* Preserve main PTX test splits

* Guard fabric half PTX overloads

* Suppress clang-tidy in generated PTX headers

---------

Co-authored-by: Federico Busato <fbusato@nvidia.com>
(cherry picked from commit f047516)
@pciolkosz
pciolkosz requested review from a team as code owners August 28, 2026 03:28
@github-project-automation github-project-automation Bot moved this to Todo in CCCL Aug 28, 2026
@cccl-authenticator-app cccl-authenticator-app Bot moved this from Todo to In Review in CCCL Aug 28, 2026
@coderabbitai

coderabbitai Bot commented Aug 28, 2026

Copy link
Copy Markdown
Contributor

Review Change Stack

📝 Walkthrough

Summary by CodeRabbit

  • New Features

    • Added support for CUDA SM 107 and SM 107a architecture detection and traits.
    • Added access to numerous PTX instructions, including fabric operations, asynchronous prefetch, matrix load/store, mbarrier operations, and priority controls.
    • Added PTX variant definitions for layouts, barrier phases, and reporting mechanisms.
  • Documentation

    • Expanded and corrected PTX instruction references, availability information, examples, and proxy-fence documentation.
  • Bug Fixes

    • Corrected an invalid variable in a mapping example and fixed SM 107 feature detection.
  • Tests

    • Added compile and architecture coverage for the new capabilities.

Walkthrough

The change adds CUDA 13.4 PTX instructions and documentation, updates PTX code generation, adds SM 107 architecture support, and introduces compile-only validation for PTX headers.

Changes

CUDA 13.4 PTX support

Layer / File(s) Summary
PTX documentation and instruction index
docs/libcudacxx/ptx/instructions.rst, docs/libcudacxx/ptx/instructions/*
Adds documentation and internal references for asynchronous bulk operations, fabric operations, mbarrier variants, ldmatrix, stmatrix, prefetch, and related instructions.
PTX headers and public variants
libcudacxx/include/cuda/__ptx/instructions/*, libcudacxx/include/cuda/__ptx/ptx_dot_variants.h, libcudacxx/include/cuda/ptx
Adds wrapper headers, umbrella-header includes, generated instruction registrations, and PTX variant enums and constants.
PTX instruction generator
libcudacxx/codegen/add_ptx_instruction.py
Adds separate C++ naming, overwrite handling, output validation, and duplicate-registration checks.
PTX compile validation
libcudacxx/test/libcudacxx/cuda/ptx/*, libcudacxx/test/utils/libcudacxx/test/format.py
Adds compile-only tests for new PTX headers and evaluates them without linking or execution.

SM 107 architecture support

Layer / File(s) Summary
Architecture definitions and detection
libcudacxx/include/cuda/__device/arch_traits.h, libcudacxx/include/cuda/std/__cccl/*, libcudacxx/include/nv/detail/__target_macros
Adds sm_107 and sm_107a traits, architecture lists, NVRTC fallbacks, SM 107 feature detection, and preprocessor support for 20 arguments.
Architecture tests
libcudacxx/test/libcudacxx/cuda/ccclrt/device/*, libcudacxx/test/nvtarget/*
Adds coverage for architecture IDs, formatting, conversions, traits, dispatch, and SM 107a/107f target-specific feature checks.

Possibly related PRs

  • NVIDIA/cccl#10909: Updates the same PTX documentation, generator, architecture support, headers, and CUDA 13.4 tests.

Suggested reviewers: ericniebler, jacobfaib

Merge Risk: 🔵 Low · up to 925f6

The PR adds CUDA 13.4 PTX declarations and compile tests. It is broadly mergeable, but owner follow-up is needed for skipped compile coverage and include-guard consistency because these could hide valid test failures or cause localized header conflicts.


Comment @coderabbitai help to get the list of available commands.

@coderabbitai coderabbitai Bot left a comment

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.

Actionable comments posted: 2

🧹 Nitpick comments (4)
libcudacxx/include/cuda/__ptx/instructions/stmatrix.h (1)

12-13: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

suggestion: Use a full-path-derived include guard. _CUDA_PTX_STMATRIX_H_ does not include the __PTX/INSTRUCTIONS path components. Rename the guard and its closing comment to the uppercase full-path form required for headers. As per coding guidelines, “Headers must use include guards derived from the uppercase full path.”

Also applies to: 52-52

Source: Coding guidelines

libcudacxx/include/cuda/__ptx/ptx_dot_variants.h (1)

251-252: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

suggestion: Use inline constexpr for the new namespace-scope constants. mbarrier_phase_primary, mbarrier_phase_conditional, layout_v0, layout_v1, and the report-mechanism objects use static constexpr. Change them to inline constexpr to follow the header rule and avoid separate definitions in each translation unit. As per coding guidelines, “All constexpr variables at namespace or global scope must also be inline.”

Also applies to: 260-261, 275-281

Source: Coding guidelines

libcudacxx/test/libcudacxx/cuda/ptx/ptx.applypriority.async.bulk.compile.pass.cpp (2)

17-17: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

suggestion: Make the generated test-header includes conform to the angle-bracket rule. Update libcudacxx/codegen/add_ptx_instruction.py and regenerate these three tests. Ensure the PTX test directory is in the angle-include search path.

Source: Coding guidelines


10-10: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

suggestion: Document the thread-related test skips. Add a reason for libcpp-has-no-threads in all three tests, or remove the directive. Update the generator template so regeneration preserves the fix.

suggestion: Use angle brackets for generated headers. Change the quoted generated/*.h includes in all three tests and in the generator template to angle-bracket includes.

Source: Coding guidelines


ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: e8134a12-5b71-4f15-8e74-6c83739ab23f

📥 Commits

Reviewing files that changed from the base of the PR and between 3af8031 and 925f6bc.

⛔ Files ignored due to path filters (182)
  • docs/libcudacxx/ptx/instructions/generated/applypriority_async_bulk.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/clusterlaunchcontrol.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/cp_async_bulk.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/cp_async_bulk_multicast.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/cp_async_bulk_prefetch.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/cp_async_bulk_prefetch_tensor.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/cp_async_bulk_tensor.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/cp_async_bulk_tensor_gather_scatter.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/cp_async_bulk_tensor_multicast.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/cp_reduce_async_bulk_tensor.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/fabric_submit.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/fabric_try_get.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/fabric_try_pullred.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/fabric_try_put.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/fabric_try_red.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/fabric_wait.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/fence_proxy_alias.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/fence_proxy_async_generic_sync_restrict.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/fence_proxy_fabric_fabric_alias.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/fence_proxy_fabric_generic_alias.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/fence_proxy_generic_fabric_alias.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/ldmatrix.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/mbarrier_arrive.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/mbarrier_arrive_drop.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/mbarrier_check_layout.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/mbarrier_complete_tx.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/mbarrier_init.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/mbarrier_pending_count.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/mbarrier_test_wait.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/mbarrier_test_wait_parity.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/mbarrier_try_wait.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/mbarrier_try_wait_parity.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/prefetch.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/setmaxnreg.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/stmatrix.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/tcgen05_alloc.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/tcgen05_commit.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/tcgen05_cp.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/tcgen05_fence.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/tcgen05_ld.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/tcgen05_mma.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/tcgen05_mma_sp.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/tcgen05_mma_ws.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/tcgen05_shift.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/tcgen05_st.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/tcgen05_wait.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/tensormap_replace.rst is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/applypriority_async_bulk.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/barrier_cluster.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/bfind.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/bmsk.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/clusterlaunchcontrol.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_commit_group.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_multicast.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_prefetch.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_prefetch_tensor.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_tensor.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_tensor_gather_scatter.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_tensor_multicast.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_wait_group.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_mbarrier_arrive.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_mbarrier_arrive_noinc.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/cp_reduce_async_bulk.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/cp_reduce_async_bulk_bf16.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/cp_reduce_async_bulk_f16.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/cp_reduce_async_bulk_tensor.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/elect_sync.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/exit.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/fabric_submit.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/fabric_try_get.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/fabric_try_pullred.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/fabric_try_put.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/fabric_try_red.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/fabric_wait.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/fence.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/fence_mbarrier_init.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/fence_proxy_alias.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/fence_proxy_async.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/fence_proxy_async_generic_sync_restrict.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/fence_proxy_fabric_fabric_alias.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/fence_proxy_fabric_generic_alias.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/fence_proxy_generic_fabric_alias.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/fence_proxy_tensormap_generic.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/fence_sync_restrict.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/get_sreg.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/getctarank.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/ld.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/ldmatrix.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_arrive.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_arrive_drop.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_arrive_expect_tx.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_arrive_no_complete.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_check_layout.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_complete_tx.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_expect_tx.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_init.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_inval.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_pending_count.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_test_wait.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_test_wait_parity.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_try_wait.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_try_wait_parity.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/multimem_ld_reduce.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/multimem_red.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/multimem_st.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/prefetch.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/prmt.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/red_async.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/setmaxnreg.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/shl.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/shr.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/st.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/st_async.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/st_bulk.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/stmatrix.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_alloc.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_commit.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_cp.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_fence.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_ld.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_mma.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_mma_sp.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_mma_ws.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_shift.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_st.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_wait.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/tensormap_cp_fenceproxy.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/tensormap_replace.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/trap.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/applypriority_async_bulk.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/clusterlaunchcontrol.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_async_bulk.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_async_bulk_multicast.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_async_bulk_prefetch.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_async_bulk_prefetch_tensor.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_async_bulk_tensor.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_async_bulk_tensor_gather_scatter.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_async_bulk_tensor_multicast.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_reduce_async_bulk_tensor.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/fabric_submit.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/fabric_try_get.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/fabric_try_pullred.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/fabric_try_put.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/fabric_try_red.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/fabric_wait.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/fence_proxy_alias.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/fence_proxy_async_generic_sync_restrict.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/fence_proxy_fabric_fabric_alias.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/fence_proxy_fabric_generic_alias.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/fence_proxy_generic_fabric_alias.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/ldmatrix.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_arrive.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_arrive_drop.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_check_layout.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_complete_tx.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_init.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_pending_count.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_test_wait.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_test_wait_parity.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_try_wait.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_try_wait_parity.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/prefetch.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/setmaxnreg.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/stmatrix.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_alloc_cta_group_1.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_alloc_cta_group_2.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_commit_cta_group_1.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_commit_cta_group_2.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_cp_cta_group_1.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_cp_cta_group_2.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_fence.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_ld.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_mma_cta_group_1.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_mma_cta_group_2.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_mma_sp.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_mma_ws.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_shift_cta_group_1.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_shift_cta_group_2.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_st.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_wait.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/tensormap_replace.h is excluded by !**/generated/**
📒 Files selected for processing (81)
  • docs/libcudacxx/ptx/instructions.rst
  • docs/libcudacxx/ptx/instructions/applypriority_async_bulk.rst
  • docs/libcudacxx/ptx/instructions/cp_async_bulk_prefetch.rst
  • docs/libcudacxx/ptx/instructions/cp_async_bulk_prefetch_tensor.rst
  • docs/libcudacxx/ptx/instructions/cp_async_mbarrier_arrive_noinc.rst
  • docs/libcudacxx/ptx/instructions/fabric_submit.rst
  • docs/libcudacxx/ptx/instructions/fabric_try_get.rst
  • docs/libcudacxx/ptx/instructions/fabric_try_pullred.rst
  • docs/libcudacxx/ptx/instructions/fabric_try_put.rst
  • docs/libcudacxx/ptx/instructions/fabric_try_red.rst
  • docs/libcudacxx/ptx/instructions/fabric_wait.rst
  • docs/libcudacxx/ptx/instructions/fence.rst
  • docs/libcudacxx/ptx/instructions/ldmatrix.rst
  • docs/libcudacxx/ptx/instructions/mapa.rst
  • docs/libcudacxx/ptx/instructions/mbarrier_arrive.rst
  • docs/libcudacxx/ptx/instructions/mbarrier_check_layout.rst
  • docs/libcudacxx/ptx/instructions/mbarrier_complete_tx.rst
  • docs/libcudacxx/ptx/instructions/mbarrier_pending_count.rst
  • docs/libcudacxx/ptx/instructions/mbarrier_test_wait.rst
  • docs/libcudacxx/ptx/instructions/mbarrier_try_wait.rst
  • docs/libcudacxx/ptx/instructions/mbarrier_wait.rst
  • docs/libcudacxx/ptx/instructions/prefetch.rst
  • docs/libcudacxx/ptx/instructions/stmatrix.rst
  • docs/libcudacxx/ptx/instructions/tcgen05_mma.rst
  • libcudacxx/codegen/add_ptx_instruction.py
  • libcudacxx/include/cuda/__device/arch_traits.h
  • libcudacxx/include/cuda/__ptx/instructions/applypriority_async_bulk.h
  • libcudacxx/include/cuda/__ptx/instructions/cp_async_bulk_prefetch.h
  • libcudacxx/include/cuda/__ptx/instructions/cp_async_bulk_prefetch_tensor.h
  • libcudacxx/include/cuda/__ptx/instructions/cp_async_mbarrier_arrive_noinc.h
  • libcudacxx/include/cuda/__ptx/instructions/fabric_submit.h
  • libcudacxx/include/cuda/__ptx/instructions/fabric_try_get.h
  • libcudacxx/include/cuda/__ptx/instructions/fabric_try_pullred.h
  • libcudacxx/include/cuda/__ptx/instructions/fabric_try_put.h
  • libcudacxx/include/cuda/__ptx/instructions/fabric_try_red.h
  • libcudacxx/include/cuda/__ptx/instructions/fabric_wait.h
  • libcudacxx/include/cuda/__ptx/instructions/fence.h
  • libcudacxx/include/cuda/__ptx/instructions/ldmatrix.h
  • libcudacxx/include/cuda/__ptx/instructions/mbarrier_arrive.h
  • libcudacxx/include/cuda/__ptx/instructions/mbarrier_check_layout.h
  • libcudacxx/include/cuda/__ptx/instructions/mbarrier_complete_tx.h
  • libcudacxx/include/cuda/__ptx/instructions/mbarrier_pending_count.h
  • libcudacxx/include/cuda/__ptx/instructions/prefetch.h
  • libcudacxx/include/cuda/__ptx/instructions/stmatrix.h
  • libcudacxx/include/cuda/__ptx/instructions/tcgen05_mma.h
  • libcudacxx/include/cuda/__ptx/ptx_dot_variants.h
  • libcudacxx/include/cuda/ptx
  • libcudacxx/include/cuda/std/__cccl/execution_space.h
  • libcudacxx/include/cuda/std/__cccl/preprocessor.h
  • libcudacxx/include/cuda/std/__cccl/ptx_isa.h
  • libcudacxx/include/nv/detail/__target_macros
  • libcudacxx/test/libcudacxx/cuda/ccclrt/device/all_arch_ids.pass.cpp
  • libcudacxx/test/libcudacxx/cuda/ccclrt/device/arch_id.pass.cpp
  • libcudacxx/test/libcudacxx/cuda/ccclrt/device/arch_id_fmt.pass.cpp
  • libcudacxx/test/libcudacxx/cuda/ccclrt/device/arch_traits.c2h.cu
  • libcudacxx/test/libcudacxx/cuda/ccclrt/device/compute_capability.pass.cpp
  • libcudacxx/test/libcudacxx/cuda/ccclrt/device/compute_capability_fmt.pass.cpp
  • libcudacxx/test/libcudacxx/cuda/ccclrt/device/is_specific_arch.pass.cpp
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.applypriority.async.bulk.compile.pass.cpp
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.async.bulk.prefetch.compile.pass.cpp
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.async.bulk.prefetch.tensor.compile.pass.cpp
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.submit.compile.pass.cpp
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try.get.compile.pass.cpp
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try.pullred.compile.pass.cpp
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try.put.compile.pass.cpp
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try.red.compile.pass.cpp
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.wait.compile.pass.cpp
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.fence.compile.pass.cpp
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.ldmatrix.compile.pass.cpp
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.mbarrier.arrive.compile.pass.cpp
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.mbarrier.check.layout.compile.pass.cpp
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.mbarrier.complete.tx.compile.pass.cpp
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.mbarrier.pending.count.compile.pass.cpp
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.prefetch.compile.pass.cpp
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.stmatrix.compile.pass.cpp
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.tcgen05.mma.sp.compile.pass.cpp
  • libcudacxx/test/nvtarget/arch_specific/CMakeLists.txt
  • libcudacxx/test/nvtarget/arch_specific/arch_specific.cu
  • libcudacxx/test/nvtarget/family_specific/CMakeLists.txt
  • libcudacxx/test/nvtarget/family_specific/family_specific.cu
  • libcudacxx/test/utils/libcudacxx/test/format.py
💤 Files with no reviewable changes (2)
  • docs/libcudacxx/ptx/instructions/mbarrier_test_wait.rst
  • docs/libcudacxx/ptx/instructions/mbarrier_try_wait.rst

Included review availability: Your plan provides up to 12 included reviews per hour; 11 remain after this review.

Comment on lines +12 to +13
#ifndef _CUDA_PTX_LDMATRIX_H_
#define _CUDA_PTX_LDMATRIX_H_

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.

📐 Maintainability & Code Quality | 🟡 Minor | ⚡ Quick win

important: Use full-path-derived include guards in the new PTX headers.

These headers use basename guards instead of guards derived from the uppercase full path. Rename each guard to comply with the shared header contract and reduce collision risk.

  • libcudacxx/include/cuda/__ptx/instructions/ldmatrix.h#L12-L13: replace _CUDA_PTX_LDMATRIX_H_.
  • libcudacxx/include/cuda/__ptx/instructions/mbarrier_check_layout.h#L12-L13: replace _CUDA_PTX_MBARRIER_CHECK_LAYOUT_H_.
  • libcudacxx/include/cuda/__ptx/instructions/mbarrier_complete_tx.h#L12-L13: replace _CUDA_PTX_MBARRIER_COMPLETE_TX_H_.
  • libcudacxx/include/cuda/__ptx/instructions/mbarrier_pending_count.h#L12-L13: replace _CUDA_PTX_MBARRIER_PENDING_COUNT_H_.
  • libcudacxx/include/cuda/__ptx/instructions/prefetch.h#L12-L13: replace _CUDA_PTX_PREFETCH_H_.

As per coding guidelines, “Headers must use include guards derived from the uppercase full path.”

📍 Affects 5 files
  • libcudacxx/include/cuda/__ptx/instructions/ldmatrix.h#L12-L13 (this comment)
  • libcudacxx/include/cuda/__ptx/instructions/mbarrier_check_layout.h#L12-L13
  • libcudacxx/include/cuda/__ptx/instructions/mbarrier_complete_tx.h#L12-L13
  • libcudacxx/include/cuda/__ptx/instructions/mbarrier_pending_count.h#L12-L13
  • libcudacxx/include/cuda/__ptx/instructions/prefetch.h#L12-L13

Source: Coding guidelines

// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
//
//===----------------------------------------------------------------------===//
// UNSUPPORTED: libcpp-has-no-threads

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.

📐 Maintainability & Code Quality | 🟡 Minor | ⚡ Quick win

important: Motivate or remove the repeated thread-support skip. The visible tests only include PTX headers and return from main, so no thread dependency is shown. These directives can skip valid PTX compile coverage. If the harness requires threads, state that requirement in a nearby comment. Otherwise remove the directives. As per coding guidelines, unsupported tests must be motivated with a comment.

  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.submit.compile.pass.cpp#L10-L10: explain or remove the gate.
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try.get.compile.pass.cpp#L10-L10: explain or remove the gate.
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try.pullred.compile.pass.cpp#L10-L10: explain or remove the gate.
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try.put.compile.pass.cpp#L10-L10: explain or remove the gate.
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try.red.compile.pass.cpp#L10-L10: explain or remove the gate.
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.wait.compile.pass.cpp#L10-L10: explain or remove the gate.
📍 Affects 6 files
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.submit.compile.pass.cpp#L10-L10 (this comment)
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try.get.compile.pass.cpp#L10-L10
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try.pullred.compile.pass.cpp#L10-L10
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try.put.compile.pass.cpp#L10-L10
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try.red.compile.pass.cpp#L10-L10
  • libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.wait.compile.pass.cpp#L10-L10

Source: Coding guidelines

@github-actions

Copy link
Copy Markdown
Contributor

⏱️ CCCL compile-time benchmark comparison: Public headers compile-time bench

Result: 1 regression row(s), 9 improvement row(s) above threshold.

Run Value
Config public-headers-gcc13
Baseline origin/main
Preset all-dev
Targets cub.headers.base, thrust.cpp.cuda.headers.base, libcudacxx.test.public_headers
GPU / launch args rtx2080 / --cuda 13.3 --host gcc13

Artifacts: reports and traces

Direct file processing

-f file-processing exclusive --sort total

🔴 Direct file processing — Regressions
Rank Regression impact Selected Δ Baseline Current Event Matched traces
1 2.009659 2.009659 0.175863 2.185522 Processing Header File: cub/cub/device/dispatch/tuning/common.cuh 131
🟢 Direct file processing — Improvements
Rank Improvement impact Selected Δ Baseline Current Event Matched traces
1 9.095768 -9.095768 19.056045 9.960277 Processing Header File: libcudacxx/include/cuda/__device/physical_device.h 93
2 1.221327 -1.221327 5.462829 4.241502 Processing Header File: libcudacxx/include/cuda/std/__cccl/prologue.h 547
3 0.336481 -0.336481 5.671880 5.335399 Processing Header File: libcudacxx/include/cuda/std/__iterator/concepts.h 436
4 0.323122 -0.323122 0.584637 0.261515 Processing Header File: cub/cub/device/dispatch/dispatch_batched_topk.cuh 3
5 0.300755 -0.300755 0.957104 0.656349 Processing Header File: libcudacxx/include/cuda/std/__functional/operations.h 278
6 0.300300 -0.300300 0.301010 0.000710 Processing Header File: cub/cub/device/device_run_length_encode.cuh 2
7 0.256844 -0.256844 0.936581 0.679737 Processing Header File: libcudacxx/include/cuda/__ptx/ptx_dot_variants.h 416
8 0.220731 -0.220731 1.711257 1.490526 Processing Header File: libcudacxx/include/cuda/std/__cccl/epilogue.h 547
9 0.210733 -0.210733 1.549608 1.338875 Processing Header File: libcudacxx/include/cuda/std/__atomic/functions/cuda_ptx_generated.h 56

@github-actions

Copy link
Copy Markdown
Contributor

🥳 CI Workflow Results

🟩 Finished in 4h 45m: Pass: 100%/115 | Total: 6d 22h | Max: 4h 45m | Hits: 37%/2764949

See results here.

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

Labels

None yet

Projects

Status: In Review

Development

Successfully merging this pull request may close these issues.

2 participants