[Mlir-commits] [llvm] [mlir] [NVPTX] Add "exclusive" intrinsic variants for tcgen05.alloc/dealloc (PR #216016)
Dharuni R Acharya
llvmlistbot at llvm.org
Mon Aug 17 21:59:29 PDT 2026
https://github.com/DharuniRAcharya updated https://github.com/llvm/llvm-project/pull/216016
>From 61fc8c54ee4cebf172c4410f02592de76ac97a4d Mon Sep 17 00:00:00 2001
From: DharuniRAcharya <dharunira at nvidia.com>
Date: Thu, 13 Aug 2026 11:27:06 +0000
Subject: [PATCH 1/6] [NVPTX] Add "exclusive" intrinsic variants for
tcgen05.alloc/dealloc
This patch adds exclusive intrinsic variants for tcgen05.alloc and tcgen05.dealloc,
which allow exclusive ownership to the allocation.
- tcgen05.alloc{.exclusive}.cta_group.sync.aligned{.shared::cta}.b32
- tcgen05.dealloc{.exclusive}.cta_group.sync.aligned.b32
PTX ISA Reference: https://docs.nvidia.com/cuda/developer-preview/13.4/parallel-thread-execution/index.html#tcgen05-instructions-tcgen05-alloc-dealloc-relinquish-alloc-permit
Signed-off-by: DharuniRAcharya <dharunira at nvidia.com>
---
llvm/docs/NVPTXUsage.md | 38 ++--
llvm/include/llvm/IR/IntrinsicsNVVM.td | 30 ++--
llvm/lib/IR/AutoUpgrade.cpp | 21 +++
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 12 ++
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 91 ++++++++--
.../Assembler/auto_upgrade_nvvm_intrinsics.ll | 18 ++
.../NVPTX/tcgen05-alloc-dealloc-exclusive.ll | 165 ++++++++++++++++++
llvm/test/CodeGen/NVPTX/tcgen05-alloc.ll | 16 +-
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 2 +-
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp | 14 +-
.../Target/LLVMIR/nvvm/tcgen05-alloc.mlir | 8 +-
11 files changed, 342 insertions(+), 73 deletions(-)
create mode 100644 llvm/test/CodeGen/NVPTX/tcgen05-alloc-dealloc-exclusive.ll
diff --git a/llvm/docs/NVPTXUsage.md b/llvm/docs/NVPTXUsage.md
index 8924a44e43a8e..98a987fcab4e6 100644
--- a/llvm/docs/NVPTXUsage.md
+++ b/llvm/docs/NVPTXUsage.md
@@ -2798,24 +2798,27 @@ For more information on tensor-memory load/store instructions, refer to [PTX ISA
##### Syntax:
```llvm
-declare void @llvm.nvvm.tcgen05.alloc.cg1(ptr %dst, i32 %ncols)
-declare void @llvm.nvvm.tcgen05.alloc.cg2(ptr %dst, i32 %ncols)
-declare void @llvm.nvvm.tcgen05.alloc.shared.cg1(ptr addrspace(3) %dst, i32 %ncols)
-declare void @llvm.nvvm.tcgen05.alloc.shared.cg2(ptr addrspace(3) %dst, i32 %ncols)
+declare void @llvm.nvvm.tcgen05.alloc.{cg1,cg2}.p0(ptr %dst, i32 %ncols)
+declare void @llvm.nvvm.tcgen05.alloc.{cg1,cg2}.p3(ptr addrspace(3) %dst, i32 %ncols)
+
+declare void @llvm.nvvm.tcgen05.alloc.exclusive.{cg1,cg2}.p0(ptr %dst, i32 %ncols)
+declare void @llvm.nvvm.tcgen05.alloc.exclusive.{cg1,cg2}.p3(ptr addrspace(3) %dst, i32 %ncols)
```
##### Overview:
The '`@llvm.nvvm.tcgen05.alloc.*`' intrinsics correspond to the
-`tcgen05.alloc.cta_group*.sync.aligned.b32` family of PTX instructions.
-The `tcgen05.alloc` is a potentially blocking instruction which dynamically
-allocates the specified number of columns in the Tensor Memory and writes the
-address of the allocated Tensor Memory into shared memory at the location
-specified by `%dst`. The 32-bit operand `%ncols` specifies the number of
-columns to be allocated and it must be a power-of-two. The `.shared` variant
-explicitly uses shared memory address space for the `%dst` operand. The
-`.cg1` and `.cg2` variants generate `cta_group::1` and `cta_group::2`
-variants of the instruction respectively.
+`tcgen05.alloc{.exclusive}.cta_group*.sync.aligned.b32` family of PTX
+instructions. The `tcgen05.alloc` is a potentially blocking instruction which
+dynamically allocates the specified number of columns in the Tensor Memory
+and writes the address of the allocated Tensor Memory into shared memory at the
+location specified by `%dst`. The 32-bit operand `%ncols` specifies the number
+of columns to be allocated and it must be a power-of-two for non-exclusive
+allocations and it must be a multiple of 32 for exclusive allocations. The
+overloaded pointer argument may use generic or shared memory address space;
+a shared pointer emits the `.shared::cta` qualifier. The `.cg1` and `.cg2`
+variants generate `cta_group::1` and `cta_group::2` variants of the instruction
+respectively. The `.exclusive` variants claim ownership of the allocation permit.
For more information, refer to the [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory-allocation-and-management-instructions).
@@ -2824,19 +2827,22 @@ For more information, refer to the [PTX ISA](https://docs.nvidia.com/cuda/parall
##### Syntax:
```llvm
-declare void @llvm.nvvm.tcgen05.dealloc.cg1(ptr addrspace(6) %tmem_addr, i32 %ncols)
-declare void @llvm.nvvm.tcgen05.dealloc.cg2(ptr addrspace(6) %tmem_addr, i32 %ncols)
+declare void @llvm.nvvm.tcgen05.dealloc.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i32 %ncols)
+
+declare void @llvm.nvvm.tcgen05.dealloc.exclusive.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i32 %ncols)
```
##### Overview:
The '`@llvm.nvvm.tcgen05.dealloc.*`' intrinsics correspond to the
-`tcgen05.dealloc.*` set of PTX instructions. The `tcgen05.dealloc`
+`tcgen05.dealloc{.exclusive}.*` set of PTX instructions. The `tcgen05.dealloc`
instructions deallocates the Tensor Memory specified by the Tensor Memory
address `%tmem_addr`. The operand `%tmem_addr` must point to a previous
Tensor Memory allocation. The 32-bit operand `%ncols` specifies the number
of columns to be de-allocated. The `.cg1` and `.cg2` variants generate
`cta_group::1` and `cta_group::2` variants of the instruction respectively.
+Memory must be deallocated with `.exclusive` if and only if it was allocated
+with `.exclusive`.
For more information, refer to the [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory-allocation-and-management-instructions).
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 4357ad367d269..a9b746e05461b 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -3210,23 +3210,19 @@ def int_nvvm_griddepcontrol_wait : Intrinsic<[], [], [IntrNoMem, IntrHasSideEffe
// Tcgen05 alloc/dealloc related intrinsics
foreach cta_group = ["cg1", "cg2"] in {
- def int_nvvm_tcgen05_alloc_ # cta_group : Intrinsic<[],
- [llvm_ptr_ty, // dst_ptr
- llvm_i32_ty] , // num_columns
- [IntrConvergent, IntrInaccessibleMemOrArgMemOnly,
- WriteOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>]>;
-
- def int_nvvm_tcgen05_alloc_shared_ # cta_group : Intrinsic<[],
- [llvm_shared_ptr_ty, // dst_ptr
- llvm_i32_ty], // num_columns
- [IntrConvergent, IntrInaccessibleMemOrArgMemOnly,
- WriteOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>]>;
-
- def int_nvvm_tcgen05_dealloc_ # cta_group : Intrinsic<[],
- [llvm_tmem_ptr_ty, // tmem_addr
- llvm_i32_ty], // num_columns
- [IntrConvergent, IntrArgMemOnly,
- NoCapture<ArgIndex<0>>]>;
+ foreach exclusive = ["", "exclusive_"] in {
+ def int_nvvm_tcgen05_alloc_ # exclusive # cta_group : Intrinsic<[],
+ [llvm_anyptr_ty, // dst_ptr
+ llvm_i32_ty], // num_columns
+ [IntrConvergent, IntrInaccessibleMemOrArgMemOnly,
+ WriteOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>]>;
+
+ def int_nvvm_tcgen05_dealloc_ # exclusive # cta_group : Intrinsic<[],
+ [llvm_tmem_ptr_ty, // tmem_addr
+ llvm_i32_ty], // num_columns
+ [IntrConvergent, IntrArgMemOnly,
+ NoCapture<ArgIndex<0>>]>;
+ }
def int_nvvm_tcgen05_relinq_alloc_permit_ # cta_group : Intrinsic<[], [],
[IntrConvergent, IntrInaccessibleMemOnly]>;
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index 5b6f50df5d4d0..22a148c5512d7 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -1279,6 +1279,17 @@ shouldUpgradeNVPTXTcgen05CommitSharedIntrinsic(Function *F, StringRef Name) {
return Intrinsic::not_intrinsic;
}
+static Intrinsic::ID
+shouldUpgradeNVPTXTcgen05AllocSharedIntrinsic(StringRef Name) {
+ if (!Name.consume_front("tcgen05.alloc.shared."))
+ return Intrinsic::not_intrinsic;
+
+ return StringSwitch<Intrinsic::ID>(Name)
+ .Case("cg1", Intrinsic::nvvm_tcgen05_alloc_cg1)
+ .Case("cg2", Intrinsic::nvvm_tcgen05_alloc_cg2)
+ .Default(Intrinsic::not_intrinsic);
+}
+
static Intrinsic::ID shouldUpgradeNVPTXBF16Intrinsic(StringRef Name) {
if (Name.consume_front("fma.rn."))
return StringSwitch<Intrinsic::ID>(Name)
@@ -1811,6 +1822,16 @@ static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn,
return true;
}
+ // Upgrade tcgen05.alloc shared variants to anyptr intrinsics.
+ IID = shouldUpgradeNVPTXTcgen05AllocSharedIntrinsic(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/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index e788b0e44041f..be3078ed835fd 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -5546,6 +5546,18 @@ void NVPTXTargetLowering::getTgtMemIntrinsic(
Infos.push_back(Info);
return;
}
+ case Intrinsic::nvvm_tcgen05_alloc_cg1:
+ case Intrinsic::nvvm_tcgen05_alloc_cg2:
+ case Intrinsic::nvvm_tcgen05_alloc_exclusive_cg1:
+ case Intrinsic::nvvm_tcgen05_alloc_exclusive_cg2:
+ Info.opc = ISD::INTRINSIC_VOID;
+ Info.memVT = MVT::i32;
+ Info.ptrVal = I.getArgOperand(0);
+ Info.offset = 0;
+ Info.flags = MachineMemOperand::MOStore;
+ Info.align = Align(4);
+ Infos.push_back(Info);
+ return;
}
}
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 02bcfcb6af926..2f9815bd64679 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -5844,30 +5844,85 @@ let Predicates = [SM90] in {
def EXIT : NullaryInst<"exit", int_nvvm_exit>;
// Tcgen05 intrinsics
-let isConvergent = true in {
-let Predicates = [hasTcgen05InstSupport] in {
-multiclass TCGEN05_ALLOC_INTR<string AS, string num, Intrinsic Intr> {
- def "" : BasicNVPTXInst<(outs),
- (ins ADDR:$dst, B32:$ncols),
- "tcgen05.alloc.cta_group::" # num # ".sync.aligned" # AS # ".b32",
- [(Intr addr:$dst, B32:$ncols)]>;
+multiclass TCGEN05_ALLOC_PATFRAGS<Intrinsic Intr> {
+ defvar frag_pat = (Intr node:$dst, node:$ncols);
+ def _GENERIC
+ : PatFrag<!setdagop(frag_pat, ops), frag_pat, AS_match.generic>;
+ def _SHARED
+ : PatFrag<!setdagop(frag_pat, ops), frag_pat, AS_match.shared>;
+}
+
+defm TCGEN05_ALLOC_CG1_PAT
+ : TCGEN05_ALLOC_PATFRAGS<int_nvvm_tcgen05_alloc_cg1>;
+defm TCGEN05_ALLOC_CG2_PAT
+ : TCGEN05_ALLOC_PATFRAGS<int_nvvm_tcgen05_alloc_cg2>;
+defm TCGEN05_ALLOC_EXCLUSIVE_CG1_PAT
+ : TCGEN05_ALLOC_PATFRAGS<int_nvvm_tcgen05_alloc_exclusive_cg1>;
+defm TCGEN05_ALLOC_EXCLUSIVE_CG2_PAT
+ : TCGEN05_ALLOC_PATFRAGS<int_nvvm_tcgen05_alloc_exclusive_cg2>;
+
+multiclass TCGEN05_ALLOC_INTR<string num, PatFrag GenericPat,
+ PatFrag SharedPat,
+ bit IsExclusive = 0> {
+ defvar exclusive = !if(IsExclusive, ".exclusive", "");
+ defvar Pred = !if(IsExclusive,
+ [hasTcgen05InstSupport, PTX94],
+ [hasTcgen05InstSupport]);
+
+ let isConvergent = true, Predicates = Pred in {
+ def _GENERIC : BasicNVPTXInst<(outs), (ins ADDR:$dst, B32:$ncols),
+ "tcgen05.alloc" # exclusive # ".cta_group::" # num #
+ ".sync.aligned.b32",
+ [(GenericPat addr:$dst, B32:$ncols)]>;
+ def _SHARED : BasicNVPTXInst<(outs), (ins ADDR:$dst, B32:$ncols),
+ "tcgen05.alloc" # exclusive # ".cta_group::" # num #
+ ".sync.aligned.shared::cta.b32",
+ [(SharedPat addr:$dst, B32:$ncols)]>;
+ }
}
-defm TCGEN05_ALLOC_CG1 : TCGEN05_ALLOC_INTR<"", "1", int_nvvm_tcgen05_alloc_cg1>;
-defm TCGEN05_ALLOC_CG2 : TCGEN05_ALLOC_INTR<"", "2", int_nvvm_tcgen05_alloc_cg2>;
-
-defm TCGEN05_ALLOC_S64_CG1 : TCGEN05_ALLOC_INTR<".shared::cta", "1", int_nvvm_tcgen05_alloc_shared_cg1>;
-defm TCGEN05_ALLOC_S64_CG2 : TCGEN05_ALLOC_INTR<".shared::cta", "2", int_nvvm_tcgen05_alloc_shared_cg2>;
-
-multiclass TCGEN05_DEALLOC_INTR<string num, Intrinsic Intr> {
- def "" : BasicNVPTXInst<(outs),
- (ins B32:$tmem_addr, B32:$ncols),
- "tcgen05.dealloc.cta_group::" # num # ".sync.aligned.b32",
- [(Intr B32:$tmem_addr, B32:$ncols)]>;
+defm TCGEN05_ALLOC_CG1
+ : TCGEN05_ALLOC_INTR<"1", TCGEN05_ALLOC_CG1_PAT_GENERIC,
+ TCGEN05_ALLOC_CG1_PAT_SHARED>;
+defm TCGEN05_ALLOC_CG2
+ : TCGEN05_ALLOC_INTR<"2", TCGEN05_ALLOC_CG2_PAT_GENERIC,
+ TCGEN05_ALLOC_CG2_PAT_SHARED>;
+
+defm TCGEN05_ALLOC_EXCLUSIVE_CG1
+ : TCGEN05_ALLOC_INTR<"1", TCGEN05_ALLOC_EXCLUSIVE_CG1_PAT_GENERIC,
+ TCGEN05_ALLOC_EXCLUSIVE_CG1_PAT_SHARED,
+ /*IsExclusive=*/1>;
+defm TCGEN05_ALLOC_EXCLUSIVE_CG2
+ : TCGEN05_ALLOC_INTR<"2", TCGEN05_ALLOC_EXCLUSIVE_CG2_PAT_GENERIC,
+ TCGEN05_ALLOC_EXCLUSIVE_CG2_PAT_SHARED,
+ /*IsExclusive=*/1>;
+
+multiclass TCGEN05_DEALLOC_INTR<string num, Intrinsic Intr,
+ bit IsExclusive = 0> {
+ defvar exclusive = !if(IsExclusive, ".exclusive", "");
+ defvar Pred = !if(IsExclusive,
+ [hasTcgen05InstSupport, PTX94],
+ [hasTcgen05InstSupport]);
+ let isConvergent = true, Predicates = Pred in {
+ def "" : BasicNVPTXInst<(outs),
+ (ins B32:$tmem_addr, B32:$ncols),
+ "tcgen05.dealloc" # exclusive # ".cta_group::" # num #
+ ".sync.aligned.b32",
+ [(Intr B32:$tmem_addr, B32:$ncols)]>;
+ }
}
defm TCGEN05_DEALLOC_CG1: TCGEN05_DEALLOC_INTR<"1", int_nvvm_tcgen05_dealloc_cg1>;
defm TCGEN05_DEALLOC_CG2: TCGEN05_DEALLOC_INTR<"2", int_nvvm_tcgen05_dealloc_cg2>;
+defm TCGEN05_DEALLOC_EXCLUSIVE_CG1
+ : TCGEN05_DEALLOC_INTR<"1", int_nvvm_tcgen05_dealloc_exclusive_cg1,
+ /*IsExclusive=*/1>;
+defm TCGEN05_DEALLOC_EXCLUSIVE_CG2
+ : TCGEN05_DEALLOC_INTR<"2", int_nvvm_tcgen05_dealloc_exclusive_cg2,
+ /*IsExclusive=*/1>;
+
+let isConvergent = true in {
+let Predicates = [hasTcgen05InstSupport] in {
multiclass TCGEN05_RELINQ_PERMIT_INTR<string num, Intrinsic Intr> {
def "" : NullaryInst<"tcgen05.relinquish_alloc_permit.cta_group::" # num # ".sync.aligned", Intr>;
}
diff --git a/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll b/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
index 244ced4825517..5813e3bda1c49 100644
--- a/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
+++ b/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
@@ -124,6 +124,11 @@ 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.tcgen05.alloc.cg1(ptr, i32)
+declare void @llvm.nvvm.tcgen05.alloc.cg2(ptr, i32)
+declare void @llvm.nvvm.tcgen05.alloc.shared.cg1(ptr addrspace(3), i32)
+declare void @llvm.nvvm.tcgen05.alloc.shared.cg2(ptr addrspace(3), i32)
+
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);
@@ -503,6 +508,19 @@ define void @nvvm_tcgen05_commit_intrinsics(ptr %bar, ptr addrspace(3) %bar_shar
ret void
}
+; CHECK-LABEL: @nvvm_tcgen05_alloc_intrinsics
+define void @nvvm_tcgen05_alloc_intrinsics(ptr %dst, ptr addrspace(3) %dst_shared, i32 %ncols) {
+; CHECK: call void @llvm.nvvm.tcgen05.alloc.cg1.p0(ptr %dst, i32 %ncols)
+; CHECK: call void @llvm.nvvm.tcgen05.alloc.cg2.p0(ptr %dst, i32 %ncols)
+; CHECK: call void @llvm.nvvm.tcgen05.alloc.cg1.p3(ptr addrspace(3) %dst_shared, i32 %ncols)
+; CHECK: call void @llvm.nvvm.tcgen05.alloc.cg2.p3(ptr addrspace(3) %dst_shared, i32 %ncols)
+ call void @llvm.nvvm.tcgen05.alloc.cg1(ptr %dst, i32 %ncols)
+ call void @llvm.nvvm.tcgen05.alloc.cg2(ptr %dst, i32 %ncols)
+ call void @llvm.nvvm.tcgen05.alloc.shared.cg1(ptr addrspace(3) %dst_shared, i32 %ncols)
+ call void @llvm.nvvm.tcgen05.alloc.shared.cg2(ptr addrspace(3) %dst_shared, i32 %ncols)
+ 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.p0.i16(ptr %bar, i16 %cta_mask)
diff --git a/llvm/test/CodeGen/NVPTX/tcgen05-alloc-dealloc-exclusive.ll b/llvm/test/CodeGen/NVPTX/tcgen05-alloc-dealloc-exclusive.ll
new file mode 100644
index 0000000000000..deb477ce41e54
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/tcgen05-alloc-dealloc-exclusive.ll
@@ -0,0 +1,165 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_100f -mattr=+ptx94 | FileCheck --check-prefixes=CHECK_PTX64 %s
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_100f -mattr=+ptx94 -target-abi=shortptr | FileCheck --check-prefixes=CHECK_PTX64_SHARED32 %s
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_110f -mattr=+ptx94 | FileCheck --check-prefixes=CHECK_PTX64 %s
+; RUN: %if ptxas-sm_100f && ptxas-isa-9.4 %{ llc < %s -march=nvptx64 -mcpu=sm_100f -mattr=+ptx94 | %ptxas-verify -arch=sm_100f %}
+; RUN: %if ptxas-sm_100f && ptxas-isa-9.4 %{ llc < %s -march=nvptx64 -mcpu=sm_100f -mattr=+ptx94 -target-abi=shortptr | %ptxas-verify -arch=sm_100f %}
+; RUN: %if ptxas-sm_110f && ptxas-isa-9.4 %{ llc < %s -march=nvptx64 -mcpu=sm_110f -mattr=+ptx94 | %ptxas-verify -arch=sm_110f %}
+
+declare void @llvm.nvvm.tcgen05.alloc.exclusive.cg1.p0(ptr %addr, i32 %ncols)
+declare void @llvm.nvvm.tcgen05.alloc.exclusive.cg2.p0(ptr %addr, i32 %ncols)
+declare void @llvm.nvvm.tcgen05.alloc.exclusive.cg1.p3(ptr addrspace(3) %addr, i32 %ncols)
+declare void @llvm.nvvm.tcgen05.alloc.exclusive.cg2.p3(ptr addrspace(3) %addr, i32 %ncols)
+
+define void @test_tcgen05_alloc_exclusive_cg1(ptr %addr, i32 %ncols) {
+; CHECK_PTX64-LABEL: test_tcgen05_alloc_exclusive_cg1(
+; 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::func.b64 %rd1, [test_tcgen05_alloc_exclusive_cg1_param_0];
+; CHECK_PTX64-NEXT: ld.param::func.b32 %r1, [test_tcgen05_alloc_exclusive_cg1_param_1];
+; CHECK_PTX64-NEXT: tcgen05.alloc.exclusive.cta_group::1.sync.aligned.b32 [%rd1], %r1;
+; CHECK_PTX64-NEXT: ret;
+;
+; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_alloc_exclusive_cg1(
+; 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::func.b64 %rd1, [test_tcgen05_alloc_exclusive_cg1_param_0];
+; CHECK_PTX64_SHARED32-NEXT: ld.param::func.b32 %r1, [test_tcgen05_alloc_exclusive_cg1_param_1];
+; CHECK_PTX64_SHARED32-NEXT: tcgen05.alloc.exclusive.cta_group::1.sync.aligned.b32 [%rd1], %r1;
+; CHECK_PTX64_SHARED32-NEXT: ret;
+ call void @llvm.nvvm.tcgen05.alloc.exclusive.cg1.p0(ptr %addr, i32 %ncols)
+ ret void
+}
+
+define void @test_tcgen05_alloc_exclusive_cg2(ptr %addr, i32 %ncols) {
+; CHECK_PTX64-LABEL: test_tcgen05_alloc_exclusive_cg2(
+; 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::func.b64 %rd1, [test_tcgen05_alloc_exclusive_cg2_param_0];
+; CHECK_PTX64-NEXT: ld.param::func.b32 %r1, [test_tcgen05_alloc_exclusive_cg2_param_1];
+; CHECK_PTX64-NEXT: tcgen05.alloc.exclusive.cta_group::2.sync.aligned.b32 [%rd1], %r1;
+; CHECK_PTX64-NEXT: ret;
+;
+; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_alloc_exclusive_cg2(
+; 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::func.b64 %rd1, [test_tcgen05_alloc_exclusive_cg2_param_0];
+; CHECK_PTX64_SHARED32-NEXT: ld.param::func.b32 %r1, [test_tcgen05_alloc_exclusive_cg2_param_1];
+; CHECK_PTX64_SHARED32-NEXT: tcgen05.alloc.exclusive.cta_group::2.sync.aligned.b32 [%rd1], %r1;
+; CHECK_PTX64_SHARED32-NEXT: ret;
+ call void @llvm.nvvm.tcgen05.alloc.exclusive.cg2.p0(ptr %addr, i32 %ncols)
+ ret void
+}
+
+define void @test_tcgen05_alloc_exclusive_shared_cg1(ptr addrspace(3) %addr, i32 %ncols) {
+; CHECK_PTX64-LABEL: test_tcgen05_alloc_exclusive_shared_cg1(
+; 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::func.b64 %rd1, [test_tcgen05_alloc_exclusive_shared_cg1_param_0];
+; CHECK_PTX64-NEXT: ld.param::func.b32 %r1, [test_tcgen05_alloc_exclusive_shared_cg1_param_1];
+; CHECK_PTX64-NEXT: tcgen05.alloc.exclusive.cta_group::1.sync.aligned.shared::cta.b32 [%rd1], %r1;
+; CHECK_PTX64-NEXT: ret;
+;
+; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_alloc_exclusive_shared_cg1(
+; 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::func.b32 %r1, [test_tcgen05_alloc_exclusive_shared_cg1_param_0];
+; CHECK_PTX64_SHARED32-NEXT: ld.param::func.b32 %r2, [test_tcgen05_alloc_exclusive_shared_cg1_param_1];
+; CHECK_PTX64_SHARED32-NEXT: tcgen05.alloc.exclusive.cta_group::1.sync.aligned.shared::cta.b32 [%r1], %r2;
+; CHECK_PTX64_SHARED32-NEXT: ret;
+ call void @llvm.nvvm.tcgen05.alloc.exclusive.cg1.p3(ptr addrspace(3) %addr, i32 %ncols)
+ ret void
+}
+
+define void @test_tcgen05_alloc_exclusive_shared_cg2(ptr addrspace(3) %addr, i32 %ncols) {
+; CHECK_PTX64-LABEL: test_tcgen05_alloc_exclusive_shared_cg2(
+; 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::func.b64 %rd1, [test_tcgen05_alloc_exclusive_shared_cg2_param_0];
+; CHECK_PTX64-NEXT: ld.param::func.b32 %r1, [test_tcgen05_alloc_exclusive_shared_cg2_param_1];
+; CHECK_PTX64-NEXT: tcgen05.alloc.exclusive.cta_group::2.sync.aligned.shared::cta.b32 [%rd1], %r1;
+; CHECK_PTX64-NEXT: ret;
+;
+; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_alloc_exclusive_shared_cg2(
+; 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::func.b32 %r1, [test_tcgen05_alloc_exclusive_shared_cg2_param_0];
+; CHECK_PTX64_SHARED32-NEXT: ld.param::func.b32 %r2, [test_tcgen05_alloc_exclusive_shared_cg2_param_1];
+; CHECK_PTX64_SHARED32-NEXT: tcgen05.alloc.exclusive.cta_group::2.sync.aligned.shared::cta.b32 [%r1], %r2;
+; CHECK_PTX64_SHARED32-NEXT: ret;
+ call void @llvm.nvvm.tcgen05.alloc.exclusive.cg2.p3(ptr addrspace(3) %addr, i32 %ncols)
+ ret void
+}
+
+declare void @llvm.nvvm.tcgen05.dealloc.exclusive.cg1(ptr addrspace(6) %tmem_addr, i32 %ncols)
+declare void @llvm.nvvm.tcgen05.dealloc.exclusive.cg2(ptr addrspace(6) %tmem_addr, i32 %ncols)
+
+define void @test_tcgen05_dealloc_exclusive_cg1(ptr addrspace(6) %tmem_addr, i32 %ncols) {
+; CHECK_PTX64-LABEL: test_tcgen05_dealloc_exclusive_cg1(
+; CHECK_PTX64: {
+; CHECK_PTX64-NEXT: .reg .b32 %r<3>;
+; CHECK_PTX64-EMPTY:
+; CHECK_PTX64-NEXT: // %bb.0:
+; CHECK_PTX64-NEXT: ld.param::func.b32 %r1, [test_tcgen05_dealloc_exclusive_cg1_param_0];
+; CHECK_PTX64-NEXT: ld.param::func.b32 %r2, [test_tcgen05_dealloc_exclusive_cg1_param_1];
+; CHECK_PTX64-NEXT: tcgen05.dealloc.exclusive.cta_group::1.sync.aligned.b32 %r1, %r2;
+; CHECK_PTX64-NEXT: ret;
+;
+; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_dealloc_exclusive_cg1(
+; 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::func.b32 %r1, [test_tcgen05_dealloc_exclusive_cg1_param_0];
+; CHECK_PTX64_SHARED32-NEXT: ld.param::func.b32 %r2, [test_tcgen05_dealloc_exclusive_cg1_param_1];
+; CHECK_PTX64_SHARED32-NEXT: tcgen05.dealloc.exclusive.cta_group::1.sync.aligned.b32 %r1, %r2;
+; CHECK_PTX64_SHARED32-NEXT: ret;
+ call void @llvm.nvvm.tcgen05.dealloc.exclusive.cg1(ptr addrspace(6) %tmem_addr, i32 %ncols)
+ ret void
+}
+
+define void @test_tcgen05_dealloc_exclusive_cg2(ptr addrspace(6) %tmem_addr, i32 %ncols) {
+; CHECK_PTX64-LABEL: test_tcgen05_dealloc_exclusive_cg2(
+; CHECK_PTX64: {
+; CHECK_PTX64-NEXT: .reg .b32 %r<3>;
+; CHECK_PTX64-EMPTY:
+; CHECK_PTX64-NEXT: // %bb.0:
+; CHECK_PTX64-NEXT: ld.param::func.b32 %r1, [test_tcgen05_dealloc_exclusive_cg2_param_0];
+; CHECK_PTX64-NEXT: ld.param::func.b32 %r2, [test_tcgen05_dealloc_exclusive_cg2_param_1];
+; CHECK_PTX64-NEXT: tcgen05.dealloc.exclusive.cta_group::2.sync.aligned.b32 %r1, %r2;
+; CHECK_PTX64-NEXT: ret;
+;
+; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_dealloc_exclusive_cg2(
+; 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::func.b32 %r1, [test_tcgen05_dealloc_exclusive_cg2_param_0];
+; CHECK_PTX64_SHARED32-NEXT: ld.param::func.b32 %r2, [test_tcgen05_dealloc_exclusive_cg2_param_1];
+; CHECK_PTX64_SHARED32-NEXT: tcgen05.dealloc.exclusive.cta_group::2.sync.aligned.b32 %r1, %r2;
+; CHECK_PTX64_SHARED32-NEXT: ret;
+ call void @llvm.nvvm.tcgen05.dealloc.exclusive.cg2(ptr addrspace(6) %tmem_addr, i32 %ncols)
+ ret void
+}
diff --git a/llvm/test/CodeGen/NVPTX/tcgen05-alloc.ll b/llvm/test/CodeGen/NVPTX/tcgen05-alloc.ll
index 43cd82615491b..65fa2b1df38c4 100644
--- a/llvm/test/CodeGen/NVPTX/tcgen05-alloc.ll
+++ b/llvm/test/CodeGen/NVPTX/tcgen05-alloc.ll
@@ -11,10 +11,10 @@
; 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.alloc.cg1(ptr %addr, i32 %ncols)
-declare void @llvm.nvvm.tcgen05.alloc.cg2(ptr %addr, i32 %ncols)
-declare void @llvm.nvvm.tcgen05.alloc.shared.cg1(ptr addrspace(3) %addr, i32 %ncols)
-declare void @llvm.nvvm.tcgen05.alloc.shared.cg2(ptr addrspace(3) %addr, i32 %ncols)
+declare void @llvm.nvvm.tcgen05.alloc.cg1.p0(ptr %addr, i32 %ncols)
+declare void @llvm.nvvm.tcgen05.alloc.cg2.p0(ptr %addr, i32 %ncols)
+declare void @llvm.nvvm.tcgen05.alloc.cg1.p3(ptr addrspace(3) %addr, i32 %ncols)
+declare void @llvm.nvvm.tcgen05.alloc.cg2.p3(ptr addrspace(3) %addr, i32 %ncols)
define void @test_tcgen05_alloc_cg1(ptr %addr, i32 %ncols) {
; CHECK_PTX64-LABEL: test_tcgen05_alloc_cg1(
@@ -38,7 +38,7 @@ define void @test_tcgen05_alloc_cg1(ptr %addr, i32 %ncols) {
; CHECK_PTX64_SHARED32-NEXT: ld.param::func.b32 %r1, [test_tcgen05_alloc_cg1_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.alloc.cta_group::1.sync.aligned.b32 [%rd1], %r1;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.alloc.cg1(ptr %addr, i32 %ncols)
+ call void @llvm.nvvm.tcgen05.alloc.cg1.p0(ptr %addr, i32 %ncols)
ret void
}
@@ -64,7 +64,7 @@ define void @test_tcgen05_alloc_cg2(ptr %addr, i32 %ncols) {
; CHECK_PTX64_SHARED32-NEXT: ld.param::func.b32 %r1, [test_tcgen05_alloc_cg2_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.alloc.cta_group::2.sync.aligned.b32 [%rd1], %r1;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.alloc.cg2(ptr %addr, i32 %ncols)
+ call void @llvm.nvvm.tcgen05.alloc.cg2.p0(ptr %addr, i32 %ncols)
ret void
}
@@ -89,7 +89,7 @@ define void @test_tcgen05_alloc_shared_cg1(ptr addrspace(3) %addr, i32 %ncols) {
; CHECK_PTX64_SHARED32-NEXT: ld.param::func.b32 %r2, [test_tcgen05_alloc_shared_cg1_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.alloc.cta_group::1.sync.aligned.shared::cta.b32 [%r1], %r2;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.alloc.shared.cg1(ptr addrspace(3) %addr, i32 %ncols)
+ call void @llvm.nvvm.tcgen05.alloc.cg1.p3(ptr addrspace(3) %addr, i32 %ncols)
ret void
}
@@ -114,7 +114,7 @@ define void @test_tcgen05_alloc_shared_cg2(ptr addrspace(3) %addr, i32 %ncols) {
; CHECK_PTX64_SHARED32-NEXT: ld.param::func.b32 %r2, [test_tcgen05_alloc_shared_cg2_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.alloc.cta_group::2.sync.aligned.shared::cta.b32 [%r1], %r2;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.alloc.shared.cg2(ptr addrspace(3) %addr, i32 %ncols)
+ call void @llvm.nvvm.tcgen05.alloc.cg2.p3(ptr addrspace(3) %addr, i32 %ncols)
ret void
}
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 6241606122f43..f4ed0c3535ddb 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -5259,7 +5259,7 @@ def NVVM_Tcgen05AllocOp : NVVM_Op<"tcgen05.alloc", [NVVMRequiresSMf<[100, 101, 1
llvm::SmallVector<llvm::Value *> args;
auto id = NVVM::Tcgen05AllocOp::getIntrinsicIDAndArgs(
*op, moduleTranslation, args);
- createIntrinsicCall(builder, id, args);
+ createIntrinsicCall(builder, id, builder.getVoidTy(), args);
}];
}
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index 3cb03e297d03c..d1d7f5c9162b6 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -5017,17 +5017,13 @@ Tcgen05AllocOp::getIntrinsicIDAndArgs(Operation &op,
auto curOp = cast<NVVM::Tcgen05AllocOp>(op);
unsigned as = llvm::cast<LLVM::LLVMPointerType>(curOp.getAddr().getType())
.getAddressSpace();
- bool isShared = as == NVVMMemorySpace::Shared;
+ assert((as == NVVMMemorySpace::Generic || as == NVVMMemorySpace::Shared) &&
+ "expected generic or shared memory address");
bool is2CTAMode = curOp.getGroup() == CTAGroupKind::CTA_2;
- llvm::Intrinsic::ID id;
- if (isShared) {
- id = is2CTAMode ? llvm::Intrinsic::nvvm_tcgen05_alloc_shared_cg2
- : llvm::Intrinsic::nvvm_tcgen05_alloc_shared_cg1;
- } else {
- id = is2CTAMode ? llvm::Intrinsic::nvvm_tcgen05_alloc_cg2
- : llvm::Intrinsic::nvvm_tcgen05_alloc_cg1;
- }
+ llvm::Intrinsic::ID id =
+ is2CTAMode ? llvm::Intrinsic::nvvm_tcgen05_alloc_cg2
+ : llvm::Intrinsic::nvvm_tcgen05_alloc_cg1;
// Fill the Intrinsic Args
args.push_back(mt.lookupValue(curOp.getAddr()));
diff --git a/mlir/test/Target/LLVMIR/nvvm/tcgen05-alloc.mlir b/mlir/test/Target/LLVMIR/nvvm/tcgen05-alloc.mlir
index a8f80296f20ae..9f46437a0ca7c 100644
--- a/mlir/test/Target/LLVMIR/nvvm/tcgen05-alloc.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/tcgen05-alloc.mlir
@@ -2,20 +2,20 @@
// CHECK-LABEL: @llvm_nvvm_tcgen05_alloc
llvm.func @llvm_nvvm_tcgen05_alloc(%addr : !llvm.ptr, %ncols : i32) {
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.alloc.cg1(ptr %{{.*}}, i32 %{{.*}})
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.alloc.cg1.p0(ptr %{{.*}}, i32 %{{.*}})
nvvm.tcgen05.alloc %addr, %ncols : !llvm.ptr, i32
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.alloc.cg2(ptr %{{.*}}, i32 %{{.*}})
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.alloc.cg2.p0(ptr %{{.*}}, i32 %{{.*}})
nvvm.tcgen05.alloc %addr, %ncols {group = #nvvm.cta_group<cta_2>} : !llvm.ptr, i32
llvm.return
}
// CHECK-LABEL: @llvm_nvvm_tcgen05_alloc_shared
llvm.func @llvm_nvvm_tcgen05_alloc_shared(%addr : !llvm.ptr<3>, %ncols : i32) {
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.alloc.shared.cg1(ptr addrspace(3) %{{.*}}, i32 %{{.*}})
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.alloc.cg1.p3(ptr addrspace(3) %{{.*}}, i32 %{{.*}})
nvvm.tcgen05.alloc %addr, %ncols : !llvm.ptr<3>, i32
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.alloc.shared.cg2(ptr addrspace(3) %{{.*}}, i32 %{{.*}})
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.alloc.cg2.p3(ptr addrspace(3) %{{.*}}, i32 %{{.*}})
nvvm.tcgen05.alloc %addr, %ncols {group = #nvvm.cta_group<cta_2>} : !llvm.ptr<3>, i32
llvm.return
}
>From 19f9454ac7e4583abb19e17c4aaf7caa11c0035e Mon Sep 17 00:00:00 2001
From: DharuniRAcharya <dharunira at nvidia.com>
Date: Thu, 13 Aug 2026 11:43:35 +0000
Subject: [PATCH 2/6] Fix formatting
---
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp | 5 ++---
1 file changed, 2 insertions(+), 3 deletions(-)
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index d1d7f5c9162b6..5e107efffbab7 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -5021,9 +5021,8 @@ Tcgen05AllocOp::getIntrinsicIDAndArgs(Operation &op,
"expected generic or shared memory address");
bool is2CTAMode = curOp.getGroup() == CTAGroupKind::CTA_2;
- llvm::Intrinsic::ID id =
- is2CTAMode ? llvm::Intrinsic::nvvm_tcgen05_alloc_cg2
- : llvm::Intrinsic::nvvm_tcgen05_alloc_cg1;
+ llvm::Intrinsic::ID id = is2CTAMode ? llvm::Intrinsic::nvvm_tcgen05_alloc_cg2
+ : llvm::Intrinsic::nvvm_tcgen05_alloc_cg1;
// Fill the Intrinsic Args
args.push_back(mt.lookupValue(curOp.getAddr()));
>From 932104b739d7b1779bfe611fd6012b5de4bfae62 Mon Sep 17 00:00:00 2001
From: DharuniRAcharya <dharunira at nvidia.com>
Date: Fri, 14 Aug 2026 11:06:51 +0000
Subject: [PATCH 3/6] Address comments
---
llvm/docs/NVPTXUsage.md | 1 +
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 79 ++++++----------------
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp | 2 -
3 files changed, 23 insertions(+), 59 deletions(-)
diff --git a/llvm/docs/NVPTXUsage.md b/llvm/docs/NVPTXUsage.md
index 98a987fcab4e6..80b6215154bb9 100644
--- a/llvm/docs/NVPTXUsage.md
+++ b/llvm/docs/NVPTXUsage.md
@@ -2819,6 +2819,7 @@ overloaded pointer argument may use generic or shared memory address space;
a shared pointer emits the `.shared::cta` qualifier. The `.cg1` and `.cg2`
variants generate `cta_group::1` and `cta_group::2` variants of the instruction
respectively. The `.exclusive` variants claim ownership of the allocation permit.
+No other allocation may exist at the same time as an exclusive allocation.
For more information, refer to the [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory-allocation-and-management-instructions).
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 2f9815bd64679..f7a1e9478295e 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -5844,82 +5844,47 @@ let Predicates = [SM90] in {
def EXIT : NullaryInst<"exit", int_nvvm_exit>;
// Tcgen05 intrinsics
-multiclass TCGEN05_ALLOC_PATFRAGS<Intrinsic Intr> {
- defvar frag_pat = (Intr node:$dst, node:$ncols);
- def _GENERIC
- : PatFrag<!setdagop(frag_pat, ops), frag_pat, AS_match.generic>;
- def _SHARED
- : PatFrag<!setdagop(frag_pat, ops), frag_pat, AS_match.shared>;
-}
-
-defm TCGEN05_ALLOC_CG1_PAT
- : TCGEN05_ALLOC_PATFRAGS<int_nvvm_tcgen05_alloc_cg1>;
-defm TCGEN05_ALLOC_CG2_PAT
- : TCGEN05_ALLOC_PATFRAGS<int_nvvm_tcgen05_alloc_cg2>;
-defm TCGEN05_ALLOC_EXCLUSIVE_CG1_PAT
- : TCGEN05_ALLOC_PATFRAGS<int_nvvm_tcgen05_alloc_exclusive_cg1>;
-defm TCGEN05_ALLOC_EXCLUSIVE_CG2_PAT
- : TCGEN05_ALLOC_PATFRAGS<int_nvvm_tcgen05_alloc_exclusive_cg2>;
-
-multiclass TCGEN05_ALLOC_INTR<string num, PatFrag GenericPat,
- PatFrag SharedPat,
- bit IsExclusive = 0> {
+multiclass TCGEN05_ALLOC_INTR<string CtaGroup, Intrinsic Intr, bit IsExclusive = 0> {
defvar exclusive = !if(IsExclusive, ".exclusive", "");
defvar Pred = !if(IsExclusive,
[hasTcgen05InstSupport, PTX94],
[hasTcgen05InstSupport]);
+ defvar frag_pat = (Intr node:$dst, node:$ncols);
+ defvar ops_dag = !setdagop(frag_pat, ops);
+
+ def _GPAT : PatFrag<ops_dag, frag_pat, AS_match.generic>;
+ def _SPAT : PatFrag<ops_dag, frag_pat, AS_match.shared>;
let isConvergent = true, Predicates = Pred in {
def _GENERIC : BasicNVPTXInst<(outs), (ins ADDR:$dst, B32:$ncols),
- "tcgen05.alloc" # exclusive # ".cta_group::" # num #
- ".sync.aligned.b32",
- [(GenericPat addr:$dst, B32:$ncols)]>;
+ "tcgen05.alloc" # exclusive # ".cta_group::" # CtaGroup # ".sync.aligned.b32",
+ [(!cast<PatFrag>(NAME # "_GPAT") addr:$dst, B32:$ncols)]>;
def _SHARED : BasicNVPTXInst<(outs), (ins ADDR:$dst, B32:$ncols),
- "tcgen05.alloc" # exclusive # ".cta_group::" # num #
- ".sync.aligned.shared::cta.b32",
- [(SharedPat addr:$dst, B32:$ncols)]>;
+ "tcgen05.alloc" # exclusive # ".cta_group::" # CtaGroup # ".sync.aligned.shared::cta.b32",
+ [(!cast<PatFrag>(NAME # "_SPAT") addr:$dst, B32:$ncols)]>;
}
}
-defm TCGEN05_ALLOC_CG1
- : TCGEN05_ALLOC_INTR<"1", TCGEN05_ALLOC_CG1_PAT_GENERIC,
- TCGEN05_ALLOC_CG1_PAT_SHARED>;
-defm TCGEN05_ALLOC_CG2
- : TCGEN05_ALLOC_INTR<"2", TCGEN05_ALLOC_CG2_PAT_GENERIC,
- TCGEN05_ALLOC_CG2_PAT_SHARED>;
-
-defm TCGEN05_ALLOC_EXCLUSIVE_CG1
- : TCGEN05_ALLOC_INTR<"1", TCGEN05_ALLOC_EXCLUSIVE_CG1_PAT_GENERIC,
- TCGEN05_ALLOC_EXCLUSIVE_CG1_PAT_SHARED,
- /*IsExclusive=*/1>;
-defm TCGEN05_ALLOC_EXCLUSIVE_CG2
- : TCGEN05_ALLOC_INTR<"2", TCGEN05_ALLOC_EXCLUSIVE_CG2_PAT_GENERIC,
- TCGEN05_ALLOC_EXCLUSIVE_CG2_PAT_SHARED,
- /*IsExclusive=*/1>;
-
-multiclass TCGEN05_DEALLOC_INTR<string num, Intrinsic Intr,
- bit IsExclusive = 0> {
+defm TCGEN05_ALLOC_CG1 : TCGEN05_ALLOC_INTR<"1", int_nvvm_tcgen05_alloc_cg1>;
+defm TCGEN05_ALLOC_CG2 : TCGEN05_ALLOC_INTR<"2", int_nvvm_tcgen05_alloc_cg2>;
+defm TCGEN05_ALLOC_EXCLUSIVE_CG1 : TCGEN05_ALLOC_INTR<"1", int_nvvm_tcgen05_alloc_exclusive_cg1, 1>;
+defm TCGEN05_ALLOC_EXCLUSIVE_CG2 : TCGEN05_ALLOC_INTR<"2", int_nvvm_tcgen05_alloc_exclusive_cg2, 1>;
+
+multiclass TCGEN05_DEALLOC_INTR<string CtaGroup, Intrinsic Intr, bit IsExclusive = 0> {
defvar exclusive = !if(IsExclusive, ".exclusive", "");
defvar Pred = !if(IsExclusive,
[hasTcgen05InstSupport, PTX94],
[hasTcgen05InstSupport]);
let isConvergent = true, Predicates = Pred in {
- def "" : BasicNVPTXInst<(outs),
- (ins B32:$tmem_addr, B32:$ncols),
- "tcgen05.dealloc" # exclusive # ".cta_group::" # num #
- ".sync.aligned.b32",
+ def "" : BasicNVPTXInst<(outs), (ins B32:$tmem_addr, B32:$ncols),
+ "tcgen05.dealloc" # exclusive # ".cta_group::" # CtaGroup # ".sync.aligned.b32",
[(Intr B32:$tmem_addr, B32:$ncols)]>;
}
}
-defm TCGEN05_DEALLOC_CG1: TCGEN05_DEALLOC_INTR<"1", int_nvvm_tcgen05_dealloc_cg1>;
-defm TCGEN05_DEALLOC_CG2: TCGEN05_DEALLOC_INTR<"2", int_nvvm_tcgen05_dealloc_cg2>;
-
-defm TCGEN05_DEALLOC_EXCLUSIVE_CG1
- : TCGEN05_DEALLOC_INTR<"1", int_nvvm_tcgen05_dealloc_exclusive_cg1,
- /*IsExclusive=*/1>;
-defm TCGEN05_DEALLOC_EXCLUSIVE_CG2
- : TCGEN05_DEALLOC_INTR<"2", int_nvvm_tcgen05_dealloc_exclusive_cg2,
- /*IsExclusive=*/1>;
+defm TCGEN05_DEALLOC_CG1 : TCGEN05_DEALLOC_INTR<"1", int_nvvm_tcgen05_dealloc_cg1>;
+defm TCGEN05_DEALLOC_CG2 : TCGEN05_DEALLOC_INTR<"2", int_nvvm_tcgen05_dealloc_cg2>;
+defm TCGEN05_DEALLOC_EXCLUSIVE_CG1 : TCGEN05_DEALLOC_INTR<"1", int_nvvm_tcgen05_dealloc_exclusive_cg1, 1>;
+defm TCGEN05_DEALLOC_EXCLUSIVE_CG2 : TCGEN05_DEALLOC_INTR<"2", int_nvvm_tcgen05_dealloc_exclusive_cg2, 1>;
let isConvergent = true in {
let Predicates = [hasTcgen05InstSupport] in {
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index 5e107efffbab7..2c339ff04d95a 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -5017,8 +5017,6 @@ Tcgen05AllocOp::getIntrinsicIDAndArgs(Operation &op,
auto curOp = cast<NVVM::Tcgen05AllocOp>(op);
unsigned as = llvm::cast<LLVM::LLVMPointerType>(curOp.getAddr().getType())
.getAddressSpace();
- assert((as == NVVMMemorySpace::Generic || as == NVVMMemorySpace::Shared) &&
- "expected generic or shared memory address");
bool is2CTAMode = curOp.getGroup() == CTAGroupKind::CTA_2;
llvm::Intrinsic::ID id = is2CTAMode ? llvm::Intrinsic::nvvm_tcgen05_alloc_cg2
>From 5c1c13bee60db386981cc6d301275c4ad97c3bc8 Mon Sep 17 00:00:00 2001
From: DharuniRAcharya <dharunira at nvidia.com>
Date: Fri, 14 Aug 2026 15:59:27 +0000
Subject: [PATCH 4/6] Fix build issue
---
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp | 2 --
1 file changed, 2 deletions(-)
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index 2c339ff04d95a..b7dadbb0c21d1 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -5015,8 +5015,6 @@ Tcgen05AllocOp::getIntrinsicIDAndArgs(Operation &op,
LLVM::ModuleTranslation &mt,
llvm::SmallVector<llvm::Value *> &args) {
auto curOp = cast<NVVM::Tcgen05AllocOp>(op);
- unsigned as = llvm::cast<LLVM::LLVMPointerType>(curOp.getAddr().getType())
- .getAddressSpace();
bool is2CTAMode = curOp.getGroup() == CTAGroupKind::CTA_2;
llvm::Intrinsic::ID id = is2CTAMode ? llvm::Intrinsic::nvvm_tcgen05_alloc_cg2
>From baf94ac021f34f7e4ba2ad932db9a0b6d44052a2 Mon Sep 17 00:00:00 2001
From: DharuniRAcharya <dharunira at nvidia.com>
Date: Mon, 17 Aug 2026 05:58:17 +0000
Subject: [PATCH 5/6] Update exclusive as an immediate boolean
---
llvm/docs/NVPTXUsage.md | 21 +++---
llvm/include/llvm/IR/IntrinsicsNVVM.td | 28 +++----
llvm/lib/IR/AutoUpgrade.cpp | 43 ++++++++---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 2 -
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 75 ++++++++++---------
.../Assembler/auto_upgrade_nvvm_intrinsics.ll | 19 ++++-
.../NVPTX/tcgen05-alloc-dealloc-exclusive.ll | 24 +++---
llvm/test/CodeGen/NVPTX/tcgen05-alloc.ll | 25 +++----
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp | 2 +
.../Target/LLVMIR/nvvm/tcgen05-alloc.mlir | 12 +--
10 files changed, 140 insertions(+), 111 deletions(-)
diff --git a/llvm/docs/NVPTXUsage.md b/llvm/docs/NVPTXUsage.md
index 80b6215154bb9..969ad9ec3cb11 100644
--- a/llvm/docs/NVPTXUsage.md
+++ b/llvm/docs/NVPTXUsage.md
@@ -2798,11 +2798,8 @@ For more information on tensor-memory load/store instructions, refer to [PTX ISA
##### Syntax:
```llvm
-declare void @llvm.nvvm.tcgen05.alloc.{cg1,cg2}.p0(ptr %dst, i32 %ncols)
-declare void @llvm.nvvm.tcgen05.alloc.{cg1,cg2}.p3(ptr addrspace(3) %dst, i32 %ncols)
-
-declare void @llvm.nvvm.tcgen05.alloc.exclusive.{cg1,cg2}.p0(ptr %dst, i32 %ncols)
-declare void @llvm.nvvm.tcgen05.alloc.exclusive.{cg1,cg2}.p3(ptr addrspace(3) %dst, i32 %ncols)
+declare void @llvm.nvvm.tcgen05.alloc.{cg1,cg2}.p0(ptr %dst, i32 %ncols, i1 %is_exclusive)
+declare void @llvm.nvvm.tcgen05.alloc.{cg1,cg2}.p3(ptr addrspace(3) %dst, i32 %ncols, i1 %is_exclusive)
```
##### Overview:
@@ -2818,8 +2815,9 @@ allocations and it must be a multiple of 32 for exclusive allocations. The
overloaded pointer argument may use generic or shared memory address space;
a shared pointer emits the `.shared::cta` qualifier. The `.cg1` and `.cg2`
variants generate `cta_group::1` and `cta_group::2` variants of the instruction
-respectively. The `.exclusive` variants claim ownership of the allocation permit.
-No other allocation may exist at the same time as an exclusive allocation.
+respectively. When `%is_exclusive` is true, the `.exclusive` variant is emitted,
+which claims ownership of the allocation permit. No other allocation may exist
+at the same time as an exclusive allocation.
For more information, refer to the [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory-allocation-and-management-instructions).
@@ -2828,9 +2826,7 @@ For more information, refer to the [PTX ISA](https://docs.nvidia.com/cuda/parall
##### Syntax:
```llvm
-declare void @llvm.nvvm.tcgen05.dealloc.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i32 %ncols)
-
-declare void @llvm.nvvm.tcgen05.dealloc.exclusive.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i32 %ncols)
+declare void @llvm.nvvm.tcgen05.dealloc.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i32 %ncols, i1 %is_exclusive)
```
##### Overview:
@@ -2842,8 +2838,9 @@ address `%tmem_addr`. The operand `%tmem_addr` must point to a previous
Tensor Memory allocation. The 32-bit operand `%ncols` specifies the number
of columns to be de-allocated. The `.cg1` and `.cg2` variants generate
`cta_group::1` and `cta_group::2` variants of the instruction respectively.
-Memory must be deallocated with `.exclusive` if and only if it was allocated
-with `.exclusive`.
+When `%is_exclusive` is true, the `.exclusive` variant is emitted. Memory must
+be deallocated with `.exclusive` if and only if it was allocated with
+`.exclusive`.
For more information, refer to the [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory-allocation-and-management-instructions).
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index a9b746e05461b..56f6a0ace1529 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -3208,21 +3208,21 @@ def int_nvvm_griddepcontrol_wait : Intrinsic<[], [], [IntrNoMem, IntrHasSideEffe
//
// Tcgen05 alloc/dealloc related intrinsics
-
foreach cta_group = ["cg1", "cg2"] in {
- foreach exclusive = ["", "exclusive_"] in {
- def int_nvvm_tcgen05_alloc_ # exclusive # cta_group : Intrinsic<[],
- [llvm_anyptr_ty, // dst_ptr
- llvm_i32_ty], // num_columns
- [IntrConvergent, IntrInaccessibleMemOrArgMemOnly,
- WriteOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>]>;
-
- def int_nvvm_tcgen05_dealloc_ # exclusive # cta_group : Intrinsic<[],
- [llvm_tmem_ptr_ty, // tmem_addr
- llvm_i32_ty], // num_columns
- [IntrConvergent, IntrArgMemOnly,
- NoCapture<ArgIndex<0>>]>;
- }
+ def int_nvvm_tcgen05_alloc_ # cta_group : Intrinsic<[],
+ [llvm_anyptr_ty, // dst_ptr
+ llvm_i32_ty, // num_columns
+ llvm_i1_ty], // is_exclusive
+ [IntrConvergent, IntrInaccessibleMemOrArgMemOnly,
+ WriteOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>,
+ ImmArg<ArgIndex<2>>]>;
+
+ def int_nvvm_tcgen05_dealloc_ # cta_group : Intrinsic<[],
+ [llvm_tmem_ptr_ty, // tmem_addr
+ llvm_i32_ty, // num_columns
+ llvm_i1_ty], // is_exclusive
+ [IntrConvergent, IntrArgMemOnly, NoCapture<ArgIndex<0>>,
+ ImmArg<ArgIndex<2>>]>;
def int_nvvm_tcgen05_relinq_alloc_permit_ # cta_group : Intrinsic<[], [],
[IntrConvergent, IntrInaccessibleMemOnly]>;
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index 22a148c5512d7..123806ec4e01f 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -1280,14 +1280,24 @@ shouldUpgradeNVPTXTcgen05CommitSharedIntrinsic(Function *F, StringRef Name) {
}
static Intrinsic::ID
-shouldUpgradeNVPTXTcgen05AllocSharedIntrinsic(StringRef Name) {
- if (!Name.consume_front("tcgen05.alloc.shared."))
+shouldUpgradeNVPTXTcgen05AllocDeallocIntrinsic(Function *F, StringRef Name) {
+ if (F->arg_size() != 2)
return Intrinsic::not_intrinsic;
- return StringSwitch<Intrinsic::ID>(Name)
- .Case("cg1", Intrinsic::nvvm_tcgen05_alloc_cg1)
- .Case("cg2", Intrinsic::nvvm_tcgen05_alloc_cg2)
- .Default(Intrinsic::not_intrinsic);
+ if (Name.consume_front("tcgen05.alloc.shared.") ||
+ Name.consume_front("tcgen05.alloc."))
+ return StringSwitch<Intrinsic::ID>(Name)
+ .Case("cg1", Intrinsic::nvvm_tcgen05_alloc_cg1)
+ .Case("cg2", Intrinsic::nvvm_tcgen05_alloc_cg2)
+ .Default(Intrinsic::not_intrinsic);
+
+ if (Name.consume_front("tcgen05.dealloc."))
+ return StringSwitch<Intrinsic::ID>(Name)
+ .Case("cg1", Intrinsic::nvvm_tcgen05_dealloc_cg1)
+ .Case("cg2", Intrinsic::nvvm_tcgen05_dealloc_cg2)
+ .Default(Intrinsic::not_intrinsic);
+
+ return Intrinsic::not_intrinsic;
}
static Intrinsic::ID shouldUpgradeNVPTXBF16Intrinsic(StringRef Name) {
@@ -1822,13 +1832,16 @@ static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn,
return true;
}
- // Upgrade tcgen05.alloc shared variants to anyptr intrinsics.
- IID = shouldUpgradeNVPTXTcgen05AllocSharedIntrinsic(Name);
+ // Upgrade tcgen05.alloc/dealloc with the is_exclusive argument and
+ // tcgen05.alloc shared variants to anyptr intrinsics.
+ IID = shouldUpgradeNVPTXTcgen05AllocDeallocIntrinsic(F, Name);
if (IID != Intrinsic::not_intrinsic) {
rename(F);
- NewFn = Intrinsic::getOrInsertDeclaration(
- F->getParent(), IID, F->getReturnType(),
- F->getFunctionType()->params());
+ if (Intrinsic::isOverloaded(IID))
+ NewFn = Intrinsic::getOrInsertDeclaration(F->getParent(), IID,
+ {F->getArg(0)->getType()});
+ else
+ NewFn = Intrinsic::getOrInsertDeclaration(F->getParent(), IID);
return true;
}
@@ -5590,6 +5603,14 @@ void llvm::UpgradeIntrinsicCall(CallBase *CI, Function *NewFn) {
case Intrinsic::ctpop:
NewCall = Builder.CreateCall(NewFn, {CI->getArgOperand(0)});
break;
+ case Intrinsic::nvvm_tcgen05_alloc_cg1:
+ case Intrinsic::nvvm_tcgen05_alloc_cg2:
+ case Intrinsic::nvvm_tcgen05_dealloc_cg1:
+ case Intrinsic::nvvm_tcgen05_dealloc_cg2:
+ NewCall =
+ Builder.CreateCall(NewFn, {CI->getArgOperand(0), CI->getArgOperand(1),
+ Builder.getFalse()});
+ break;
case Intrinsic::dbg_value: {
StringRef Name = F->getName();
Name = Name.substr(5); // Strip llvm.
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index be3078ed835fd..2b01aab98d54a 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -5548,8 +5548,6 @@ void NVPTXTargetLowering::getTgtMemIntrinsic(
}
case Intrinsic::nvvm_tcgen05_alloc_cg1:
case Intrinsic::nvvm_tcgen05_alloc_cg2:
- case Intrinsic::nvvm_tcgen05_alloc_exclusive_cg1:
- case Intrinsic::nvvm_tcgen05_alloc_exclusive_cg2:
Info.opc = ISD::INTRINSIC_VOID;
Info.memVT = MVT::i32;
Info.ptrVal = I.getArgOperand(0);
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index f7a1e9478295e..e24399e6e2929 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -5844,47 +5844,48 @@ let Predicates = [SM90] in {
def EXIT : NullaryInst<"exit", int_nvvm_exit>;
// Tcgen05 intrinsics
-multiclass TCGEN05_ALLOC_INTR<string CtaGroup, Intrinsic Intr, bit IsExclusive = 0> {
- defvar exclusive = !if(IsExclusive, ".exclusive", "");
- defvar Pred = !if(IsExclusive,
- [hasTcgen05InstSupport, PTX94],
- [hasTcgen05InstSupport]);
- defvar frag_pat = (Intr node:$dst, node:$ncols);
- defvar ops_dag = !setdagop(frag_pat, ops);
-
- def _GPAT : PatFrag<ops_dag, frag_pat, AS_match.generic>;
- def _SPAT : PatFrag<ops_dag, frag_pat, AS_match.shared>;
-
- let isConvergent = true, Predicates = Pred in {
- def _GENERIC : BasicNVPTXInst<(outs), (ins ADDR:$dst, B32:$ncols),
- "tcgen05.alloc" # exclusive # ".cta_group::" # CtaGroup # ".sync.aligned.b32",
- [(!cast<PatFrag>(NAME # "_GPAT") addr:$dst, B32:$ncols)]>;
- def _SHARED : BasicNVPTXInst<(outs), (ins ADDR:$dst, B32:$ncols),
- "tcgen05.alloc" # exclusive # ".cta_group::" # CtaGroup # ".sync.aligned.shared::cta.b32",
- [(!cast<PatFrag>(NAME # "_SPAT") addr:$dst, B32:$ncols)]>;
+multiclass TCGEN05_ALLOC_INTR<string CtaGroup, Intrinsic Intr> {
+ foreach IsExclusive = [0, -1] in {
+ defvar exclusive = !if(IsExclusive, ".exclusive", "");
+ defvar Pred = !if(IsExclusive,
+ [hasTcgen05InstSupport, PTX94],
+ [hasTcgen05InstSupport]);
+ defvar suffix = !if(IsExclusive, "_EXCLUSIVE", "");
+
+ let isConvergent = true, Predicates = Pred in {
+ def suffix # "_GENERIC" : BasicNVPTXInst<(outs), (ins ADDR:$dst, B32:$ncols),
+ "tcgen05.alloc" # exclusive # ".cta_group::" # CtaGroup # ".sync.aligned.b32",
+ [(PatFrag<(ops node:$dst, node:$ncols),
+ (Intr node:$dst, node:$ncols, IsExclusive),
+ AS_match.generic> addr:$dst, B32:$ncols)]>;
+ def suffix # "_SHARED" : BasicNVPTXInst<(outs), (ins ADDR:$dst, B32:$ncols),
+ "tcgen05.alloc" # exclusive # ".cta_group::" # CtaGroup # ".sync.aligned.shared::cta.b32",
+ [(PatFrag<(ops node:$dst, node:$ncols),
+ (Intr node:$dst, node:$ncols, IsExclusive),
+ AS_match.shared> addr:$dst, B32:$ncols)]>;
+ }
}
}
-
-defm TCGEN05_ALLOC_CG1 : TCGEN05_ALLOC_INTR<"1", int_nvvm_tcgen05_alloc_cg1>;
-defm TCGEN05_ALLOC_CG2 : TCGEN05_ALLOC_INTR<"2", int_nvvm_tcgen05_alloc_cg2>;
-defm TCGEN05_ALLOC_EXCLUSIVE_CG1 : TCGEN05_ALLOC_INTR<"1", int_nvvm_tcgen05_alloc_exclusive_cg1, 1>;
-defm TCGEN05_ALLOC_EXCLUSIVE_CG2 : TCGEN05_ALLOC_INTR<"2", int_nvvm_tcgen05_alloc_exclusive_cg2, 1>;
-
-multiclass TCGEN05_DEALLOC_INTR<string CtaGroup, Intrinsic Intr, bit IsExclusive = 0> {
- defvar exclusive = !if(IsExclusive, ".exclusive", "");
- defvar Pred = !if(IsExclusive,
- [hasTcgen05InstSupport, PTX94],
- [hasTcgen05InstSupport]);
- let isConvergent = true, Predicates = Pred in {
- def "" : BasicNVPTXInst<(outs), (ins B32:$tmem_addr, B32:$ncols),
- "tcgen05.dealloc" # exclusive # ".cta_group::" # CtaGroup # ".sync.aligned.b32",
- [(Intr B32:$tmem_addr, B32:$ncols)]>;
+defm TCGEN05_ALLOC_CG1 : TCGEN05_ALLOC_INTR<"1", int_nvvm_tcgen05_alloc_cg1>;
+defm TCGEN05_ALLOC_CG2 : TCGEN05_ALLOC_INTR<"2", int_nvvm_tcgen05_alloc_cg2>;
+
+multiclass TCGEN05_DEALLOC_INTR<string CtaGroup, Intrinsic Intr> {
+ foreach IsExclusive = [0, -1] in {
+ defvar exclusive = !if(IsExclusive, ".exclusive", "");
+ defvar Pred = !if(IsExclusive,
+ [hasTcgen05InstSupport, PTX94],
+ [hasTcgen05InstSupport]);
+ defvar suffix = !if(IsExclusive, "_EXCLUSIVE", "");
+
+ let isConvergent = true, Predicates = Pred in {
+ def suffix : BasicNVPTXInst<(outs), (ins B32:$tmem_addr, B32:$ncols),
+ "tcgen05.dealloc" # exclusive # ".cta_group::" # CtaGroup # ".sync.aligned.b32",
+ [(Intr B32:$tmem_addr, B32:$ncols, IsExclusive)]>;
+ }
}
}
-defm TCGEN05_DEALLOC_CG1 : TCGEN05_DEALLOC_INTR<"1", int_nvvm_tcgen05_dealloc_cg1>;
-defm TCGEN05_DEALLOC_CG2 : TCGEN05_DEALLOC_INTR<"2", int_nvvm_tcgen05_dealloc_cg2>;
-defm TCGEN05_DEALLOC_EXCLUSIVE_CG1 : TCGEN05_DEALLOC_INTR<"1", int_nvvm_tcgen05_dealloc_exclusive_cg1, 1>;
-defm TCGEN05_DEALLOC_EXCLUSIVE_CG2 : TCGEN05_DEALLOC_INTR<"2", int_nvvm_tcgen05_dealloc_exclusive_cg2, 1>;
+defm TCGEN05_DEALLOC_CG1 : TCGEN05_DEALLOC_INTR<"1", int_nvvm_tcgen05_dealloc_cg1>;
+defm TCGEN05_DEALLOC_CG2 : TCGEN05_DEALLOC_INTR<"2", int_nvvm_tcgen05_dealloc_cg2>;
let isConvergent = true in {
let Predicates = [hasTcgen05InstSupport] in {
diff --git a/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll b/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
index 5813e3bda1c49..7228e286ad510 100644
--- a/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
+++ b/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
@@ -128,6 +128,8 @@ declare void @llvm.nvvm.tcgen05.alloc.cg1(ptr, i32)
declare void @llvm.nvvm.tcgen05.alloc.cg2(ptr, i32)
declare void @llvm.nvvm.tcgen05.alloc.shared.cg1(ptr addrspace(3), i32)
declare void @llvm.nvvm.tcgen05.alloc.shared.cg2(ptr addrspace(3), i32)
+declare void @llvm.nvvm.tcgen05.dealloc.cg1(ptr addrspace(6), i32)
+declare void @llvm.nvvm.tcgen05.dealloc.cg2(ptr addrspace(6), i32)
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);
@@ -510,10 +512,10 @@ define void @nvvm_tcgen05_commit_intrinsics(ptr %bar, ptr addrspace(3) %bar_shar
; CHECK-LABEL: @nvvm_tcgen05_alloc_intrinsics
define void @nvvm_tcgen05_alloc_intrinsics(ptr %dst, ptr addrspace(3) %dst_shared, i32 %ncols) {
-; CHECK: call void @llvm.nvvm.tcgen05.alloc.cg1.p0(ptr %dst, i32 %ncols)
-; CHECK: call void @llvm.nvvm.tcgen05.alloc.cg2.p0(ptr %dst, i32 %ncols)
-; CHECK: call void @llvm.nvvm.tcgen05.alloc.cg1.p3(ptr addrspace(3) %dst_shared, i32 %ncols)
-; CHECK: call void @llvm.nvvm.tcgen05.alloc.cg2.p3(ptr addrspace(3) %dst_shared, i32 %ncols)
+; CHECK: call void @llvm.nvvm.tcgen05.alloc.cg1.p0(ptr %dst, i32 %ncols, i1 false)
+; CHECK: call void @llvm.nvvm.tcgen05.alloc.cg2.p0(ptr %dst, i32 %ncols, i1 false)
+; CHECK: call void @llvm.nvvm.tcgen05.alloc.cg1.p3(ptr addrspace(3) %dst_shared, i32 %ncols, i1 false)
+; CHECK: call void @llvm.nvvm.tcgen05.alloc.cg2.p3(ptr addrspace(3) %dst_shared, i32 %ncols, i1 false)
call void @llvm.nvvm.tcgen05.alloc.cg1(ptr %dst, i32 %ncols)
call void @llvm.nvvm.tcgen05.alloc.cg2(ptr %dst, i32 %ncols)
call void @llvm.nvvm.tcgen05.alloc.shared.cg1(ptr addrspace(3) %dst_shared, i32 %ncols)
@@ -521,6 +523,15 @@ define void @nvvm_tcgen05_alloc_intrinsics(ptr %dst, ptr addrspace(3) %dst_share
ret void
}
+; CHECK-LABEL: @nvvm_tcgen05_dealloc_intrinsics
+define void @nvvm_tcgen05_dealloc_intrinsics(ptr addrspace(6) %tmem_addr, i32 %ncols) {
+; CHECK: call void @llvm.nvvm.tcgen05.dealloc.cg1(ptr addrspace(6) %tmem_addr, i32 %ncols, i1 false)
+; CHECK: call void @llvm.nvvm.tcgen05.dealloc.cg2(ptr addrspace(6) %tmem_addr, i32 %ncols, i1 false)
+ call void @llvm.nvvm.tcgen05.dealloc.cg1(ptr addrspace(6) %tmem_addr, i32 %ncols)
+ call void @llvm.nvvm.tcgen05.dealloc.cg2(ptr addrspace(6) %tmem_addr, i32 %ncols)
+ 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.p0.i16(ptr %bar, i16 %cta_mask)
diff --git a/llvm/test/CodeGen/NVPTX/tcgen05-alloc-dealloc-exclusive.ll b/llvm/test/CodeGen/NVPTX/tcgen05-alloc-dealloc-exclusive.ll
index deb477ce41e54..4c46204e39dff 100644
--- a/llvm/test/CodeGen/NVPTX/tcgen05-alloc-dealloc-exclusive.ll
+++ b/llvm/test/CodeGen/NVPTX/tcgen05-alloc-dealloc-exclusive.ll
@@ -6,10 +6,10 @@
; RUN: %if ptxas-sm_100f && ptxas-isa-9.4 %{ llc < %s -march=nvptx64 -mcpu=sm_100f -mattr=+ptx94 -target-abi=shortptr | %ptxas-verify -arch=sm_100f %}
; RUN: %if ptxas-sm_110f && ptxas-isa-9.4 %{ llc < %s -march=nvptx64 -mcpu=sm_110f -mattr=+ptx94 | %ptxas-verify -arch=sm_110f %}
-declare void @llvm.nvvm.tcgen05.alloc.exclusive.cg1.p0(ptr %addr, i32 %ncols)
-declare void @llvm.nvvm.tcgen05.alloc.exclusive.cg2.p0(ptr %addr, i32 %ncols)
-declare void @llvm.nvvm.tcgen05.alloc.exclusive.cg1.p3(ptr addrspace(3) %addr, i32 %ncols)
-declare void @llvm.nvvm.tcgen05.alloc.exclusive.cg2.p3(ptr addrspace(3) %addr, i32 %ncols)
+declare void @llvm.nvvm.tcgen05.alloc.cg1.p0(ptr %addr, i32 %ncols, i1 %is_exclusive)
+declare void @llvm.nvvm.tcgen05.alloc.cg2.p0(ptr %addr, i32 %ncols, i1 %is_exclusive)
+declare void @llvm.nvvm.tcgen05.alloc.cg1.p3(ptr addrspace(3) %addr, i32 %ncols, i1 %is_exclusive)
+declare void @llvm.nvvm.tcgen05.alloc.cg2.p3(ptr addrspace(3) %addr, i32 %ncols, i1 %is_exclusive)
define void @test_tcgen05_alloc_exclusive_cg1(ptr %addr, i32 %ncols) {
; CHECK_PTX64-LABEL: test_tcgen05_alloc_exclusive_cg1(
@@ -33,7 +33,7 @@ define void @test_tcgen05_alloc_exclusive_cg1(ptr %addr, i32 %ncols) {
; CHECK_PTX64_SHARED32-NEXT: ld.param::func.b32 %r1, [test_tcgen05_alloc_exclusive_cg1_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.alloc.exclusive.cta_group::1.sync.aligned.b32 [%rd1], %r1;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.alloc.exclusive.cg1.p0(ptr %addr, i32 %ncols)
+ call void @llvm.nvvm.tcgen05.alloc.cg1.p0(ptr %addr, i32 %ncols, i1 true)
ret void
}
@@ -59,7 +59,7 @@ define void @test_tcgen05_alloc_exclusive_cg2(ptr %addr, i32 %ncols) {
; CHECK_PTX64_SHARED32-NEXT: ld.param::func.b32 %r1, [test_tcgen05_alloc_exclusive_cg2_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.alloc.exclusive.cta_group::2.sync.aligned.b32 [%rd1], %r1;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.alloc.exclusive.cg2.p0(ptr %addr, i32 %ncols)
+ call void @llvm.nvvm.tcgen05.alloc.cg2.p0(ptr %addr, i32 %ncols, i1 true)
ret void
}
@@ -84,7 +84,7 @@ define void @test_tcgen05_alloc_exclusive_shared_cg1(ptr addrspace(3) %addr, i32
; CHECK_PTX64_SHARED32-NEXT: ld.param::func.b32 %r2, [test_tcgen05_alloc_exclusive_shared_cg1_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.alloc.exclusive.cta_group::1.sync.aligned.shared::cta.b32 [%r1], %r2;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.alloc.exclusive.cg1.p3(ptr addrspace(3) %addr, i32 %ncols)
+ call void @llvm.nvvm.tcgen05.alloc.cg1.p3(ptr addrspace(3) %addr, i32 %ncols, i1 true)
ret void
}
@@ -109,12 +109,12 @@ define void @test_tcgen05_alloc_exclusive_shared_cg2(ptr addrspace(3) %addr, i32
; CHECK_PTX64_SHARED32-NEXT: ld.param::func.b32 %r2, [test_tcgen05_alloc_exclusive_shared_cg2_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.alloc.exclusive.cta_group::2.sync.aligned.shared::cta.b32 [%r1], %r2;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.alloc.exclusive.cg2.p3(ptr addrspace(3) %addr, i32 %ncols)
+ call void @llvm.nvvm.tcgen05.alloc.cg2.p3(ptr addrspace(3) %addr, i32 %ncols, i1 true)
ret void
}
-declare void @llvm.nvvm.tcgen05.dealloc.exclusive.cg1(ptr addrspace(6) %tmem_addr, i32 %ncols)
-declare void @llvm.nvvm.tcgen05.dealloc.exclusive.cg2(ptr addrspace(6) %tmem_addr, i32 %ncols)
+declare void @llvm.nvvm.tcgen05.dealloc.cg1(ptr addrspace(6) %tmem_addr, i32 %ncols, i1 %is_exclusive)
+declare void @llvm.nvvm.tcgen05.dealloc.cg2(ptr addrspace(6) %tmem_addr, i32 %ncols, i1 %is_exclusive)
define void @test_tcgen05_dealloc_exclusive_cg1(ptr addrspace(6) %tmem_addr, i32 %ncols) {
; CHECK_PTX64-LABEL: test_tcgen05_dealloc_exclusive_cg1(
@@ -136,7 +136,7 @@ define void @test_tcgen05_dealloc_exclusive_cg1(ptr addrspace(6) %tmem_addr, i32
; CHECK_PTX64_SHARED32-NEXT: ld.param::func.b32 %r2, [test_tcgen05_dealloc_exclusive_cg1_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.dealloc.exclusive.cta_group::1.sync.aligned.b32 %r1, %r2;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.dealloc.exclusive.cg1(ptr addrspace(6) %tmem_addr, i32 %ncols)
+ call void @llvm.nvvm.tcgen05.dealloc.cg1(ptr addrspace(6) %tmem_addr, i32 %ncols, i1 true)
ret void
}
@@ -160,6 +160,6 @@ define void @test_tcgen05_dealloc_exclusive_cg2(ptr addrspace(6) %tmem_addr, i32
; CHECK_PTX64_SHARED32-NEXT: ld.param::func.b32 %r2, [test_tcgen05_dealloc_exclusive_cg2_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.dealloc.exclusive.cta_group::2.sync.aligned.b32 %r1, %r2;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.dealloc.exclusive.cg2(ptr addrspace(6) %tmem_addr, i32 %ncols)
+ call void @llvm.nvvm.tcgen05.dealloc.cg2(ptr addrspace(6) %tmem_addr, i32 %ncols, i1 true)
ret void
}
diff --git a/llvm/test/CodeGen/NVPTX/tcgen05-alloc.ll b/llvm/test/CodeGen/NVPTX/tcgen05-alloc.ll
index 65fa2b1df38c4..201ea4da53942 100644
--- a/llvm/test/CodeGen/NVPTX/tcgen05-alloc.ll
+++ b/llvm/test/CodeGen/NVPTX/tcgen05-alloc.ll
@@ -10,11 +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.alloc.cg1.p0(ptr %addr, i32 %ncols)
-declare void @llvm.nvvm.tcgen05.alloc.cg2.p0(ptr %addr, i32 %ncols)
-declare void @llvm.nvvm.tcgen05.alloc.cg1.p3(ptr addrspace(3) %addr, i32 %ncols)
-declare void @llvm.nvvm.tcgen05.alloc.cg2.p3(ptr addrspace(3) %addr, i32 %ncols)
+declare void @llvm.nvvm.tcgen05.alloc.cg1.p0(ptr %addr, i32 %ncols, i1 %is_exclusive)
+declare void @llvm.nvvm.tcgen05.alloc.cg2.p0(ptr %addr, i32 %ncols, i1 %is_exclusive)
+declare void @llvm.nvvm.tcgen05.alloc.cg1.p3(ptr addrspace(3) %addr, i32 %ncols, i1 %is_exclusive)
+declare void @llvm.nvvm.tcgen05.alloc.cg2.p3(ptr addrspace(3) %addr, i32 %ncols, i1 %is_exclusive)
define void @test_tcgen05_alloc_cg1(ptr %addr, i32 %ncols) {
; CHECK_PTX64-LABEL: test_tcgen05_alloc_cg1(
@@ -38,7 +37,7 @@ define void @test_tcgen05_alloc_cg1(ptr %addr, i32 %ncols) {
; CHECK_PTX64_SHARED32-NEXT: ld.param::func.b32 %r1, [test_tcgen05_alloc_cg1_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.alloc.cta_group::1.sync.aligned.b32 [%rd1], %r1;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.alloc.cg1.p0(ptr %addr, i32 %ncols)
+ call void @llvm.nvvm.tcgen05.alloc.cg1.p0(ptr %addr, i32 %ncols, i1 false)
ret void
}
@@ -64,7 +63,7 @@ define void @test_tcgen05_alloc_cg2(ptr %addr, i32 %ncols) {
; CHECK_PTX64_SHARED32-NEXT: ld.param::func.b32 %r1, [test_tcgen05_alloc_cg2_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.alloc.cta_group::2.sync.aligned.b32 [%rd1], %r1;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.alloc.cg2.p0(ptr %addr, i32 %ncols)
+ call void @llvm.nvvm.tcgen05.alloc.cg2.p0(ptr %addr, i32 %ncols, i1 false)
ret void
}
@@ -89,7 +88,7 @@ define void @test_tcgen05_alloc_shared_cg1(ptr addrspace(3) %addr, i32 %ncols) {
; CHECK_PTX64_SHARED32-NEXT: ld.param::func.b32 %r2, [test_tcgen05_alloc_shared_cg1_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.alloc.cta_group::1.sync.aligned.shared::cta.b32 [%r1], %r2;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.alloc.cg1.p3(ptr addrspace(3) %addr, i32 %ncols)
+ call void @llvm.nvvm.tcgen05.alloc.cg1.p3(ptr addrspace(3) %addr, i32 %ncols, i1 false)
ret void
}
@@ -114,12 +113,12 @@ define void @test_tcgen05_alloc_shared_cg2(ptr addrspace(3) %addr, i32 %ncols) {
; CHECK_PTX64_SHARED32-NEXT: ld.param::func.b32 %r2, [test_tcgen05_alloc_shared_cg2_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.alloc.cta_group::2.sync.aligned.shared::cta.b32 [%r1], %r2;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.alloc.cg2.p3(ptr addrspace(3) %addr, i32 %ncols)
+ call void @llvm.nvvm.tcgen05.alloc.cg2.p3(ptr addrspace(3) %addr, i32 %ncols, i1 false)
ret void
}
-declare void @llvm.nvvm.tcgen05.dealloc.cg1(ptr addrspace(6) %tmem_addr, i32 %ncols)
-declare void @llvm.nvvm.tcgen05.dealloc.cg2(ptr addrspace(6) %tmem_addr, i32 %ncols)
+declare void @llvm.nvvm.tcgen05.dealloc.cg1(ptr addrspace(6) %tmem_addr, i32 %ncols, i1 %is_exclusive)
+declare void @llvm.nvvm.tcgen05.dealloc.cg2(ptr addrspace(6) %tmem_addr, i32 %ncols, i1 %is_exclusive)
define void @test_tcgen05_dealloc_cg1(ptr addrspace(6) %tmem_addr, i32 %ncols) {
; CHECK_PTX64-LABEL: test_tcgen05_dealloc_cg1(
@@ -141,7 +140,7 @@ define void @test_tcgen05_dealloc_cg1(ptr addrspace(6) %tmem_addr, i32 %ncols) {
; CHECK_PTX64_SHARED32-NEXT: ld.param::func.b32 %r2, [test_tcgen05_dealloc_cg1_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.dealloc.cta_group::1.sync.aligned.b32 %r1, %r2;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.dealloc.cg1(ptr addrspace(6) %tmem_addr, i32 %ncols)
+ call void @llvm.nvvm.tcgen05.dealloc.cg1(ptr addrspace(6) %tmem_addr, i32 %ncols, i1 false)
ret void
}
@@ -165,7 +164,7 @@ define void @test_tcgen05_dealloc_cg2(ptr addrspace(6) %tmem_addr, i32 %ncols) {
; CHECK_PTX64_SHARED32-NEXT: ld.param::func.b32 %r2, [test_tcgen05_dealloc_cg2_param_1];
; CHECK_PTX64_SHARED32-NEXT: tcgen05.dealloc.cta_group::2.sync.aligned.b32 %r1, %r2;
; CHECK_PTX64_SHARED32-NEXT: ret;
- call void @llvm.nvvm.tcgen05.dealloc.cg2(ptr addrspace(6) %tmem_addr, i32 %ncols)
+ call void @llvm.nvvm.tcgen05.dealloc.cg2(ptr addrspace(6) %tmem_addr, i32 %ncols, i1 false)
ret void
}
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index b7dadbb0c21d1..61f7296a63543 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -5023,6 +5023,7 @@ Tcgen05AllocOp::getIntrinsicIDAndArgs(Operation &op,
// Fill the Intrinsic Args
args.push_back(mt.lookupValue(curOp.getAddr()));
args.push_back(mt.lookupValue(curOp.getNCols()));
+ args.push_back(llvm::ConstantInt::getFalse(mt.getLLVMContext()));
return id;
}
@@ -5038,6 +5039,7 @@ llvm::Intrinsic::ID Tcgen05DeallocOp::getIntrinsicIDAndArgs(
// Fill the Intrinsic Args
args.push_back(mt.lookupValue(curOp.getTaddr()));
args.push_back(mt.lookupValue(curOp.getNCols()));
+ args.push_back(llvm::ConstantInt::getFalse(mt.getLLVMContext()));
return id;
}
diff --git a/mlir/test/Target/LLVMIR/nvvm/tcgen05-alloc.mlir b/mlir/test/Target/LLVMIR/nvvm/tcgen05-alloc.mlir
index 9f46437a0ca7c..7d53e951e6127 100644
--- a/mlir/test/Target/LLVMIR/nvvm/tcgen05-alloc.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/tcgen05-alloc.mlir
@@ -2,30 +2,30 @@
// CHECK-LABEL: @llvm_nvvm_tcgen05_alloc
llvm.func @llvm_nvvm_tcgen05_alloc(%addr : !llvm.ptr, %ncols : i32) {
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.alloc.cg1.p0(ptr %{{.*}}, i32 %{{.*}})
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.alloc.cg1.p0(ptr %{{.*}}, i32 %{{.*}}, i1 false)
nvvm.tcgen05.alloc %addr, %ncols : !llvm.ptr, i32
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.alloc.cg2.p0(ptr %{{.*}}, i32 %{{.*}})
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.alloc.cg2.p0(ptr %{{.*}}, i32 %{{.*}}, i1 false)
nvvm.tcgen05.alloc %addr, %ncols {group = #nvvm.cta_group<cta_2>} : !llvm.ptr, i32
llvm.return
}
// CHECK-LABEL: @llvm_nvvm_tcgen05_alloc_shared
llvm.func @llvm_nvvm_tcgen05_alloc_shared(%addr : !llvm.ptr<3>, %ncols : i32) {
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.alloc.cg1.p3(ptr addrspace(3) %{{.*}}, i32 %{{.*}})
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.alloc.cg1.p3(ptr addrspace(3) %{{.*}}, i32 %{{.*}}, i1 false)
nvvm.tcgen05.alloc %addr, %ncols : !llvm.ptr<3>, i32
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.alloc.cg2.p3(ptr addrspace(3) %{{.*}}, i32 %{{.*}})
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.alloc.cg2.p3(ptr addrspace(3) %{{.*}}, i32 %{{.*}}, i1 false)
nvvm.tcgen05.alloc %addr, %ncols {group = #nvvm.cta_group<cta_2>} : !llvm.ptr<3>, i32
llvm.return
}
// CHECK-LABEL: @llvm_nvvm_tcgen05_dealloc
llvm.func @llvm_nvvm_tcgen05_dealloc(%addr : !llvm.ptr<6>, %ncols : i32) {
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.dealloc.cg1(ptr addrspace(6) %{{.*}}, i32 %{{.*}})
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.dealloc.cg1(ptr addrspace(6) %{{.*}}, i32 %{{.*}}, i1 false)
nvvm.tcgen05.dealloc %addr, %ncols : !llvm.ptr<6>, i32
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.dealloc.cg2(ptr addrspace(6) %{{.*}}, i32 %{{.*}})
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.dealloc.cg2(ptr addrspace(6) %{{.*}}, i32 %{{.*}}, i1 false)
nvvm.tcgen05.dealloc %addr, %ncols {group = #nvvm.cta_group<cta_2>} : !llvm.ptr<6>, i32
llvm.return
}
>From 98c8c942207d8489f8cdab047742836c44885c8d Mon Sep 17 00:00:00 2001
From: DharuniRAcharya <dharunira at nvidia.com>
Date: Tue, 18 Aug 2026 04:58:58 +0000
Subject: [PATCH 6/6] Simplify backend lowering
---
llvm/lib/IR/AutoUpgrade.cpp | 16 ++++++++--------
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 8 ++------
2 files changed, 10 insertions(+), 14 deletions(-)
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index 123806ec4e01f..6c66d4e5f583f 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -5603,14 +5603,6 @@ void llvm::UpgradeIntrinsicCall(CallBase *CI, Function *NewFn) {
case Intrinsic::ctpop:
NewCall = Builder.CreateCall(NewFn, {CI->getArgOperand(0)});
break;
- case Intrinsic::nvvm_tcgen05_alloc_cg1:
- case Intrinsic::nvvm_tcgen05_alloc_cg2:
- case Intrinsic::nvvm_tcgen05_dealloc_cg1:
- case Intrinsic::nvvm_tcgen05_dealloc_cg2:
- NewCall =
- Builder.CreateCall(NewFn, {CI->getArgOperand(0), CI->getArgOperand(1),
- Builder.getFalse()});
- break;
case Intrinsic::dbg_value: {
StringRef Name = F->getName();
Name = Name.substr(5); // Strip llvm.
@@ -5845,6 +5837,14 @@ void llvm::UpgradeIntrinsicCall(CallBase *CI, Function *NewFn) {
NewCall = Builder.CreateCall(NewFn, Args);
break;
}
+ case Intrinsic::nvvm_tcgen05_alloc_cg1:
+ case Intrinsic::nvvm_tcgen05_alloc_cg2:
+ case Intrinsic::nvvm_tcgen05_dealloc_cg1:
+ case Intrinsic::nvvm_tcgen05_dealloc_cg2:
+ NewCall =
+ Builder.CreateCall(NewFn, {CI->getArgOperand(0), CI->getArgOperand(1),
+ Builder.getFalse()});
+ break;
case Intrinsic::riscv_sha256sig0:
case Intrinsic::riscv_sha256sig1:
case Intrinsic::riscv_sha256sum0:
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index e24399e6e2929..16885d0e3880d 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -5855,14 +5855,10 @@ multiclass TCGEN05_ALLOC_INTR<string CtaGroup, Intrinsic Intr> {
let isConvergent = true, Predicates = Pred in {
def suffix # "_GENERIC" : BasicNVPTXInst<(outs), (ins ADDR:$dst, B32:$ncols),
"tcgen05.alloc" # exclusive # ".cta_group::" # CtaGroup # ".sync.aligned.b32",
- [(PatFrag<(ops node:$dst, node:$ncols),
- (Intr node:$dst, node:$ncols, IsExclusive),
- AS_match.generic> addr:$dst, B32:$ncols)]>;
+ [(IntrinsicInAS<Intr, AddrSpaceGeneric> addr:$dst, B32:$ncols, IsExclusive)]>;
def suffix # "_SHARED" : BasicNVPTXInst<(outs), (ins ADDR:$dst, B32:$ncols),
"tcgen05.alloc" # exclusive # ".cta_group::" # CtaGroup # ".sync.aligned.shared::cta.b32",
- [(PatFrag<(ops node:$dst, node:$ncols),
- (Intr node:$dst, node:$ncols, IsExclusive),
- AS_match.shared> addr:$dst, B32:$ncols)]>;
+ [(IntrinsicInAS<Intr, AddrSpaceShared> addr:$dst, B32:$ncols, IsExclusive)]>;
}
}
}
More information about the Mlir-commits
mailing list