[Mlir-commits] [llvm] [mlir] [NVPTX] Add Rubin extensions to tcgen05.commit (PR #211577)
Rajat Bajpai
llvmlistbot at llvm.org
Thu Jul 30 23:04:46 PDT 2026
https://github.com/rajatbajpai updated https://github.com/llvm/llvm-project/pull/211577
>From 1c16e3c6b1bf1ae8d5822d777cf42b4b6562cc38 Mon Sep 17 00:00:00 2001
From: Rajat Bajpai <rbajpai at nvidia.com>
Date: Thu, 23 Jul 2026 10:40:42 +0000
Subject: [PATCH 1/3] [NVPTX] Add Rubin extensions to tcgen05.commit
The Rubin architecture extends `tcgen05.commit` operations with two
additional features: support for 32-bit CTA multicast masks and the
ability to track completion of Matrix A reads from shared memory for
all prior `tcgen05.mma` operations.
This change adds support for these features to the `tcgen05.commit`
intrinsics.
---
llvm/docs/NVPTXUsage.md | 16 +-
llvm/include/llvm/IR/IntrinsicsNVVM.td | 40 +-
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 41 +-
llvm/lib/Target/NVPTX/NVPTXSubtarget.h | 6 +
.../Assembler/auto_upgrade_nvvm_intrinsics.ll | 18 +
.../CodeGen/NVPTX/tcgen05-commit-rubin.ll | 438 ++++++++++++++++++
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 2 +-
.../Target/LLVMIR/nvvm/tcgen05-commit.mlir | 8 +-
8 files changed, 532 insertions(+), 37 deletions(-)
create mode 100644 llvm/test/CodeGen/NVPTX/tcgen05-commit-rubin.ll
diff --git a/llvm/docs/NVPTXUsage.md b/llvm/docs/NVPTXUsage.md
index 810938303c591..c9a6d619716a4 100644
--- a/llvm/docs/NVPTXUsage.md
+++ b/llvm/docs/NVPTXUsage.md
@@ -2630,8 +2630,17 @@ For more information, refer to the [PTX ISA](https://docs.nvidia.com/cuda/parall
```llvm
declare void @llvm.nvvm.tcgen05.commit.{cg1,cg2}(ptr %mbar)
declare void @llvm.nvvm.tcgen05.commit.shared.{cg1,cg2}(ptr addrspace(3) %mbar)
-declare void @llvm.nvvm.tcgen05.commit.mc.{cg1,cg2}(ptr %mbar, i16 %mc)
-declare void @llvm.nvvm.tcgen05.commit.mc.shared.{cg1,cg2}(ptr addrspace(3) %mbar, i16 %mc)
+declare void @llvm.nvvm.tcgen05.commit.mc.{cg1,cg2}.i16(ptr %mbar, i16 %mc)
+declare void @llvm.nvvm.tcgen05.commit.mc.shared.{cg1,cg2}.i16(ptr addrspace(3) %mbar, i16 %mc)
+declare void @llvm.nvvm.tcgen05.commit.mc.{cg1,cg2}.i32(ptr %mbar, i32 %mc)
+declare void @llvm.nvvm.tcgen05.commit.mc.shared.{cg1,cg2}.i32(ptr addrspace(3) %mbar, i32 %mc)
+
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.{cg1,cg2}(ptr %bar_addr)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.shared.{cg1,cg2}(ptr addrspace(3) %bar_addr)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.{cg1,cg2}.i16(ptr %bar_addr, i16 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.shared.{cg1,cg2}.i16(ptr addrspace(3) %bar_addr, i16 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.{cg1,cg2}.i32(ptr %bar_addr, i32 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.shared.{cg1,cg2}.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
```
##### Overview:
@@ -2643,7 +2652,8 @@ object (`%mbar`) track the completion of all prior asynchronous tcgen05
operations. The `.mc` variants allow signaling on the mbarrier objects of
multiple CTAs (specified by `%mc`) in the cluster. The `.cg1` and `.cg2`
variants generate `cta_group::1` and `cta_group::2` flavors of the
-instruction respectively.
+instruction, respectively. The `smem.a.read` variants allow tracking the completion
+of reads of Matrix A from shared memory for all prior `tcgen05.mma` operations.
For more information, refer to the [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen-async-sync-operations-commit).
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 40f3c40a76bc6..26be050d4b22a 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -3085,25 +3085,27 @@ foreach cta_group = ["cg1", "cg2"] in {
def int_nvvm_tcgen05_relinq_alloc_permit_ # cta_group : Intrinsic<[], [],
[IntrConvergent, IntrInaccessibleMemOnly]>;
- def int_nvvm_tcgen05_commit_ # cta_group : Intrinsic<[],
- [llvm_ptr_ty], // mbar_ptr
- [IntrConvergent, IntrInaccessibleMemOrArgMemOnly,
- NoCapture<ArgIndex<0>>]>;
-
- def int_nvvm_tcgen05_commit_shared_ # cta_group : Intrinsic<[],
- [llvm_shared_ptr_ty], // mbar_ptr
- [IntrConvergent, IntrInaccessibleMemOrArgMemOnly,
- NoCapture<ArgIndex<0>>]>;
-
- def int_nvvm_tcgen05_commit_mc_ # cta_group : Intrinsic<[],
- [llvm_ptr_ty, llvm_i16_ty], // mbar_ptr, cta_mask
- [IntrConvergent, IntrInaccessibleMemOrArgMemOnly,
- NoCapture<ArgIndex<0>>]>;
-
- def int_nvvm_tcgen05_commit_mc_shared_ # cta_group : Intrinsic<[],
- [llvm_shared_ptr_ty, llvm_i16_ty], // mbar_ptr, cta_mask
- [IntrConvergent, IntrInaccessibleMemOrArgMemOnly,
- NoCapture<ArgIndex<0>>]>;
+ foreach smem_a_read = ["", "smem_a_read_"] in {
+ def int_nvvm_tcgen05_commit_ # smem_a_read # cta_group : Intrinsic<[],
+ [llvm_ptr_ty], // mbar_ptr
+ [IntrConvergent, IntrInaccessibleMemOrArgMemOnly,
+ NoCapture<ArgIndex<0>>]>;
+
+ def int_nvvm_tcgen05_commit_ # smem_a_read # shared_ # cta_group : Intrinsic<[],
+ [llvm_shared_ptr_ty], // mbar_ptr
+ [IntrConvergent, IntrInaccessibleMemOrArgMemOnly,
+ NoCapture<ArgIndex<0>>]>;
+
+ def int_nvvm_tcgen05_commit_ # smem_a_read # mc_ # cta_group : Intrinsic<[],
+ [llvm_ptr_ty, llvm_anyint_ty], // mbar_ptr, cta_mask[16 or 32 bits]
+ [IntrConvergent, IntrInaccessibleMemOrArgMemOnly,
+ NoCapture<ArgIndex<0>>]>;
+
+ def int_nvvm_tcgen05_commit_ # smem_a_read # mc_shared_ # cta_group : Intrinsic<[],
+ [llvm_shared_ptr_ty, llvm_anyint_ty], // mbar_ptr, cta_mask[16 or 32 bits]
+ [IntrConvergent, IntrInaccessibleMemOrArgMemOnly,
+ NoCapture<ArgIndex<0>>]>;
+ }
def int_nvvm_tcgen05_shift_down_ # cta_group : Intrinsic<[],
[llvm_tmem_ptr_ty], // tmem_addr
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 5230e01261fd6..ea2147e720da4 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -5711,25 +5711,41 @@ def tcgen05_wait_ld: NullaryInst<"tcgen05.wait::ld.sync.aligned", int_nvvm_tcgen
def tcgen05_wait_st: NullaryInst<"tcgen05.wait::st.sync.aligned", int_nvvm_tcgen05_wait_st>;
multiclass TCGEN05_COMMIT_INTR<string AS, string num> {
- defvar prefix = "tcgen05.commit.cta_group::" # num #".mbarrier::arrive::one.shared::cluster";
+ defvar prefix = "tcgen05.commit.cta_group::" # num #
+ ".mbarrier::arrive::one.shared::cluster";
+ defvar prefix_smem_a_read = prefix # ".sync_restrict::shared::read::mma::a";
+ defvar intr_prefix = "int_nvvm_tcgen05_commit";
defvar intr_suffix = !if(!eq(AS, "shared"), "_shared", "") # "_cg" # num;
- defvar Intr = !cast<Intrinsic>("int_nvvm_tcgen05_commit" # intr_suffix);
- defvar IntrMC = !cast<Intrinsic>("int_nvvm_tcgen05_commit_mc" # intr_suffix);
+ defvar Intr = !cast<Intrinsic>(intr_prefix # intr_suffix);
+ defvar IntrMC = !cast<Intrinsic>(intr_prefix # "_mc" # intr_suffix);
+ defvar IntrSmemARead =
+ !cast<Intrinsic>(intr_prefix # "_smem_a_read" # intr_suffix);
+ defvar IntrSmemAReadMC =
+ !cast<Intrinsic>(intr_prefix # "_smem_a_read_mc" # intr_suffix);
def "" : BasicNVPTXInst<(outs), (ins ADDR:$mbar),
prefix # ".b64",
[(Intr addr:$mbar)]>;
def _MC : BasicNVPTXInst<(outs), (ins ADDR:$mbar, B16:$mc),
- prefix # ".multicast::cluster.b64",
- [(IntrMC addr:$mbar, B16:$mc)]>;
+ prefix # ".multicast::cluster.b64",
+ [(IntrMC addr:$mbar, i16:$mc)]>;
+ let Predicates = [callSubtarget<"hasTcgen05RubinFamilySupport">] in {
+ def _MC_32B : BasicNVPTXInst<(outs), (ins ADDR:$mbar, B32:$mc),
+ prefix # ".multicast::cluster::32b.b64",
+ [(IntrMC addr:$mbar, i32:$mc)]>;
+ def _SMRA : BasicNVPTXInst<(outs), (ins ADDR:$mbar),
+ prefix_smem_a_read # ".b64",
+ [(IntrSmemARead addr:$mbar)]>;
+ def _SMRA_MC : BasicNVPTXInst<(outs), (ins ADDR:$mbar, B16:$mc),
+ prefix_smem_a_read # ".multicast::cluster.b64",
+ [(IntrSmemAReadMC addr:$mbar, i16:$mc)]>;
+ def _SMRA_MC_32B : BasicNVPTXInst<(outs), (ins ADDR:$mbar, B32:$mc),
+ prefix_smem_a_read # ".multicast::cluster::32b.b64",
+ [(IntrSmemAReadMC addr:$mbar, i32:$mc)]>;
+ }
}
-defm TCGEN05_COMMIT_CG1 : TCGEN05_COMMIT_INTR<"", "1">;
-defm TCGEN05_COMMIT_CG2 : TCGEN05_COMMIT_INTR<"", "2">;
-defm TCGEN05_COMMIT_S64_CG1 : TCGEN05_COMMIT_INTR<"shared", "1">;
-defm TCGEN05_COMMIT_S64_CG2 : TCGEN05_COMMIT_INTR<"shared", "2">;
-
multiclass TCGEN05_CP_INTR<string shape, string src_fmt, string mc = ""> {
defvar dst_fmt = !if(!eq(src_fmt, ""), "", ".b8x16");
defvar fmt_asm = StrJoin<".", [dst_fmt, src_fmt]>.ret;
@@ -5762,6 +5778,11 @@ foreach src_fmt = ["", "b6x16_p32", "b4x16_p64"] in {
}
} // Predicates
+defm TCGEN05_COMMIT_CG1 : TCGEN05_COMMIT_INTR<"", "1">;
+defm TCGEN05_COMMIT_CG2 : TCGEN05_COMMIT_INTR<"", "2">;
+defm TCGEN05_COMMIT_S64_CG1 : TCGEN05_COMMIT_INTR<"shared", "1">;
+defm TCGEN05_COMMIT_S64_CG2 : TCGEN05_COMMIT_INTR<"shared", "2">;
+
let Predicates = [callSubtarget<"hasTcgen05ShiftSupport">] in {
multiclass TCGEN05_SHIFT_INTR<string num, Intrinsic Intr> {
def "" : BasicNVPTXInst<(outs),
diff --git a/llvm/lib/Target/NVPTX/NVPTXSubtarget.h b/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
index 1df5d326f63a6..9561e00c84023 100644
--- a/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
+++ b/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
@@ -140,6 +140,12 @@ class NVPTXSubtarget : public NVPTXGenSubtargetInfo {
hasPTXWithAccelSMs(86, {100, 101});
}
+ // Checks Rubin family extensions support.
+ // - tcgen05.commit
+ bool hasTcgen05RubinFamilySupport() const {
+ return hasPTXWithFamilySMs(94, {107});
+ }
+
// Checks tcgen05.shift instruction support.
bool hasTcgen05ShiftSupport() const {
// sm_101 renamed to sm_110 in PTX 9.0
diff --git a/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll b/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
index c59c9bda203b8..d3b815a1ef708 100644
--- a/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
+++ b/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
@@ -115,6 +115,11 @@ declare ptr addrspace(3) @llvm.nvvm.mapa.shared.cluster(ptr addrspace(3), i32)
declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster(ptr addrspace(3), ptr addrspace(3), ptr addrspace(1), i32, i16, i64, i1, i1)
declare void @llvm.nvvm.cp.async.bulk.shared.cta.to.cluster(ptr addrspace(3), ptr addrspace(3), ptr addrspace(3), i32)
+declare void @llvm.nvvm.tcgen05.commit.mc.cg1(ptr, i16)
+declare void @llvm.nvvm.tcgen05.commit.mc.cg2(ptr, i16)
+declare void @llvm.nvvm.tcgen05.commit.mc.shared.cg1(ptr addrspace(3), i16)
+declare void @llvm.nvvm.tcgen05.commit.mc.shared.cg2(ptr addrspace(3), i16)
+
declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.1d(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr %tm, i32 %d0, i16 %mc, i64 %ch, i1 %f1, i1 %f2);
declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.2d(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr %tm, i32 %d0, i32 %d1, i16 %mc, i64 %ch, i1 %f1, i1 %f2);
declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.3d(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr %tm, i32 %d0, i32 %d1, i32 %d2, i16 %mc, i64 %ch, i1 %f1, i1 %f2);
@@ -456,6 +461,19 @@ define void @nvvm_cp_async_bulk_intrinsics(ptr addrspace(3) %dst, ptr addrspace(
ret void
}
+; CHECK-LABEL: @nvvm_tcgen05_commit_mc_intrinsics
+define void @nvvm_tcgen05_commit_mc_intrinsics(ptr %bar, ptr addrspace(3) %bar_shared, i16 %cta_mask) {
+; CHECK: call void @llvm.nvvm.tcgen05.commit.mc.cg1.i16(ptr %bar, i16 %cta_mask)
+; CHECK: call void @llvm.nvvm.tcgen05.commit.mc.cg2.i16(ptr %bar, i16 %cta_mask)
+; CHECK: call void @llvm.nvvm.tcgen05.commit.mc.shared.cg1.i16(ptr addrspace(3) %bar_shared, i16 %cta_mask)
+; CHECK: call void @llvm.nvvm.tcgen05.commit.mc.shared.cg2.i16(ptr addrspace(3) %bar_shared, i16 %cta_mask)
+ call void @llvm.nvvm.tcgen05.commit.mc.cg1(ptr %bar, i16 %cta_mask)
+ call void @llvm.nvvm.tcgen05.commit.mc.cg2(ptr %bar, i16 %cta_mask)
+ call void @llvm.nvvm.tcgen05.commit.mc.shared.cg1(ptr addrspace(3) %bar_shared, i16 %cta_mask)
+ call void @llvm.nvvm.tcgen05.commit.mc.shared.cg2(ptr addrspace(3) %bar_shared, i16 %cta_mask)
+ ret void
+}
+
; CHECK-LABEL: @nvvm_cp_async_bulk_tensor_g2s_im2col
define void @nvvm_cp_async_bulk_tensor_g2s_im2col(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i16 %im2col0, i16 %im2col1, i16 %im2col2, i16 %mc, i64 %ch) {
; CHECK: call void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.3d(ptr addrspace(7) %1, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i16 %im2col0, i16 0, i64 0, i1 false, i1 false, i32 0)
diff --git a/llvm/test/CodeGen/NVPTX/tcgen05-commit-rubin.ll b/llvm/test/CodeGen/NVPTX/tcgen05-commit-rubin.ll
new file mode 100644
index 0000000000000..0b54e735ef3a9
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/tcgen05-commit-rubin.ll
@@ -0,0 +1,438 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_107f | FileCheck --check-prefixes=CHECK_PTX64 %s
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_107f --nvptx-short-ptr | FileCheck --check-prefixes=CHECK_PTX64_SHARED32 %s
+
+; RUN: %if ptxas-sm_107f && ptxas-isa-9.4 %{ llc < %s -march=nvptx64 -mcpu=sm_107f | %ptxas-verify -arch=sm_107f %}
+; RUN: %if ptxas-sm_107f && ptxas-isa-9.4 %{ llc < %s -march=nvptx64 -mcpu=sm_107f --nvptx-short-ptr | %ptxas-verify -arch=sm_107f %}
+
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.cg1(ptr %bar_addr)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.cg2(ptr %bar_addr)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.shared.cg1(ptr addrspace(3) %bar_addr)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.shared.cg2(ptr addrspace(3) %bar_addr)
+
+define void @test_tcgen05_commit_smem_a_read_cg1(ptr %bar_addr) {
+; CHECK_PTX64-LABEL: test_tcgen05_commit_smem_a_read_cg1(
+; CHECK_PTX64: {
+; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;
+; CHECK_PTX64-EMPTY:
+; CHECK_PTX64-NEXT: // %bb.0:
+; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_smem_a_read_cg1_param_0];
+; CHECK_PTX64-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.b64 [%rd1];
+; CHECK_PTX64-NEXT: ret;
+;
+; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_commit_smem_a_read_cg1(
+; CHECK_PTX64_SHARED32: {
+; CHECK_PTX64_SHARED32-NEXT: .reg .b64 %rd<2>;
+; CHECK_PTX64_SHARED32-EMPTY:
+; CHECK_PTX64_SHARED32-NEXT: // %bb.0:
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_smem_a_read_cg1_param_0];
+; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.b64 [%rd1];
+; CHECK_PTX64_SHARED32-NEXT: ret;
+ call void @llvm.nvvm.tcgen05.commit.smem.a.read.cg1(ptr %bar_addr)
+
+ ret void
+}
+
+define void @test_tcgen05_commit_smem_a_read_cg2(ptr %bar_addr) {
+; CHECK_PTX64-LABEL: test_tcgen05_commit_smem_a_read_cg2(
+; CHECK_PTX64: {
+; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;
+; CHECK_PTX64-EMPTY:
+; CHECK_PTX64-NEXT: // %bb.0:
+; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_smem_a_read_cg2_param_0];
+; CHECK_PTX64-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.b64 [%rd1];
+; CHECK_PTX64-NEXT: ret;
+;
+; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_commit_smem_a_read_cg2(
+; CHECK_PTX64_SHARED32: {
+; CHECK_PTX64_SHARED32-NEXT: .reg .b64 %rd<2>;
+; CHECK_PTX64_SHARED32-EMPTY:
+; CHECK_PTX64_SHARED32-NEXT: // %bb.0:
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_smem_a_read_cg2_param_0];
+; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.b64 [%rd1];
+; CHECK_PTX64_SHARED32-NEXT: ret;
+ call void @llvm.nvvm.tcgen05.commit.smem.a.read.cg2(ptr %bar_addr)
+
+ ret void
+}
+
+define void @test_tcgen05_commit_smem_a_read_shared_cg1(ptr addrspace(3) %bar_addr) {
+; CHECK_PTX64-LABEL: test_tcgen05_commit_smem_a_read_shared_cg1(
+; CHECK_PTX64: {
+; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;
+; CHECK_PTX64-EMPTY:
+; CHECK_PTX64-NEXT: // %bb.0:
+; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_smem_a_read_shared_cg1_param_0];
+; CHECK_PTX64-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.b64 [%rd1];
+; CHECK_PTX64-NEXT: ret;
+;
+; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_commit_smem_a_read_shared_cg1(
+; CHECK_PTX64_SHARED32: {
+; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<2>;
+; CHECK_PTX64_SHARED32-EMPTY:
+; CHECK_PTX64_SHARED32-NEXT: // %bb.0:
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_smem_a_read_shared_cg1_param_0];
+; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.b64 [%r1];
+; CHECK_PTX64_SHARED32-NEXT: ret;
+ call void @llvm.nvvm.tcgen05.commit.smem.a.read.shared.cg1(ptr addrspace(3) %bar_addr)
+
+ ret void
+}
+
+define void @test_tcgen05_commit_smem_a_read_shared_cg2(ptr addrspace(3) %bar_addr) {
+; CHECK_PTX64-LABEL: test_tcgen05_commit_smem_a_read_shared_cg2(
+; CHECK_PTX64: {
+; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;
+; CHECK_PTX64-EMPTY:
+; CHECK_PTX64-NEXT: // %bb.0:
+; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_smem_a_read_shared_cg2_param_0];
+; CHECK_PTX64-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.b64 [%rd1];
+; CHECK_PTX64-NEXT: ret;
+;
+; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_commit_smem_a_read_shared_cg2(
+; CHECK_PTX64_SHARED32: {
+; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<2>;
+; CHECK_PTX64_SHARED32-EMPTY:
+; CHECK_PTX64_SHARED32-NEXT: // %bb.0:
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_smem_a_read_shared_cg2_param_0];
+; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.b64 [%r1];
+; CHECK_PTX64_SHARED32-NEXT: ret;
+ call void @llvm.nvvm.tcgen05.commit.smem.a.read.shared.cg2(ptr addrspace(3) %bar_addr)
+
+ ret void
+}
+
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg1.i16(ptr %bar_addr, i16 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg2.i16(ptr %bar_addr, i16 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.shared.cg1.i16(ptr addrspace(3) %bar_addr, i16 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.shared.cg2.i16(ptr addrspace(3) %bar_addr, i16 %cta_mask)
+
+define void @test_tcgen05_commit_smem_a_read_mc_cg1(ptr %bar_addr, i16 %cta_mask) {
+; CHECK_PTX64-LABEL: test_tcgen05_commit_smem_a_read_mc_cg1(
+; CHECK_PTX64: {
+; CHECK_PTX64-NEXT: .reg .b16 %rs<2>;
+; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;
+; CHECK_PTX64-EMPTY:
+; CHECK_PTX64-NEXT: // %bb.0:
+; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_smem_a_read_mc_cg1_param_0];
+; CHECK_PTX64-NEXT: ld.param.b16 %rs1, [test_tcgen05_commit_smem_a_read_mc_cg1_param_1];
+; CHECK_PTX64-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.multicast::cluster.b64 [%rd1], %rs1;
+; CHECK_PTX64-NEXT: ret;
+;
+; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_commit_smem_a_read_mc_cg1(
+; CHECK_PTX64_SHARED32: {
+; CHECK_PTX64_SHARED32-NEXT: .reg .b16 %rs<2>;
+; CHECK_PTX64_SHARED32-NEXT: .reg .b64 %rd<2>;
+; CHECK_PTX64_SHARED32-EMPTY:
+; CHECK_PTX64_SHARED32-NEXT: // %bb.0:
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_smem_a_read_mc_cg1_param_0];
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b16 %rs1, [test_tcgen05_commit_smem_a_read_mc_cg1_param_1];
+; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.multicast::cluster.b64 [%rd1], %rs1;
+; CHECK_PTX64_SHARED32-NEXT: ret;
+ call void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg1.i16(ptr %bar_addr, i16 %cta_mask)
+
+ ret void
+}
+
+define void @test_tcgen05_commit_smem_a_read_mc_cg2(ptr %bar_addr, i16 %cta_mask) {
+; CHECK_PTX64-LABEL: test_tcgen05_commit_smem_a_read_mc_cg2(
+; CHECK_PTX64: {
+; CHECK_PTX64-NEXT: .reg .b16 %rs<2>;
+; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;
+; CHECK_PTX64-EMPTY:
+; CHECK_PTX64-NEXT: // %bb.0:
+; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_smem_a_read_mc_cg2_param_0];
+; CHECK_PTX64-NEXT: ld.param.b16 %rs1, [test_tcgen05_commit_smem_a_read_mc_cg2_param_1];
+; CHECK_PTX64-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.multicast::cluster.b64 [%rd1], %rs1;
+; CHECK_PTX64-NEXT: ret;
+;
+; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_commit_smem_a_read_mc_cg2(
+; CHECK_PTX64_SHARED32: {
+; CHECK_PTX64_SHARED32-NEXT: .reg .b16 %rs<2>;
+; CHECK_PTX64_SHARED32-NEXT: .reg .b64 %rd<2>;
+; CHECK_PTX64_SHARED32-EMPTY:
+; CHECK_PTX64_SHARED32-NEXT: // %bb.0:
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_smem_a_read_mc_cg2_param_0];
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b16 %rs1, [test_tcgen05_commit_smem_a_read_mc_cg2_param_1];
+; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.multicast::cluster.b64 [%rd1], %rs1;
+; CHECK_PTX64_SHARED32-NEXT: ret;
+ call void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg2.i16(ptr %bar_addr, i16 %cta_mask)
+
+ ret void
+}
+
+define void @test_tcgen05_commit_smem_a_read_mc_shared_cg1(ptr addrspace(3) %bar_addr, i16 %cta_mask) {
+; CHECK_PTX64-LABEL: test_tcgen05_commit_smem_a_read_mc_shared_cg1(
+; CHECK_PTX64: {
+; CHECK_PTX64-NEXT: .reg .b16 %rs<2>;
+; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;
+; CHECK_PTX64-EMPTY:
+; CHECK_PTX64-NEXT: // %bb.0:
+; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_smem_a_read_mc_shared_cg1_param_0];
+; CHECK_PTX64-NEXT: ld.param.b16 %rs1, [test_tcgen05_commit_smem_a_read_mc_shared_cg1_param_1];
+; CHECK_PTX64-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.multicast::cluster.b64 [%rd1], %rs1;
+; CHECK_PTX64-NEXT: ret;
+;
+; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_commit_smem_a_read_mc_shared_cg1(
+; CHECK_PTX64_SHARED32: {
+; CHECK_PTX64_SHARED32-NEXT: .reg .b16 %rs<2>;
+; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<2>;
+; CHECK_PTX64_SHARED32-EMPTY:
+; CHECK_PTX64_SHARED32-NEXT: // %bb.0:
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_smem_a_read_mc_shared_cg1_param_0];
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b16 %rs1, [test_tcgen05_commit_smem_a_read_mc_shared_cg1_param_1];
+; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.multicast::cluster.b64 [%r1], %rs1;
+; CHECK_PTX64_SHARED32-NEXT: ret;
+ call void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.shared.cg1.i16(ptr addrspace(3) %bar_addr, i16 %cta_mask)
+
+ ret void
+}
+
+define void @test_tcgen05_commit_smem_a_read_mc_shared_cg2(ptr addrspace(3) %bar_addr, i16 %cta_mask) {
+; CHECK_PTX64-LABEL: test_tcgen05_commit_smem_a_read_mc_shared_cg2(
+; CHECK_PTX64: {
+; CHECK_PTX64-NEXT: .reg .b16 %rs<2>;
+; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;
+; CHECK_PTX64-EMPTY:
+; CHECK_PTX64-NEXT: // %bb.0:
+; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_smem_a_read_mc_shared_cg2_param_0];
+; CHECK_PTX64-NEXT: ld.param.b16 %rs1, [test_tcgen05_commit_smem_a_read_mc_shared_cg2_param_1];
+; CHECK_PTX64-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.multicast::cluster.b64 [%rd1], %rs1;
+; CHECK_PTX64-NEXT: ret;
+;
+; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_commit_smem_a_read_mc_shared_cg2(
+; CHECK_PTX64_SHARED32: {
+; CHECK_PTX64_SHARED32-NEXT: .reg .b16 %rs<2>;
+; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<2>;
+; CHECK_PTX64_SHARED32-EMPTY:
+; CHECK_PTX64_SHARED32-NEXT: // %bb.0:
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_smem_a_read_mc_shared_cg2_param_0];
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b16 %rs1, [test_tcgen05_commit_smem_a_read_mc_shared_cg2_param_1];
+; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.multicast::cluster.b64 [%r1], %rs1;
+; CHECK_PTX64_SHARED32-NEXT: ret;
+ call void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.shared.cg2.i16(ptr addrspace(3) %bar_addr, i16 %cta_mask)
+
+ ret void
+}
+
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg1.i32(ptr %bar_addr, i32 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg2.i32(ptr %bar_addr, i32 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.shared.cg1.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.shared.cg2.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
+
+define void @test_tcgen05_commit_smem_a_read_mc_cg1_i32(ptr %bar_addr, i32 %cta_mask) {
+; CHECK_PTX64-LABEL: test_tcgen05_commit_smem_a_read_mc_cg1_i32(
+; CHECK_PTX64: {
+; CHECK_PTX64-NEXT: .reg .b32 %r<2>;
+; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;
+; CHECK_PTX64-EMPTY:
+; CHECK_PTX64-NEXT: // %bb.0:
+; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_smem_a_read_mc_cg1_i32_param_0];
+; CHECK_PTX64-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_smem_a_read_mc_cg1_i32_param_1];
+; CHECK_PTX64-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.multicast::cluster::32b.b64 [%rd1], %r1;
+; CHECK_PTX64-NEXT: ret;
+;
+; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_commit_smem_a_read_mc_cg1_i32(
+; CHECK_PTX64_SHARED32: {
+; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<2>;
+; CHECK_PTX64_SHARED32-NEXT: .reg .b64 %rd<2>;
+; CHECK_PTX64_SHARED32-EMPTY:
+; CHECK_PTX64_SHARED32-NEXT: // %bb.0:
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_smem_a_read_mc_cg1_i32_param_0];
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_smem_a_read_mc_cg1_i32_param_1];
+; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.multicast::cluster::32b.b64 [%rd1], %r1;
+; CHECK_PTX64_SHARED32-NEXT: ret;
+ call void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg1.i32(ptr %bar_addr, i32 %cta_mask)
+
+ ret void
+}
+
+define void @test_tcgen05_commit_smem_a_read_mc_cg2_i32(ptr %bar_addr, i32 %cta_mask) {
+; CHECK_PTX64-LABEL: test_tcgen05_commit_smem_a_read_mc_cg2_i32(
+; CHECK_PTX64: {
+; CHECK_PTX64-NEXT: .reg .b32 %r<2>;
+; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;
+; CHECK_PTX64-EMPTY:
+; CHECK_PTX64-NEXT: // %bb.0:
+; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_smem_a_read_mc_cg2_i32_param_0];
+; CHECK_PTX64-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_smem_a_read_mc_cg2_i32_param_1];
+; CHECK_PTX64-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.multicast::cluster::32b.b64 [%rd1], %r1;
+; CHECK_PTX64-NEXT: ret;
+;
+; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_commit_smem_a_read_mc_cg2_i32(
+; CHECK_PTX64_SHARED32: {
+; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<2>;
+; CHECK_PTX64_SHARED32-NEXT: .reg .b64 %rd<2>;
+; CHECK_PTX64_SHARED32-EMPTY:
+; CHECK_PTX64_SHARED32-NEXT: // %bb.0:
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_smem_a_read_mc_cg2_i32_param_0];
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_smem_a_read_mc_cg2_i32_param_1];
+; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.multicast::cluster::32b.b64 [%rd1], %r1;
+; CHECK_PTX64_SHARED32-NEXT: ret;
+ call void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg2.i32(ptr %bar_addr, i32 %cta_mask)
+
+ ret void
+}
+
+define void @test_tcgen05_commit_smem_a_read_mc_shared_cg1_i32(ptr addrspace(3) %bar_addr, i32 %cta_mask) {
+; CHECK_PTX64-LABEL: test_tcgen05_commit_smem_a_read_mc_shared_cg1_i32(
+; CHECK_PTX64: {
+; CHECK_PTX64-NEXT: .reg .b32 %r<2>;
+; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;
+; CHECK_PTX64-EMPTY:
+; CHECK_PTX64-NEXT: // %bb.0:
+; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_smem_a_read_mc_shared_cg1_i32_param_0];
+; CHECK_PTX64-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_smem_a_read_mc_shared_cg1_i32_param_1];
+; CHECK_PTX64-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.multicast::cluster::32b.b64 [%rd1], %r1;
+; CHECK_PTX64-NEXT: ret;
+;
+; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_commit_smem_a_read_mc_shared_cg1_i32(
+; CHECK_PTX64_SHARED32: {
+; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<3>;
+; CHECK_PTX64_SHARED32-EMPTY:
+; CHECK_PTX64_SHARED32-NEXT: // %bb.0:
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_smem_a_read_mc_shared_cg1_i32_param_0];
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r2, [test_tcgen05_commit_smem_a_read_mc_shared_cg1_i32_param_1];
+; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.multicast::cluster::32b.b64 [%r1], %r2;
+; CHECK_PTX64_SHARED32-NEXT: ret;
+ call void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.shared.cg1.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
+
+ ret void
+}
+
+define void @test_tcgen05_commit_smem_a_read_mc_shared_cg2_i32(ptr addrspace(3) %bar_addr, i32 %cta_mask) {
+; CHECK_PTX64-LABEL: test_tcgen05_commit_smem_a_read_mc_shared_cg2_i32(
+; CHECK_PTX64: {
+; CHECK_PTX64-NEXT: .reg .b32 %r<2>;
+; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;
+; CHECK_PTX64-EMPTY:
+; CHECK_PTX64-NEXT: // %bb.0:
+; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_smem_a_read_mc_shared_cg2_i32_param_0];
+; CHECK_PTX64-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_smem_a_read_mc_shared_cg2_i32_param_1];
+; CHECK_PTX64-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.multicast::cluster::32b.b64 [%rd1], %r1;
+; CHECK_PTX64-NEXT: ret;
+;
+; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_commit_smem_a_read_mc_shared_cg2_i32(
+; CHECK_PTX64_SHARED32: {
+; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<3>;
+; CHECK_PTX64_SHARED32-EMPTY:
+; CHECK_PTX64_SHARED32-NEXT: // %bb.0:
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_smem_a_read_mc_shared_cg2_i32_param_0];
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r2, [test_tcgen05_commit_smem_a_read_mc_shared_cg2_i32_param_1];
+; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.multicast::cluster::32b.b64 [%r1], %r2;
+; CHECK_PTX64_SHARED32-NEXT: ret;
+ call void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.shared.cg2.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
+
+ ret void
+}
+
+declare void @llvm.nvvm.tcgen05.commit.mc.cg1.i32(ptr %bar_addr, i32 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.mc.cg2.i32(ptr %bar_addr, i32 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.mc.shared.cg1.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.mc.shared.cg2.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
+
+define void @test_tcgen05_commit_mc_cg1_i32(ptr %bar_addr, i32 %cta_mask) {
+; CHECK_PTX64-LABEL: test_tcgen05_commit_mc_cg1_i32(
+; CHECK_PTX64: {
+; CHECK_PTX64-NEXT: .reg .b32 %r<2>;
+; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;
+; CHECK_PTX64-EMPTY:
+; CHECK_PTX64-NEXT: // %bb.0:
+; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_mc_cg1_i32_param_0];
+; CHECK_PTX64-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_mc_cg1_i32_param_1];
+; CHECK_PTX64-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.multicast::cluster::32b.b64 [%rd1], %r1;
+; CHECK_PTX64-NEXT: ret;
+;
+; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_commit_mc_cg1_i32(
+; CHECK_PTX64_SHARED32: {
+; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<2>;
+; CHECK_PTX64_SHARED32-NEXT: .reg .b64 %rd<2>;
+; CHECK_PTX64_SHARED32-EMPTY:
+; CHECK_PTX64_SHARED32-NEXT: // %bb.0:
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_mc_cg1_i32_param_0];
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_mc_cg1_i32_param_1];
+; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.multicast::cluster::32b.b64 [%rd1], %r1;
+; CHECK_PTX64_SHARED32-NEXT: ret;
+ call void @llvm.nvvm.tcgen05.commit.mc.cg1.i32(ptr %bar_addr, i32 %cta_mask)
+
+ ret void
+}
+
+define void @test_tcgen05_commit_mc_cg2_i32(ptr %bar_addr, i32 %cta_mask) {
+; CHECK_PTX64-LABEL: test_tcgen05_commit_mc_cg2_i32(
+; CHECK_PTX64: {
+; CHECK_PTX64-NEXT: .reg .b32 %r<2>;
+; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;
+; CHECK_PTX64-EMPTY:
+; CHECK_PTX64-NEXT: // %bb.0:
+; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_mc_cg2_i32_param_0];
+; CHECK_PTX64-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_mc_cg2_i32_param_1];
+; CHECK_PTX64-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.multicast::cluster::32b.b64 [%rd1], %r1;
+; CHECK_PTX64-NEXT: ret;
+;
+; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_commit_mc_cg2_i32(
+; CHECK_PTX64_SHARED32: {
+; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<2>;
+; CHECK_PTX64_SHARED32-NEXT: .reg .b64 %rd<2>;
+; CHECK_PTX64_SHARED32-EMPTY:
+; CHECK_PTX64_SHARED32-NEXT: // %bb.0:
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_mc_cg2_i32_param_0];
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_mc_cg2_i32_param_1];
+; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.multicast::cluster::32b.b64 [%rd1], %r1;
+; CHECK_PTX64_SHARED32-NEXT: ret;
+ call void @llvm.nvvm.tcgen05.commit.mc.cg2.i32(ptr %bar_addr, i32 %cta_mask)
+
+ ret void
+}
+
+define void @test_tcgen05_commit_mc_shared_cg1_i32(ptr addrspace(3) %bar_addr, i32 %cta_mask) {
+; CHECK_PTX64-LABEL: test_tcgen05_commit_mc_shared_cg1_i32(
+; CHECK_PTX64: {
+; CHECK_PTX64-NEXT: .reg .b32 %r<2>;
+; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;
+; CHECK_PTX64-EMPTY:
+; CHECK_PTX64-NEXT: // %bb.0:
+; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_mc_shared_cg1_i32_param_0];
+; CHECK_PTX64-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_mc_shared_cg1_i32_param_1];
+; CHECK_PTX64-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.multicast::cluster::32b.b64 [%rd1], %r1;
+; CHECK_PTX64-NEXT: ret;
+;
+; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_commit_mc_shared_cg1_i32(
+; CHECK_PTX64_SHARED32: {
+; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<3>;
+; CHECK_PTX64_SHARED32-EMPTY:
+; CHECK_PTX64_SHARED32-NEXT: // %bb.0:
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_mc_shared_cg1_i32_param_0];
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r2, [test_tcgen05_commit_mc_shared_cg1_i32_param_1];
+; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.multicast::cluster::32b.b64 [%r1], %r2;
+; CHECK_PTX64_SHARED32-NEXT: ret;
+ call void @llvm.nvvm.tcgen05.commit.mc.shared.cg1.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
+
+ ret void
+}
+
+define void @test_tcgen05_commit_mc_shared_cg2_i32(ptr addrspace(3) %bar_addr, i32 %cta_mask) {
+; CHECK_PTX64-LABEL: test_tcgen05_commit_mc_shared_cg2_i32(
+; CHECK_PTX64: {
+; CHECK_PTX64-NEXT: .reg .b32 %r<2>;
+; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;
+; CHECK_PTX64-EMPTY:
+; CHECK_PTX64-NEXT: // %bb.0:
+; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_mc_shared_cg2_i32_param_0];
+; CHECK_PTX64-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_mc_shared_cg2_i32_param_1];
+; CHECK_PTX64-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.multicast::cluster::32b.b64 [%rd1], %r1;
+; CHECK_PTX64-NEXT: ret;
+;
+; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_commit_mc_shared_cg2_i32(
+; CHECK_PTX64_SHARED32: {
+; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<3>;
+; CHECK_PTX64_SHARED32-EMPTY:
+; CHECK_PTX64_SHARED32-NEXT: // %bb.0:
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_mc_shared_cg2_i32_param_0];
+; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r2, [test_tcgen05_commit_mc_shared_cg2_i32_param_1];
+; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.multicast::cluster::32b.b64 [%r1], %r2;
+; CHECK_PTX64_SHARED32-NEXT: ret;
+ call void @llvm.nvvm.tcgen05.commit.mc.shared.cg2.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
+
+ ret void
+}
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 1f76218936e16..2bf21db6b6e6f 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -5331,7 +5331,7 @@ def NVVM_Tcgen05CommitOp : NVVM_Op<"tcgen05.commit", [NVVMRequiresSMf<[100, 101,
llvm::SmallVector<llvm::Value *> args;
auto id = NVVM::Tcgen05CommitOp::getIntrinsicIDAndArgs(
*op, moduleTranslation, args);
- createIntrinsicCall(builder, id, args);
+ createIntrinsicCall(builder, id, builder.getVoidTy(), args);
}];
}
diff --git a/mlir/test/Target/LLVMIR/nvvm/tcgen05-commit.mlir b/mlir/test/Target/LLVMIR/nvvm/tcgen05-commit.mlir
index 60475bf64ae7a..6997f8a98b2fc 100644
--- a/mlir/test/Target/LLVMIR/nvvm/tcgen05-commit.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/tcgen05-commit.mlir
@@ -8,10 +8,10 @@ llvm.func @llvm_nvvm_tcgen05_commit_generic(%barrier : !llvm.ptr, %cta_mask : i1
// CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.cg2(ptr %{{.*}})
nvvm.tcgen05.commit %barrier {group = #nvvm.cta_group<cta_2>} : !llvm.ptr
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.mc.cg1(ptr %{{.*}}, i16 %{{.*}})
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.mc.cg1.i16(ptr %{{.*}}, i16 %{{.*}})
nvvm.tcgen05.commit %barrier, multicast_mask = %cta_mask : !llvm.ptr, i16
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.mc.cg2(ptr %{{.*}}, i16 %{{.*}})
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.mc.cg2.i16(ptr %{{.*}}, i16 %{{.*}})
nvvm.tcgen05.commit %barrier, multicast_mask = %cta_mask {group = #nvvm.cta_group<cta_2>} : !llvm.ptr, i16
llvm.return
}
@@ -24,10 +24,10 @@ llvm.func @llvm_nvvm_tcgen05_commit_shared(%barrier : !llvm.ptr<3>, %cta_mask :
// CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.shared.cg2(ptr addrspace(3) %{{.*}})
nvvm.tcgen05.commit %barrier {group = #nvvm.cta_group<cta_2>} : !llvm.ptr<3>
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.mc.shared.cg1(ptr addrspace(3) %{{.*}}, i16 %{{.*}})
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.mc.shared.cg1.i16(ptr addrspace(3) %{{.*}}, i16 %{{.*}})
nvvm.tcgen05.commit %barrier, multicast_mask = %cta_mask : !llvm.ptr<3>, i16
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.mc.shared.cg2(ptr addrspace(3) %{{.*}}, i16 %{{.*}})
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.mc.shared.cg2.i16(ptr addrspace(3) %{{.*}}, i16 %{{.*}})
nvvm.tcgen05.commit %barrier, multicast_mask = %cta_mask {group = #nvvm.cta_group<cta_2>} : !llvm.ptr<3>, i16
llvm.return
}
>From d2fa4a882fd07421a2e7e5e3f6f2fd9b12fec1a8 Mon Sep 17 00:00:00 2001
From: Rajat Bajpai <rbajpai at nvidia.com>
Date: Fri, 24 Jul 2026 13:47:42 +0000
Subject: [PATCH 2/3] Addressed review comments
---
llvm/docs/NVPTXUsage.md | 4 ++--
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 6 +++---
llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll | 4 ++--
.../{tcgen05-commit-rubin.ll => tcgen05-commit-sm107.ll} | 0
4 files changed, 7 insertions(+), 7 deletions(-)
rename llvm/test/CodeGen/NVPTX/{tcgen05-commit-rubin.ll => tcgen05-commit-sm107.ll} (100%)
diff --git a/llvm/docs/NVPTXUsage.md b/llvm/docs/NVPTXUsage.md
index c9a6d619716a4..dd75e1f2d16f8 100644
--- a/llvm/docs/NVPTXUsage.md
+++ b/llvm/docs/NVPTXUsage.md
@@ -2652,8 +2652,8 @@ object (`%mbar`) track the completion of all prior asynchronous tcgen05
operations. The `.mc` variants allow signaling on the mbarrier objects of
multiple CTAs (specified by `%mc`) in the cluster. The `.cg1` and `.cg2`
variants generate `cta_group::1` and `cta_group::2` flavors of the
-instruction, respectively. The `smem.a.read` variants allow tracking the completion
-of reads of Matrix A from shared memory for all prior `tcgen05.mma` operations.
+instruction, respectively. The `smem.a.read` variants track the completion
+of reads of A-matrix from shared memory for all prior `tcgen05.mma` operations.
For more information, refer to the [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen-async-sync-operations-commit).
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index ea2147e720da4..05f9a5592cde4 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -5734,13 +5734,13 @@ multiclass TCGEN05_COMMIT_INTR<string AS, string num> {
def _MC_32B : BasicNVPTXInst<(outs), (ins ADDR:$mbar, B32:$mc),
prefix # ".multicast::cluster::32b.b64",
[(IntrMC addr:$mbar, i32:$mc)]>;
- def _SMRA : BasicNVPTXInst<(outs), (ins ADDR:$mbar),
+ def _SMEM_A_READ : BasicNVPTXInst<(outs), (ins ADDR:$mbar),
prefix_smem_a_read # ".b64",
[(IntrSmemARead addr:$mbar)]>;
- def _SMRA_MC : BasicNVPTXInst<(outs), (ins ADDR:$mbar, B16:$mc),
+ def _SMEM_A_READ_MC : BasicNVPTXInst<(outs), (ins ADDR:$mbar, B16:$mc),
prefix_smem_a_read # ".multicast::cluster.b64",
[(IntrSmemAReadMC addr:$mbar, i16:$mc)]>;
- def _SMRA_MC_32B : BasicNVPTXInst<(outs), (ins ADDR:$mbar, B32:$mc),
+ def _SMEM_A_READ_MC_32B : BasicNVPTXInst<(outs), (ins ADDR:$mbar, B32:$mc),
prefix_smem_a_read # ".multicast::cluster::32b.b64",
[(IntrSmemAReadMC addr:$mbar, i32:$mc)]>;
}
diff --git a/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll b/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
index d3b815a1ef708..5ffb1be055450 100644
--- a/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
+++ b/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
@@ -461,8 +461,8 @@ define void @nvvm_cp_async_bulk_intrinsics(ptr addrspace(3) %dst, ptr addrspace(
ret void
}
-; CHECK-LABEL: @nvvm_tcgen05_commit_mc_intrinsics
-define void @nvvm_tcgen05_commit_mc_intrinsics(ptr %bar, ptr addrspace(3) %bar_shared, i16 %cta_mask) {
+; CHECK-LABEL: @nvvm_tcgen05_commit_mc_16b_intrinsics
+define void @nvvm_tcgen05_commit_mc_16b_intrinsics(ptr %bar, ptr addrspace(3) %bar_shared, i16 %cta_mask) {
; CHECK: call void @llvm.nvvm.tcgen05.commit.mc.cg1.i16(ptr %bar, i16 %cta_mask)
; CHECK: call void @llvm.nvvm.tcgen05.commit.mc.cg2.i16(ptr %bar, i16 %cta_mask)
; CHECK: call void @llvm.nvvm.tcgen05.commit.mc.shared.cg1.i16(ptr addrspace(3) %bar_shared, i16 %cta_mask)
diff --git a/llvm/test/CodeGen/NVPTX/tcgen05-commit-rubin.ll b/llvm/test/CodeGen/NVPTX/tcgen05-commit-sm107.ll
similarity index 100%
rename from llvm/test/CodeGen/NVPTX/tcgen05-commit-rubin.ll
rename to llvm/test/CodeGen/NVPTX/tcgen05-commit-sm107.ll
>From 01aab9a8b4c89a5a76dc7ecc85c84461e6cfacf1 Mon Sep 17 00:00:00 2001
From: Rajat Bajpai <rbajpai at nvidia.com>
Date: Mon, 27 Jul 2026 09:13:37 +0000
Subject: [PATCH 3/3] Replaced generic/shared variants with overloaded form
---
llvm/docs/NVPTXUsage.md | 28 ++++----
llvm/include/llvm/IR/IntrinsicsNVVM.td | 14 +---
llvm/lib/IR/AutoUpgrade.cpp | 35 ++++++++++
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 10 ++-
.../Assembler/auto_upgrade_nvvm_intrinsics.ll | 25 ++++++--
.../CodeGen/NVPTX/tcgen05-commit-sm107.ll | 64 +++++++++----------
llvm/test/CodeGen/NVPTX/tcgen05-commit.ll | 32 +++++-----
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp | 20 +++---
.../Target/LLVMIR/nvvm/tcgen05-commit.mlir | 16 ++---
9 files changed, 141 insertions(+), 103 deletions(-)
diff --git a/llvm/docs/NVPTXUsage.md b/llvm/docs/NVPTXUsage.md
index dd75e1f2d16f8..98208f4a96dba 100644
--- a/llvm/docs/NVPTXUsage.md
+++ b/llvm/docs/NVPTXUsage.md
@@ -2628,19 +2628,21 @@ For more information, refer to the [PTX ISA](https://docs.nvidia.com/cuda/parall
##### Syntax:
```llvm
-declare void @llvm.nvvm.tcgen05.commit.{cg1,cg2}(ptr %mbar)
-declare void @llvm.nvvm.tcgen05.commit.shared.{cg1,cg2}(ptr addrspace(3) %mbar)
-declare void @llvm.nvvm.tcgen05.commit.mc.{cg1,cg2}.i16(ptr %mbar, i16 %mc)
-declare void @llvm.nvvm.tcgen05.commit.mc.shared.{cg1,cg2}.i16(ptr addrspace(3) %mbar, i16 %mc)
-declare void @llvm.nvvm.tcgen05.commit.mc.{cg1,cg2}.i32(ptr %mbar, i32 %mc)
-declare void @llvm.nvvm.tcgen05.commit.mc.shared.{cg1,cg2}.i32(ptr addrspace(3) %mbar, i32 %mc)
-
-declare void @llvm.nvvm.tcgen05.commit.smem.a.read.{cg1,cg2}(ptr %bar_addr)
-declare void @llvm.nvvm.tcgen05.commit.smem.a.read.shared.{cg1,cg2}(ptr addrspace(3) %bar_addr)
-declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.{cg1,cg2}.i16(ptr %bar_addr, i16 %cta_mask)
-declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.shared.{cg1,cg2}.i16(ptr addrspace(3) %bar_addr, i16 %cta_mask)
-declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.{cg1,cg2}.i32(ptr %bar_addr, i32 %cta_mask)
-declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.shared.{cg1,cg2}.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.{cg1,cg2}.p0(ptr %mbar)
+declare void @llvm.nvvm.tcgen05.commit.{cg1,cg2}.p3(ptr addrspace(3) %mbar)
+
+declare void @llvm.nvvm.tcgen05.commit.mc.{cg1,cg2}.p0.i16(ptr %mbar, i16 %mc)
+declare void @llvm.nvvm.tcgen05.commit.mc.{cg1,cg2}.p0.i32(ptr %mbar, i32 %mc)
+declare void @llvm.nvvm.tcgen05.commit.mc.{cg1,cg2}.p3.i16(ptr addrspace(3) %mbar, i16 %mc)
+declare void @llvm.nvvm.tcgen05.commit.mc.{cg1,cg2}.p3.i32(ptr addrspace(3) %mbar, i32 %mc)
+
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.{cg1,cg2}.p0(ptr %bar_addr)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.{cg1,cg2}.p3(ptr addrspace(3) %bar_addr)
+
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.{cg1,cg2}.p0.i16(ptr %bar_addr, i16 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.{cg1,cg2}.p0.i32(ptr %bar_addr, i32 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.{cg1,cg2}.p3.i16(ptr addrspace(3) %bar_addr, i16 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.{cg1,cg2}.p3.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
```
##### Overview:
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 26be050d4b22a..0a39e60f121e1 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -3087,22 +3087,12 @@ foreach cta_group = ["cg1", "cg2"] in {
foreach smem_a_read = ["", "smem_a_read_"] in {
def int_nvvm_tcgen05_commit_ # smem_a_read # cta_group : Intrinsic<[],
- [llvm_ptr_ty], // mbar_ptr
- [IntrConvergent, IntrInaccessibleMemOrArgMemOnly,
- NoCapture<ArgIndex<0>>]>;
-
- def int_nvvm_tcgen05_commit_ # smem_a_read # shared_ # cta_group : Intrinsic<[],
- [llvm_shared_ptr_ty], // mbar_ptr
+ [llvm_anyptr_ty], // mbar_ptr
[IntrConvergent, IntrInaccessibleMemOrArgMemOnly,
NoCapture<ArgIndex<0>>]>;
def int_nvvm_tcgen05_commit_ # smem_a_read # mc_ # cta_group : Intrinsic<[],
- [llvm_ptr_ty, llvm_anyint_ty], // mbar_ptr, cta_mask[16 or 32 bits]
- [IntrConvergent, IntrInaccessibleMemOrArgMemOnly,
- NoCapture<ArgIndex<0>>]>;
-
- def int_nvvm_tcgen05_commit_ # smem_a_read # mc_shared_ # cta_group : Intrinsic<[],
- [llvm_shared_ptr_ty, llvm_anyint_ty], // mbar_ptr, cta_mask[16 or 32 bits]
+ [llvm_anyptr_ty, llvm_anyint_ty], // mbar_ptr, cta_mask[16 or 32 bits]
[IntrConvergent, IntrInaccessibleMemOrArgMemOnly,
NoCapture<ArgIndex<0>>]>;
}
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index 4502759417c5a..b97785335d20f 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -1217,6 +1217,31 @@ static Intrinsic::ID shouldUpgradeNVPTXSharedClusterIntrinsic(Function *F,
return Intrinsic::not_intrinsic;
}
+static Intrinsic::ID
+shouldUpgradeNVPTXTcgen05CommitSharedIntrinsic(Function *F, StringRef Name) {
+ if (!Name.consume_front("tcgen05.commit."))
+ return Intrinsic::not_intrinsic;
+
+ if (Name.consume_front("shared."))
+ return StringSwitch<Intrinsic::ID>(Name)
+ .Case("cg1", Intrinsic::nvvm_tcgen05_commit_cg1)
+ .Case("cg2", Intrinsic::nvvm_tcgen05_commit_cg2)
+ .Default(Intrinsic::not_intrinsic);
+
+ if (Name.consume_front("mc.shared.")) {
+ // Only upgrade older i16 mc variants.
+ if (!F->getArg(1)->getType()->isIntegerTy(16))
+ return Intrinsic::not_intrinsic;
+
+ return StringSwitch<Intrinsic::ID>(Name)
+ .Case("cg1", Intrinsic::nvvm_tcgen05_commit_mc_cg1)
+ .Case("cg2", Intrinsic::nvvm_tcgen05_commit_mc_cg2)
+ .Default(Intrinsic::not_intrinsic);
+ }
+
+ return Intrinsic::not_intrinsic;
+}
+
static Intrinsic::ID shouldUpgradeNVPTXBF16Intrinsic(StringRef Name) {
if (Name.consume_front("fma.rn."))
return StringSwitch<Intrinsic::ID>(Name)
@@ -1718,6 +1743,16 @@ static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn,
return true;
}
+ // Upgrade tcgen05.commit shared variants to anyptr intrinsics.
+ IID = shouldUpgradeNVPTXTcgen05CommitSharedIntrinsic(F, Name);
+ if (IID != Intrinsic::not_intrinsic) {
+ rename(F);
+ NewFn = Intrinsic::getOrInsertDeclaration(
+ F->getParent(), IID, F->getReturnType(),
+ F->getFunctionType()->params());
+ return true;
+ }
+
// Upgrade TMA copy G2S Intrinsics
IID = shouldUpgradeNVPTXTMAG2SIntrinsics(F, Name);
if (IID != Intrinsic::not_intrinsic) {
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 05f9a5592cde4..0205b964016ab 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -5710,13 +5710,13 @@ defm TCGEN05_RELINQ_CG2: TCGEN05_RELINQ_PERMIT_INTR<"2", int_nvvm_tcgen05_relinq
def tcgen05_wait_ld: NullaryInst<"tcgen05.wait::ld.sync.aligned", int_nvvm_tcgen05_wait_ld>;
def tcgen05_wait_st: NullaryInst<"tcgen05.wait::st.sync.aligned", int_nvvm_tcgen05_wait_st>;
-multiclass TCGEN05_COMMIT_INTR<string AS, string num> {
+multiclass TCGEN05_COMMIT_INTR<string num> {
defvar prefix = "tcgen05.commit.cta_group::" # num #
".mbarrier::arrive::one.shared::cluster";
defvar prefix_smem_a_read = prefix # ".sync_restrict::shared::read::mma::a";
defvar intr_prefix = "int_nvvm_tcgen05_commit";
- defvar intr_suffix = !if(!eq(AS, "shared"), "_shared", "") # "_cg" # num;
+ defvar intr_suffix = "_cg" # num;
defvar Intr = !cast<Intrinsic>(intr_prefix # intr_suffix);
defvar IntrMC = !cast<Intrinsic>(intr_prefix # "_mc" # intr_suffix);
defvar IntrSmemARead =
@@ -5778,10 +5778,8 @@ foreach src_fmt = ["", "b6x16_p32", "b4x16_p64"] in {
}
} // Predicates
-defm TCGEN05_COMMIT_CG1 : TCGEN05_COMMIT_INTR<"", "1">;
-defm TCGEN05_COMMIT_CG2 : TCGEN05_COMMIT_INTR<"", "2">;
-defm TCGEN05_COMMIT_S64_CG1 : TCGEN05_COMMIT_INTR<"shared", "1">;
-defm TCGEN05_COMMIT_S64_CG2 : TCGEN05_COMMIT_INTR<"shared", "2">;
+defm TCGEN05_COMMIT_CG1 : TCGEN05_COMMIT_INTR<"1">;
+defm TCGEN05_COMMIT_CG2 : TCGEN05_COMMIT_INTR<"2">;
let Predicates = [callSubtarget<"hasTcgen05ShiftSupport">] in {
multiclass TCGEN05_SHIFT_INTR<string num, Intrinsic Intr> {
diff --git a/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll b/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
index 5ffb1be055450..9db89f79bab71 100644
--- a/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
+++ b/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
@@ -115,6 +115,10 @@ declare ptr addrspace(3) @llvm.nvvm.mapa.shared.cluster(ptr addrspace(3), i32)
declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster(ptr addrspace(3), ptr addrspace(3), ptr addrspace(1), i32, i16, i64, i1, i1)
declare void @llvm.nvvm.cp.async.bulk.shared.cta.to.cluster(ptr addrspace(3), ptr addrspace(3), ptr addrspace(3), i32)
+declare void @llvm.nvvm.tcgen05.commit.cg1(ptr)
+declare void @llvm.nvvm.tcgen05.commit.cg2(ptr)
+declare void @llvm.nvvm.tcgen05.commit.shared.cg1(ptr addrspace(3))
+declare void @llvm.nvvm.tcgen05.commit.shared.cg2(ptr addrspace(3))
declare void @llvm.nvvm.tcgen05.commit.mc.cg1(ptr, i16)
declare void @llvm.nvvm.tcgen05.commit.mc.cg2(ptr, i16)
declare void @llvm.nvvm.tcgen05.commit.mc.shared.cg1(ptr addrspace(3), i16)
@@ -461,12 +465,25 @@ define void @nvvm_cp_async_bulk_intrinsics(ptr addrspace(3) %dst, ptr addrspace(
ret void
}
+; CHECK-LABEL: @nvvm_tcgen05_commit_intrinsics
+define void @nvvm_tcgen05_commit_intrinsics(ptr %bar, ptr addrspace(3) %bar_shared) {
+; CHECK: call void @llvm.nvvm.tcgen05.commit.cg1.p0(ptr %bar)
+; CHECK: call void @llvm.nvvm.tcgen05.commit.cg2.p0(ptr %bar)
+; CHECK: call void @llvm.nvvm.tcgen05.commit.cg1.p3(ptr addrspace(3) %bar_shared)
+; CHECK: call void @llvm.nvvm.tcgen05.commit.cg2.p3(ptr addrspace(3) %bar_shared)
+ call void @llvm.nvvm.tcgen05.commit.cg1(ptr %bar)
+ call void @llvm.nvvm.tcgen05.commit.cg2(ptr %bar)
+ call void @llvm.nvvm.tcgen05.commit.shared.cg1(ptr addrspace(3) %bar_shared)
+ call void @llvm.nvvm.tcgen05.commit.shared.cg2(ptr addrspace(3) %bar_shared)
+ ret void
+}
+
; CHECK-LABEL: @nvvm_tcgen05_commit_mc_16b_intrinsics
define void @nvvm_tcgen05_commit_mc_16b_intrinsics(ptr %bar, ptr addrspace(3) %bar_shared, i16 %cta_mask) {
-; CHECK: call void @llvm.nvvm.tcgen05.commit.mc.cg1.i16(ptr %bar, i16 %cta_mask)
-; CHECK: call void @llvm.nvvm.tcgen05.commit.mc.cg2.i16(ptr %bar, i16 %cta_mask)
-; CHECK: call void @llvm.nvvm.tcgen05.commit.mc.shared.cg1.i16(ptr addrspace(3) %bar_shared, i16 %cta_mask)
-; CHECK: call void @llvm.nvvm.tcgen05.commit.mc.shared.cg2.i16(ptr addrspace(3) %bar_shared, i16 %cta_mask)
+; CHECK: call void @llvm.nvvm.tcgen05.commit.mc.cg1.p0.i16(ptr %bar, i16 %cta_mask)
+; CHECK: call void @llvm.nvvm.tcgen05.commit.mc.cg2.p0.i16(ptr %bar, i16 %cta_mask)
+; CHECK: call void @llvm.nvvm.tcgen05.commit.mc.cg1.p3.i16(ptr addrspace(3) %bar_shared, i16 %cta_mask)
+; CHECK: call void @llvm.nvvm.tcgen05.commit.mc.cg2.p3.i16(ptr addrspace(3) %bar_shared, i16 %cta_mask)
call void @llvm.nvvm.tcgen05.commit.mc.cg1(ptr %bar, i16 %cta_mask)
call void @llvm.nvvm.tcgen05.commit.mc.cg2(ptr %bar, i16 %cta_mask)
call void @llvm.nvvm.tcgen05.commit.mc.shared.cg1(ptr addrspace(3) %bar_shared, i16 %cta_mask)
diff --git a/llvm/test/CodeGen/NVPTX/tcgen05-commit-sm107.ll b/llvm/test/CodeGen/NVPTX/tcgen05-commit-sm107.ll
index 0b54e735ef3a9..8750e9e006865 100644
--- a/llvm/test/CodeGen/NVPTX/tcgen05-commit-sm107.ll
+++ b/llvm/test/CodeGen/NVPTX/tcgen05-commit-sm107.ll
@@ -5,10 +5,10 @@
; RUN: %if ptxas-sm_107f && ptxas-isa-9.4 %{ llc < %s -march=nvptx64 -mcpu=sm_107f | %ptxas-verify -arch=sm_107f %}
; RUN: %if ptxas-sm_107f && ptxas-isa-9.4 %{ llc < %s -march=nvptx64 -mcpu=sm_107f --nvptx-short-ptr | %ptxas-verify -arch=sm_107f %}
-declare void @llvm.nvvm.tcgen05.commit.smem.a.read.cg1(ptr %bar_addr)
-declare void @llvm.nvvm.tcgen05.commit.smem.a.read.cg2(ptr %bar_addr)
-declare void @llvm.nvvm.tcgen05.commit.smem.a.read.shared.cg1(ptr addrspace(3) %bar_addr)
-declare void @llvm.nvvm.tcgen05.commit.smem.a.read.shared.cg2(ptr addrspace(3) %bar_addr)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.cg1.p0(ptr %bar_addr)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.cg2.p0(ptr %bar_addr)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.cg1.p3(ptr addrspace(3) %bar_addr)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.cg2.p3(ptr addrspace(3) %bar_addr)
define void @test_tcgen05_commit_smem_a_read_cg1(ptr %bar_addr) {
; CHECK_PTX64-LABEL: test_tcgen05_commit_smem_a_read_cg1(
@@ -28,7 +28,7 @@ define void @test_tcgen05_commit_smem_a_read_cg1(ptr %bar_addr) {
; CHECK_PTX64_SHARED32-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_smem_a_read_cg1_param_0];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.b64 [%rd1];
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.commit.smem.a.read.cg1(ptr %bar_addr)
+ call void @llvm.nvvm.tcgen05.commit.smem.a.read.cg1.p0(ptr %bar_addr)
ret void
}
@@ -51,7 +51,7 @@ define void @test_tcgen05_commit_smem_a_read_cg2(ptr %bar_addr) {
; CHECK_PTX64_SHARED32-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_smem_a_read_cg2_param_0];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.b64 [%rd1];
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.commit.smem.a.read.cg2(ptr %bar_addr)
+ call void @llvm.nvvm.tcgen05.commit.smem.a.read.cg2.p0(ptr %bar_addr)
ret void
}
@@ -74,7 +74,7 @@ define void @test_tcgen05_commit_smem_a_read_shared_cg1(ptr addrspace(3) %bar_ad
; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_smem_a_read_shared_cg1_param_0];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.b64 [%r1];
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.commit.smem.a.read.shared.cg1(ptr addrspace(3) %bar_addr)
+ call void @llvm.nvvm.tcgen05.commit.smem.a.read.cg1.p3(ptr addrspace(3) %bar_addr)
ret void
}
@@ -97,15 +97,15 @@ define void @test_tcgen05_commit_smem_a_read_shared_cg2(ptr addrspace(3) %bar_ad
; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_smem_a_read_shared_cg2_param_0];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.b64 [%r1];
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.commit.smem.a.read.shared.cg2(ptr addrspace(3) %bar_addr)
+ call void @llvm.nvvm.tcgen05.commit.smem.a.read.cg2.p3(ptr addrspace(3) %bar_addr)
ret void
}
-declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg1.i16(ptr %bar_addr, i16 %cta_mask)
-declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg2.i16(ptr %bar_addr, i16 %cta_mask)
-declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.shared.cg1.i16(ptr addrspace(3) %bar_addr, i16 %cta_mask)
-declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.shared.cg2.i16(ptr addrspace(3) %bar_addr, i16 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg1.p0.i16(ptr %bar_addr, i16 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg2.p0.i16(ptr %bar_addr, i16 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg1.p3.i16(ptr addrspace(3) %bar_addr, i16 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg2.p3.i16(ptr addrspace(3) %bar_addr, i16 %cta_mask)
define void @test_tcgen05_commit_smem_a_read_mc_cg1(ptr %bar_addr, i16 %cta_mask) {
; CHECK_PTX64-LABEL: test_tcgen05_commit_smem_a_read_mc_cg1(
@@ -129,7 +129,7 @@ define void @test_tcgen05_commit_smem_a_read_mc_cg1(ptr %bar_addr, i16 %cta_mask
; CHECK_PTX64_SHARED32-NEXT: ld.param.b16 %rs1, [test_tcgen05_commit_smem_a_read_mc_cg1_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.multicast::cluster.b64 [%rd1], %rs1;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg1.i16(ptr %bar_addr, i16 %cta_mask)
+ call void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg1.p0.i16(ptr %bar_addr, i16 %cta_mask)
ret void
}
@@ -156,7 +156,7 @@ define void @test_tcgen05_commit_smem_a_read_mc_cg2(ptr %bar_addr, i16 %cta_mask
; CHECK_PTX64_SHARED32-NEXT: ld.param.b16 %rs1, [test_tcgen05_commit_smem_a_read_mc_cg2_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.multicast::cluster.b64 [%rd1], %rs1;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg2.i16(ptr %bar_addr, i16 %cta_mask)
+ call void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg2.p0.i16(ptr %bar_addr, i16 %cta_mask)
ret void
}
@@ -183,7 +183,7 @@ define void @test_tcgen05_commit_smem_a_read_mc_shared_cg1(ptr addrspace(3) %bar
; CHECK_PTX64_SHARED32-NEXT: ld.param.b16 %rs1, [test_tcgen05_commit_smem_a_read_mc_shared_cg1_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.multicast::cluster.b64 [%r1], %rs1;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.shared.cg1.i16(ptr addrspace(3) %bar_addr, i16 %cta_mask)
+ call void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg1.p3.i16(ptr addrspace(3) %bar_addr, i16 %cta_mask)
ret void
}
@@ -210,15 +210,15 @@ define void @test_tcgen05_commit_smem_a_read_mc_shared_cg2(ptr addrspace(3) %bar
; CHECK_PTX64_SHARED32-NEXT: ld.param.b16 %rs1, [test_tcgen05_commit_smem_a_read_mc_shared_cg2_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.multicast::cluster.b64 [%r1], %rs1;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.shared.cg2.i16(ptr addrspace(3) %bar_addr, i16 %cta_mask)
+ call void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg2.p3.i16(ptr addrspace(3) %bar_addr, i16 %cta_mask)
ret void
}
-declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg1.i32(ptr %bar_addr, i32 %cta_mask)
-declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg2.i32(ptr %bar_addr, i32 %cta_mask)
-declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.shared.cg1.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
-declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.shared.cg2.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg1.p0.i32(ptr %bar_addr, i32 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg2.p0.i32(ptr %bar_addr, i32 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg1.p3.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg2.p3.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
define void @test_tcgen05_commit_smem_a_read_mc_cg1_i32(ptr %bar_addr, i32 %cta_mask) {
; CHECK_PTX64-LABEL: test_tcgen05_commit_smem_a_read_mc_cg1_i32(
@@ -242,7 +242,7 @@ define void @test_tcgen05_commit_smem_a_read_mc_cg1_i32(ptr %bar_addr, i32 %cta_
; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_smem_a_read_mc_cg1_i32_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.multicast::cluster::32b.b64 [%rd1], %r1;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg1.i32(ptr %bar_addr, i32 %cta_mask)
+ call void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg1.p0.i32(ptr %bar_addr, i32 %cta_mask)
ret void
}
@@ -269,7 +269,7 @@ define void @test_tcgen05_commit_smem_a_read_mc_cg2_i32(ptr %bar_addr, i32 %cta_
; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_smem_a_read_mc_cg2_i32_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.multicast::cluster::32b.b64 [%rd1], %r1;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg2.i32(ptr %bar_addr, i32 %cta_mask)
+ call void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg2.p0.i32(ptr %bar_addr, i32 %cta_mask)
ret void
}
@@ -295,7 +295,7 @@ define void @test_tcgen05_commit_smem_a_read_mc_shared_cg1_i32(ptr addrspace(3)
; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r2, [test_tcgen05_commit_smem_a_read_mc_shared_cg1_i32_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.multicast::cluster::32b.b64 [%r1], %r2;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.shared.cg1.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
+ call void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg1.p3.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
ret void
}
@@ -321,15 +321,15 @@ define void @test_tcgen05_commit_smem_a_read_mc_shared_cg2_i32(ptr addrspace(3)
; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r2, [test_tcgen05_commit_smem_a_read_mc_shared_cg2_i32_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.sync_restrict::shared::read::mma::a.multicast::cluster::32b.b64 [%r1], %r2;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.shared.cg2.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
+ call void @llvm.nvvm.tcgen05.commit.smem.a.read.mc.cg2.p3.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
ret void
}
-declare void @llvm.nvvm.tcgen05.commit.mc.cg1.i32(ptr %bar_addr, i32 %cta_mask)
-declare void @llvm.nvvm.tcgen05.commit.mc.cg2.i32(ptr %bar_addr, i32 %cta_mask)
-declare void @llvm.nvvm.tcgen05.commit.mc.shared.cg1.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
-declare void @llvm.nvvm.tcgen05.commit.mc.shared.cg2.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.mc.cg1.p0.i32(ptr %bar_addr, i32 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.mc.cg2.p0.i32(ptr %bar_addr, i32 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.mc.cg1.p3.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.mc.cg2.p3.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
define void @test_tcgen05_commit_mc_cg1_i32(ptr %bar_addr, i32 %cta_mask) {
; CHECK_PTX64-LABEL: test_tcgen05_commit_mc_cg1_i32(
@@ -353,7 +353,7 @@ define void @test_tcgen05_commit_mc_cg1_i32(ptr %bar_addr, i32 %cta_mask) {
; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_mc_cg1_i32_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.multicast::cluster::32b.b64 [%rd1], %r1;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.commit.mc.cg1.i32(ptr %bar_addr, i32 %cta_mask)
+ call void @llvm.nvvm.tcgen05.commit.mc.cg1.p0.i32(ptr %bar_addr, i32 %cta_mask)
ret void
}
@@ -380,7 +380,7 @@ define void @test_tcgen05_commit_mc_cg2_i32(ptr %bar_addr, i32 %cta_mask) {
; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_mc_cg2_i32_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.multicast::cluster::32b.b64 [%rd1], %r1;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.commit.mc.cg2.i32(ptr %bar_addr, i32 %cta_mask)
+ call void @llvm.nvvm.tcgen05.commit.mc.cg2.p0.i32(ptr %bar_addr, i32 %cta_mask)
ret void
}
@@ -406,7 +406,7 @@ define void @test_tcgen05_commit_mc_shared_cg1_i32(ptr addrspace(3) %bar_addr, i
; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r2, [test_tcgen05_commit_mc_shared_cg1_i32_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.multicast::cluster::32b.b64 [%r1], %r2;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.commit.mc.shared.cg1.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
+ call void @llvm.nvvm.tcgen05.commit.mc.cg1.p3.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
ret void
}
@@ -432,7 +432,7 @@ define void @test_tcgen05_commit_mc_shared_cg2_i32(ptr addrspace(3) %bar_addr, i
; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r2, [test_tcgen05_commit_mc_shared_cg2_i32_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.multicast::cluster::32b.b64 [%r1], %r2;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.commit.mc.shared.cg2.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
+ call void @llvm.nvvm.tcgen05.commit.mc.cg2.p3.i32(ptr addrspace(3) %bar_addr, i32 %cta_mask)
ret void
}
diff --git a/llvm/test/CodeGen/NVPTX/tcgen05-commit.ll b/llvm/test/CodeGen/NVPTX/tcgen05-commit.ll
index 29b130f8cf7c3..254533fd09d18 100644
--- a/llvm/test/CodeGen/NVPTX/tcgen05-commit.ll
+++ b/llvm/test/CodeGen/NVPTX/tcgen05-commit.ll
@@ -10,10 +10,10 @@
; RUN: %if ptxas-sm_100f && ptxas-isa-8.8 %{ llc < %s -march=nvptx64 -mcpu=sm_100f -mattr=+ptx88 | %ptxas-verify -arch=sm_100f %}
; RUN: %if ptxas-sm_110f && ptxas-isa-9.0 %{ llc < %s -march=nvptx64 -mcpu=sm_110f -mattr=+ptx90 | %ptxas-verify -arch=sm_110f %}
-declare void @llvm.nvvm.tcgen05.commit.cg1(ptr %bar_addr)
-declare void @llvm.nvvm.tcgen05.commit.cg2(ptr %bar_addr)
-declare void @llvm.nvvm.tcgen05.commit.shared.cg1(ptr addrspace(3) %bar_addr)
-declare void @llvm.nvvm.tcgen05.commit.shared.cg2(ptr addrspace(3) %bar_addr)
+declare void @llvm.nvvm.tcgen05.commit.cg1.p0(ptr %bar_addr)
+declare void @llvm.nvvm.tcgen05.commit.cg2.p0(ptr %bar_addr)
+declare void @llvm.nvvm.tcgen05.commit.cg1.p3(ptr addrspace(3) %bar_addr)
+declare void @llvm.nvvm.tcgen05.commit.cg2.p3(ptr addrspace(3) %bar_addr)
define void @test_tcgen05_commit_cg1(ptr %bar_addr) {
; CHECK_PTX64-LABEL: test_tcgen05_commit_cg1(
@@ -33,7 +33,7 @@ define void @test_tcgen05_commit_cg1(ptr %bar_addr) {
; CHECK_PTX64_SHARED32-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_cg1_param_0];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.b64 [%rd1];
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.commit.cg1(ptr %bar_addr)
+ call void @llvm.nvvm.tcgen05.commit.cg1.p0(ptr %bar_addr)
ret void
}
@@ -56,7 +56,7 @@ define void @test_tcgen05_commit_cg2(ptr %bar_addr) {
; CHECK_PTX64_SHARED32-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_cg2_param_0];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.b64 [%rd1];
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.commit.cg2(ptr %bar_addr)
+ call void @llvm.nvvm.tcgen05.commit.cg2.p0(ptr %bar_addr)
ret void
}
@@ -79,7 +79,7 @@ define void @test_tcgen05_commit_shared_cg1(ptr addrspace(3) %bar_addr) {
; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_shared_cg1_param_0];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.b64 [%r1];
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.commit.shared.cg1(ptr addrspace(3) %bar_addr)
+ call void @llvm.nvvm.tcgen05.commit.cg1.p3(ptr addrspace(3) %bar_addr)
ret void
}
@@ -102,15 +102,15 @@ define void @test_tcgen05_commit_shared_cg2(ptr addrspace(3) %bar_addr) {
; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_shared_cg2_param_0];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.b64 [%r1];
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.commit.shared.cg2(ptr addrspace(3) %bar_addr)
+ call void @llvm.nvvm.tcgen05.commit.cg2.p3(ptr addrspace(3) %bar_addr)
ret void
}
-declare void @llvm.nvvm.tcgen05.commit.mc.cg1(ptr %bar_addr, i16 %cta_mask)
-declare void @llvm.nvvm.tcgen05.commit.mc.cg2(ptr %bar_addr, i16 %cta_mask)
-declare void @llvm.nvvm.tcgen05.commit.mc.shared.cg1(ptr addrspace(3) %bar_addr, i16 %cta_mask)
-declare void @llvm.nvvm.tcgen05.commit.mc.shared.cg2(ptr addrspace(3) %bar_addr, i16 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.mc.cg1.p0.i16(ptr %bar_addr, i16 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.mc.cg2.p0.i16(ptr %bar_addr, i16 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.mc.cg1.p3.i16(ptr addrspace(3) %bar_addr, i16 %cta_mask)
+declare void @llvm.nvvm.tcgen05.commit.mc.cg2.p3.i16(ptr addrspace(3) %bar_addr, i16 %cta_mask)
define void @test_tcgen05_commit_mc_cg1(ptr %bar_addr, i16 %cta_mask) {
; CHECK_PTX64-LABEL: test_tcgen05_commit_mc_cg1(
@@ -134,7 +134,7 @@ define void @test_tcgen05_commit_mc_cg1(ptr %bar_addr, i16 %cta_mask) {
; CHECK_PTX64_SHARED32-NEXT: ld.param.b16 %rs1, [test_tcgen05_commit_mc_cg1_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.multicast::cluster.b64 [%rd1], %rs1;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.commit.mc.cg1(ptr %bar_addr, i16 %cta_mask)
+ call void @llvm.nvvm.tcgen05.commit.mc.cg1.p0.i16(ptr %bar_addr, i16 %cta_mask)
ret void
}
@@ -160,7 +160,7 @@ define void @test_tcgen05_commit_mc_cg2(ptr %bar_addr, i16 %cta_mask) {
; CHECK_PTX64_SHARED32-NEXT: ld.param.b16 %rs1, [test_tcgen05_commit_mc_cg2_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.multicast::cluster.b64 [%rd1], %rs1;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.commit.mc.cg2(ptr %bar_addr, i16 %cta_mask)
+ call void @llvm.nvvm.tcgen05.commit.mc.cg2.p0.i16(ptr %bar_addr, i16 %cta_mask)
ret void
}
@@ -186,7 +186,7 @@ define void @test_tcgen05_commit_mc_shared_cg1(ptr addrspace(3) %bar_addr, i16 %
; CHECK_PTX64_SHARED32-NEXT: ld.param.b16 %rs1, [test_tcgen05_commit_mc_shared_cg1_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.multicast::cluster.b64 [%r1], %rs1;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.commit.mc.shared.cg1(ptr addrspace(3) %bar_addr, i16 %cta_mask)
+ call void @llvm.nvvm.tcgen05.commit.mc.cg1.p3.i16(ptr addrspace(3) %bar_addr, i16 %cta_mask)
ret void
}
@@ -212,6 +212,6 @@ define void @test_tcgen05_commit_mc_shared_cg2(ptr addrspace(3) %bar_addr, i16 %
; CHECK_PTX64_SHARED32-NEXT: ld.param.b16 %rs1, [test_tcgen05_commit_mc_shared_cg2_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.multicast::cluster.b64 [%r1], %rs1;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.commit.mc.shared.cg2(ptr addrspace(3) %bar_addr, i16 %cta_mask)
+ call void @llvm.nvvm.tcgen05.commit.mc.cg2.p3.i16(ptr addrspace(3) %bar_addr, i16 %cta_mask)
ret void
}
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index b29cba96a410f..6e98e681c651c 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -5103,28 +5103,24 @@ llvm::Intrinsic::ID Tcgen05DeallocOp::getIntrinsicIDAndArgs(
return id;
}
-#define TCGEN05_COMMIT_IMPL(cg, is_shared, mc) \
- is_shared ? llvm::Intrinsic::nvvm_tcgen05_commit##mc##_shared##_##cg \
- : llvm::Intrinsic::nvvm_tcgen05_commit##mc##_##cg
+#define TCGEN05_COMMIT_IMPL(cg, mc) \
+ llvm::Intrinsic::nvvm_tcgen05_commit##mc##_##cg
-#define GET_TCGEN05_COMMIT_ID(cta_group, is_shared, has_mc) \
- has_mc ? TCGEN05_COMMIT_IMPL(cta_group, is_shared, _mc) \
- : TCGEN05_COMMIT_IMPL(cta_group, is_shared, )
+#define GET_TCGEN05_COMMIT_ID(cta_group, has_mc) \
+ has_mc ? TCGEN05_COMMIT_IMPL(cta_group, _mc) \
+ : TCGEN05_COMMIT_IMPL(cta_group, )
llvm::Intrinsic::ID
Tcgen05CommitOp::getIntrinsicIDAndArgs(Operation &op,
LLVM::ModuleTranslation &mt,
llvm::SmallVector<llvm::Value *> &args) {
auto curOp = cast<NVVM::Tcgen05CommitOp>(op);
- unsigned as = llvm::cast<LLVM::LLVMPointerType>(curOp.getAddr().getType())
- .getAddressSpace();
- bool isShared = as == NVVMMemorySpace::Shared;
bool hasMulticast = static_cast<bool>(curOp.getMulticastMask());
bool is2CTAMode = curOp.getGroup() == CTAGroupKind::CTA_2;
- llvm::Intrinsic::ID id =
- is2CTAMode ? GET_TCGEN05_COMMIT_ID(cg2, isShared, hasMulticast)
- : GET_TCGEN05_COMMIT_ID(cg1, isShared, hasMulticast);
+ llvm::Intrinsic::ID id = is2CTAMode
+ ? GET_TCGEN05_COMMIT_ID(cg2, hasMulticast)
+ : GET_TCGEN05_COMMIT_ID(cg1, hasMulticast);
// Fill the Intrinsic Args
args.push_back(mt.lookupValue(curOp.getAddr()));
diff --git a/mlir/test/Target/LLVMIR/nvvm/tcgen05-commit.mlir b/mlir/test/Target/LLVMIR/nvvm/tcgen05-commit.mlir
index 6997f8a98b2fc..6ef6f9914ffb4 100644
--- a/mlir/test/Target/LLVMIR/nvvm/tcgen05-commit.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/tcgen05-commit.mlir
@@ -2,32 +2,32 @@
// CHECK-LABEL: @llvm_nvvm_tcgen05_commit_generic
llvm.func @llvm_nvvm_tcgen05_commit_generic(%barrier : !llvm.ptr, %cta_mask : i16) {
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.cg1(ptr %{{.*}})
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.cg1.p0(ptr %{{.*}})
nvvm.tcgen05.commit %barrier : !llvm.ptr
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.cg2(ptr %{{.*}})
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.cg2.p0(ptr %{{.*}})
nvvm.tcgen05.commit %barrier {group = #nvvm.cta_group<cta_2>} : !llvm.ptr
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.mc.cg1.i16(ptr %{{.*}}, i16 %{{.*}})
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.mc.cg1.p0.i16(ptr %{{.*}}, i16 %{{.*}})
nvvm.tcgen05.commit %barrier, multicast_mask = %cta_mask : !llvm.ptr, i16
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.mc.cg2.i16(ptr %{{.*}}, i16 %{{.*}})
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.mc.cg2.p0.i16(ptr %{{.*}}, i16 %{{.*}})
nvvm.tcgen05.commit %barrier, multicast_mask = %cta_mask {group = #nvvm.cta_group<cta_2>} : !llvm.ptr, i16
llvm.return
}
// CHECK-LABEL: @llvm_nvvm_tcgen05_commit_shared
llvm.func @llvm_nvvm_tcgen05_commit_shared(%barrier : !llvm.ptr<3>, %cta_mask : i16) {
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.shared.cg1(ptr addrspace(3) %{{.*}})
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.cg1.p3(ptr addrspace(3) %{{.*}})
nvvm.tcgen05.commit %barrier : !llvm.ptr<3>
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.shared.cg2(ptr addrspace(3) %{{.*}})
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.cg2.p3(ptr addrspace(3) %{{.*}})
nvvm.tcgen05.commit %barrier {group = #nvvm.cta_group<cta_2>} : !llvm.ptr<3>
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.mc.shared.cg1.i16(ptr addrspace(3) %{{.*}}, i16 %{{.*}})
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.mc.cg1.p3.i16(ptr addrspace(3) %{{.*}}, i16 %{{.*}})
nvvm.tcgen05.commit %barrier, multicast_mask = %cta_mask : !llvm.ptr<3>, i16
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.mc.shared.cg2.i16(ptr addrspace(3) %{{.*}}, i16 %{{.*}})
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.commit.mc.cg2.p3.i16(ptr addrspace(3) %{{.*}}, i16 %{{.*}})
nvvm.tcgen05.commit %barrier, multicast_mask = %cta_mask {group = #nvvm.cta_group<cta_2>} : !llvm.ptr<3>, i16
llvm.return
}
More information about the Mlir-commits
mailing list