---
id: pr-cccl-cub-3568
repo: NVIDIA/cccl
pr: 3568
title: 'PTX: Update generated files with Blackwell instructions '
author: bernhardmgruber
date: '2025-01-29'
url: https://github.com/NVIDIA/cccl/pull/3568
source_category: upstream-code
architectures:
- sm100
tags:
- blackwell
techniques: []
hardware_features:
- blackwell
kernel_types: []
languages:
- ptx
captured_at: '2026-05-16'
status: merged
merge_sha: d21e0c9804ad63d23950c8b0a2462e5b7ebc8701
inclusion_reason: Search-matched CUDA kernel optimization PR; materialize diff/source before synthesis.
changed_paths:
- docs/libcudacxx/ptx/instructions/generated/barrier_cluster_aligned.rst
- docs/libcudacxx/ptx/instructions/generated/clusterlaunchcontrol.rst
- docs/libcudacxx/ptx/instructions/generated/cp_async_bulk.rst
- docs/libcudacxx/ptx/instructions/generated/cp_async_bulk_multicast.rst
- docs/libcudacxx/ptx/instructions/generated/cp_async_bulk_tensor.rst
- docs/libcudacxx/ptx/instructions/generated/cp_async_bulk_tensor_gather_scatter.rst
- docs/libcudacxx/ptx/instructions/generated/cp_async_bulk_tensor_multicast.rst
- docs/libcudacxx/ptx/instructions/generated/cp_async_mbarrier_arrive.rst
- docs/libcudacxx/ptx/instructions/generated/cp_async_mbarrier_arrive_noinc.rst
- docs/libcudacxx/ptx/instructions/generated/cp_reduce_async_bulk.rst
- docs/libcudacxx/ptx/instructions/generated/elect_sync.rst
- docs/libcudacxx/ptx/instructions/generated/fence.rst
- docs/libcudacxx/ptx/instructions/generated/fence_proxy_async.rst
- docs/libcudacxx/ptx/instructions/generated/fence_proxy_async_generic_sync_restrict.rst
- docs/libcudacxx/ptx/instructions/generated/fence_sync_restrict.rst
- docs/libcudacxx/ptx/instructions/generated/get_sreg.rst
- docs/libcudacxx/ptx/instructions/generated/getctarank.rst
- docs/libcudacxx/ptx/instructions/generated/mapa.rst
- docs/libcudacxx/ptx/instructions/generated/mbarrier_arrive.rst
- docs/libcudacxx/ptx/instructions/generated/mbarrier_arrive_expect_tx.rst
- docs/libcudacxx/ptx/instructions/generated/mbarrier_test_wait.rst
- docs/libcudacxx/ptx/instructions/generated/mbarrier_test_wait_parity.rst
- docs/libcudacxx/ptx/instructions/generated/mbarrier_try_wait.rst
- docs/libcudacxx/ptx/instructions/generated/mbarrier_try_wait_parity.rst
- docs/libcudacxx/ptx/instructions/generated/multimem_ld_reduce.rst
- docs/libcudacxx/ptx/instructions/generated/multimem_red.rst
- docs/libcudacxx/ptx/instructions/generated/multimem_st.rst
- docs/libcudacxx/ptx/instructions/generated/red_async.rst
- docs/libcudacxx/ptx/instructions/generated/st_async.rst
- docs/libcudacxx/ptx/instructions/generated/st_bulk.rst
- docs/libcudacxx/ptx/instructions/generated/tcgen05_alloc.rst
- docs/libcudacxx/ptx/instructions/generated/tcgen05_commit.rst
- docs/libcudacxx/ptx/instructions/generated/tcgen05_cp.rst
- docs/libcudacxx/ptx/instructions/generated/tcgen05_fence.rst
- docs/libcudacxx/ptx/instructions/generated/tcgen05_ld.rst
- docs/libcudacxx/ptx/instructions/generated/tcgen05_mma.rst
- docs/libcudacxx/ptx/instructions/generated/tcgen05_mma_ws.rst
- docs/libcudacxx/ptx/instructions/generated/tcgen05_shift.rst
- docs/libcudacxx/ptx/instructions/generated/tcgen05_st.rst
- docs/libcudacxx/ptx/instructions/generated/tcgen05_wait.rst
- docs/libcudacxx/ptx/instructions/generated/tensormap_replace.rst
- docs/libcudacxx/ptx/instructions/special_registers.rst
- docs/repo.toml
- libcudacxx/include/cuda/__ptx/instructions/generated/barrier_cluster.h
- libcudacxx/include/cuda/__ptx/instructions/generated/barrier_cluster_aligned.h
- libcudacxx/include/cuda/__ptx/instructions/generated/clusterlaunchcontrol.h
- libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk.h
- libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_commit_group.h
- libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_multicast.h
- libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_tensor.h
- libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_tensor_gather_scatter.h
- libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_tensor_multicast.h
- libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_wait_group.h
- libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_mbarrier_arrive.h
- libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_mbarrier_arrive_noinc.h
- libcudacxx/include/cuda/__ptx/instructions/generated/cp_reduce_async_bulk.h
- libcudacxx/include/cuda/__ptx/instructions/generated/cp_reduce_async_bulk_bf16.h
- libcudacxx/include/cuda/__ptx/instructions/generated/cp_reduce_async_bulk_f16.h
- libcudacxx/include/cuda/__ptx/instructions/generated/cp_reduce_async_bulk_tensor.h
- libcudacxx/include/cuda/__ptx/instructions/generated/elect_sync.h
- libcudacxx/include/cuda/__ptx/instructions/generated/fence.h
- libcudacxx/include/cuda/__ptx/instructions/generated/fence_mbarrier_init.h
- libcudacxx/include/cuda/__ptx/instructions/generated/fence_proxy_alias.h
- libcudacxx/include/cuda/__ptx/instructions/generated/fence_proxy_async.h
- libcudacxx/include/cuda/__ptx/instructions/generated/fence_proxy_async_generic_sync_restrict.h
- libcudacxx/include/cuda/__ptx/instructions/generated/fence_proxy_tensormap_generic.h
- libcudacxx/include/cuda/__ptx/instructions/generated/fence_sync_restrict.h
- libcudacxx/include/cuda/__ptx/instructions/generated/get_sreg.h
- libcudacxx/include/cuda/__ptx/instructions/generated/getctarank.h
- libcudacxx/include/cuda/__ptx/instructions/generated/mapa.h
- libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_arrive.h
- libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_arrive_expect_tx.h
- libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_arrive_no_complete.h
- libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_expect_tx.h
- libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_init.h
- libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_test_wait.h
- libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_test_wait_parity.h
- libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_try_wait.h
- libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_try_wait_parity.h
- libcudacxx/include/cuda/__ptx/instructions/generated/multimem_ld_reduce.h
- libcudacxx/include/cuda/__ptx/instructions/generated/multimem_red.h
- libcudacxx/include/cuda/__ptx/instructions/generated/multimem_st.h
- libcudacxx/include/cuda/__ptx/instructions/generated/red_async.h
- libcudacxx/include/cuda/__ptx/instructions/generated/st_async.h
- libcudacxx/include/cuda/__ptx/instructions/generated/st_bulk.h
- libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_alloc.h
- libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_commit.h
- libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_cp.h
- libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_fence.h
- libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_ld.h
- libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_mma.h
- libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_mma_ws.h
- libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_shift.h
- libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_st.h
- libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_wait.h
- libcudacxx/include/cuda/__ptx/instructions/generated/tensormap_cp_fenceproxy.h
- libcudacxx/include/cuda/__ptx/instructions/generated/tensormap_replace.h
- libcudacxx/test/libcudacxx/cuda/ptx/generated/barrier_cluster_aligned.h
- libcudacxx/test/libcudacxx/cuda/ptx/generated/clusterlaunchcontrol.h
- libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_async_bulk.h
- libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_async_bulk_multicast.h
- libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_async_bulk_tensor.h
- libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_async_bulk_tensor_gather_scatter.h
- libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_async_bulk_tensor_multicast.h
- libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_async_mbarrier_arrive.h
- libcudacxx/test/libcudacxx/cuda/ptx/generated/cp_async_mbarrier_arrive_noinc.h
- libcudacxx/test/libcudacxx/cuda/ptx/generated/elect_sync.h
- libcudacxx/test/libcudacxx/cuda/ptx/generated/fence.h
- libcudacxx/test/libcudacxx/cuda/ptx/generated/fence_proxy_async_generic_sync_restrict.h
- libcudacxx/test/libcudacxx/cuda/ptx/generated/fence_sync_restrict.h
- libcudacxx/test/libcudacxx/cuda/ptx/generated/mapa.h
- libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_arrive.h
- libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_arrive_expect_tx.h
- libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_expect_tx.h
- libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_test_wait.h
- libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_test_wait_parity.h
- libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_try_wait.h
- libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_try_wait_parity.h
- libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_wait.h
- libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_wait_parity.h
artifact_dir: evidence/pull-bundles/cccl-cub/gh-3568
head_ref_oid: deebc024508aac9b20012ac1e972afa4437e92f5
---
# PTX: Update generated files with Blackwell instructions 

- Repository: `NVIDIA/cccl`
- PR: [3568](https://github.com/NVIDIA/cccl/pull/3568)
- Decision: `Search-matched CUDA kernel optimization PR; materialize diff/source before synthesis.`
- Reason: Search-matched CUDA kernel optimization PR; materialize diff/source before synthesis.

## Changed Paths

- `docs/libcudacxx/ptx/instructions/generated/barrier_cluster_aligned.rst`
- `docs/libcudacxx/ptx/instructions/generated/clusterlaunchcontrol.rst`
- `docs/libcudacxx/ptx/instructions/generated/cp_async_bulk.rst`
- `docs/libcudacxx/ptx/instructions/generated/cp_async_bulk_multicast.rst`
- `docs/libcudacxx/ptx/instructions/generated/cp_async_bulk_tensor.rst`
- `docs/libcudacxx/ptx/instructions/generated/cp_async_bulk_tensor_gather_scatter.rst`
- `docs/libcudacxx/ptx/instructions/generated/cp_async_bulk_tensor_multicast.rst`
- `docs/libcudacxx/ptx/instructions/generated/cp_async_mbarrier_arrive.rst`
- `docs/libcudacxx/ptx/instructions/generated/cp_async_mbarrier_arrive_noinc.rst`
- `docs/libcudacxx/ptx/instructions/generated/cp_reduce_async_bulk.rst`
- `docs/libcudacxx/ptx/instructions/generated/elect_sync.rst`
- `docs/libcudacxx/ptx/instructions/generated/fence.rst`
- `docs/libcudacxx/ptx/instructions/generated/fence_proxy_async.rst`
- `docs/libcudacxx/ptx/instructions/generated/fence_proxy_async_generic_sync_restrict.rst`
- `docs/libcudacxx/ptx/instructions/generated/fence_sync_restrict.rst`
- `docs/libcudacxx/ptx/instructions/generated/get_sreg.rst`
- `docs/libcudacxx/ptx/instructions/generated/getctarank.rst`
- `docs/libcudacxx/ptx/instructions/generated/mapa.rst`
- `docs/libcudacxx/ptx/instructions/generated/mbarrier_arrive.rst`
- `docs/libcudacxx/ptx/instructions/generated/mbarrier_arrive_expect_tx.rst`
- `docs/libcudacxx/ptx/instructions/generated/mbarrier_test_wait.rst`
- `docs/libcudacxx/ptx/instructions/generated/mbarrier_test_wait_parity.rst`
- `docs/libcudacxx/ptx/instructions/generated/mbarrier_try_wait.rst`
- `docs/libcudacxx/ptx/instructions/generated/mbarrier_try_wait_parity.rst`
- `docs/libcudacxx/ptx/instructions/generated/multimem_ld_reduce.rst`
- `docs/libcudacxx/ptx/instructions/generated/multimem_red.rst`
- `docs/libcudacxx/ptx/instructions/generated/multimem_st.rst`
- `docs/libcudacxx/ptx/instructions/generated/red_async.rst`
- `docs/libcudacxx/ptx/instructions/generated/st_async.rst`
- `docs/libcudacxx/ptx/instructions/generated/st_bulk.rst`
- `docs/libcudacxx/ptx/instructions/generated/tcgen05_alloc.rst`
- `docs/libcudacxx/ptx/instructions/generated/tcgen05_commit.rst`
- `docs/libcudacxx/ptx/instructions/generated/tcgen05_cp.rst`
- `docs/libcudacxx/ptx/instructions/generated/tcgen05_fence.rst`
- `docs/libcudacxx/ptx/instructions/generated/tcgen05_ld.rst`
- `docs/libcudacxx/ptx/instructions/generated/tcgen05_mma.rst`
- `docs/libcudacxx/ptx/instructions/generated/tcgen05_mma_ws.rst`
- `docs/libcudacxx/ptx/instructions/generated/tcgen05_shift.rst`
- `docs/libcudacxx/ptx/instructions/generated/tcgen05_st.rst`
- `docs/libcudacxx/ptx/instructions/generated/tcgen05_wait.rst`
- `docs/libcudacxx/ptx/instructions/generated/tensormap_replace.rst`
- `docs/libcudacxx/ptx/instructions/special_registers.rst`
- `docs/repo.toml`
- `libcudacxx/include/cuda/__ptx/instructions/generated/barrier_cluster.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/barrier_cluster_aligned.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/clusterlaunchcontrol.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_commit_group.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_multicast.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_tensor.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_tensor_gather_scatter.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_tensor_multicast.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_wait_group.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_mbarrier_arrive.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_mbarrier_arrive_noinc.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/cp_reduce_async_bulk.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/cp_reduce_async_bulk_bf16.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/cp_reduce_async_bulk_f16.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/cp_reduce_async_bulk_tensor.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/elect_sync.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/fence.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/fence_mbarrier_init.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/fence_proxy_alias.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/fence_proxy_async.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/fence_proxy_async_generic_sync_restrict.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/fence_proxy_tensormap_generic.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/fence_sync_restrict.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/get_sreg.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/getctarank.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/mapa.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_arrive.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_arrive_expect_tx.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_arrive_no_complete.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_expect_tx.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_init.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_test_wait.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_test_wait_parity.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_try_wait.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_try_wait_parity.h`
- `libcudacxx/include/cuda/__ptx/instructions/generated/multimem_ld_reduce.h`

## Evidence Bundle

`evidence/pull-bundles/cccl-cub/gh-3568`
