[CI Only][Do not review][Do not merge] Test cuda::ptx CTK-13.4#10002
Open
fbusato wants to merge 1 commit into
Open
[CI Only][Do not review][Do not merge] Test cuda::ptx CTK-13.4#10002fbusato wants to merge 1 commit into
cuda::ptx CTK-13.4#10002fbusato wants to merge 1 commit into
Conversation
Contributor
📝 WalkthroughSummary by CodeRabbit
WalkthroughAdds CUDA PTX instruction wrappers and exports, dot-variant types, SM_107 target support, compile-pass coverage for new generated instructions, and compile-only execution handling for CUDA PTX tests. ChangesCUDA PTX support
Suggested reviewers: Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
Actionable comments posted: 2
🧹 Nitpick comments (6)
libcudacxx/include/nv/detail/__target_macros (1)
39-39: 🎯 Functional Correctness | 🔵 Trivial | ⚡ Quick winsuggestion: Add focused compile-time coverage for the complete SM_107 dispatch chain.
The change duplicates selector, integer, exact/provides, and SM_107a/SM_107f feature mappings. Assert the new predicates under representative host, device, architecture-specific, and family-specific configurations so a table typo cannot silently omit downstream PTX registrations.
As per coding guidelines, validate changes with relevant targeted builds and tests. The supplied downstream wrapper consumes
NV_HAS_FEATURE_SM_107aandNV_HAS_FEATURE_SM_107f.Also applies to: 63-63, 90-90, 131-131, 184-184, 220-220, 244-244, 302-309, 446-451, 506-514, 571-579, 734-739
Source: Coding guidelines
libcudacxx/include/cuda/__ptx/instructions/fabric_submit.h (2)
12-14: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Apply the full-path include-guard rule consistently.
All ten guards omit the
__ptx/instructionspath components. Rename each guard to encode the full path and update its closing#endifcomment.
libcudacxx/include/cuda/__ptx/instructions/fabric_submit.h#L12-L14: update_CUDA_PTX_FABRIC_SUBMIT_H_.libcudacxx/include/cuda/__ptx/instructions/fabric_try_get.h#L12-L14: update_CUDA_PTX_FABRIC_TRY_GET_H_.libcudacxx/include/cuda/__ptx/instructions/fabric_try_pullred.h#L12-L14: update_CUDA_PTX_FABRIC_TRY_PULLRED_H_.libcudacxx/include/cuda/__ptx/instructions/fabric_try_put.h#L12-L14: update_CUDA_PTX_FABRIC_TRY_PUT_H_.libcudacxx/include/cuda/__ptx/instructions/fabric_try_red.h#L12-L14: update_CUDA_PTX_FABRIC_TRY_RED_H_.libcudacxx/include/cuda/__ptx/instructions/fabric_wait.h#L12-L14: update_CUDA_PTX_FABRIC_WAIT_H_.libcudacxx/include/cuda/__ptx/instructions/ldmatrix.h#L12-L14: update_CUDA_PTX_LDMATRIX_H_.libcudacxx/include/cuda/__ptx/instructions/mbarrier_check_layout.h#L12-L14: update_CUDA_PTX_MBARRIER_CHECK_LAYOUT_H_.libcudacxx/include/cuda/__ptx/instructions/mbarrier_complete_tx.h#L12-L14: update_CUDA_PTX_MBARRIER_COMPLETE_TX_H_.libcudacxx/include/cuda/__ptx/instructions/mbarrier_pending_count.h#L12-L14: update_CUDA_PTX_MBARRIER_PENDING_COUNT_H_.As per coding guidelines, headers must use include guards derived from the uppercase full path.
Source: Coding guidelines
17-23: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Make the preprocessor closing annotations verbatim.
Each
#endif // no system headershould either repeat the initial#ifcondition exactly or omit the comment.
libcudacxx/include/cuda/__ptx/instructions/fabric_submit.h#L17-L23: fix the system-header pragma annotation.libcudacxx/include/cuda/__ptx/instructions/fabric_try_get.h#L17-L23: fix the system-header pragma annotation.libcudacxx/include/cuda/__ptx/instructions/fabric_try_pullred.h#L17-L23: fix the system-header pragma annotation.libcudacxx/include/cuda/__ptx/instructions/fabric_try_put.h#L17-L23: fix the system-header pragma annotation.libcudacxx/include/cuda/__ptx/instructions/fabric_try_red.h#L17-L23: fix the system-header pragma annotation.libcudacxx/include/cuda/__ptx/instructions/fabric_wait.h#L17-L23: fix the system-header pragma annotation.libcudacxx/include/cuda/__ptx/instructions/ldmatrix.h#L17-L23: fix the system-header pragma annotation.libcudacxx/include/cuda/__ptx/instructions/mbarrier_check_layout.h#L17-L23: fix the system-header pragma annotation.libcudacxx/include/cuda/__ptx/instructions/mbarrier_complete_tx.h#L17-L23: fix the system-header pragma annotation.libcudacxx/include/cuda/__ptx/instructions/mbarrier_pending_count.h#L17-L23: fix the system-header pragma annotation.Based on learnings, annotated preprocessor branch comments must repeat the corresponding
#ifcondition verbatim.Source: Learnings
libcudacxx/test/libcudacxx/cuda/ptx/ptx.applypriority.async.bulk.compile.pass.cpp (1)
10-10: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Add a rationale comment for each
UNSUPPORTED: libcpp-has-no-threadsdirective. The test guidance requires skipped or unsupported tests to explain why they are unavailable.
libcudacxx/test/libcudacxx/cuda/ptx/ptx.applypriority.async.bulk.compile.pass.cpp#L10-L10: explain the thread dependency.libcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.async.bulk.prefetch.compile.pass.cpp#L10-L10: explain the thread dependency.libcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.async.bulk.prefetch.tensor.compile.pass.cpp#L10-L10: explain the thread dependency.libcudacxx/test/libcudacxx/cuda/ptx/ptx.prefetch.compile.pass.cpp#L10-L10: explain the thread dependency.libcudacxx/test/libcudacxx/cuda/ptx/ptx.stmatrix.compile.pass.cpp#L10-L10: explain the thread dependency.libcudacxx/test/libcudacxx/cuda/ptx/ptx.tcgen05.alloc.compile.pass.cpp#L11-L11: explain the thread dependency.libcudacxx/test/libcudacxx/cuda/ptx/ptx.tcgen05.commit.compile.pass.cpp#L11-L11: explain the thread dependency.libcudacxx/test/libcudacxx/cuda/ptx/ptx.tcgen05.cp.compile.pass.cpp#L11-L11: explain the thread dependency.As per path instructions, unsupported features must be motivated with a comment.
Source: Path instructions
libcudacxx/test/utils/libcudacxx/test/format.py (1)
274-282: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Add a regression fixture that cannot link.
The current PTX fixtures contain
main, so they would still pass if compile-only dispatch regressed to the normal compile-link-run path. Add a compile-pass case without a linkable entry point, or a focused harness test, to prove this branch compiles without linking or executing.libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try.red.compile.pass.cpp (1)
17-17: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: These generated-header includes use quotes, while the repository rule requires angle-bracket syntax. Apply the same include-style fix at each site, or document a scoped exception if quote lookup is intentional.
libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try.red.compile.pass.cpp#L17-L17: change the generatedfabric_try_redinclude.libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.wait.compile.pass.cpp#L17-L17: change the generatedfabric_waitinclude.libcudacxx/test/libcudacxx/cuda/ptx/ptx.fence.compile.pass.cpp#L23-L25: change all three generated fence includes.libcudacxx/test/libcudacxx/cuda/ptx/ptx.ldmatrix.compile.pass.cpp#L17-L17: change the generatedldmatrixinclude.libcudacxx/test/libcudacxx/cuda/ptx/ptx.mbarrier.arrive.compile.pass.cpp#L22-L22: change the generatedmbarrier_arrive_dropinclude.libcudacxx/test/libcudacxx/cuda/ptx/ptx.mbarrier.check.layout.compile.pass.cpp#L17-L17: change the generated layout include.libcudacxx/test/libcudacxx/cuda/ptx/ptx.mbarrier.complete.tx.compile.pass.cpp#L17-L17: change the generated transaction-completion include.libcudacxx/test/libcudacxx/cuda/ptx/ptx.mbarrier.pending.count.compile.pass.cpp#L17-L17: change the generated pending-count include.Source: Coding guidelines
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: adcf099f-a250-4b42-a8b1-631f6082096d
⛔ Files ignored due to path filters (138)
libcudacxx/include/cuda/__ptx/instructions/generated/applypriority_async_bulk.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/barrier_cluster.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/bfind.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/bmsk.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/clusterlaunchcontrol.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_commit_group.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_multicast.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_prefetch.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_prefetch_tensor.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_tensor.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_tensor_gather_scatter.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_tensor_multicast.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_wait_group.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_mbarrier_arrive.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_mbarrier_arrive_noinc.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/cp_reduce_async_bulk.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/cp_reduce_async_bulk_bf16.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/cp_reduce_async_bulk_f16.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/cp_reduce_async_bulk_tensor.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/elect_sync.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/exit.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/fabric_submit.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/fabric_try_get.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/fabric_try_pullred.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/fabric_try_put.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/fabric_try_red.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/fabric_wait.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/fence.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/fence_mbarrier_init.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/fence_proxy_alias.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/fence_proxy_async.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/fence_proxy_async_generic_sync_restrict.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/fence_proxy_fabric_fabric_alias.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/fence_proxy_fabric_generic_alias.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/fence_proxy_generic_fabric_alias.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/fence_proxy_tensormap_generic.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/fence_sync_restrict.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/get_sreg.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/getctarank.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/ld.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/ldmatrix.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_arrive.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_arrive_drop.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_arrive_expect_tx.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_arrive_no_complete.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_check_layout.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_complete_tx.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_expect_tx.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_init.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_inval.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_pending_count.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_test_wait.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_test_wait_parity.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_try_wait.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_try_wait_parity.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/multimem_ld_reduce.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/multimem_red.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/multimem_st.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/prefetch.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/prmt.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/red_async.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/setmaxnreg.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/shl.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/shr.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/st.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/st_async.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/st_bulk.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/stmatrix.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_alloc.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_commit.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_cp.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_fence.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_ld.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_mma.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_mma_sp.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_mma_ws.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_shift.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_st.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_wait.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/tensormap_cp_fenceproxy.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/tensormap_replace.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/trap.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/applypriority_async_bulk.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/clusterlaunchcontrol.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_async_bulk.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_async_bulk_multicast.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_async_bulk_prefetch.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_async_bulk_prefetch_tensor.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_async_bulk_tensor.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_async_bulk_tensor_gather_scatter.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_async_bulk_tensor_multicast.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_reduce_async_bulk_tensor.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/fabric_submit.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/fabric_try_get.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/fabric_try_pullred.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/fabric_try_put.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/fabric_try_red.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/fabric_wait.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/fence_proxy_alias.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/fence_proxy_async_generic_sync_restrict.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/fence_proxy_fabric_fabric_alias.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/fence_proxy_fabric_generic_alias.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/fence_proxy_generic_fabric_alias.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/ldmatrix.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_arrive.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_arrive_drop.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_check_layout.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_complete_tx.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_init.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_pending_count.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_test_wait.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_test_wait_parity.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_try_wait.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_try_wait_parity.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/prefetch.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/setmaxnreg.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/stmatrix.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_alloc.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_alloc_cta_group_1.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_alloc_cta_group_2.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_commit.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_commit_cta_group_1.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_commit_cta_group_2.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_cp.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_cp_cta_group_1.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_cp_cta_group_2.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_fence.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_ld.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_mma.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_mma_cta_group_2.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_mma_sp.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_mma_ws.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_shift.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_shift_cta_group_2.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_st.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/tcgen05_wait.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/tensormap_replace.his excluded by!**/generated/**
📒 Files selected for processing (54)
libcudacxx/include/cuda/__ptx/instructions/applypriority_async_bulk.hlibcudacxx/include/cuda/__ptx/instructions/cp_async_bulk_prefetch.hlibcudacxx/include/cuda/__ptx/instructions/cp_async_bulk_prefetch_tensor.hlibcudacxx/include/cuda/__ptx/instructions/cp_async_mbarrier_arrive_noinc.hlibcudacxx/include/cuda/__ptx/instructions/fabric_submit.hlibcudacxx/include/cuda/__ptx/instructions/fabric_try_get.hlibcudacxx/include/cuda/__ptx/instructions/fabric_try_pullred.hlibcudacxx/include/cuda/__ptx/instructions/fabric_try_put.hlibcudacxx/include/cuda/__ptx/instructions/fabric_try_red.hlibcudacxx/include/cuda/__ptx/instructions/fabric_wait.hlibcudacxx/include/cuda/__ptx/instructions/fence.hlibcudacxx/include/cuda/__ptx/instructions/ldmatrix.hlibcudacxx/include/cuda/__ptx/instructions/mbarrier_arrive.hlibcudacxx/include/cuda/__ptx/instructions/mbarrier_check_layout.hlibcudacxx/include/cuda/__ptx/instructions/mbarrier_complete_tx.hlibcudacxx/include/cuda/__ptx/instructions/mbarrier_pending_count.hlibcudacxx/include/cuda/__ptx/instructions/prefetch.hlibcudacxx/include/cuda/__ptx/instructions/stmatrix.hlibcudacxx/include/cuda/__ptx/instructions/tcgen05_mma.hlibcudacxx/include/cuda/__ptx/ptx_dot_variants.hlibcudacxx/include/cuda/ptxlibcudacxx/include/nv/detail/__target_macroslibcudacxx/include/nv/targetlibcudacxx/test/libcudacxx/cuda/ptx/ptx.applypriority.async.bulk.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.async.bulk.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.async.bulk.prefetch.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.async.bulk.prefetch.tensor.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.async.bulk.tensor.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.submit.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try.get.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try.pullred.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try.put.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try.red.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.wait.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.fence.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.ldmatrix.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.mbarrier.arrive.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.mbarrier.check.layout.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.mbarrier.complete.tx.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.mbarrier.pending.count.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.prefetch.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.stmatrix.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.tcgen05.alloc.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.tcgen05.alloc.cta_group_1.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.tcgen05.alloc.cta_group_2.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.tcgen05.commit.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.tcgen05.commit.cta_group_2.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.tcgen05.cp.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.tcgen05.mma.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.tcgen05.mma.cta_group_2.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.tcgen05.shift.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.tcgen05.shift.cta_group_1.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.tcgen05.shift.cta_group_2.compile.pass.cpplibcudacxx/test/utils/libcudacxx/test/format.py
Comment on lines
+245
to
+282
| template <dot_phase_type __phase> | ||
| using mbarrier_phase_t = ::cuda::std::integral_constant<dot_phase_type, __phase>; | ||
|
|
||
| using mbarrier_phase_primary_t = mbarrier_phase_t<dot_phase_type::primary>; | ||
| using mbarrier_phase_conditional_t = mbarrier_phase_t<dot_phase_type::conditional>; | ||
|
|
||
| [[maybe_unused]] static constexpr mbarrier_phase_primary_t mbarrier_phase_primary{}; | ||
| [[maybe_unused]] static constexpr mbarrier_phase_conditional_t mbarrier_phase_conditional{}; | ||
|
|
||
| template <dot_layout __layout> | ||
| using layout_t = ::cuda::std::integral_constant<dot_layout, __layout>; | ||
|
|
||
| using layout_v0_t = layout_t<dot_layout::v0>; | ||
| using layout_v1_t = layout_t<dot_layout::v1>; | ||
|
|
||
| [[maybe_unused]] static constexpr layout_v0_t layout_v0{}; | ||
| [[maybe_unused]] static constexpr layout_v1_t layout_v1{}; | ||
|
|
||
| template <int n> | ||
| using n32_t = ::cuda::std::integral_constant<int, n>; | ||
|
|
||
| template <dot_report_mechanism __report_mechanism> | ||
| using report_mechanism_t = ::cuda::std::integral_constant<dot_report_mechanism, __report_mechanism>; | ||
| using mbarrier_report_disabled_t = report_mechanism_t<dot_report_mechanism::disabled>; | ||
| using mbarrier_report_valid_per_16bytes_80000000_t = report_mechanism_t<dot_report_mechanism::per_16bytes_80000000>; | ||
| using mbarrier_report_valid_per_16bytes_8000_t = report_mechanism_t<dot_report_mechanism::per_16bytes_8000>; | ||
| using mbarrier_report_valid_per_16bytes_80_t = report_mechanism_t<dot_report_mechanism::per_16bytes_80>; | ||
| using mbarrier_report_valid_per_16bytes_8_t = report_mechanism_t<dot_report_mechanism::per_16bytes_8>; | ||
| using mbarrier_report_valid_per_element_ff_t = report_mechanism_t<dot_report_mechanism::per_element_ff>; | ||
|
|
||
| [[maybe_unused]] static constexpr mbarrier_report_disabled_t mbarrier_report_disabled{}; | ||
| [[maybe_unused]] static constexpr mbarrier_report_valid_per_16bytes_80000000_t | ||
| mbarrier_report_valid_per_16bytes_80000000{}; | ||
| [[maybe_unused]] static constexpr mbarrier_report_valid_per_16bytes_8000_t mbarrier_report_valid_per_16bytes_8000{}; | ||
| [[maybe_unused]] static constexpr mbarrier_report_valid_per_16bytes_80_t mbarrier_report_valid_per_16bytes_80{}; | ||
| [[maybe_unused]] static constexpr mbarrier_report_valid_per_16bytes_8_t mbarrier_report_valid_per_16bytes_8{}; | ||
| [[maybe_unused]] static constexpr mbarrier_report_valid_per_element_ff_t mbarrier_report_valid_per_element_ff{}; | ||
|
|
Contributor
There was a problem hiding this comment.
📐 Maintainability & Code Quality | 🟠 Major | ⚡ Quick win
suggestion: Replace the namespace-scope static constexpr instances with inline constexpr. This applies to the phase, layout, and report-mechanism objects at Line 251-252, Line 260-261, and Line 275-281.
As per coding guidelines, all namespace- or global-scope constexpr variables must also be inline.
Source: Coding guidelines
Comment on lines
+72
to
+110
| #include <cuda/__ptx/instructions/applypriority_async_bulk.h> | ||
| #include <cuda/__ptx/instructions/barrier_cluster.h> | ||
| #include <cuda/__ptx/instructions/bfind.h> | ||
| #include <cuda/__ptx/instructions/bmsk.h> | ||
| #include <cuda/__ptx/instructions/clusterlaunchcontrol.h> | ||
| #include <cuda/__ptx/instructions/cp_async_bulk.h> | ||
| #include <cuda/__ptx/instructions/cp_async_bulk_commit_group.h> | ||
| #include <cuda/__ptx/instructions/cp_async_bulk_prefetch.h> | ||
| #include <cuda/__ptx/instructions/cp_async_bulk_prefetch_tensor.h> | ||
| #include <cuda/__ptx/instructions/cp_async_bulk_tensor.h> | ||
| #include <cuda/__ptx/instructions/cp_async_bulk_wait_group.h> | ||
| #include <cuda/__ptx/instructions/cp_async_mbarrier_arrive.h> | ||
| #include <cuda/__ptx/instructions/cp_reduce_async_bulk.h> | ||
| #include <cuda/__ptx/instructions/cp_reduce_async_bulk_tensor.h> | ||
| #include <cuda/__ptx/instructions/elect_sync.h> | ||
| #include <cuda/__ptx/instructions/exit.h> | ||
| #include <cuda/__ptx/instructions/fabric_submit.h> | ||
| #include <cuda/__ptx/instructions/fabric_try_get.h> | ||
| #include <cuda/__ptx/instructions/fabric_try_pullred.h> | ||
| #include <cuda/__ptx/instructions/fabric_try_put.h> | ||
| #include <cuda/__ptx/instructions/fabric_try_red.h> | ||
| #include <cuda/__ptx/instructions/fabric_wait.h> | ||
| #include <cuda/__ptx/instructions/fence.h> | ||
| #include <cuda/__ptx/instructions/get_sreg.h> | ||
| #include <cuda/__ptx/instructions/getctarank.h> | ||
| #include <cuda/__ptx/instructions/ld.h> | ||
| #include <cuda/__ptx/instructions/ldmatrix.h> | ||
| #include <cuda/__ptx/instructions/mbarrier_arrive.h> | ||
| #include <cuda/__ptx/instructions/mbarrier_check_layout.h> | ||
| #include <cuda/__ptx/instructions/mbarrier_complete_tx.h> | ||
| #include <cuda/__ptx/instructions/mbarrier_expect_tx.h> | ||
| #include <cuda/__ptx/instructions/mbarrier_init.h> | ||
| #include <cuda/__ptx/instructions/mbarrier_inval.h> | ||
| #include <cuda/__ptx/instructions/mbarrier_pending_count.h> | ||
| #include <cuda/__ptx/instructions/mbarrier_wait.h> | ||
| #include <cuda/__ptx/instructions/multimem_ld_reduce.h> | ||
| #include <cuda/__ptx/instructions/multimem_red.h> | ||
| #include <cuda/__ptx/instructions/multimem_st.h> | ||
| #include <cuda/__ptx/instructions/prefetch.h> |
Contributor
There was a problem hiding this comment.
🎯 Functional Correctness | 🟠 Major | ⚡ Quick win
important: Export cp_async_mbarrier_arrive_noinc through <cuda/ptx>. The new wrapper at libcudacxx/include/cuda/__ptx/instructions/cp_async_mbarrier_arrive_noinc.h is not included in this aggregate list, so users must include the lower-level wrapper directly to access it. Add #include <cuda/__ptx/instructions/cp_async_mbarrier_arrive_noinc.h> next to the existing cp_async_mbarrier_arrive.h include.
Contributor
😬 CI Workflow Results🟥 Finished in 4h 42m: Pass: 96%/120 | Total: 7d 03h | Max: 4h 41m | Hits: 34%/3583257See results here. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Test PTX for CTK 13.4 preview https://docs.nvidia.com/cuda/developer-preview/13.4/cuda-toolkit-release-notes/index.html