[libcu++] Update cuda::ptx:: for CUDA 13.5 - #10909
Open
pciolkosz wants to merge 1 commit into
Open
Conversation
Contributor
|
Important Review skippedReview was skipped as selected files did not have any reviewable changes. 💤 Files selected but had no reviewable changes (1)
⛔ Files ignored due to path filters (141)
⚙️ Run configurationConfiguration used: Path: .coderabbit.yaml Review profile: CHILL Plan: Enterprise Run ID: ⛔ Files ignored due to path filters (141)
📒 Files selected for processing (1)
You can disable this status message by setting the Use the checkbox below for a quick retry:
📝 WalkthroughSummary by CodeRabbit
WalkthroughChangesThe PR adds PTX instruction wrapper headers, variant enums and selectors, umbrella-header wiring, generated-content documentation pages, availability updates, and compile-pass tests for asynchronous bulk, Fabric, matrix, mbarrier, prefetch, and related instructions. PTX instruction support
Possibly related PRs
Suggested reviewers: ✨ Finishing Touches 💡 1🛠️ Fix failing CI checks 💡
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: 4
🧹 Nitpick comments (3)
docs/libcudacxx/ptx/instructions.rst (1)
299-301: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Replace generic
Yesvalues with the exact CCCL and CUDA versions used by the surrounding table.Yesdoes not tell users which release provides each API. Verify the release metadata before merging. As per path instructions, documentation changes must cover API/version consistency.Also applies to: 315-315, 336-336, 342-342, 357-357, 359-359, 361-361, 363-363, 365-365, 367-367, 468-468, 472-472, 482-482, 484-484, 507-507, 509-509, 551-551
Source: Path instructions
libcudacxx/include/cuda/__ptx/instructions/fabric_try_get.h (1)
17-23: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Use one consistent closing annotation for the system-header preprocessor chain. The current
#endif // no system headerdoes not repeat the exact controlling condition text.
libcudacxx/include/cuda/__ptx/instructions/fabric_try_get.h#L17-L23: remove or correct the Line 23 annotation.libcudacxx/include/cuda/__ptx/instructions/fabric_try_pullred.h#L17-L23: remove or correct the Line 23 annotation.libcudacxx/include/cuda/__ptx/instructions/fabric_try_put.h#L17-L23: remove or correct the Line 23 annotation.libcudacxx/include/cuda/__ptx/instructions/fabric_try_red.h#L17-L23: remove or correct the Line 23 annotation.libcudacxx/include/cuda/__ptx/instructions/fabric_wait.h#L17-L23: remove or correct the Line 23 annotation.libcudacxx/include/cuda/__ptx/instructions/ldmatrix.h#L17-L23: remove or correct the Line 23 annotation.libcudacxx/include/cuda/__ptx/instructions/ldmatrix_m16n16_trans.h#L17-L23: remove or correct the Line 23 annotation.libcudacxx/include/cuda/__ptx/instructions/mapa.h#L17-L23: remove or correct the Line 23 annotation.libcudacxx/include/cuda/__ptx/instructions/mbarrier_check_layout.h#L17-L23: remove or correct the Line 23 annotation.
Based on learnings, annotated#endifcomments must repeat the exact condition text.Source: Learnings
libcudacxx/include/cuda/__ptx/instructions/fabric_try_pullred.h (1)
31-42: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winimportant: Move the half and bfloat16 forward declarations after
<cuda/std/__cccl/prologue.h>. The current order places code before the required prologue include.
libcudacxx/include/cuda/__ptx/instructions/fabric_try_pullred.h#L31-L42: move the forward declarations below the prologue include and keep them outside_CCCL_BEGIN_NAMESPACE_CUDA_PTX.libcudacxx/include/cuda/__ptx/instructions/fabric_try_red.h#L31-L42: move the forward declarations below the prologue include and keep them outside_CCCL_BEGIN_NAMESPACE_CUDA_PTX.
As per coding guidelines, the last header included before code must be<cuda/std/__cccl/prologue.h>.Source: Coding guidelines
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 0414e4f9-c00b-4078-9447-e0ac4202eab5
⛔ Files ignored due to path filters (226)
docs/libcudacxx/ptx/instructions/generated/applypriority_async_bulk.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/bfind.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/clusterlaunchcontrol.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/cp_async_bulk.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/cp_async_bulk_multicast.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/cp_async_bulk_prefetch.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/cp_async_bulk_prefetch_tensor.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/cp_async_bulk_tensor.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/cp_async_bulk_tensor_gather_scatter.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/cp_async_bulk_tensor_multicast.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/cp_reduce_async_bulk_tensor.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/fabric_submit.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/fabric_try_get.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/fabric_try_pullred.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/fabric_try_put.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/fabric_try_red.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/fabric_wait.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/fence_proxy_alias.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/fence_proxy_async_generic_sync_restrict.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/fence_proxy_fabric_fabric_alias.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/fence_proxy_fabric_generic_alias.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/fence_proxy_generic_fabric_alias.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/ldmatrix.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/ldmatrix_m16n16_trans.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/mapa.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/mbarrier_arrive.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/mbarrier_arrive_drop.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/mbarrier_check_layout.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/mbarrier_complete_tx.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/mbarrier_init.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/mbarrier_pending_count.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/mbarrier_test_wait.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/mbarrier_test_wait_parity.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/mbarrier_try_wait.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/mbarrier_try_wait_parity.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/prefetch.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/setmaxnreg.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/shr.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/stmatrix.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/tcgen05_alloc.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/tcgen05_commit.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/tcgen05_cp.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/tcgen05_fence.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/tcgen05_ld.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/tcgen05_mma.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/tcgen05_mma_sp.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/tcgen05_mma_ws.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/tcgen05_shift.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/tcgen05_st.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/tcgen05_wait.rstis excluded by!**/generated/**docs/libcudacxx/ptx/instructions/generated/tensormap_replace.rstis excluded by!**/generated/**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/ldmatrix_m16n16_trans.his excluded by!**/generated/**libcudacxx/include/cuda/__ptx/instructions/generated/mapa.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/barrier_cluster.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/bfind.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/bmsk.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_commit_group.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_async_bulk_wait_group.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_async_mbarrier_arrive.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_async_mbarrier_arrive_noinc.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_reduce_async_bulk.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_reduce_async_bulk_bf16.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_reduce_async_bulk_f16.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/elect_sync.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/exit.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.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/fence_mbarrier_init.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.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/fence_proxy_tensormap_generic.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/fence_sync_restrict.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/get_sreg.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/getctarank.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/ld.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/ldmatrix.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/ldmatrix_m16n16_trans.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/mapa.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_arrive_expect_tx.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_arrive_no_complete.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_expect_tx.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_init.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_inval.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/multimem_ld_reduce.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/multimem_red.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/multimem_st.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/prefetch.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/prmt.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/red_async.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/setmaxnreg.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/shl.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/shr.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/st.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/st_async.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/st_bulk.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/stmatrix.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_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_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_cta_group_1.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_cta_group_1.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_cp_fenceproxy.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/tensormap_replace.his excluded by!**/generated/**libcudacxx/test/libcudacxx/cuda/ptx/generated/trap.his excluded by!**/generated/**
📒 Files selected for processing (63)
docs/libcudacxx/ptx/instructions.rstdocs/libcudacxx/ptx/instructions/applypriority_async_bulk.rstdocs/libcudacxx/ptx/instructions/cp_async_bulk_prefetch.rstdocs/libcudacxx/ptx/instructions/cp_async_bulk_prefetch_tensor.rstdocs/libcudacxx/ptx/instructions/fabric_submit.rstdocs/libcudacxx/ptx/instructions/fabric_try_get.rstdocs/libcudacxx/ptx/instructions/fabric_try_pullred.rstdocs/libcudacxx/ptx/instructions/fabric_try_put.rstdocs/libcudacxx/ptx/instructions/fabric_try_red.rstdocs/libcudacxx/ptx/instructions/fabric_wait.rstdocs/libcudacxx/ptx/instructions/fence.rstdocs/libcudacxx/ptx/instructions/ldmatrix.rstdocs/libcudacxx/ptx/instructions/ldmatrix_m16n16_trans.rstdocs/libcudacxx/ptx/instructions/mapa.rstdocs/libcudacxx/ptx/instructions/mbarrier_arrive.rstdocs/libcudacxx/ptx/instructions/mbarrier_check_layout.rstdocs/libcudacxx/ptx/instructions/mbarrier_complete_tx.rstdocs/libcudacxx/ptx/instructions/mbarrier_pending_count.rstdocs/libcudacxx/ptx/instructions/prefetch.rstdocs/libcudacxx/ptx/instructions/stmatrix.rstdocs/libcudacxx/ptx/instructions/tcgen05_mma.rstlibcudacxx/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/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/ldmatrix_m16n16_trans.hlibcudacxx/include/cuda/__ptx/instructions/mapa.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/test/libcudacxx/cuda/ptx/ptx.applypriority.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.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.ldmatrix.m16n16.trans.compile.pass.cpplibcudacxx/test/libcudacxx/cuda/ptx/ptx.mapa.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.mma.sp.compile.pass.cpp
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_APPLYPRIORITY_ASYNC_BULK_H_ | ||
| #define _CUDA_PTX_APPLYPRIORITY_ASYNC_BULK_H_ |
Contributor
There was a problem hiding this comment.
📐 Maintainability & Code Quality | 🟠 Major | ⚡ Quick win
important: Apply one full-path include-guard convention to all new instruction headers.
libcudacxx/include/cuda/__ptx/instructions/applypriority_async_bulk.h#L12-L13: include the__ptx/instructionspath in the guard.libcudacxx/include/cuda/__ptx/instructions/cp_async_bulk_prefetch.h#L12-L13: include the__ptx/instructionspath in the guard.libcudacxx/include/cuda/__ptx/instructions/cp_async_bulk_prefetch_tensor.h#L12-L13: include the__ptx/instructionspath in the guard.libcudacxx/include/cuda/__ptx/instructions/fabric_submit.h#L12-L13: include the__ptx/instructionspath in the guard.libcudacxx/include/cuda/__ptx/instructions/mbarrier_complete_tx.h#L12-L13: include the__ptx/instructionspath in the guard.libcudacxx/include/cuda/__ptx/instructions/mbarrier_pending_count.h#L12-L13: include the__ptx/instructionspath in the guard.libcudacxx/include/cuda/__ptx/instructions/prefetch.h#L12-L13: include the__ptx/instructionspath in the guard.libcudacxx/include/cuda/__ptx/instructions/stmatrix.h#L12-L13: include the__ptx/instructionspath in the guard.
As per coding guidelines, “Headers must use include guards derived from the uppercase full path.”
📍 Affects 8 files
libcudacxx/include/cuda/__ptx/instructions/applypriority_async_bulk.h#L12-L13(this comment)libcudacxx/include/cuda/__ptx/instructions/cp_async_bulk_prefetch.h#L12-L13libcudacxx/include/cuda/__ptx/instructions/cp_async_bulk_prefetch_tensor.h#L12-L13libcudacxx/include/cuda/__ptx/instructions/fabric_submit.h#L12-L13libcudacxx/include/cuda/__ptx/instructions/mbarrier_complete_tx.h#L12-L13libcudacxx/include/cuda/__ptx/instructions/mbarrier_pending_count.h#L12-L13libcudacxx/include/cuda/__ptx/instructions/prefetch.h#L12-L13libcudacxx/include/cuda/__ptx/instructions/stmatrix.h#L12-L13
Source: Coding guidelines
Comment on lines
+251
to
+252
| [[maybe_unused]] static constexpr mbarrier_phase_primary_t mbarrier_phase_primary{}; | ||
| [[maybe_unused]] static constexpr mbarrier_phase_conditional_t mbarrier_phase_conditional{}; |
Contributor
There was a problem hiding this comment.
📐 Maintainability & Code Quality | 🟠 Major | ⚡ Quick win
important: Use inline constexpr for all namespace-scope tag objects in this header. The declarations around lines 251–283 currently use static constexpr; update the mbarrier_phase_*, layout_*, and mbarrier_report_* objects to match libcudacxx conventions.
📍 Affects 1 file
libcudacxx/include/cuda/__ptx/ptx_dot_variants.h#L251-L252(this comment)libcudacxx/include/cuda/__ptx/ptx_dot_variants.h#L135-L155
Source: Coding guidelines
| // SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. | ||
| // | ||
| //===----------------------------------------------------------------------===// | ||
| // UNSUPPORTED: libcpp-has-no-threads |
Contributor
There was a problem hiding this comment.
📐 Maintainability & Code Quality | 🟡 Minor | ⚡ Quick win
important: Document the reason for the common UNSUPPORTED: libcpp-has-no-threads directive in every test. The libcudacxx test guidance requires a motivation for unsupported platforms.
libcudacxx/test/libcudacxx/cuda/ptx/ptx.applypriority.async.bulk.compile.pass.cpp#L10-L10: add the rationale comment.libcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.async.bulk.prefetch.compile.pass.cpp#L10-L10: add the rationale comment.libcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.async.bulk.prefetch.tensor.compile.pass.cpp#L10-L10: add the rationale comment.libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.submit.compile.pass.cpp#L10-L10: add the rationale comment.libcudacxx/test/libcudacxx/cuda/ptx/ptx.mbarrier.check_layout.compile.pass.cpp#L10-L10: add the rationale comment.libcudacxx/test/libcudacxx/cuda/ptx/ptx.mbarrier.complete_tx.compile.pass.cpp#L10-L10: add the rationale comment.libcudacxx/test/libcudacxx/cuda/ptx/ptx.mbarrier.pending_count.compile.pass.cpp#L10-L10: add the rationale comment.libcudacxx/test/libcudacxx/cuda/ptx/ptx.prefetch.compile.pass.cpp#L10-L10: add the rationale comment.libcudacxx/test/libcudacxx/cuda/ptx/ptx.stmatrix.compile.pass.cpp#L10-L10: add the rationale comment.libcudacxx/test/libcudacxx/cuda/ptx/ptx.tcgen05.mma.sp.compile.pass.cpp#L10-L10: add the rationale comment.
📍 Affects 10 files
libcudacxx/test/libcudacxx/cuda/ptx/ptx.applypriority.async.bulk.compile.pass.cpp#L10-L10(this comment)libcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.async.bulk.prefetch.compile.pass.cpp#L10-L10libcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.async.bulk.prefetch.tensor.compile.pass.cpp#L10-L10libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.submit.compile.pass.cpp#L10-L10libcudacxx/test/libcudacxx/cuda/ptx/ptx.mbarrier.check_layout.compile.pass.cpp#L10-L10libcudacxx/test/libcudacxx/cuda/ptx/ptx.mbarrier.complete_tx.compile.pass.cpp#L10-L10libcudacxx/test/libcudacxx/cuda/ptx/ptx.mbarrier.pending_count.compile.pass.cpp#L10-L10libcudacxx/test/libcudacxx/cuda/ptx/ptx.prefetch.compile.pass.cpp#L10-L10libcudacxx/test/libcudacxx/cuda/ptx/ptx.stmatrix.compile.pass.cpp#L10-L10libcudacxx/test/libcudacxx/cuda/ptx/ptx.tcgen05.mma.sp.compile.pass.cpp#L10-L10
Source: Coding guidelines
| // SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. | ||
| // | ||
| //===----------------------------------------------------------------------===// | ||
| // UNSUPPORTED: libcpp-has-no-threads |
Contributor
There was a problem hiding this comment.
📐 Maintainability & Code Quality | 🟡 Minor | ⚡ Quick win
suggestion: Document or remove the repeated UNSUPPORTED: libcpp-has-no-threads directive. The changed tests contain no direct thread-dependent code. If the harness requires threads, state that dependency; otherwise remove the directive to preserve compile coverage on no-thread configurations. As per coding guidelines, unsupported or skipped tests must state why they are skipped.
- libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try_get.compile.pass.cpp#L10-L10: Add the skip rationale or remove the directive.
- libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try_pullred.compile.pass.cpp#L10-L10: Add the skip rationale or remove the directive.
- libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try_put.compile.pass.cpp#L10-L10: Add the skip rationale or remove the directive.
- libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try_red.compile.pass.cpp#L10-L10: Add the skip rationale or remove the directive.
- libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.wait.compile.pass.cpp#L10-L10: Add the skip rationale or remove the directive.
- libcudacxx/test/libcudacxx/cuda/ptx/ptx.ldmatrix.compile.pass.cpp#L10-L10: Add the skip rationale or remove the directive.
- libcudacxx/test/libcudacxx/cuda/ptx/ptx.ldmatrix.m16n16.trans.compile.pass.cpp#L10-L10: Add the skip rationale or remove the directive.
- libcudacxx/test/libcudacxx/cuda/ptx/ptx.mapa.compile.pass.cpp#L10-L10: Add the skip rationale or remove the directive.
📍 Affects 8 files
libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try_get.compile.pass.cpp#L10-L10(this comment)libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try_pullred.compile.pass.cpp#L10-L10libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try_put.compile.pass.cpp#L10-L10libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.try_red.compile.pass.cpp#L10-L10libcudacxx/test/libcudacxx/cuda/ptx/ptx.fabric.wait.compile.pass.cpp#L10-L10libcudacxx/test/libcudacxx/cuda/ptx/ptx.ldmatrix.compile.pass.cpp#L10-L10libcudacxx/test/libcudacxx/cuda/ptx/ptx.ldmatrix.m16n16.trans.compile.pass.cpp#L10-L10libcudacxx/test/libcudacxx/cuda/ptx/ptx.mapa.compile.pass.cpp#L10-L10
Source: Coding guidelines
pciolkosz
force-pushed
the
ptx-sync-libcuda-ptx
branch
from
August 20, 2026 01:33
b0a306d to
4fcb94b
Compare
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.
This PR updates cuda/ptx in main with 13.4 updates plus a few extra additions (
ldmatrix.m16n16.transandmapavariants)