[llvm] [NVPTX] Add asynchronous store intrinsics (PR #200768)
Srinivasa Ravi via llvm-commits
llvm-commits at lists.llvm.org
Tue Jun 9 06:27:49 PDT 2026
https://github.com/Wolfram70 updated https://github.com/llvm/llvm-project/pull/200768
>From 8ca9f580e2175f24843c64e5028edbc9a4f1c924 Mon Sep 17 00:00:00 2001
From: Srinivasa Ravi <srinivasar at nvidia.com>
Date: Fri, 29 May 2026 12:30:40 +0000
Subject: [PATCH 1/7] [NVPTX] Add asynchronous store intrinsics
Adds the following intrinsics for asynchronous store operations:
- `st.async.space.cluster`
- `st.async.scope.sys.space.global`
- `st.async.scope.gpu.space.global`
- `st.async.mmio.scope.sys.space.global`
Tests verified through `ptxas-13.2`.
---
llvm/include/llvm/IR/IntrinsicsNVVM.td | 20 ++
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 135 ++++++++++++-
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 118 ++++++++++++
llvm/test/CodeGen/NVPTX/st_async_mbarrier.ll | 140 ++++++++++++++
llvm/test/CodeGen/NVPTX/st_async_release.ll | 192 +++++++++++++++++++
5 files changed, 603 insertions(+), 2 deletions(-)
create mode 100644 llvm/test/CodeGen/NVPTX/st_async_mbarrier.ll
create mode 100644 llvm/test/CodeGen/NVPTX/st_async_release.ll
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 647b65cf7714a..9e72aa6d8afe2 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -3272,6 +3272,26 @@ let IntrProperties = [IntrArgMemOnly, IntrWriteMem, WriteOnly<ArgIndex<0>>,
DefaultAttrsIntrinsic<[], [llvm_shared_ptr_ty, llvm_i64_ty, llvm_i64_ty]>;
}
+//
+// Asynchronous store intrinsics
+//
+
+def int_nvvm_st_async_space_cluster :
+ DefaultAttrsIntrinsic<[],
+ [llvm_shared_cluster_ptr_ty, llvm_anyint_ty, llvm_shared_cluster_ptr_ty],
+ [WriteOnly<ArgIndex<0>>]>;
+
+foreach scope = ["_sys", "_gpu"] in {
+ def int_nvvm_st_async_scope # scope # _space_global :
+ DefaultAttrsIntrinsic<[],
+ [llvm_global_ptr_ty, llvm_anyint_ty],
+ [WriteOnly<ArgIndex<0>>]>;
+}
+
+def int_nvvm_st_async_mmio_scope_sys_space_global :
+ DefaultAttrsIntrinsic<[],
+ [llvm_global_ptr_ty, llvm_anyint_ty], [WriteOnly<ArgIndex<0>>]>;
+
//
// clusterlaunchcontorl Intrinsics
//
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index d1d01089aa49e..348006d320e19 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -1118,8 +1118,9 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM,
// Custom lowering for tcgen05.st vector operands
setOperationAction(ISD::INTRINSIC_VOID,
- {MVT::v2i32, MVT::v4i32, MVT::v8i32, MVT::v16i32,
- MVT::v32i32, MVT::v64i32, MVT::v128i32, MVT::Other},
+ {MVT::i8, MVT::v2i32, MVT::v4i32, MVT::v8i32, MVT::v16i32,
+ MVT::v32i32, MVT::v64i32, MVT::v128i32, MVT::v2i64,
+ MVT::Other},
Custom);
// Enable custom lowering for the following:
@@ -2640,6 +2641,130 @@ static SDValue lowerBSWAP(SDValue Op, SelectionDAG &DAG) {
}
}
+static SDValue lowerStAsyncWithMbarrier(SDValue Op, SelectionDAG &DAG) {
+ const Function &Fn = DAG.getMachineFunction().getFunction();
+ SDNode *N = Op.getNode();
+ SDLoc DL(N);
+ Intrinsic::ID IntrinsicID = N->getConstantOperandVal(1);
+ SDValue DestAddr = N->getOperand(2);
+ SDValue Value = N->getOperand(3);
+ SDValue MbarAddr = N->getOperand(4);
+
+ MVT ValueVT = Value.getSimpleValueType();
+
+ if (ValueVT.getSizeInBits() > 128) {
+ DAG.getContext()->diagnose(DiagnosticInfoUnsupported(
+ Fn, "total bit-width of the value to be stored must be <= 128",
+ DiagnosticLocation(DL.getDebugLoc())));
+ return Op.getOperand(0); // Return only the chain
+ }
+
+ auto OpCode = [&]() -> std::optional<unsigned> {
+ switch (ValueVT.SimpleTy) {
+ case MVT::i32:
+ return NVPTXISD::ST_ASYNC_MBARRIER_B32;
+ case MVT::i64:
+ return NVPTXISD::ST_ASYNC_MBARRIER_B64;
+ case MVT::v2i32:
+ return NVPTXISD::ST_ASYNC_MBARRIER_V2B32;
+ case MVT::v2i64:
+ return NVPTXISD::ST_ASYNC_MBARRIER_V2B64;
+ case MVT::v4i32:
+ return NVPTXISD::ST_ASYNC_MBARRIER_V4B32;
+ default:
+ DAG.getContext()->diagnose(DiagnosticInfoUnsupported(
+ Fn,
+ Twine("unsupported argument type ") +
+ llvm::EVT(ValueVT).getEVTString() + " for " +
+ llvm::Intrinsic::getName(IntrinsicID) + " intrinsic",
+ DiagnosticLocation(DL.getDebugLoc())));
+ return {};
+ }
+ }();
+
+ if (!OpCode)
+ return Op.getOperand(0); // Return only the chain
+
+ SmallVector<SDValue, 6> Ops;
+
+ Ops.push_back(N->getOperand(0)); // Chain
+ Ops.push_back(DestAddr);
+ if (ValueVT.isVector()) {
+ for (unsigned i = 0; i < ValueVT.getVectorNumElements(); ++i) {
+ Ops.push_back(DAG.getNode(ISD::EXTRACT_VECTOR_ELT, DL,
+ ValueVT.getVectorElementType(), Value,
+ DAG.getIntPtrConstant(i, DL)));
+ }
+ } else {
+ Ops.push_back(Value);
+ }
+ Ops.push_back(MbarAddr);
+
+ return DAG.getNode(OpCode.value(), DL, MVT::Other, Ops);
+}
+
+static SDValue lowerStAsyncRelease(SDValue Op, SelectionDAG &DAG) {
+ const Function &Fn = DAG.getMachineFunction().getFunction();
+ SDNode *N = Op.getNode();
+ SDLoc DL(N);
+ Intrinsic::ID IntrinsicID = N->getConstantOperandVal(1);
+ SDValue DestAddr = N->getOperand(2);
+ SDValue Value = N->getOperand(3);
+
+ MVT ValueVT = Value.getSimpleValueType();
+
+ auto OpCode = [&]() -> std::optional<unsigned> {
+ switch (IntrinsicID) {
+ case Intrinsic::nvvm_st_async_scope_sys_space_global:
+ switch (ValueVT.SimpleTy) {
+ case MVT::i8: return NVPTXISD::ST_ASYNC_SYS_B8;
+ case MVT::i16: return NVPTXISD::ST_ASYNC_SYS_B16;
+ case MVT::i32: return NVPTXISD::ST_ASYNC_SYS_B32;
+ case MVT::i64: return NVPTXISD::ST_ASYNC_SYS_B64;
+ default: break;
+ }
+ break;
+ case Intrinsic::nvvm_st_async_scope_gpu_space_global:
+ switch (ValueVT.SimpleTy) {
+ case MVT::i8: return NVPTXISD::ST_ASYNC_GPU_B8;
+ case MVT::i16: return NVPTXISD::ST_ASYNC_GPU_B16;
+ case MVT::i32: return NVPTXISD::ST_ASYNC_GPU_B32;
+ case MVT::i64: return NVPTXISD::ST_ASYNC_GPU_B64;
+ default: break;
+ }
+ break;
+ case Intrinsic::nvvm_st_async_mmio_scope_sys_space_global:
+ switch (ValueVT.SimpleTy) {
+ case MVT::i8: return NVPTXISD::ST_ASYNC_MMIO_SYS_B8;
+ case MVT::i16: return NVPTXISD::ST_ASYNC_MMIO_SYS_B16;
+ case MVT::i32: return NVPTXISD::ST_ASYNC_MMIO_SYS_B32;
+ case MVT::i64: return NVPTXISD::ST_ASYNC_MMIO_SYS_B64;
+ default: break;
+ }
+ break;
+ }
+ return std::nullopt;
+ }();
+
+ if (!OpCode) {
+ DAG.getContext()->diagnose(DiagnosticInfoUnsupported(
+ Fn,
+ Twine("unsupported argument type ") +
+ llvm::EVT(ValueVT).getEVTString() + " for " +
+ llvm::Intrinsic::getName(IntrinsicID) + " intrinsic",
+ DiagnosticLocation(DL.getDebugLoc())));
+ return Op.getOperand(0); // Return only the chain
+ }
+
+ // NVPTX has no i8 register class; widen i8 to i16. The selected ST_ASYNC_*_B8
+ // node still emits the `.b8` qualifier in PTX.
+ if (ValueVT == MVT::i8)
+ Value = DAG.getNode(ISD::ZERO_EXTEND, DL, MVT::i16, Value);
+
+ SDValue Ops[] = {N->getOperand(0), DestAddr, Value};
+ return DAG.getNode(OpCode.value(), DL, MVT::Other, Ops);
+}
+
static unsigned getTcgen05MMADisableOutputLane(unsigned IID) {
switch (IID) {
case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg1:
@@ -2834,6 +2959,12 @@ static SDValue lowerIntrinsicVoid(SDValue Op, SelectionDAG &DAG) {
switch (IntrinNo) {
default:
break;
+ case Intrinsic::nvvm_st_async_space_cluster:
+ return lowerStAsyncWithMbarrier(Op, DAG);
+ case Intrinsic::nvvm_st_async_scope_sys_space_global:
+ case Intrinsic::nvvm_st_async_scope_gpu_space_global:
+ case Intrinsic::nvvm_st_async_mmio_scope_sys_space_global:
+ return lowerStAsyncRelease(Op, DAG);
case Intrinsic::nvvm_tcgen05_st_16x64b_x2:
case Intrinsic::nvvm_tcgen05_st_16x64b_x4:
case Intrinsic::nvvm_tcgen05_st_16x64b_x8:
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 1a3420ac6a7c7..1e3fd12ace8bf 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -6083,6 +6083,124 @@ let Predicates = [hasSM<100>, hasPTX<86>] in {
[(int_nvvm_st_bulk_shared_cta addr:$dest_addr, i64:$size, st_bulk_imm:$value)]>;
}
+//
+// Asynchronous store instructions
+//
+
+def SDTAsyncStoreMbarrier : SDTypeProfile<0, 3,
+ [SDTCisPtrTy<0>, SDTCisPtrTy<2>]>;
+def SDTAsyncStoreMbarrierV2 : SDTypeProfile<0, 4,
+ [SDTCisPtrTy<0>, SDTCisPtrTy<3>]>;
+def SDTAsyncStoreMbarrierV4 : SDTypeProfile<0, 6,
+ [SDTCisPtrTy<0>, SDTCisPtrTy<5>]>;
+
+class StAsyncMbarrierNode<string vec, string type> :
+ SDNode<"NVPTXISD::ST_ASYNC_MBARRIER_" # vec # type,
+ !cast<SDTypeProfile>("SDTAsyncStoreMbarrier" # vec),
+ [SDNPHasChain, SDNPMayStore]>;
+
+
+class ST_ASYNC_MBARRIER<string type, NVPTXRegClass reg> :
+ NVPTXInst<(outs),
+ (ins ADDR:$dest_addr, reg:$value, ADDR:$mbar),
+ "st.async.shared::cluster.mbarrier::complete_tx::bytes." # type # "\t [$dest_addr], $value, [$mbar];">;
+
+class ST_ASYNC_MBARRIER_V2<string type, NVPTXRegClass reg> :
+ NVPTXInst<(outs),
+ (ins ADDR:$dest_addr, reg:$value1, reg:$value2, ADDR:$mbar),
+ "st.async.shared::cluster.mbarrier::complete_tx::bytes.v2." # type # "\t [$dest_addr], {{$value1, $value2}}, [$mbar];">;
+
+class ST_ASYNC_MBARRIER_V4<string type, NVPTXRegClass reg> :
+ NVPTXInst<(outs),
+ (ins ADDR:$dest_addr, reg:$value1, reg:$value2, reg:$value3,
+ reg:$value4, ADDR:$mbar),
+ "st.async.shared::cluster.mbarrier::complete_tx::bytes.v4." # type # "\t [$dest_addr], {{$value1, $value2, $value3, $value4}}, [$mbar];">;
+
+
+let Predicates = [hasSM<90>, hasPTX<81>] in {
+ foreach size = ["32", "64"] in {
+ defvar valueType = !cast<ValueType>("i" # size);
+ defvar reg = !cast<NVPTXRegClass>("B" # size);
+ defvar type_str = "B" # size;
+
+ def : Pat<(StAsyncMbarrierNode<"", type_str>
+ addr:$dest_addr, valueType:$value, addr:$mbar),
+ (ST_ASYNC_MBARRIER<!tolower(type_str), reg>
+ addr:$dest_addr, valueType:$value, addr:$mbar)>;
+ def : Pat<(StAsyncMbarrierNode<"V2", type_str>
+ addr:$dest_addr, valueType:$value1, valueType:$value2,
+ addr:$mbar),
+ (ST_ASYNC_MBARRIER_V2<!tolower(type_str), reg>
+ addr:$dest_addr, valueType:$value1, valueType:$value2,
+ addr:$mbar)>;
+ }
+
+ def : Pat<(StAsyncMbarrierNode<"V4", "B32">
+ addr:$dest_addr, i32:$value1, i32:$value2,
+ i32:$value3, i32:$value4, addr:$mbar),
+ (ST_ASYNC_MBARRIER_V4<"b32", B32>
+ addr:$dest_addr, i32:$value1, i32:$value2,
+ i32:$value3, i32:$value4, addr:$mbar)>;
+}
+
+def SDTAsyncStoreRelease : SDTypeProfile<0, 2, [SDTCisPtrTy<0>]>;
+
+class StAsyncReleaseNode<string variant, string type> :
+ SDNode<"NVPTXISD::ST_ASYNC_" # variant # "_" # type,
+ SDTAsyncStoreRelease,
+ [SDNPHasChain, SDNPMayStore]>;
+
+class ST_ASYNC_RELEASE_BODY<string instr_prefix, string type> {
+ string body = !if(!eq(type, "b8"),
+ // The minimum supported register size to pass in arguments is 16-bits, so
+ // the 8-bit value is passed in using a 16-bit register and truncated to
+ // 8-bits with a cvt instruction.
+ "{{ \n\t"
+ ".reg .b8 \t%st_async_val; \n\t"
+ "cvt.u8.u16 \t%st_async_val, $value; \n\t" #
+ instr_prefix # ".b8\t [$dest_addr], %st_async_val; \n\t"
+ "}}",
+ instr_prefix # "." # type # "\t [$dest_addr], $value;");
+}
+
+class ST_ASYNC_SYS<string type, NVPTXRegClass reg> :
+ NVPTXInst<(outs),
+ (ins ADDR:$dest_addr, reg:$value),
+ ST_ASYNC_RELEASE_BODY<"st.async.release.sys.global", type>.body>;
+
+class ST_ASYNC_GPU<string type, NVPTXRegClass reg> :
+ NVPTXInst<(outs),
+ (ins ADDR:$dest_addr, reg:$value),
+ ST_ASYNC_RELEASE_BODY<"st.async.release.gpu.global", type>.body>;
+
+class ST_ASYNC_MMIO_SYS<string type, NVPTXRegClass reg> :
+ NVPTXInst<(outs),
+ (ins ADDR:$dest_addr, reg:$value),
+ ST_ASYNC_RELEASE_BODY<"st.async.mmio.release.sys.global", type>.body>;
+
+
+let Predicates = [hasSM<100>, hasPTX<87>] in {
+ foreach i = !range(4) in {
+ defvar size = ["8", "16", "32", "64"][i];
+ defvar reg = [B16, B16, B32, B64 ][i];
+ defvar valueType = [i16, i16, i32, i64 ][i];
+ defvar width = "b" # size;
+
+ def : Pat<(StAsyncReleaseNode<"SYS", !toupper(width)>
+ addr:$dest_addr, valueType:$value),
+ (ST_ASYNC_SYS<width, reg>
+ addr:$dest_addr, valueType:$value)>;
+ def : Pat<(StAsyncReleaseNode<"GPU", !toupper(width)>
+ addr:$dest_addr, valueType:$value),
+ (ST_ASYNC_GPU<width, reg>
+ addr:$dest_addr, valueType:$value)>;
+ def : Pat<(StAsyncReleaseNode<"MMIO_SYS", !toupper(width)>
+ addr:$dest_addr, valueType:$value),
+ (ST_ASYNC_MMIO_SYS<width, reg>
+ addr:$dest_addr, valueType:$value)>;
+ }
+}
+
//
// clusterlaunchcontorl Instructions
//
diff --git a/llvm/test/CodeGen/NVPTX/st_async_mbarrier.ll b/llvm/test/CodeGen/NVPTX/st_async_mbarrier.ll
new file mode 100644
index 0000000000000..8ccb0e3e69a6c
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/st_async_mbarrier.ll
@@ -0,0 +1,140 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx81 | FileCheck --check-prefixes=CHECK-PTX64 %s
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx81 --nvptx-short-ptr | FileCheck --check-prefixes=CHECK-SHARED32 %s
+; RUN: %if ptxas-sm_90 && ptxas-isa-8.1 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx81 | %ptxas-verify -arch=sm_90 %}
+; RUN: %if ptxas-sm_90 && ptxas-isa-8.1 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx81 --nvptx-short-ptr | %ptxas-verify -arch=sm_90 %}
+
+define void @test_st_async_mbarrier_b32(ptr addrspace(7) %addr, i32 %value, ptr addrspace(7) %mbar) {
+; CHECK-PTX64-LABEL: test_st_async_mbarrier_b32(
+; CHECK-PTX64: {
+; CHECK-PTX64-NEXT: .reg .b32 %r<2>;
+; CHECK-PTX64-NEXT: .reg .b64 %rd<3>;
+; CHECK-PTX64-EMPTY:
+; CHECK-PTX64-NEXT: // %bb.0:
+; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [test_st_async_mbarrier_b32_param_0];
+; CHECK-PTX64-NEXT: ld.param.b32 %r1, [test_st_async_mbarrier_b32_param_1];
+; CHECK-PTX64-NEXT: ld.param.b64 %rd2, [test_st_async_mbarrier_b32_param_2];
+; CHECK-PTX64-NEXT: st.async.shared::cluster.mbarrier::complete_tx::bytes.b32 [%rd1], %r1, [%rd2];
+; CHECK-PTX64-NEXT: ret;
+;
+; CHECK-SHARED32-LABEL: test_st_async_mbarrier_b32(
+; CHECK-SHARED32: {
+; CHECK-SHARED32-NEXT: .reg .b32 %r<4>;
+; CHECK-SHARED32-EMPTY:
+; CHECK-SHARED32-NEXT: // %bb.0:
+; CHECK-SHARED32-NEXT: ld.param.b32 %r1, [test_st_async_mbarrier_b32_param_0];
+; CHECK-SHARED32-NEXT: ld.param.b32 %r2, [test_st_async_mbarrier_b32_param_1];
+; CHECK-SHARED32-NEXT: ld.param.b32 %r3, [test_st_async_mbarrier_b32_param_2];
+; CHECK-SHARED32-NEXT: st.async.shared::cluster.mbarrier::complete_tx::bytes.b32 [%r1], %r2, [%r3];
+; CHECK-SHARED32-NEXT: ret;
+ call void @llvm.nvvm.st.async.space.cluster.i32(ptr addrspace(7) %addr, i32 %value, ptr addrspace(7) %mbar)
+ ret void
+}
+
+define void @test_st_async_mbarrier_b64(ptr addrspace(7) %addr, i64 %value, ptr addrspace(7) %mbar) {
+; CHECK-PTX64-LABEL: test_st_async_mbarrier_b64(
+; CHECK-PTX64: {
+; CHECK-PTX64-NEXT: .reg .b64 %rd<4>;
+; CHECK-PTX64-EMPTY:
+; CHECK-PTX64-NEXT: // %bb.0:
+; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [test_st_async_mbarrier_b64_param_0];
+; CHECK-PTX64-NEXT: ld.param.b64 %rd2, [test_st_async_mbarrier_b64_param_1];
+; CHECK-PTX64-NEXT: ld.param.b64 %rd3, [test_st_async_mbarrier_b64_param_2];
+; CHECK-PTX64-NEXT: st.async.shared::cluster.mbarrier::complete_tx::bytes.b64 [%rd1], %rd2, [%rd3];
+; CHECK-PTX64-NEXT: ret;
+;
+; CHECK-SHARED32-LABEL: test_st_async_mbarrier_b64(
+; CHECK-SHARED32: {
+; CHECK-SHARED32-NEXT: .reg .b32 %r<3>;
+; CHECK-SHARED32-NEXT: .reg .b64 %rd<2>;
+; CHECK-SHARED32-EMPTY:
+; CHECK-SHARED32-NEXT: // %bb.0:
+; CHECK-SHARED32-NEXT: ld.param.b32 %r1, [test_st_async_mbarrier_b64_param_0];
+; CHECK-SHARED32-NEXT: ld.param.b64 %rd1, [test_st_async_mbarrier_b64_param_1];
+; CHECK-SHARED32-NEXT: ld.param.b32 %r2, [test_st_async_mbarrier_b64_param_2];
+; CHECK-SHARED32-NEXT: st.async.shared::cluster.mbarrier::complete_tx::bytes.b64 [%r1], %rd1, [%r2];
+; CHECK-SHARED32-NEXT: ret;
+ call void @llvm.nvvm.st.async.space.cluster.i64(ptr addrspace(7) %addr, i64 %value, ptr addrspace(7) %mbar)
+ ret void
+}
+
+define void @test_st_async_mbarrier_v2b32(ptr addrspace(7) %addr, <2 x i32> %value, ptr addrspace(7) %mbar) {
+; CHECK-PTX64-LABEL: test_st_async_mbarrier_v2b32(
+; CHECK-PTX64: {
+; CHECK-PTX64-NEXT: .reg .b32 %r<3>;
+; CHECK-PTX64-NEXT: .reg .b64 %rd<3>;
+; CHECK-PTX64-EMPTY:
+; CHECK-PTX64-NEXT: // %bb.0:
+; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [test_st_async_mbarrier_v2b32_param_0];
+; CHECK-PTX64-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_st_async_mbarrier_v2b32_param_1];
+; CHECK-PTX64-NEXT: ld.param.b64 %rd2, [test_st_async_mbarrier_v2b32_param_2];
+; CHECK-PTX64-NEXT: st.async.shared::cluster.mbarrier::complete_tx::bytes.v2.b32 [%rd1], {%r1, %r2}, [%rd2];
+; CHECK-PTX64-NEXT: ret;
+;
+; CHECK-SHARED32-LABEL: test_st_async_mbarrier_v2b32(
+; CHECK-SHARED32: {
+; CHECK-SHARED32-NEXT: .reg .b32 %r<5>;
+; CHECK-SHARED32-EMPTY:
+; CHECK-SHARED32-NEXT: // %bb.0:
+; CHECK-SHARED32-NEXT: ld.param.b32 %r1, [test_st_async_mbarrier_v2b32_param_0];
+; CHECK-SHARED32-NEXT: ld.param.v2.b32 {%r2, %r3}, [test_st_async_mbarrier_v2b32_param_1];
+; CHECK-SHARED32-NEXT: ld.param.b32 %r4, [test_st_async_mbarrier_v2b32_param_2];
+; CHECK-SHARED32-NEXT: st.async.shared::cluster.mbarrier::complete_tx::bytes.v2.b32 [%r1], {%r2, %r3}, [%r4];
+; CHECK-SHARED32-NEXT: ret;
+ call void @llvm.nvvm.st.async.space.cluster.v2i32(ptr addrspace(7) %addr, <2 x i32> %value, ptr addrspace(7) %mbar)
+ ret void
+}
+
+define void @test_st_async_mbarrier_v2b64(ptr addrspace(7) %addr, <2 x i64> %value, ptr addrspace(7) %mbar) {
+; CHECK-PTX64-LABEL: test_st_async_mbarrier_v2b64(
+; CHECK-PTX64: {
+; CHECK-PTX64-NEXT: .reg .b64 %rd<5>;
+; CHECK-PTX64-EMPTY:
+; CHECK-PTX64-NEXT: // %bb.0:
+; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [test_st_async_mbarrier_v2b64_param_0];
+; CHECK-PTX64-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_st_async_mbarrier_v2b64_param_1];
+; CHECK-PTX64-NEXT: ld.param.b64 %rd4, [test_st_async_mbarrier_v2b64_param_2];
+; CHECK-PTX64-NEXT: st.async.shared::cluster.mbarrier::complete_tx::bytes.v2.b64 [%rd1], {%rd2, %rd3}, [%rd4];
+; CHECK-PTX64-NEXT: ret;
+;
+; CHECK-SHARED32-LABEL: test_st_async_mbarrier_v2b64(
+; CHECK-SHARED32: {
+; CHECK-SHARED32-NEXT: .reg .b32 %r<3>;
+; CHECK-SHARED32-NEXT: .reg .b64 %rd<3>;
+; CHECK-SHARED32-EMPTY:
+; CHECK-SHARED32-NEXT: // %bb.0:
+; CHECK-SHARED32-NEXT: ld.param.b32 %r1, [test_st_async_mbarrier_v2b64_param_0];
+; CHECK-SHARED32-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_st_async_mbarrier_v2b64_param_1];
+; CHECK-SHARED32-NEXT: ld.param.b32 %r2, [test_st_async_mbarrier_v2b64_param_2];
+; CHECK-SHARED32-NEXT: st.async.shared::cluster.mbarrier::complete_tx::bytes.v2.b64 [%r1], {%rd1, %rd2}, [%r2];
+; CHECK-SHARED32-NEXT: ret;
+ call void @llvm.nvvm.st.async.space.cluster.v2i64(ptr addrspace(7) %addr, <2 x i64> %value, ptr addrspace(7) %mbar)
+ ret void
+}
+
+define void @test_st_async_mbarrier_v4b32(ptr addrspace(7) %addr, <4 x i32> %value, ptr addrspace(7) %mbar) {
+; CHECK-PTX64-LABEL: test_st_async_mbarrier_v4b32(
+; CHECK-PTX64: {
+; CHECK-PTX64-NEXT: .reg .b32 %r<5>;
+; CHECK-PTX64-NEXT: .reg .b64 %rd<3>;
+; CHECK-PTX64-EMPTY:
+; CHECK-PTX64-NEXT: // %bb.0:
+; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [test_st_async_mbarrier_v4b32_param_0];
+; CHECK-PTX64-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_st_async_mbarrier_v4b32_param_1];
+; CHECK-PTX64-NEXT: ld.param.b64 %rd2, [test_st_async_mbarrier_v4b32_param_2];
+; CHECK-PTX64-NEXT: st.async.shared::cluster.mbarrier::complete_tx::bytes.v4.b32 [%rd1], {%r1, %r2, %r3, %r4}, [%rd2];
+; CHECK-PTX64-NEXT: ret;
+;
+; CHECK-SHARED32-LABEL: test_st_async_mbarrier_v4b32(
+; CHECK-SHARED32: {
+; CHECK-SHARED32-NEXT: .reg .b32 %r<7>;
+; CHECK-SHARED32-EMPTY:
+; CHECK-SHARED32-NEXT: // %bb.0:
+; CHECK-SHARED32-NEXT: ld.param.b32 %r1, [test_st_async_mbarrier_v4b32_param_0];
+; CHECK-SHARED32-NEXT: ld.param.v4.b32 {%r2, %r3, %r4, %r5}, [test_st_async_mbarrier_v4b32_param_1];
+; CHECK-SHARED32-NEXT: ld.param.b32 %r6, [test_st_async_mbarrier_v4b32_param_2];
+; CHECK-SHARED32-NEXT: st.async.shared::cluster.mbarrier::complete_tx::bytes.v4.b32 [%r1], {%r2, %r3, %r4, %r5}, [%r6];
+; CHECK-SHARED32-NEXT: ret;
+ call void @llvm.nvvm.st.async.space.cluster.v4i32(ptr addrspace(7) %addr, <4 x i32> %value, ptr addrspace(7) %mbar)
+ ret void
+}
diff --git a/llvm/test/CodeGen/NVPTX/st_async_release.ll b/llvm/test/CodeGen/NVPTX/st_async_release.ll
new file mode 100644
index 0000000000000..091b6bcd71f2d
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/st_async_release.ll
@@ -0,0 +1,192 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx87 | FileCheck %s
+; RUN: %if ptxas-sm_100 && ptxas-isa-8.7 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx87 | %ptxas-verify -arch=sm_100 %}
+
+define void @test_st_async_release_sys_b8(ptr addrspace(1) %addr, i8 %value) {
+; CHECK-LABEL: test_st_async_release_sys_b8(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_st_async_release_sys_b8_param_0];
+; CHECK-NEXT: ld.param.b8 %rs1, [test_st_async_release_sys_b8_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %st_async_val;
+; CHECK-NEXT: cvt.u8.u16 %st_async_val, %rs1;
+; CHECK-NEXT: st.async.release.sys.global.b8 [%rd1], %st_async_val;
+; CHECK-NEXT: }
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.st.async.scope.sys.space.global.i8(ptr addrspace(1) %addr, i8 %value)
+ ret void
+}
+
+define void @test_st_async_release_sys_b16(ptr addrspace(1) %addr, i16 %value) {
+; CHECK-LABEL: test_st_async_release_sys_b16(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_st_async_release_sys_b16_param_0];
+; CHECK-NEXT: ld.param.b16 %rs1, [test_st_async_release_sys_b16_param_1];
+; CHECK-NEXT: st.async.release.sys.global.b16 [%rd1], %rs1;
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.st.async.scope.sys.space.global.i16(ptr addrspace(1) %addr, i16 %value)
+ ret void
+}
+
+define void @test_st_async_release_sys_b32(ptr addrspace(1) %addr, i32 %value) {
+; CHECK-LABEL: test_st_async_release_sys_b32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_st_async_release_sys_b32_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_st_async_release_sys_b32_param_1];
+; CHECK-NEXT: st.async.release.sys.global.b32 [%rd1], %r1;
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.st.async.scope.sys.space.global.i32(ptr addrspace(1) %addr, i32 %value)
+ ret void
+}
+
+define void @test_st_async_release_sys_b64(ptr addrspace(1) %addr, i64 %value) {
+; CHECK-LABEL: test_st_async_release_sys_b64(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_st_async_release_sys_b64_param_0];
+; CHECK-NEXT: ld.param.b64 %rd2, [test_st_async_release_sys_b64_param_1];
+; CHECK-NEXT: st.async.release.sys.global.b64 [%rd1], %rd2;
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.st.async.scope.sys.space.global.i64(ptr addrspace(1) %addr, i64 %value)
+ ret void
+}
+
+define void @test_st_async_release_gpu_b8(ptr addrspace(1) %addr, i8 %value) {
+; CHECK-LABEL: test_st_async_release_gpu_b8(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_st_async_release_gpu_b8_param_0];
+; CHECK-NEXT: ld.param.b8 %rs1, [test_st_async_release_gpu_b8_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %st_async_val;
+; CHECK-NEXT: cvt.u8.u16 %st_async_val, %rs1;
+; CHECK-NEXT: st.async.release.gpu.global.b8 [%rd1], %st_async_val;
+; CHECK-NEXT: }
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.st.async.scope.gpu.space.global.i8(ptr addrspace(1) %addr, i8 %value)
+ ret void
+}
+
+define void @test_st_async_release_gpu_b16(ptr addrspace(1) %addr, i16 %value) {
+; CHECK-LABEL: test_st_async_release_gpu_b16(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_st_async_release_gpu_b16_param_0];
+; CHECK-NEXT: ld.param.b16 %rs1, [test_st_async_release_gpu_b16_param_1];
+; CHECK-NEXT: st.async.release.gpu.global.b16 [%rd1], %rs1;
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.st.async.scope.gpu.space.global.i16(ptr addrspace(1) %addr, i16 %value)
+ ret void
+}
+
+define void @test_st_async_release_gpu_b32(ptr addrspace(1) %addr, i32 %value) {
+; CHECK-LABEL: test_st_async_release_gpu_b32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_st_async_release_gpu_b32_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_st_async_release_gpu_b32_param_1];
+; CHECK-NEXT: st.async.release.gpu.global.b32 [%rd1], %r1;
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.st.async.scope.gpu.space.global.i32(ptr addrspace(1) %addr, i32 %value)
+ ret void
+}
+
+define void @test_st_async_release_gpu_b64(ptr addrspace(1) %addr, i64 %value) {
+; CHECK-LABEL: test_st_async_release_gpu_b64(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_st_async_release_gpu_b64_param_0];
+; CHECK-NEXT: ld.param.b64 %rd2, [test_st_async_release_gpu_b64_param_1];
+; CHECK-NEXT: st.async.release.gpu.global.b64 [%rd1], %rd2;
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.st.async.scope.gpu.space.global.i64(ptr addrspace(1) %addr, i64 %value)
+ ret void
+}
+
+define void @test_st_async_mmio_release_sys_b8(ptr addrspace(1) %addr, i8 %value) {
+; CHECK-LABEL: test_st_async_mmio_release_sys_b8(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_st_async_mmio_release_sys_b8_param_0];
+; CHECK-NEXT: ld.param.b8 %rs1, [test_st_async_mmio_release_sys_b8_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %st_async_val;
+; CHECK-NEXT: cvt.u8.u16 %st_async_val, %rs1;
+; CHECK-NEXT: st.async.mmio.release.sys.global.b8 [%rd1], %st_async_val;
+; CHECK-NEXT: }
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.st.async.mmio.scope.sys.space.global.i8(ptr addrspace(1) %addr, i8 %value)
+ ret void
+}
+
+define void @test_st_async_mmio_release_sys_b16(ptr addrspace(1) %addr, i16 %value) {
+; CHECK-LABEL: test_st_async_mmio_release_sys_b16(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_st_async_mmio_release_sys_b16_param_0];
+; CHECK-NEXT: ld.param.b16 %rs1, [test_st_async_mmio_release_sys_b16_param_1];
+; CHECK-NEXT: st.async.mmio.release.sys.global.b16 [%rd1], %rs1;
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.st.async.mmio.scope.sys.space.global.i16(ptr addrspace(1) %addr, i16 %value)
+ ret void
+}
+
+define void @test_st_async_mmio_release_sys_b32(ptr addrspace(1) %addr, i32 %value) {
+; CHECK-LABEL: test_st_async_mmio_release_sys_b32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_st_async_mmio_release_sys_b32_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_st_async_mmio_release_sys_b32_param_1];
+; CHECK-NEXT: st.async.mmio.release.sys.global.b32 [%rd1], %r1;
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.st.async.mmio.scope.sys.space.global.i32(ptr addrspace(1) %addr, i32 %value)
+ ret void
+}
+
+define void @test_st_async_mmio_release_sys_b64(ptr addrspace(1) %addr, i64 %value) {
+; CHECK-LABEL: test_st_async_mmio_release_sys_b64(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_st_async_mmio_release_sys_b64_param_0];
+; CHECK-NEXT: ld.param.b64 %rd2, [test_st_async_mmio_release_sys_b64_param_1];
+; CHECK-NEXT: st.async.mmio.release.sys.global.b64 [%rd1], %rd2;
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.st.async.mmio.scope.sys.space.global.i64(ptr addrspace(1) %addr, i64 %value)
+ ret void
+}
>From 031dc99a1cde67183194cfc696aabd8784e87bf1 Mon Sep 17 00:00:00 2001
From: Srinivasa Ravi <srinivasar at nvidia.com>
Date: Mon, 1 Jun 2026 09:55:39 +0000
Subject: [PATCH 2/7] add docs
---
llvm/docs/NVPTXUsage.rst | 90 ++++++++++++++++++++++++++++++++++++++++
1 file changed, 90 insertions(+)
diff --git a/llvm/docs/NVPTXUsage.rst b/llvm/docs/NVPTXUsage.rst
index 507786e2e076c..b9ba22e2fb592 100644
--- a/llvm/docs/NVPTXUsage.rst
+++ b/llvm/docs/NVPTXUsage.rst
@@ -3831,6 +3831,96 @@ similar but the latter uses generic addressing (see `Generic Addressing
For more information, refer `PTX ISA
<https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-st-bulk>`__.
+'``llvm.nvvm.st.async.space.cluster``'
+^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+Syntax:
+"""""""
+
+.. code-block:: llvm
+
+ declare void @llvm.nvvm.st.async.space.cluster.i32(ptr addrspace(7) %dest_addr, i32 %value, ptr addrspace(7) %mbarrier_addr)
+ declare void @llvm.nvvm.st.async.space.cluster.i64(ptr addrspace(7) %dest_addr, i64 %value, ptr addrspace(7) %mbarrier_addr)
+ declare void @llvm.nvvm.st.async.space.cluster.v2i32(ptr addrspace(7) %dest_addr, <2 x i32> %value, ptr addrspace(7) %mbarrier_addr)
+ declare void @llvm.nvvm.st.async.space.cluster.v2i64(ptr addrspace(7) %dest_addr, <2 x i64> %value, ptr addrspace(7) %mbarrier_addr)
+ declare void @llvm.nvvm.st.async.space.cluster.v4i32(ptr addrspace(7) %dest_addr, <4 x i32> %value, ptr addrspace(7) %mbarrier_addr)
+
+Overview:
+"""""""""
+
+The '``llvm.nvvm.st.async.space.cluster``' intrinsic initiates a weak
+asynchronous store operation to shared memory that stores the value specified
+by the `%value` operand to the destination address specified by the
+`%dest_addr` operand.
+
+The store operation is treated as a weak memory operation. The effects of this
+operation become visible to other threads only when synchronization is
+established by other means.
+
+The operation is performed asynchronously and the completion is signalled using
+the mbarrier object specified by the `%mbarrier_addr` operand. Upon completion,
+a `complete-tx <https://docs.nvidia.com/cuda/parallel-thread-execution/#parallel-synchronization-and-communication-instructions-mbarrier-complete-tx-operation>`__ operation is performed on the mbarrier object, with the
+`completeCount` argument equal to the amount of data stored in bytes.
+
+For more information, refer `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-st-async>`__.
+
+'``llvm.nvvm.st.async.scope.{sys,gpu}.*``'
+^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+Syntax:
+"""""""
+
+.. code-block:: llvm
+
+ ; sys scope
+ declare void @llvm.nvvm.st.async.scope.sys.space.global.i8(ptr addrspace(1) %dest_addr, i8 %value)
+ declare void @llvm.nvvm.st.async.scope.sys.space.global.i16(ptr addrspace(1) %dest_addr, i16 %value)
+ declare void @llvm.nvvm.st.async.scope.sys.space.global.i32(ptr addrspace(1) %dest_addr, i32 %value)
+ declare void @llvm.nvvm.st.async.scope.sys.space.global.i64(ptr addrspace(1) %dest_addr, i64 %value)
+
+ ; gpu scope
+ declare void @llvm.nvvm.st.async.scope.gpu.space.global.i8(ptr addrspace(1) %dest_addr, i8 %value)
+ declare void @llvm.nvvm.st.async.scope.gpu.space.global.i16(ptr addrspace(1) %dest_addr, i16 %value)
+ declare void @llvm.nvvm.st.async.scope.gpu.space.global.i32(ptr addrspace(1) %dest_addr, i32 %value)
+ declare void @llvm.nvvm.st.async.scope.gpu.space.global.i64(ptr addrspace(1) %dest_addr, i64 %value)
+
+Overview:
+"""""""""
+
+The '``llvm.nvvm.st.async.scope.sys.space.global``' and
+'``llvm.nvvm.st.async.scope.gpu.space.global``' intrinsics initiate an
+asynchronous release store to global memory that stores the value specified by
+the `%value` operand to the destination address specified by the `%dest_addr`
+operand.
+
+The store operation is treated as a strong memory operation with ``.release``
+semantics. The effects of prior stores from the current thread are made visible
+to some operations in other threads in the scope.
+
+The scope of this operation can be ``sys`` or ``gpu``.
+
+For more information, refer `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-st-async>`__.
+
+'``llvm.nvvm.st.async.mmio.scope.sys.*``'
+^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+Syntax:
+"""""""
+
+.. code-block:: llvm
+
+ declare void @llvm.nvvm.st.async.mmio.scope.sys.space.global.i8(ptr addrspace(1) %dest_addr, i8 %value)
+ declare void @llvm.nvvm.st.async.mmio.scope.sys.space.global.i16(ptr addrspace(1) %dest_addr, i16 %value)
+ declare void @llvm.nvvm.st.async.mmio.scope.sys.space.global.i32(ptr addrspace(1) %dest_addr, i32 %value)
+ declare void @llvm.nvvm.st.async.mmio.scope.sys.space.global.i64(ptr addrspace(1) %dest_addr, i64 %value)
+
+Overview:
+"""""""""
+
+The '``llvm.nvvm.st.async.mmio.scope.sys.space.global``' intrinsic performs an `MMIO <https://docs.nvidia.com/cuda/parallel-thread-execution/#mmio-operation>`__
+store to global memory with ``.release`` semantics at the ``sys`` scope.
+
+For more information, refer `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-st-async>`__.
clusterlaunchcontrol Intrinsics
-------------------------------
>From 99c47e7340a551a73ada468f785a9334f77e0a47 Mon Sep 17 00:00:00 2001
From: Srinivasa Ravi <srinivasar at nvidia.com>
Date: Mon, 1 Jun 2026 11:09:06 +0000
Subject: [PATCH 3/7] fix formatting
---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 45 ++++++++++++++-------
1 file changed, 30 insertions(+), 15 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 348006d320e19..2c8a3361a123a 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -2717,29 +2717,44 @@ static SDValue lowerStAsyncRelease(SDValue Op, SelectionDAG &DAG) {
switch (IntrinsicID) {
case Intrinsic::nvvm_st_async_scope_sys_space_global:
switch (ValueVT.SimpleTy) {
- case MVT::i8: return NVPTXISD::ST_ASYNC_SYS_B8;
- case MVT::i16: return NVPTXISD::ST_ASYNC_SYS_B16;
- case MVT::i32: return NVPTXISD::ST_ASYNC_SYS_B32;
- case MVT::i64: return NVPTXISD::ST_ASYNC_SYS_B64;
- default: break;
+ case MVT::i8:
+ return NVPTXISD::ST_ASYNC_SYS_B8;
+ case MVT::i16:
+ return NVPTXISD::ST_ASYNC_SYS_B16;
+ case MVT::i32:
+ return NVPTXISD::ST_ASYNC_SYS_B32;
+ case MVT::i64:
+ return NVPTXISD::ST_ASYNC_SYS_B64;
+ default:
+ break;
}
break;
case Intrinsic::nvvm_st_async_scope_gpu_space_global:
switch (ValueVT.SimpleTy) {
- case MVT::i8: return NVPTXISD::ST_ASYNC_GPU_B8;
- case MVT::i16: return NVPTXISD::ST_ASYNC_GPU_B16;
- case MVT::i32: return NVPTXISD::ST_ASYNC_GPU_B32;
- case MVT::i64: return NVPTXISD::ST_ASYNC_GPU_B64;
- default: break;
+ case MVT::i8:
+ return NVPTXISD::ST_ASYNC_GPU_B8;
+ case MVT::i16:
+ return NVPTXISD::ST_ASYNC_GPU_B16;
+ case MVT::i32:
+ return NVPTXISD::ST_ASYNC_GPU_B32;
+ case MVT::i64:
+ return NVPTXISD::ST_ASYNC_GPU_B64;
+ default:
+ break;
}
break;
case Intrinsic::nvvm_st_async_mmio_scope_sys_space_global:
switch (ValueVT.SimpleTy) {
- case MVT::i8: return NVPTXISD::ST_ASYNC_MMIO_SYS_B8;
- case MVT::i16: return NVPTXISD::ST_ASYNC_MMIO_SYS_B16;
- case MVT::i32: return NVPTXISD::ST_ASYNC_MMIO_SYS_B32;
- case MVT::i64: return NVPTXISD::ST_ASYNC_MMIO_SYS_B64;
- default: break;
+ case MVT::i8:
+ return NVPTXISD::ST_ASYNC_MMIO_SYS_B8;
+ case MVT::i16:
+ return NVPTXISD::ST_ASYNC_MMIO_SYS_B16;
+ case MVT::i32:
+ return NVPTXISD::ST_ASYNC_MMIO_SYS_B32;
+ case MVT::i64:
+ return NVPTXISD::ST_ASYNC_MMIO_SYS_B64;
+ default:
+ break;
}
break;
}
>From 104b4f78bb6e536372d8ebab7c447c89ed23f43b Mon Sep 17 00:00:00 2001
From: Srinivasa Ravi <srinivasar at nvidia.com>
Date: Wed, 3 Jun 2026 12:18:08 +0000
Subject: [PATCH 4/7] add multimem support
---
llvm/docs/NVPTXUsage.rst | 24 ++--
llvm/include/llvm/IR/IntrinsicsNVVM.td | 5 +-
.../NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp | 7 +
.../NVPTX/MCTargetDesc/NVPTXInstPrinter.h | 2 +
llvm/lib/Target/NVPTX/NVPTX.td | 2 +-
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 11 +-
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 43 +++---
llvm/test/CodeGen/NVPTX/st_async_release.ll | 16 +--
.../NVPTX/st_async_release_multimem.ll | 129 ++++++++++++++++++
9 files changed, 201 insertions(+), 38 deletions(-)
create mode 100644 llvm/test/CodeGen/NVPTX/st_async_release_multimem.ll
diff --git a/llvm/docs/NVPTXUsage.rst b/llvm/docs/NVPTXUsage.rst
index b9ba22e2fb592..eb6d9a627e77e 100644
--- a/llvm/docs/NVPTXUsage.rst
+++ b/llvm/docs/NVPTXUsage.rst
@@ -3873,16 +3873,16 @@ Syntax:
.. code-block:: llvm
; sys scope
- declare void @llvm.nvvm.st.async.scope.sys.space.global.i8(ptr addrspace(1) %dest_addr, i8 %value)
- declare void @llvm.nvvm.st.async.scope.sys.space.global.i16(ptr addrspace(1) %dest_addr, i16 %value)
- declare void @llvm.nvvm.st.async.scope.sys.space.global.i32(ptr addrspace(1) %dest_addr, i32 %value)
- declare void @llvm.nvvm.st.async.scope.sys.space.global.i64(ptr addrspace(1) %dest_addr, i64 %value)
+ declare void @llvm.nvvm.st.async.scope.sys.space.global.i8(ptr addrspace(1) %dest_addr, i8 %value, i1 immarg %is_multimem)
+ declare void @llvm.nvvm.st.async.scope.sys.space.global.i16(ptr addrspace(1) %dest_addr, i16 %value, i1 immarg %is_multimem)
+ declare void @llvm.nvvm.st.async.scope.sys.space.global.i32(ptr addrspace(1) %dest_addr, i32 %value, i1 immarg %is_multimem)
+ declare void @llvm.nvvm.st.async.scope.sys.space.global.i64(ptr addrspace(1) %dest_addr, i64 %value, i1 immarg %is_multimem)
; gpu scope
- declare void @llvm.nvvm.st.async.scope.gpu.space.global.i8(ptr addrspace(1) %dest_addr, i8 %value)
- declare void @llvm.nvvm.st.async.scope.gpu.space.global.i16(ptr addrspace(1) %dest_addr, i16 %value)
- declare void @llvm.nvvm.st.async.scope.gpu.space.global.i32(ptr addrspace(1) %dest_addr, i32 %value)
- declare void @llvm.nvvm.st.async.scope.gpu.space.global.i64(ptr addrspace(1) %dest_addr, i64 %value)
+ declare void @llvm.nvvm.st.async.scope.gpu.space.global.i8(ptr addrspace(1) %dest_addr, i8 %value, i1 immarg %is_multimem)
+ declare void @llvm.nvvm.st.async.scope.gpu.space.global.i16(ptr addrspace(1) %dest_addr, i16 %value, i1 immarg %is_multimem)
+ declare void @llvm.nvvm.st.async.scope.gpu.space.global.i32(ptr addrspace(1) %dest_addr, i32 %value, i1 immarg %is_multimem)
+ declare void @llvm.nvvm.st.async.scope.gpu.space.global.i64(ptr addrspace(1) %dest_addr, i64 %value, i1 immarg %is_multimem)
Overview:
"""""""""
@@ -3893,12 +3893,20 @@ asynchronous release store to global memory that stores the value specified by
the `%value` operand to the destination address specified by the `%dest_addr`
operand.
+The `%is_multimem` immediate argument selects the variant of the store. When it
+is `0`, a regular ``st.async`` is emitted. When it is `1`, a
+``multimem.st.async`` is emitted and the `%dest_addr` operand must be a
+multimem address.
+
The store operation is treated as a strong memory operation with ``.release``
semantics. The effects of prior stores from the current thread are made visible
to some operations in other threads in the scope.
The scope of this operation can be ``sys`` or ``gpu``.
+These intrinsics lower to the `.b8`, `.b16`, `.b32`, and `.b64` variants of the
+``st.async`` instruction.
+
For more information, refer `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-st-async>`__.
'``llvm.nvvm.st.async.mmio.scope.sys.*``'
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 9e72aa6d8afe2..1b840bdf41cec 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -3284,8 +3284,9 @@ def int_nvvm_st_async_space_cluster :
foreach scope = ["_sys", "_gpu"] in {
def int_nvvm_st_async_scope # scope # _space_global :
DefaultAttrsIntrinsic<[],
- [llvm_global_ptr_ty, llvm_anyint_ty],
- [WriteOnly<ArgIndex<0>>]>;
+ [llvm_global_ptr_ty, llvm_anyint_ty, llvm_i1_ty],
+ [WriteOnly<ArgIndex<0>>, ImmArg<ArgIndex<2>>,
+ ArgInfo<ArgIndex<2>, [ArgName<"isMultimem">]>]>;
}
def int_nvvm_st_async_mmio_scope_sys_space_global :
diff --git a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
index 5af9e75bca399..6e123fecf93d1 100644
--- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
+++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
@@ -178,6 +178,13 @@ void NVPTXInstPrinter::printFTZFlag(const MCInst *MI, int OpNum,
O << ".ftz";
}
+void NVPTXInstPrinter::printMultimem(const MCInst *MI, int OpNum,
+ const MCSubtargetInfo &, raw_ostream &O) {
+ const MCOperand &MO = MI->getOperand(OpNum);
+ if (MO.getImm())
+ O << "multimem.";
+}
+
void NVPTXInstPrinter::printNegatedPredicate(const MCInst *MI, int OpNum,
const MCSubtargetInfo &,
raw_ostream &O) {
diff --git a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.h b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.h
index 3f5dd99a2be05..d9a6c3c85a89b 100644
--- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.h
+++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.h
@@ -66,6 +66,8 @@ class NVPTXInstPrinter : public MCInstPrinter {
raw_ostream &O, StringRef Modifier = {});
void printFTZFlag(const MCInst *MI, int OpNum, const MCSubtargetInfo &STI,
raw_ostream &O);
+ void printMultimem(const MCInst *MI, int OpNum, const MCSubtargetInfo &STI,
+ raw_ostream &O);
void printNegatedPredicate(const MCInst *MI, int OpNum,
const MCSubtargetInfo &STI, raw_ostream &O);
diff --git a/llvm/lib/Target/NVPTX/NVPTX.td b/llvm/lib/Target/NVPTX/NVPTX.td
index 1c169b05841c3..0f02a81e3cefc 100644
--- a/llvm/lib/Target/NVPTX/NVPTX.td
+++ b/llvm/lib/Target/NVPTX/NVPTX.td
@@ -106,7 +106,7 @@ foreach sm = [20, 21, 30, 32, 35, 37, 50, 52, 53, 60,
foreach version = [32, 40, 41, 42, 43, 50, 60, 61, 62, 63, 64, 65, 70, 71, 72,
73, 74, 75, 76, 77, 78, 80, 81, 82, 83, 84, 85, 86, 87, 88,
- 90, 91, 92] in
+ 90, 91, 92, 93] in
def PTX#version : FeaturePTX<version>;
def Is64Bit : Predicate<"Subtarget->getTargetTriple().getArch() == Triple::nvptx64">;
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 2c8a3361a123a..4a1977800afcb 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -2776,7 +2776,16 @@ static SDValue lowerStAsyncRelease(SDValue Op, SelectionDAG &DAG) {
if (ValueVT == MVT::i8)
Value = DAG.getNode(ISD::ZERO_EXTEND, DL, MVT::i16, Value);
- SDValue Ops[] = {N->getOperand(0), DestAddr, Value};
+ // The `.mmio` variant has no multimem form and therefore no `isMultimem`
+ // operand.
+ if (IntrinsicID == Intrinsic::nvvm_st_async_mmio_scope_sys_space_global) {
+ SDValue Ops[] = {N->getOperand(0), DestAddr, Value};
+ return DAG.getNode(OpCode.value(), DL, MVT::Other, Ops);
+ }
+
+ SDValue IsMultimem =
+ DAG.getTargetConstant(N->getConstantOperandVal(4), DL, MVT::i32);
+ SDValue Ops[] = {N->getOperand(0), DestAddr, Value, IsMultimem};
return DAG.getNode(OpCode.value(), DL, MVT::Other, Ops);
}
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 1e3fd12ace8bf..5f5f5b6a1ebc3 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -6143,14 +6143,21 @@ let Predicates = [hasSM<90>, hasPTX<81>] in {
i32:$value3, i32:$value4, addr:$mbar)>;
}
-def SDTAsyncStoreRelease : SDTypeProfile<0, 2, [SDTCisPtrTy<0>]>;
+def SDTAsyncStoreRelease : SDTypeProfile<0, 3,
+ [SDTCisPtrTy<0>, SDTCisVT<2, i32>]>;
+def SDTAsyncStoreMMIO : SDTypeProfile<0, 2, [SDTCisPtrTy<0>]>;
-class StAsyncReleaseNode<string variant, string type> :
+class StAsyncReleaseNode<string variant, string type, SDTypeProfile sdt> :
SDNode<"NVPTXISD::ST_ASYNC_" # variant # "_" # type,
- SDTAsyncStoreRelease,
+ sdt,
[SDNPHasChain, SDNPMayStore]>;
-class ST_ASYNC_RELEASE_BODY<string instr_prefix, string type> {
+def MultimemFlag : Operand<i32> {
+ let PrintMethod = "printMultimem";
+}
+
+class ST_ASYNC_RELEASE_BODY<string instr_prefix, string type, bit is_mmio> {
+ string prefix = !if(is_mmio, "", "${is_multimem}");
string body = !if(!eq(type, "b8"),
// The minimum supported register size to pass in arguments is 16-bits, so
// the 8-bit value is passed in using a 16-bit register and truncated to
@@ -6158,25 +6165,25 @@ class ST_ASYNC_RELEASE_BODY<string instr_prefix, string type> {
"{{ \n\t"
".reg .b8 \t%st_async_val; \n\t"
"cvt.u8.u16 \t%st_async_val, $value; \n\t" #
- instr_prefix # ".b8\t [$dest_addr], %st_async_val; \n\t"
+ prefix # instr_prefix # ".b8\t [$dest_addr], %st_async_val; \n\t"
"}}",
- instr_prefix # "." # type # "\t [$dest_addr], $value;");
+ prefix # instr_prefix # "." # type # "\t [$dest_addr], $value;");
}
class ST_ASYNC_SYS<string type, NVPTXRegClass reg> :
NVPTXInst<(outs),
- (ins ADDR:$dest_addr, reg:$value),
- ST_ASYNC_RELEASE_BODY<"st.async.release.sys.global", type>.body>;
+ (ins ADDR:$dest_addr, reg:$value, MultimemFlag:$is_multimem),
+ ST_ASYNC_RELEASE_BODY<"st.async.release.sys.global", type, 0>.body>;
class ST_ASYNC_GPU<string type, NVPTXRegClass reg> :
NVPTXInst<(outs),
- (ins ADDR:$dest_addr, reg:$value),
- ST_ASYNC_RELEASE_BODY<"st.async.release.gpu.global", type>.body>;
+ (ins ADDR:$dest_addr, reg:$value, MultimemFlag:$is_multimem),
+ ST_ASYNC_RELEASE_BODY<"st.async.release.gpu.global", type, 0>.body>;
class ST_ASYNC_MMIO_SYS<string type, NVPTXRegClass reg> :
NVPTXInst<(outs),
(ins ADDR:$dest_addr, reg:$value),
- ST_ASYNC_RELEASE_BODY<"st.async.mmio.release.sys.global", type>.body>;
+ ST_ASYNC_RELEASE_BODY<"st.async.mmio.release.sys.global", type, 1>.body>;
let Predicates = [hasSM<100>, hasPTX<87>] in {
@@ -6186,15 +6193,15 @@ let Predicates = [hasSM<100>, hasPTX<87>] in {
defvar valueType = [i16, i16, i32, i64 ][i];
defvar width = "b" # size;
- def : Pat<(StAsyncReleaseNode<"SYS", !toupper(width)>
- addr:$dest_addr, valueType:$value),
+ def : Pat<(StAsyncReleaseNode<"SYS", !toupper(width), SDTAsyncStoreRelease>
+ addr:$dest_addr, valueType:$value, timm:$is_multimem),
(ST_ASYNC_SYS<width, reg>
- addr:$dest_addr, valueType:$value)>;
- def : Pat<(StAsyncReleaseNode<"GPU", !toupper(width)>
- addr:$dest_addr, valueType:$value),
+ addr:$dest_addr, valueType:$value, timm:$is_multimem)>;
+ def : Pat<(StAsyncReleaseNode<"GPU", !toupper(width), SDTAsyncStoreRelease>
+ addr:$dest_addr, valueType:$value, timm:$is_multimem),
(ST_ASYNC_GPU<width, reg>
- addr:$dest_addr, valueType:$value)>;
- def : Pat<(StAsyncReleaseNode<"MMIO_SYS", !toupper(width)>
+ addr:$dest_addr, valueType:$value, timm:$is_multimem)>;
+ def : Pat<(StAsyncReleaseNode<"MMIO_SYS", !toupper(width), SDTAsyncStoreMMIO>
addr:$dest_addr, valueType:$value),
(ST_ASYNC_MMIO_SYS<width, reg>
addr:$dest_addr, valueType:$value)>;
diff --git a/llvm/test/CodeGen/NVPTX/st_async_release.ll b/llvm/test/CodeGen/NVPTX/st_async_release.ll
index 091b6bcd71f2d..9bf4d3c8f789f 100644
--- a/llvm/test/CodeGen/NVPTX/st_async_release.ll
+++ b/llvm/test/CodeGen/NVPTX/st_async_release.ll
@@ -17,7 +17,7 @@ define void @test_st_async_release_sys_b8(ptr addrspace(1) %addr, i8 %value) {
; CHECK-NEXT: st.async.release.sys.global.b8 [%rd1], %st_async_val;
; CHECK-NEXT: }
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.scope.sys.space.global.i8(ptr addrspace(1) %addr, i8 %value)
+ call void @llvm.nvvm.st.async.scope.sys.space.global.i8(ptr addrspace(1) %addr, i8 %value, /* isMultimem= */ i1 false)
ret void
}
@@ -32,7 +32,7 @@ define void @test_st_async_release_sys_b16(ptr addrspace(1) %addr, i16 %value) {
; CHECK-NEXT: ld.param.b16 %rs1, [test_st_async_release_sys_b16_param_1];
; CHECK-NEXT: st.async.release.sys.global.b16 [%rd1], %rs1;
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.scope.sys.space.global.i16(ptr addrspace(1) %addr, i16 %value)
+ call void @llvm.nvvm.st.async.scope.sys.space.global.i16(ptr addrspace(1) %addr, i16 %value, /* isMultimem= */ i1 false)
ret void
}
@@ -47,7 +47,7 @@ define void @test_st_async_release_sys_b32(ptr addrspace(1) %addr, i32 %value) {
; CHECK-NEXT: ld.param.b32 %r1, [test_st_async_release_sys_b32_param_1];
; CHECK-NEXT: st.async.release.sys.global.b32 [%rd1], %r1;
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.scope.sys.space.global.i32(ptr addrspace(1) %addr, i32 %value)
+ call void @llvm.nvvm.st.async.scope.sys.space.global.i32(ptr addrspace(1) %addr, i32 %value, /* isMultimem= */ i1 false)
ret void
}
@@ -61,7 +61,7 @@ define void @test_st_async_release_sys_b64(ptr addrspace(1) %addr, i64 %value) {
; CHECK-NEXT: ld.param.b64 %rd2, [test_st_async_release_sys_b64_param_1];
; CHECK-NEXT: st.async.release.sys.global.b64 [%rd1], %rd2;
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.scope.sys.space.global.i64(ptr addrspace(1) %addr, i64 %value)
+ call void @llvm.nvvm.st.async.scope.sys.space.global.i64(ptr addrspace(1) %addr, i64 %value, /* isMultimem= */ i1 false)
ret void
}
@@ -80,7 +80,7 @@ define void @test_st_async_release_gpu_b8(ptr addrspace(1) %addr, i8 %value) {
; CHECK-NEXT: st.async.release.gpu.global.b8 [%rd1], %st_async_val;
; CHECK-NEXT: }
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.scope.gpu.space.global.i8(ptr addrspace(1) %addr, i8 %value)
+ call void @llvm.nvvm.st.async.scope.gpu.space.global.i8(ptr addrspace(1) %addr, i8 %value, /* isMultimem= */ i1 false)
ret void
}
@@ -95,7 +95,7 @@ define void @test_st_async_release_gpu_b16(ptr addrspace(1) %addr, i16 %value) {
; CHECK-NEXT: ld.param.b16 %rs1, [test_st_async_release_gpu_b16_param_1];
; CHECK-NEXT: st.async.release.gpu.global.b16 [%rd1], %rs1;
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.scope.gpu.space.global.i16(ptr addrspace(1) %addr, i16 %value)
+ call void @llvm.nvvm.st.async.scope.gpu.space.global.i16(ptr addrspace(1) %addr, i16 %value, /* isMultimem= */ i1 false)
ret void
}
@@ -110,7 +110,7 @@ define void @test_st_async_release_gpu_b32(ptr addrspace(1) %addr, i32 %value) {
; CHECK-NEXT: ld.param.b32 %r1, [test_st_async_release_gpu_b32_param_1];
; CHECK-NEXT: st.async.release.gpu.global.b32 [%rd1], %r1;
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.scope.gpu.space.global.i32(ptr addrspace(1) %addr, i32 %value)
+ call void @llvm.nvvm.st.async.scope.gpu.space.global.i32(ptr addrspace(1) %addr, i32 %value, /* isMultimem= */ i1 false)
ret void
}
@@ -124,7 +124,7 @@ define void @test_st_async_release_gpu_b64(ptr addrspace(1) %addr, i64 %value) {
; CHECK-NEXT: ld.param.b64 %rd2, [test_st_async_release_gpu_b64_param_1];
; CHECK-NEXT: st.async.release.gpu.global.b64 [%rd1], %rd2;
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.scope.gpu.space.global.i64(ptr addrspace(1) %addr, i64 %value)
+ call void @llvm.nvvm.st.async.scope.gpu.space.global.i64(ptr addrspace(1) %addr, i64 %value, /* isMultimem= */ i1 false)
ret void
}
diff --git a/llvm/test/CodeGen/NVPTX/st_async_release_multimem.ll b/llvm/test/CodeGen/NVPTX/st_async_release_multimem.ll
new file mode 100644
index 0000000000000..0631a08d3e15b
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/st_async_release_multimem.ll
@@ -0,0 +1,129 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx93 | FileCheck %s
+; RUN: %if ptxas-sm_100 && ptxas-isa-9.3 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx93 | %ptxas-verify -arch=sm_100 %}
+
+define void @test_multimem_st_async_release_sys_b8(ptr addrspace(1) %addr, i8 %value) {
+; CHECK-LABEL: test_multimem_st_async_release_sys_b8(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_multimem_st_async_release_sys_b8_param_0];
+; CHECK-NEXT: ld.param.b8 %rs1, [test_multimem_st_async_release_sys_b8_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %st_async_val;
+; CHECK-NEXT: cvt.u8.u16 %st_async_val, %rs1;
+; CHECK-NEXT: multimem.st.async.release.sys.global.b8 [%rd1], %st_async_val;
+; CHECK-NEXT: }
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.st.async.scope.sys.space.global.i8(ptr addrspace(1) %addr, i8 %value, /* isMultimem= */ i1 true)
+ ret void
+}
+
+define void @test_multimem_st_async_release_sys_b16(ptr addrspace(1) %addr, i16 %value) {
+; CHECK-LABEL: test_multimem_st_async_release_sys_b16(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_multimem_st_async_release_sys_b16_param_0];
+; CHECK-NEXT: ld.param.b16 %rs1, [test_multimem_st_async_release_sys_b16_param_1];
+; CHECK-NEXT: multimem.st.async.release.sys.global.b16 [%rd1], %rs1;
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.st.async.scope.sys.space.global.i16(ptr addrspace(1) %addr, i16 %value, /* isMultimem= */ i1 true)
+ ret void
+}
+
+define void @test_multimem_st_async_release_sys_b32(ptr addrspace(1) %addr, i32 %value) {
+; CHECK-LABEL: test_multimem_st_async_release_sys_b32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_multimem_st_async_release_sys_b32_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_multimem_st_async_release_sys_b32_param_1];
+; CHECK-NEXT: multimem.st.async.release.sys.global.b32 [%rd1], %r1;
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.st.async.scope.sys.space.global.i32(ptr addrspace(1) %addr, i32 %value, /* isMultimem= */ i1 true)
+ ret void
+}
+
+define void @test_multimem_st_async_release_sys_b64(ptr addrspace(1) %addr, i64 %value) {
+; CHECK-LABEL: test_multimem_st_async_release_sys_b64(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_multimem_st_async_release_sys_b64_param_0];
+; CHECK-NEXT: ld.param.b64 %rd2, [test_multimem_st_async_release_sys_b64_param_1];
+; CHECK-NEXT: multimem.st.async.release.sys.global.b64 [%rd1], %rd2;
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.st.async.scope.sys.space.global.i64(ptr addrspace(1) %addr, i64 %value, /* isMultimem= */ i1 true)
+ ret void
+}
+
+define void @test_multimem_st_async_release_gpu_b8(ptr addrspace(1) %addr, i8 %value) {
+; CHECK-LABEL: test_multimem_st_async_release_gpu_b8(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_multimem_st_async_release_gpu_b8_param_0];
+; CHECK-NEXT: ld.param.b8 %rs1, [test_multimem_st_async_release_gpu_b8_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %st_async_val;
+; CHECK-NEXT: cvt.u8.u16 %st_async_val, %rs1;
+; CHECK-NEXT: multimem.st.async.release.gpu.global.b8 [%rd1], %st_async_val;
+; CHECK-NEXT: }
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.st.async.scope.gpu.space.global.i8(ptr addrspace(1) %addr, i8 %value, /* isMultimem= */ i1 true)
+ ret void
+}
+
+define void @test_multimem_st_async_release_gpu_b16(ptr addrspace(1) %addr, i16 %value) {
+; CHECK-LABEL: test_multimem_st_async_release_gpu_b16(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_multimem_st_async_release_gpu_b16_param_0];
+; CHECK-NEXT: ld.param.b16 %rs1, [test_multimem_st_async_release_gpu_b16_param_1];
+; CHECK-NEXT: multimem.st.async.release.gpu.global.b16 [%rd1], %rs1;
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.st.async.scope.gpu.space.global.i16(ptr addrspace(1) %addr, i16 %value, /* isMultimem= */ i1 true)
+ ret void
+}
+
+define void @test_multimem_st_async_release_gpu_b32(ptr addrspace(1) %addr, i32 %value) {
+; CHECK-LABEL: test_multimem_st_async_release_gpu_b32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_multimem_st_async_release_gpu_b32_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_multimem_st_async_release_gpu_b32_param_1];
+; CHECK-NEXT: multimem.st.async.release.gpu.global.b32 [%rd1], %r1;
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.st.async.scope.gpu.space.global.i32(ptr addrspace(1) %addr, i32 %value, /* isMultimem= */ i1 true)
+ ret void
+}
+
+define void @test_multimem_st_async_release_gpu_b64(ptr addrspace(1) %addr, i64 %value) {
+; CHECK-LABEL: test_multimem_st_async_release_gpu_b64(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_multimem_st_async_release_gpu_b64_param_0];
+; CHECK-NEXT: ld.param.b64 %rd2, [test_multimem_st_async_release_gpu_b64_param_1];
+; CHECK-NEXT: multimem.st.async.release.gpu.global.b64 [%rd1], %rd2;
+; CHECK-NEXT: ret;
+ call void @llvm.nvvm.st.async.scope.gpu.space.global.i64(ptr addrspace(1) %addr, i64 %value, /* isMultimem= */ i1 true)
+ ret void
+}
>From d3c18c05a0b23ae9ef4e92f15e95474c2edf5878 Mon Sep 17 00:00:00 2001
From: Srinivasa Ravi <srinivasar at nvidia.com>
Date: Wed, 3 Jun 2026 13:05:49 +0000
Subject: [PATCH 5/7] add support for i128 for st.async.space.cluster instead
of vector types
---
llvm/docs/NVPTXUsage.rst | 8 +-
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 68 +++++-----------
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 57 +++++--------
llvm/test/CodeGen/NVPTX/st_async_mbarrier.ll | 81 -------------------
.../CodeGen/NVPTX/st_async_mbarrier_b128.ll | 40 +++++++++
5 files changed, 84 insertions(+), 170 deletions(-)
create mode 100644 llvm/test/CodeGen/NVPTX/st_async_mbarrier_b128.ll
diff --git a/llvm/docs/NVPTXUsage.rst b/llvm/docs/NVPTXUsage.rst
index eb6d9a627e77e..41c77f33b9611 100644
--- a/llvm/docs/NVPTXUsage.rst
+++ b/llvm/docs/NVPTXUsage.rst
@@ -3841,9 +3841,7 @@ Syntax:
declare void @llvm.nvvm.st.async.space.cluster.i32(ptr addrspace(7) %dest_addr, i32 %value, ptr addrspace(7) %mbarrier_addr)
declare void @llvm.nvvm.st.async.space.cluster.i64(ptr addrspace(7) %dest_addr, i64 %value, ptr addrspace(7) %mbarrier_addr)
- declare void @llvm.nvvm.st.async.space.cluster.v2i32(ptr addrspace(7) %dest_addr, <2 x i32> %value, ptr addrspace(7) %mbarrier_addr)
- declare void @llvm.nvvm.st.async.space.cluster.v2i64(ptr addrspace(7) %dest_addr, <2 x i64> %value, ptr addrspace(7) %mbarrier_addr)
- declare void @llvm.nvvm.st.async.space.cluster.v4i32(ptr addrspace(7) %dest_addr, <4 x i32> %value, ptr addrspace(7) %mbarrier_addr)
+ declare void @llvm.nvvm.st.async.space.cluster.i128(ptr addrspace(7) %dest_addr, i128 %value, ptr addrspace(7) %mbarrier_addr)
Overview:
"""""""""
@@ -3851,7 +3849,9 @@ Overview:
The '``llvm.nvvm.st.async.space.cluster``' intrinsic initiates a weak
asynchronous store operation to shared memory that stores the value specified
by the `%value` operand to the destination address specified by the
-`%dest_addr` operand.
+`%dest_addr` operand. The `%value` operand must be an ``i32``, ``i64`` or
+``i128``, lowering to the ``.b32``, ``.b64`` and ``.b128`` variants
+of the ``st.async`` PTX instruction respectively.
The store operation is treated as a weak memory operation. The effects of this
operation become visible to other threads only when synchronization is
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 4a1977800afcb..1a5a79a224995 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -1116,10 +1116,12 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM,
MVT::v64f32, MVT::v128f32},
Custom);
- // Custom lowering for tcgen05.st vector operands
+ // Custom lowering for tcgen05.st vector operands and the st.async mbarrier
+ // i128 (.b128) operand. MVT::i8 is needed for the st.async.release b8
+ // variant.
setOperationAction(ISD::INTRINSIC_VOID,
{MVT::i8, MVT::v2i32, MVT::v4i32, MVT::v8i32, MVT::v16i32,
- MVT::v32i32, MVT::v64i32, MVT::v128i32, MVT::v2i64,
+ MVT::v32i32, MVT::v64i32, MVT::v128i32, MVT::i128,
MVT::Other},
Custom);
@@ -2652,55 +2654,25 @@ static SDValue lowerStAsyncWithMbarrier(SDValue Op, SelectionDAG &DAG) {
MVT ValueVT = Value.getSimpleValueType();
- if (ValueVT.getSizeInBits() > 128) {
- DAG.getContext()->diagnose(DiagnosticInfoUnsupported(
- Fn, "total bit-width of the value to be stored must be <= 128",
- DiagnosticLocation(DL.getDebugLoc())));
- return Op.getOperand(0); // Return only the chain
- }
-
- auto OpCode = [&]() -> std::optional<unsigned> {
- switch (ValueVT.SimpleTy) {
- case MVT::i32:
- return NVPTXISD::ST_ASYNC_MBARRIER_B32;
- case MVT::i64:
- return NVPTXISD::ST_ASYNC_MBARRIER_B64;
- case MVT::v2i32:
- return NVPTXISD::ST_ASYNC_MBARRIER_V2B32;
- case MVT::v2i64:
- return NVPTXISD::ST_ASYNC_MBARRIER_V2B64;
- case MVT::v4i32:
- return NVPTXISD::ST_ASYNC_MBARRIER_V4B32;
- default:
- DAG.getContext()->diagnose(DiagnosticInfoUnsupported(
- Fn,
- Twine("unsupported argument type ") +
- llvm::EVT(ValueVT).getEVTString() + " for " +
- llvm::Intrinsic::getName(IntrinsicID) + " intrinsic",
- DiagnosticLocation(DL.getDebugLoc())));
- return {};
- }
- }();
-
- if (!OpCode)
- return Op.getOperand(0); // Return only the chain
-
- SmallVector<SDValue, 6> Ops;
+ if (ValueVT == MVT::i32 || ValueVT == MVT::i64)
+ return Op;
- Ops.push_back(N->getOperand(0)); // Chain
- Ops.push_back(DestAddr);
- if (ValueVT.isVector()) {
- for (unsigned i = 0; i < ValueVT.getVectorNumElements(); ++i) {
- Ops.push_back(DAG.getNode(ISD::EXTRACT_VECTOR_ELT, DL,
- ValueVT.getVectorElementType(), Value,
- DAG.getIntPtrConstant(i, DL)));
- }
- } else {
- Ops.push_back(Value);
+ if (ValueVT == MVT::i128) {
+ SDValue Cast = DAG.getNode(ISD::BITCAST, DL, MVT::v2i64, Value);
+ SDValue ValueLo = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, DL, MVT::i64, Cast,
+ DAG.getIntPtrConstant(0, DL));
+ SDValue ValueHi = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, DL, MVT::i64, Cast,
+ DAG.getIntPtrConstant(1, DL));
+ SDValue Ops[] = {N->getOperand(0), DestAddr, ValueLo, ValueHi, MbarAddr};
+ return DAG.getNode(NVPTXISD::ST_ASYNC_MBARRIER_B128, DL, MVT::Other, Ops);
}
- Ops.push_back(MbarAddr);
- return DAG.getNode(OpCode.value(), DL, MVT::Other, Ops);
+ DAG.getContext()->diagnose(DiagnosticInfoUnsupported(
+ Fn,
+ Twine("unsupported argument type ") + llvm::EVT(ValueVT).getEVTString() +
+ " for " + llvm::Intrinsic::getName(IntrinsicID) + " intrinsic",
+ DiagnosticLocation(DL.getDebugLoc())));
+ return Op.getOperand(0); // Return only the chain
}
static SDValue lowerStAsyncRelease(SDValue Op, SelectionDAG &DAG) {
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 5f5f5b6a1ebc3..355f3cbbe55ce 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -6087,60 +6087,43 @@ let Predicates = [hasSM<100>, hasPTX<86>] in {
// Asynchronous store instructions
//
-def SDTAsyncStoreMbarrier : SDTypeProfile<0, 3,
- [SDTCisPtrTy<0>, SDTCisPtrTy<2>]>;
-def SDTAsyncStoreMbarrierV2 : SDTypeProfile<0, 4,
- [SDTCisPtrTy<0>, SDTCisPtrTy<3>]>;
-def SDTAsyncStoreMbarrierV4 : SDTypeProfile<0, 6,
- [SDTCisPtrTy<0>, SDTCisPtrTy<5>]>;
-
-class StAsyncMbarrierNode<string vec, string type> :
- SDNode<"NVPTXISD::ST_ASYNC_MBARRIER_" # vec # type,
- !cast<SDTypeProfile>("SDTAsyncStoreMbarrier" # vec),
+def SDTAsyncStoreMbarrierB128 : SDTypeProfile<0, 4,
+ [SDTCisPtrTy<0>, SDTCisVT<1, i64>, SDTCisVT<2, i64>, SDTCisPtrTy<3>]>;
+def st_async_mbarrier_b128 :
+ SDNode<"NVPTXISD::ST_ASYNC_MBARRIER_B128", SDTAsyncStoreMbarrierB128,
[SDNPHasChain, SDNPMayStore]>;
-
class ST_ASYNC_MBARRIER<string type, NVPTXRegClass reg> :
NVPTXInst<(outs),
(ins ADDR:$dest_addr, reg:$value, ADDR:$mbar),
"st.async.shared::cluster.mbarrier::complete_tx::bytes." # type # "\t [$dest_addr], $value, [$mbar];">;
-class ST_ASYNC_MBARRIER_V2<string type, NVPTXRegClass reg> :
- NVPTXInst<(outs),
- (ins ADDR:$dest_addr, reg:$value1, reg:$value2, ADDR:$mbar),
- "st.async.shared::cluster.mbarrier::complete_tx::bytes.v2." # type # "\t [$dest_addr], {{$value1, $value2}}, [$mbar];">;
-
-class ST_ASYNC_MBARRIER_V4<string type, NVPTXRegClass reg> :
+def ST_ASYNC_MBARRIER_B128 :
NVPTXInst<(outs),
- (ins ADDR:$dest_addr, reg:$value1, reg:$value2, reg:$value3,
- reg:$value4, ADDR:$mbar),
- "st.async.shared::cluster.mbarrier::complete_tx::bytes.v4." # type # "\t [$dest_addr], {{$value1, $value2, $value3, $value4}}, [$mbar];">;
-
+ (ins ADDR:$dest_addr, B64:$value0, B64:$value1, ADDR:$mbar),
+ "{{\n\t"
+ ".reg .b128 %in_128;\n\t"
+ "mov.b128 %in_128, {$value0, $value1};\n\t"
+ "st.async.shared::cluster.mbarrier::complete_tx::bytes.b128 \t[$dest_addr], %in_128, [$mbar];\n\t"
+ "}}">;
let Predicates = [hasSM<90>, hasPTX<81>] in {
foreach size = ["32", "64"] in {
defvar valueType = !cast<ValueType>("i" # size);
defvar reg = !cast<NVPTXRegClass>("B" # size);
- defvar type_str = "B" # size;
- def : Pat<(StAsyncMbarrierNode<"", type_str>
- addr:$dest_addr, valueType:$value, addr:$mbar),
- (ST_ASYNC_MBARRIER<!tolower(type_str), reg>
+ def : Pat<(int_nvvm_st_async_space_cluster
+ addr:$dest_addr, valueType:$value, addr:$mbar),
+ (ST_ASYNC_MBARRIER<"b" # size, reg>
addr:$dest_addr, valueType:$value, addr:$mbar)>;
- def : Pat<(StAsyncMbarrierNode<"V2", type_str>
- addr:$dest_addr, valueType:$value1, valueType:$value2,
- addr:$mbar),
- (ST_ASYNC_MBARRIER_V2<!tolower(type_str), reg>
- addr:$dest_addr, valueType:$value1, valueType:$value2,
- addr:$mbar)>;
}
+}
- def : Pat<(StAsyncMbarrierNode<"V4", "B32">
- addr:$dest_addr, i32:$value1, i32:$value2,
- i32:$value3, i32:$value4, addr:$mbar),
- (ST_ASYNC_MBARRIER_V4<"b32", B32>
- addr:$dest_addr, i32:$value1, i32:$value2,
- i32:$value3, i32:$value4, addr:$mbar)>;
+let Predicates = [hasSM<90>, hasPTX<92>] in {
+ def : Pat<(st_async_mbarrier_b128
+ addr:$dest_addr, i64:$value0, i64:$value1, addr:$mbar),
+ (ST_ASYNC_MBARRIER_B128
+ addr:$dest_addr, B64:$value0, B64:$value1, addr:$mbar)>;
}
def SDTAsyncStoreRelease : SDTypeProfile<0, 3,
diff --git a/llvm/test/CodeGen/NVPTX/st_async_mbarrier.ll b/llvm/test/CodeGen/NVPTX/st_async_mbarrier.ll
index 8ccb0e3e69a6c..76c41d505d5d8 100644
--- a/llvm/test/CodeGen/NVPTX/st_async_mbarrier.ll
+++ b/llvm/test/CodeGen/NVPTX/st_async_mbarrier.ll
@@ -57,84 +57,3 @@ define void @test_st_async_mbarrier_b64(ptr addrspace(7) %addr, i64 %value, ptr
call void @llvm.nvvm.st.async.space.cluster.i64(ptr addrspace(7) %addr, i64 %value, ptr addrspace(7) %mbar)
ret void
}
-
-define void @test_st_async_mbarrier_v2b32(ptr addrspace(7) %addr, <2 x i32> %value, ptr addrspace(7) %mbar) {
-; CHECK-PTX64-LABEL: test_st_async_mbarrier_v2b32(
-; CHECK-PTX64: {
-; CHECK-PTX64-NEXT: .reg .b32 %r<3>;
-; CHECK-PTX64-NEXT: .reg .b64 %rd<3>;
-; CHECK-PTX64-EMPTY:
-; CHECK-PTX64-NEXT: // %bb.0:
-; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [test_st_async_mbarrier_v2b32_param_0];
-; CHECK-PTX64-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_st_async_mbarrier_v2b32_param_1];
-; CHECK-PTX64-NEXT: ld.param.b64 %rd2, [test_st_async_mbarrier_v2b32_param_2];
-; CHECK-PTX64-NEXT: st.async.shared::cluster.mbarrier::complete_tx::bytes.v2.b32 [%rd1], {%r1, %r2}, [%rd2];
-; CHECK-PTX64-NEXT: ret;
-;
-; CHECK-SHARED32-LABEL: test_st_async_mbarrier_v2b32(
-; CHECK-SHARED32: {
-; CHECK-SHARED32-NEXT: .reg .b32 %r<5>;
-; CHECK-SHARED32-EMPTY:
-; CHECK-SHARED32-NEXT: // %bb.0:
-; CHECK-SHARED32-NEXT: ld.param.b32 %r1, [test_st_async_mbarrier_v2b32_param_0];
-; CHECK-SHARED32-NEXT: ld.param.v2.b32 {%r2, %r3}, [test_st_async_mbarrier_v2b32_param_1];
-; CHECK-SHARED32-NEXT: ld.param.b32 %r4, [test_st_async_mbarrier_v2b32_param_2];
-; CHECK-SHARED32-NEXT: st.async.shared::cluster.mbarrier::complete_tx::bytes.v2.b32 [%r1], {%r2, %r3}, [%r4];
-; CHECK-SHARED32-NEXT: ret;
- call void @llvm.nvvm.st.async.space.cluster.v2i32(ptr addrspace(7) %addr, <2 x i32> %value, ptr addrspace(7) %mbar)
- ret void
-}
-
-define void @test_st_async_mbarrier_v2b64(ptr addrspace(7) %addr, <2 x i64> %value, ptr addrspace(7) %mbar) {
-; CHECK-PTX64-LABEL: test_st_async_mbarrier_v2b64(
-; CHECK-PTX64: {
-; CHECK-PTX64-NEXT: .reg .b64 %rd<5>;
-; CHECK-PTX64-EMPTY:
-; CHECK-PTX64-NEXT: // %bb.0:
-; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [test_st_async_mbarrier_v2b64_param_0];
-; CHECK-PTX64-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_st_async_mbarrier_v2b64_param_1];
-; CHECK-PTX64-NEXT: ld.param.b64 %rd4, [test_st_async_mbarrier_v2b64_param_2];
-; CHECK-PTX64-NEXT: st.async.shared::cluster.mbarrier::complete_tx::bytes.v2.b64 [%rd1], {%rd2, %rd3}, [%rd4];
-; CHECK-PTX64-NEXT: ret;
-;
-; CHECK-SHARED32-LABEL: test_st_async_mbarrier_v2b64(
-; CHECK-SHARED32: {
-; CHECK-SHARED32-NEXT: .reg .b32 %r<3>;
-; CHECK-SHARED32-NEXT: .reg .b64 %rd<3>;
-; CHECK-SHARED32-EMPTY:
-; CHECK-SHARED32-NEXT: // %bb.0:
-; CHECK-SHARED32-NEXT: ld.param.b32 %r1, [test_st_async_mbarrier_v2b64_param_0];
-; CHECK-SHARED32-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_st_async_mbarrier_v2b64_param_1];
-; CHECK-SHARED32-NEXT: ld.param.b32 %r2, [test_st_async_mbarrier_v2b64_param_2];
-; CHECK-SHARED32-NEXT: st.async.shared::cluster.mbarrier::complete_tx::bytes.v2.b64 [%r1], {%rd1, %rd2}, [%r2];
-; CHECK-SHARED32-NEXT: ret;
- call void @llvm.nvvm.st.async.space.cluster.v2i64(ptr addrspace(7) %addr, <2 x i64> %value, ptr addrspace(7) %mbar)
- ret void
-}
-
-define void @test_st_async_mbarrier_v4b32(ptr addrspace(7) %addr, <4 x i32> %value, ptr addrspace(7) %mbar) {
-; CHECK-PTX64-LABEL: test_st_async_mbarrier_v4b32(
-; CHECK-PTX64: {
-; CHECK-PTX64-NEXT: .reg .b32 %r<5>;
-; CHECK-PTX64-NEXT: .reg .b64 %rd<3>;
-; CHECK-PTX64-EMPTY:
-; CHECK-PTX64-NEXT: // %bb.0:
-; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [test_st_async_mbarrier_v4b32_param_0];
-; CHECK-PTX64-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_st_async_mbarrier_v4b32_param_1];
-; CHECK-PTX64-NEXT: ld.param.b64 %rd2, [test_st_async_mbarrier_v4b32_param_2];
-; CHECK-PTX64-NEXT: st.async.shared::cluster.mbarrier::complete_tx::bytes.v4.b32 [%rd1], {%r1, %r2, %r3, %r4}, [%rd2];
-; CHECK-PTX64-NEXT: ret;
-;
-; CHECK-SHARED32-LABEL: test_st_async_mbarrier_v4b32(
-; CHECK-SHARED32: {
-; CHECK-SHARED32-NEXT: .reg .b32 %r<7>;
-; CHECK-SHARED32-EMPTY:
-; CHECK-SHARED32-NEXT: // %bb.0:
-; CHECK-SHARED32-NEXT: ld.param.b32 %r1, [test_st_async_mbarrier_v4b32_param_0];
-; CHECK-SHARED32-NEXT: ld.param.v4.b32 {%r2, %r3, %r4, %r5}, [test_st_async_mbarrier_v4b32_param_1];
-; CHECK-SHARED32-NEXT: ld.param.b32 %r6, [test_st_async_mbarrier_v4b32_param_2];
-; CHECK-SHARED32-NEXT: st.async.shared::cluster.mbarrier::complete_tx::bytes.v4.b32 [%r1], {%r2, %r3, %r4, %r5}, [%r6];
-; CHECK-SHARED32-NEXT: ret;
- call void @llvm.nvvm.st.async.space.cluster.v4i32(ptr addrspace(7) %addr, <4 x i32> %value, ptr addrspace(7) %mbar)
- ret void
-}
diff --git a/llvm/test/CodeGen/NVPTX/st_async_mbarrier_b128.ll b/llvm/test/CodeGen/NVPTX/st_async_mbarrier_b128.ll
new file mode 100644
index 0000000000000..0753597ad5302
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/st_async_mbarrier_b128.ll
@@ -0,0 +1,40 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx92 | FileCheck --check-prefixes=CHECK-PTX64 %s
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx92 --nvptx-short-ptr | FileCheck --check-prefixes=CHECK-SHARED32 %s
+; RUN: %if ptxas-sm_90 && ptxas-isa-9.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx92 | %ptxas-verify -arch=sm_90 %}
+; RUN: %if ptxas-sm_90 && ptxas-isa-9.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx92 --nvptx-short-ptr | %ptxas-verify -arch=sm_90 %}
+
+define void @test_st_async_mbarrier_b128(ptr addrspace(7) %addr, i128 %value, ptr addrspace(7) %mbar) {
+; CHECK-PTX64-LABEL: test_st_async_mbarrier_b128(
+; CHECK-PTX64: {
+; CHECK-PTX64-NEXT: .reg .b64 %rd<5>;
+; CHECK-PTX64-EMPTY:
+; CHECK-PTX64-NEXT: // %bb.0:
+; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [test_st_async_mbarrier_b128_param_0];
+; CHECK-PTX64-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_st_async_mbarrier_b128_param_1];
+; CHECK-PTX64-NEXT: ld.param.b64 %rd4, [test_st_async_mbarrier_b128_param_2];
+; CHECK-PTX64-NEXT: {
+; CHECK-PTX64-NEXT: .reg .b128 %in_128;
+; CHECK-PTX64-NEXT: mov.b128 %in_128, {%rd2, %rd3};
+; CHECK-PTX64-NEXT: st.async.shared::cluster.mbarrier::complete_tx::bytes.b128 [%rd1], %in_128, [%rd4];
+; CHECK-PTX64-NEXT: }
+; CHECK-PTX64-NEXT: ret;
+;
+; CHECK-SHARED32-LABEL: test_st_async_mbarrier_b128(
+; CHECK-SHARED32: {
+; CHECK-SHARED32-NEXT: .reg .b32 %r<3>;
+; CHECK-SHARED32-NEXT: .reg .b64 %rd<3>;
+; CHECK-SHARED32-EMPTY:
+; CHECK-SHARED32-NEXT: // %bb.0:
+; CHECK-SHARED32-NEXT: ld.param.b32 %r1, [test_st_async_mbarrier_b128_param_0];
+; CHECK-SHARED32-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_st_async_mbarrier_b128_param_1];
+; CHECK-SHARED32-NEXT: ld.param.b32 %r2, [test_st_async_mbarrier_b128_param_2];
+; CHECK-SHARED32-NEXT: {
+; CHECK-SHARED32-NEXT: .reg .b128 %in_128;
+; CHECK-SHARED32-NEXT: mov.b128 %in_128, {%rd1, %rd2};
+; CHECK-SHARED32-NEXT: st.async.shared::cluster.mbarrier::complete_tx::bytes.b128 [%r1], %in_128, [%r2];
+; CHECK-SHARED32-NEXT: }
+; CHECK-SHARED32-NEXT: ret;
+ call void @llvm.nvvm.st.async.space.cluster.i128(ptr addrspace(7) %addr, i128 %value, ptr addrspace(7) %mbar)
+ ret void
+}
>From 2c72048e8e23a5b475f20d45bedf5a819e2e44e7 Mon Sep 17 00:00:00 2001
From: Srinivasa Ravi <srinivasar at nvidia.com>
Date: Thu, 4 Jun 2026 15:57:51 +0000
Subject: [PATCH 6/7] address comments
---
llvm/docs/NVPTXUsage.rst | 44 ++++----
llvm/include/llvm/IR/IntrinsicsNVVM.td | 8 +-
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 102 ++++++-----------
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 106 ++++++++++--------
llvm/test/CodeGen/NVPTX/st_async_mbarrier.ll | 4 +-
.../CodeGen/NVPTX/st_async_mbarrier_b128.ll | 2 +-
llvm/test/CodeGen/NVPTX/st_async_release.ll | 24 ++--
.../NVPTX/st_async_release_multimem.ll | 16 +--
8 files changed, 142 insertions(+), 164 deletions(-)
diff --git a/llvm/docs/NVPTXUsage.rst b/llvm/docs/NVPTXUsage.rst
index 41c77f33b9611..c1b2d7ccfbda8 100644
--- a/llvm/docs/NVPTXUsage.rst
+++ b/llvm/docs/NVPTXUsage.rst
@@ -3831,7 +3831,7 @@ similar but the latter uses generic addressing (see `Generic Addressing
For more information, refer `PTX ISA
<https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-st-bulk>`__.
-'``llvm.nvvm.st.async.space.cluster``'
+'``llvm.nvvm.st.async``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
Syntax:
@@ -3839,14 +3839,14 @@ Syntax:
.. code-block:: llvm
- declare void @llvm.nvvm.st.async.space.cluster.i32(ptr addrspace(7) %dest_addr, i32 %value, ptr addrspace(7) %mbarrier_addr)
- declare void @llvm.nvvm.st.async.space.cluster.i64(ptr addrspace(7) %dest_addr, i64 %value, ptr addrspace(7) %mbarrier_addr)
- declare void @llvm.nvvm.st.async.space.cluster.i128(ptr addrspace(7) %dest_addr, i128 %value, ptr addrspace(7) %mbarrier_addr)
+ declare void @llvm.nvvm.st.async.i32(ptr addrspace(7) %dest_addr, i32 %value, ptr addrspace(7) %mbarrier_addr)
+ declare void @llvm.nvvm.st.async.i64(ptr addrspace(7) %dest_addr, i64 %value, ptr addrspace(7) %mbarrier_addr)
+ declare void @llvm.nvvm.st.async.i128(ptr addrspace(7) %dest_addr, i128 %value, ptr addrspace(7) %mbarrier_addr)
Overview:
"""""""""
-The '``llvm.nvvm.st.async.space.cluster``' intrinsic initiates a weak
+The '``llvm.nvvm.st.async``' intrinsic initiates a weak
asynchronous store operation to shared memory that stores the value specified
by the `%value` operand to the destination address specified by the
`%dest_addr` operand. The `%value` operand must be an ``i32``, ``i64`` or
@@ -3864,7 +3864,7 @@ a `complete-tx <https://docs.nvidia.com/cuda/parallel-thread-execution/#parallel
For more information, refer `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-st-async>`__.
-'``llvm.nvvm.st.async.scope.{sys,gpu}.*``'
+'``llvm.nvvm.st.async.{sys,gpu}``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
Syntax:
@@ -3873,22 +3873,22 @@ Syntax:
.. code-block:: llvm
; sys scope
- declare void @llvm.nvvm.st.async.scope.sys.space.global.i8(ptr addrspace(1) %dest_addr, i8 %value, i1 immarg %is_multimem)
- declare void @llvm.nvvm.st.async.scope.sys.space.global.i16(ptr addrspace(1) %dest_addr, i16 %value, i1 immarg %is_multimem)
- declare void @llvm.nvvm.st.async.scope.sys.space.global.i32(ptr addrspace(1) %dest_addr, i32 %value, i1 immarg %is_multimem)
- declare void @llvm.nvvm.st.async.scope.sys.space.global.i64(ptr addrspace(1) %dest_addr, i64 %value, i1 immarg %is_multimem)
+ declare void @llvm.nvvm.st.async.sys.i8(ptr addrspace(1) %dest_addr, i8 %value, i1 immarg %is_multimem)
+ declare void @llvm.nvvm.st.async.sys.i16(ptr addrspace(1) %dest_addr, i16 %value, i1 immarg %is_multimem)
+ declare void @llvm.nvvm.st.async.sys.i32(ptr addrspace(1) %dest_addr, i32 %value, i1 immarg %is_multimem)
+ declare void @llvm.nvvm.st.async.sys.i64(ptr addrspace(1) %dest_addr, i64 %value, i1 immarg %is_multimem)
; gpu scope
- declare void @llvm.nvvm.st.async.scope.gpu.space.global.i8(ptr addrspace(1) %dest_addr, i8 %value, i1 immarg %is_multimem)
- declare void @llvm.nvvm.st.async.scope.gpu.space.global.i16(ptr addrspace(1) %dest_addr, i16 %value, i1 immarg %is_multimem)
- declare void @llvm.nvvm.st.async.scope.gpu.space.global.i32(ptr addrspace(1) %dest_addr, i32 %value, i1 immarg %is_multimem)
- declare void @llvm.nvvm.st.async.scope.gpu.space.global.i64(ptr addrspace(1) %dest_addr, i64 %value, i1 immarg %is_multimem)
+ declare void @llvm.nvvm.st.async.gpu.i8(ptr addrspace(1) %dest_addr, i8 %value, i1 immarg %is_multimem)
+ declare void @llvm.nvvm.st.async.gpu.i16(ptr addrspace(1) %dest_addr, i16 %value, i1 immarg %is_multimem)
+ declare void @llvm.nvvm.st.async.gpu.i32(ptr addrspace(1) %dest_addr, i32 %value, i1 immarg %is_multimem)
+ declare void @llvm.nvvm.st.async.gpu.i64(ptr addrspace(1) %dest_addr, i64 %value, i1 immarg %is_multimem)
Overview:
"""""""""
-The '``llvm.nvvm.st.async.scope.sys.space.global``' and
-'``llvm.nvvm.st.async.scope.gpu.space.global``' intrinsics initiate an
+The '``llvm.nvvm.st.async.sys``' and
+'``llvm.nvvm.st.async.gpu``' intrinsics initiate an
asynchronous release store to global memory that stores the value specified by
the `%value` operand to the destination address specified by the `%dest_addr`
operand.
@@ -3909,7 +3909,7 @@ These intrinsics lower to the `.b8`, `.b16`, `.b32`, and `.b64` variants of the
For more information, refer `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-st-async>`__.
-'``llvm.nvvm.st.async.mmio.scope.sys.*``'
+'``llvm.nvvm.st.async.mmio.sys``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
Syntax:
@@ -3917,15 +3917,15 @@ Syntax:
.. code-block:: llvm
- declare void @llvm.nvvm.st.async.mmio.scope.sys.space.global.i8(ptr addrspace(1) %dest_addr, i8 %value)
- declare void @llvm.nvvm.st.async.mmio.scope.sys.space.global.i16(ptr addrspace(1) %dest_addr, i16 %value)
- declare void @llvm.nvvm.st.async.mmio.scope.sys.space.global.i32(ptr addrspace(1) %dest_addr, i32 %value)
- declare void @llvm.nvvm.st.async.mmio.scope.sys.space.global.i64(ptr addrspace(1) %dest_addr, i64 %value)
+ declare void @llvm.nvvm.st.async.mmio.sys.i8(ptr addrspace(1) %dest_addr, i8 %value)
+ declare void @llvm.nvvm.st.async.mmio.sys.i16(ptr addrspace(1) %dest_addr, i16 %value)
+ declare void @llvm.nvvm.st.async.mmio.sys.i32(ptr addrspace(1) %dest_addr, i32 %value)
+ declare void @llvm.nvvm.st.async.mmio.sys.i64(ptr addrspace(1) %dest_addr, i64 %value)
Overview:
"""""""""
-The '``llvm.nvvm.st.async.mmio.scope.sys.space.global``' intrinsic performs an `MMIO <https://docs.nvidia.com/cuda/parallel-thread-execution/#mmio-operation>`__
+The '``llvm.nvvm.st.async.mmio.sys``' intrinsic performs an `MMIO <https://docs.nvidia.com/cuda/parallel-thread-execution/#mmio-operation>`__
store to global memory with ``.release`` semantics at the ``sys`` scope.
For more information, refer `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-st-async>`__.
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 1b840bdf41cec..581b84eda9e74 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -3276,20 +3276,20 @@ let IntrProperties = [IntrArgMemOnly, IntrWriteMem, WriteOnly<ArgIndex<0>>,
// Asynchronous store intrinsics
//
-def int_nvvm_st_async_space_cluster :
+def int_nvvm_st_async :
DefaultAttrsIntrinsic<[],
[llvm_shared_cluster_ptr_ty, llvm_anyint_ty, llvm_shared_cluster_ptr_ty],
- [WriteOnly<ArgIndex<0>>]>;
+ [WriteOnly<ArgIndex<0>>, IntrArgMemOnly]>;
foreach scope = ["_sys", "_gpu"] in {
- def int_nvvm_st_async_scope # scope # _space_global :
+ def int_nvvm_st_async # scope :
DefaultAttrsIntrinsic<[],
[llvm_global_ptr_ty, llvm_anyint_ty, llvm_i1_ty],
[WriteOnly<ArgIndex<0>>, ImmArg<ArgIndex<2>>,
ArgInfo<ArgIndex<2>, [ArgName<"isMultimem">]>]>;
}
-def int_nvvm_st_async_mmio_scope_sys_space_global :
+def int_nvvm_st_async_mmio_sys :
DefaultAttrsIntrinsic<[],
[llvm_global_ptr_ty, llvm_anyint_ty], [WriteOnly<ArgIndex<0>>]>;
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 1a5a79a224995..33b2b132aca1e 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -2685,80 +2685,46 @@ static SDValue lowerStAsyncRelease(SDValue Op, SelectionDAG &DAG) {
MVT ValueVT = Value.getSimpleValueType();
- auto OpCode = [&]() -> std::optional<unsigned> {
+ if (ValueVT == MVT::i16 || ValueVT == MVT::i32 || ValueVT == MVT::i64)
+ return Op;
+
+ if (ValueVT == MVT::i8) {
+ unsigned OpCode;
switch (IntrinsicID) {
- case Intrinsic::nvvm_st_async_scope_sys_space_global:
- switch (ValueVT.SimpleTy) {
- case MVT::i8:
- return NVPTXISD::ST_ASYNC_SYS_B8;
- case MVT::i16:
- return NVPTXISD::ST_ASYNC_SYS_B16;
- case MVT::i32:
- return NVPTXISD::ST_ASYNC_SYS_B32;
- case MVT::i64:
- return NVPTXISD::ST_ASYNC_SYS_B64;
- default:
- break;
- }
+ case Intrinsic::nvvm_st_async_sys:
+ OpCode = NVPTXISD::ST_ASYNC_SYS_B8;
break;
- case Intrinsic::nvvm_st_async_scope_gpu_space_global:
- switch (ValueVT.SimpleTy) {
- case MVT::i8:
- return NVPTXISD::ST_ASYNC_GPU_B8;
- case MVT::i16:
- return NVPTXISD::ST_ASYNC_GPU_B16;
- case MVT::i32:
- return NVPTXISD::ST_ASYNC_GPU_B32;
- case MVT::i64:
- return NVPTXISD::ST_ASYNC_GPU_B64;
- default:
- break;
- }
+ case Intrinsic::nvvm_st_async_gpu:
+ OpCode = NVPTXISD::ST_ASYNC_GPU_B8;
break;
- case Intrinsic::nvvm_st_async_mmio_scope_sys_space_global:
- switch (ValueVT.SimpleTy) {
- case MVT::i8:
- return NVPTXISD::ST_ASYNC_MMIO_SYS_B8;
- case MVT::i16:
- return NVPTXISD::ST_ASYNC_MMIO_SYS_B16;
- case MVT::i32:
- return NVPTXISD::ST_ASYNC_MMIO_SYS_B32;
- case MVT::i64:
- return NVPTXISD::ST_ASYNC_MMIO_SYS_B64;
- default:
- break;
- }
+ case Intrinsic::nvvm_st_async_mmio_sys:
+ OpCode = NVPTXISD::ST_ASYNC_MMIO_SYS_B8;
break;
+ default:
+ llvm_unreachable("unexpected intrinsic ID for st.async.release");
}
- return std::nullopt;
- }();
- if (!OpCode) {
- DAG.getContext()->diagnose(DiagnosticInfoUnsupported(
- Fn,
- Twine("unsupported argument type ") +
- llvm::EVT(ValueVT).getEVTString() + " for " +
- llvm::Intrinsic::getName(IntrinsicID) + " intrinsic",
- DiagnosticLocation(DL.getDebugLoc())));
- return Op.getOperand(0); // Return only the chain
- }
-
- // NVPTX has no i8 register class; widen i8 to i16. The selected ST_ASYNC_*_B8
- // node still emits the `.b8` qualifier in PTX.
- if (ValueVT == MVT::i8)
Value = DAG.getNode(ISD::ZERO_EXTEND, DL, MVT::i16, Value);
- // The `.mmio` variant has no multimem form and therefore no `isMultimem`
- // operand.
- if (IntrinsicID == Intrinsic::nvvm_st_async_mmio_scope_sys_space_global) {
- SDValue Ops[] = {N->getOperand(0), DestAddr, Value};
- return DAG.getNode(OpCode.value(), DL, MVT::Other, Ops);
+ // The `.mmio` variant has no multimem form and therefore no `isMultimem`
+ // operand.
+ if (IntrinsicID == Intrinsic::nvvm_st_async_mmio_sys) {
+ SDValue Ops[] = {N->getOperand(0), DestAddr, Value};
+ return DAG.getNode(OpCode, DL, MVT::Other, Ops);
+ }
+
+ SDValue IsMultimem =
+ DAG.getTargetConstant(N->getConstantOperandVal(4), DL, MVT::i1);
+ SDValue Ops[] = {N->getOperand(0), DestAddr, Value, IsMultimem};
+ return DAG.getNode(OpCode, DL, MVT::Other, Ops);
}
- SDValue IsMultimem =
- DAG.getTargetConstant(N->getConstantOperandVal(4), DL, MVT::i32);
- SDValue Ops[] = {N->getOperand(0), DestAddr, Value, IsMultimem};
- return DAG.getNode(OpCode.value(), DL, MVT::Other, Ops);
+ DAG.getContext()->diagnose(DiagnosticInfoUnsupported(
+ Fn,
+ Twine("unsupported argument type ") + llvm::EVT(ValueVT).getEVTString() +
+ " for " + llvm::Intrinsic::getName(IntrinsicID) + " intrinsic",
+ DiagnosticLocation(DL.getDebugLoc())));
+ return Op.getOperand(0); // Return only the chain
}
static unsigned getTcgen05MMADisableOutputLane(unsigned IID) {
@@ -2955,11 +2921,11 @@ static SDValue lowerIntrinsicVoid(SDValue Op, SelectionDAG &DAG) {
switch (IntrinNo) {
default:
break;
- case Intrinsic::nvvm_st_async_space_cluster:
+ case Intrinsic::nvvm_st_async:
return lowerStAsyncWithMbarrier(Op, DAG);
- case Intrinsic::nvvm_st_async_scope_sys_space_global:
- case Intrinsic::nvvm_st_async_scope_gpu_space_global:
- case Intrinsic::nvvm_st_async_mmio_scope_sys_space_global:
+ case Intrinsic::nvvm_st_async_sys:
+ case Intrinsic::nvvm_st_async_gpu:
+ case Intrinsic::nvvm_st_async_mmio_sys:
return lowerStAsyncRelease(Op, DAG);
case Intrinsic::nvvm_tcgen05_st_16x64b_x2:
case Intrinsic::nvvm_tcgen05_st_16x64b_x4:
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 355f3cbbe55ce..25610cd8f4c4e 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -6094,9 +6094,9 @@ def st_async_mbarrier_b128 :
[SDNPHasChain, SDNPMayStore]>;
class ST_ASYNC_MBARRIER<string type, NVPTXRegClass reg> :
- NVPTXInst<(outs),
+ BasicNVPTXInst<(outs),
(ins ADDR:$dest_addr, reg:$value, ADDR:$mbar),
- "st.async.shared::cluster.mbarrier::complete_tx::bytes." # type # "\t [$dest_addr], $value, [$mbar];">;
+ "st.async.shared::cluster.mbarrier::complete_tx::bytes." # type>;
def ST_ASYNC_MBARRIER_B128 :
NVPTXInst<(outs),
@@ -6112,10 +6112,10 @@ let Predicates = [hasSM<90>, hasPTX<81>] in {
defvar valueType = !cast<ValueType>("i" # size);
defvar reg = !cast<NVPTXRegClass>("B" # size);
- def : Pat<(int_nvvm_st_async_space_cluster
+ def : Pat<(int_nvvm_st_async
addr:$dest_addr, valueType:$value, addr:$mbar),
(ST_ASYNC_MBARRIER<"b" # size, reg>
- addr:$dest_addr, valueType:$value, addr:$mbar)>;
+ addr:$dest_addr, $value, addr:$mbar)>;
}
}
@@ -6123,11 +6123,11 @@ let Predicates = [hasSM<90>, hasPTX<92>] in {
def : Pat<(st_async_mbarrier_b128
addr:$dest_addr, i64:$value0, i64:$value1, addr:$mbar),
(ST_ASYNC_MBARRIER_B128
- addr:$dest_addr, B64:$value0, B64:$value1, addr:$mbar)>;
+ addr:$dest_addr, $value0, $value1, addr:$mbar)>;
}
def SDTAsyncStoreRelease : SDTypeProfile<0, 3,
- [SDTCisPtrTy<0>, SDTCisVT<2, i32>]>;
+ [SDTCisPtrTy<0>, SDTCisVT<2, i1>]>;
def SDTAsyncStoreMMIO : SDTypeProfile<0, 2, [SDTCisPtrTy<0>]>;
class StAsyncReleaseNode<string variant, string type, SDTypeProfile sdt> :
@@ -6135,60 +6135,72 @@ class StAsyncReleaseNode<string variant, string type, SDTypeProfile sdt> :
sdt,
[SDNPHasChain, SDNPMayStore]>;
-def MultimemFlag : Operand<i32> {
+def MultimemFlag : Operand<i1> {
let PrintMethod = "printMultimem";
}
-class ST_ASYNC_RELEASE_BODY<string instr_prefix, string type, bit is_mmio> {
- string prefix = !if(is_mmio, "", "${is_multimem}");
- string body = !if(!eq(type, "b8"),
- // The minimum supported register size to pass in arguments is 16-bits, so
- // the 8-bit value is passed in using a 16-bit register and truncated to
- // 8-bits with a cvt instruction.
+class ST_ASYNC_SYS<string type, NVPTXRegClass reg> :
+ BasicFlagsNVPTXInst<(outs), (ins ADDR:$dest_addr, reg:$value),
+ (ins MultimemFlag:$is_multimem),
+ "${is_multimem}st.async.release.sys.global." # type>;
+
+class ST_ASYNC_GPU<string type, NVPTXRegClass reg> :
+ BasicFlagsNVPTXInst<(outs), (ins ADDR:$dest_addr, reg:$value),
+ (ins MultimemFlag:$is_multimem),
+ "${is_multimem}st.async.release.gpu.global." # type>;
+
+class ST_ASYNC_MMIO_SYS<string type, NVPTXRegClass reg> :
+ BasicNVPTXInst<(outs), (ins ADDR:$dest_addr, reg:$value),
+ "st.async.mmio.release.sys.global." # type>;
+
+class ST_ASYNC_RELEASE_B8_BODY<string instr_prefix, string prefix> {
+ // The minimum supported register size to pass in arguments is 16-bits, so
+ // the 8-bit value is passed in using a 16-bit register and truncated to
+ // 8-bits with a cvt instruction.
+ string body =
"{{ \n\t"
".reg .b8 \t%st_async_val; \n\t"
"cvt.u8.u16 \t%st_async_val, $value; \n\t" #
prefix # instr_prefix # ".b8\t [$dest_addr], %st_async_val; \n\t"
- "}}",
- prefix # instr_prefix # "." # type # "\t [$dest_addr], $value;");
+ "}}";
}
-class ST_ASYNC_SYS<string type, NVPTXRegClass reg> :
- NVPTXInst<(outs),
- (ins ADDR:$dest_addr, reg:$value, MultimemFlag:$is_multimem),
- ST_ASYNC_RELEASE_BODY<"st.async.release.sys.global", type, 0>.body>;
-
-class ST_ASYNC_GPU<string type, NVPTXRegClass reg> :
- NVPTXInst<(outs),
- (ins ADDR:$dest_addr, reg:$value, MultimemFlag:$is_multimem),
- ST_ASYNC_RELEASE_BODY<"st.async.release.gpu.global", type, 0>.body>;
-
-class ST_ASYNC_MMIO_SYS<string type, NVPTXRegClass reg> :
- NVPTXInst<(outs),
- (ins ADDR:$dest_addr, reg:$value),
- ST_ASYNC_RELEASE_BODY<"st.async.mmio.release.sys.global", type, 1>.body>;
+def ST_ASYNC_SYS_B8 :
+ NVPTXInst<(outs), (ins ADDR:$dest_addr, B16:$value, MultimemFlag:$is_multimem),
+ ST_ASYNC_RELEASE_B8_BODY<"st.async.release.sys.global", "${is_multimem}">.body>;
+def ST_ASYNC_GPU_B8 :
+ NVPTXInst<(outs), (ins ADDR:$dest_addr, B16:$value, MultimemFlag:$is_multimem),
+ ST_ASYNC_RELEASE_B8_BODY<"st.async.release.gpu.global", "${is_multimem}">.body>;
+def ST_ASYNC_MMIO_SYS_B8 :
+ NVPTXInst<(outs), (ins ADDR:$dest_addr, B16:$value),
+ ST_ASYNC_RELEASE_B8_BODY<"st.async.mmio.release.sys.global", "">.body>;
let Predicates = [hasSM<100>, hasPTX<87>] in {
- foreach i = !range(4) in {
- defvar size = ["8", "16", "32", "64"][i];
- defvar reg = [B16, B16, B32, B64 ][i];
- defvar valueType = [i16, i16, i32, i64 ][i];
- defvar width = "b" # size;
-
- def : Pat<(StAsyncReleaseNode<"SYS", !toupper(width), SDTAsyncStoreRelease>
- addr:$dest_addr, valueType:$value, timm:$is_multimem),
- (ST_ASYNC_SYS<width, reg>
- addr:$dest_addr, valueType:$value, timm:$is_multimem)>;
- def : Pat<(StAsyncReleaseNode<"GPU", !toupper(width), SDTAsyncStoreRelease>
- addr:$dest_addr, valueType:$value, timm:$is_multimem),
- (ST_ASYNC_GPU<width, reg>
- addr:$dest_addr, valueType:$value, timm:$is_multimem)>;
- def : Pat<(StAsyncReleaseNode<"MMIO_SYS", !toupper(width), SDTAsyncStoreMMIO>
- addr:$dest_addr, valueType:$value),
- (ST_ASYNC_MMIO_SYS<width, reg>
- addr:$dest_addr, valueType:$value)>;
+ foreach rti = [I16RT, I32RT, I64RT] in {
+ def : Pat<(int_nvvm_st_async_sys
+ addr:$dest_addr, rti.Ty:$value, timm:$is_multimem),
+ (ST_ASYNC_SYS<rti.PtxType, rti.RC>
+ addr:$dest_addr, $value, $is_multimem)>;
+ def : Pat<(int_nvvm_st_async_gpu
+ addr:$dest_addr, rti.Ty:$value, timm:$is_multimem),
+ (ST_ASYNC_GPU<rti.PtxType, rti.RC>
+ addr:$dest_addr, $value, $is_multimem)>;
+ def : Pat<(int_nvvm_st_async_mmio_sys
+ addr:$dest_addr, rti.Ty:$value),
+ (ST_ASYNC_MMIO_SYS<rti.PtxType, rti.RC>
+ addr:$dest_addr, $value)>;
}
+
+ def : Pat<(StAsyncReleaseNode<"SYS", "B8", SDTAsyncStoreRelease>
+ addr:$dest_addr, i16:$value, timm:$is_multimem),
+ (ST_ASYNC_SYS_B8 addr:$dest_addr, $value, $is_multimem)>;
+ def : Pat<(StAsyncReleaseNode<"GPU", "B8", SDTAsyncStoreRelease>
+ addr:$dest_addr, i16:$value, timm:$is_multimem),
+ (ST_ASYNC_GPU_B8 addr:$dest_addr, $value, $is_multimem)>;
+ def : Pat<(StAsyncReleaseNode<"MMIO_SYS", "B8", SDTAsyncStoreMMIO>
+ addr:$dest_addr, i16:$value),
+ (ST_ASYNC_MMIO_SYS_B8 addr:$dest_addr, $value)>;
}
//
diff --git a/llvm/test/CodeGen/NVPTX/st_async_mbarrier.ll b/llvm/test/CodeGen/NVPTX/st_async_mbarrier.ll
index 76c41d505d5d8..5d3ed95565a3d 100644
--- a/llvm/test/CodeGen/NVPTX/st_async_mbarrier.ll
+++ b/llvm/test/CodeGen/NVPTX/st_async_mbarrier.ll
@@ -27,7 +27,7 @@ define void @test_st_async_mbarrier_b32(ptr addrspace(7) %addr, i32 %value, ptr
; CHECK-SHARED32-NEXT: ld.param.b32 %r3, [test_st_async_mbarrier_b32_param_2];
; CHECK-SHARED32-NEXT: st.async.shared::cluster.mbarrier::complete_tx::bytes.b32 [%r1], %r2, [%r3];
; CHECK-SHARED32-NEXT: ret;
- call void @llvm.nvvm.st.async.space.cluster.i32(ptr addrspace(7) %addr, i32 %value, ptr addrspace(7) %mbar)
+ call void @llvm.nvvm.st.async.i32(ptr addrspace(7) %addr, i32 %value, ptr addrspace(7) %mbar)
ret void
}
@@ -54,6 +54,6 @@ define void @test_st_async_mbarrier_b64(ptr addrspace(7) %addr, i64 %value, ptr
; CHECK-SHARED32-NEXT: ld.param.b32 %r2, [test_st_async_mbarrier_b64_param_2];
; CHECK-SHARED32-NEXT: st.async.shared::cluster.mbarrier::complete_tx::bytes.b64 [%r1], %rd1, [%r2];
; CHECK-SHARED32-NEXT: ret;
- call void @llvm.nvvm.st.async.space.cluster.i64(ptr addrspace(7) %addr, i64 %value, ptr addrspace(7) %mbar)
+ call void @llvm.nvvm.st.async.i64(ptr addrspace(7) %addr, i64 %value, ptr addrspace(7) %mbar)
ret void
}
diff --git a/llvm/test/CodeGen/NVPTX/st_async_mbarrier_b128.ll b/llvm/test/CodeGen/NVPTX/st_async_mbarrier_b128.ll
index 0753597ad5302..6927770c53cea 100644
--- a/llvm/test/CodeGen/NVPTX/st_async_mbarrier_b128.ll
+++ b/llvm/test/CodeGen/NVPTX/st_async_mbarrier_b128.ll
@@ -35,6 +35,6 @@ define void @test_st_async_mbarrier_b128(ptr addrspace(7) %addr, i128 %value, pt
; CHECK-SHARED32-NEXT: st.async.shared::cluster.mbarrier::complete_tx::bytes.b128 [%r1], %in_128, [%r2];
; CHECK-SHARED32-NEXT: }
; CHECK-SHARED32-NEXT: ret;
- call void @llvm.nvvm.st.async.space.cluster.i128(ptr addrspace(7) %addr, i128 %value, ptr addrspace(7) %mbar)
+ call void @llvm.nvvm.st.async.i128(ptr addrspace(7) %addr, i128 %value, ptr addrspace(7) %mbar)
ret void
}
diff --git a/llvm/test/CodeGen/NVPTX/st_async_release.ll b/llvm/test/CodeGen/NVPTX/st_async_release.ll
index 9bf4d3c8f789f..3df7642d5ed1a 100644
--- a/llvm/test/CodeGen/NVPTX/st_async_release.ll
+++ b/llvm/test/CodeGen/NVPTX/st_async_release.ll
@@ -17,7 +17,7 @@ define void @test_st_async_release_sys_b8(ptr addrspace(1) %addr, i8 %value) {
; CHECK-NEXT: st.async.release.sys.global.b8 [%rd1], %st_async_val;
; CHECK-NEXT: }
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.scope.sys.space.global.i8(ptr addrspace(1) %addr, i8 %value, /* isMultimem= */ i1 false)
+ call void @llvm.nvvm.st.async.sys.i8(ptr addrspace(1) %addr, i8 %value, /* isMultimem= */ i1 false)
ret void
}
@@ -32,7 +32,7 @@ define void @test_st_async_release_sys_b16(ptr addrspace(1) %addr, i16 %value) {
; CHECK-NEXT: ld.param.b16 %rs1, [test_st_async_release_sys_b16_param_1];
; CHECK-NEXT: st.async.release.sys.global.b16 [%rd1], %rs1;
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.scope.sys.space.global.i16(ptr addrspace(1) %addr, i16 %value, /* isMultimem= */ i1 false)
+ call void @llvm.nvvm.st.async.sys.i16(ptr addrspace(1) %addr, i16 %value, /* isMultimem= */ i1 false)
ret void
}
@@ -47,7 +47,7 @@ define void @test_st_async_release_sys_b32(ptr addrspace(1) %addr, i32 %value) {
; CHECK-NEXT: ld.param.b32 %r1, [test_st_async_release_sys_b32_param_1];
; CHECK-NEXT: st.async.release.sys.global.b32 [%rd1], %r1;
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.scope.sys.space.global.i32(ptr addrspace(1) %addr, i32 %value, /* isMultimem= */ i1 false)
+ call void @llvm.nvvm.st.async.sys.i32(ptr addrspace(1) %addr, i32 %value, /* isMultimem= */ i1 false)
ret void
}
@@ -61,7 +61,7 @@ define void @test_st_async_release_sys_b64(ptr addrspace(1) %addr, i64 %value) {
; CHECK-NEXT: ld.param.b64 %rd2, [test_st_async_release_sys_b64_param_1];
; CHECK-NEXT: st.async.release.sys.global.b64 [%rd1], %rd2;
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.scope.sys.space.global.i64(ptr addrspace(1) %addr, i64 %value, /* isMultimem= */ i1 false)
+ call void @llvm.nvvm.st.async.sys.i64(ptr addrspace(1) %addr, i64 %value, /* isMultimem= */ i1 false)
ret void
}
@@ -80,7 +80,7 @@ define void @test_st_async_release_gpu_b8(ptr addrspace(1) %addr, i8 %value) {
; CHECK-NEXT: st.async.release.gpu.global.b8 [%rd1], %st_async_val;
; CHECK-NEXT: }
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.scope.gpu.space.global.i8(ptr addrspace(1) %addr, i8 %value, /* isMultimem= */ i1 false)
+ call void @llvm.nvvm.st.async.gpu.i8(ptr addrspace(1) %addr, i8 %value, /* isMultimem= */ i1 false)
ret void
}
@@ -95,7 +95,7 @@ define void @test_st_async_release_gpu_b16(ptr addrspace(1) %addr, i16 %value) {
; CHECK-NEXT: ld.param.b16 %rs1, [test_st_async_release_gpu_b16_param_1];
; CHECK-NEXT: st.async.release.gpu.global.b16 [%rd1], %rs1;
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.scope.gpu.space.global.i16(ptr addrspace(1) %addr, i16 %value, /* isMultimem= */ i1 false)
+ call void @llvm.nvvm.st.async.gpu.i16(ptr addrspace(1) %addr, i16 %value, /* isMultimem= */ i1 false)
ret void
}
@@ -110,7 +110,7 @@ define void @test_st_async_release_gpu_b32(ptr addrspace(1) %addr, i32 %value) {
; CHECK-NEXT: ld.param.b32 %r1, [test_st_async_release_gpu_b32_param_1];
; CHECK-NEXT: st.async.release.gpu.global.b32 [%rd1], %r1;
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.scope.gpu.space.global.i32(ptr addrspace(1) %addr, i32 %value, /* isMultimem= */ i1 false)
+ call void @llvm.nvvm.st.async.gpu.i32(ptr addrspace(1) %addr, i32 %value, /* isMultimem= */ i1 false)
ret void
}
@@ -124,7 +124,7 @@ define void @test_st_async_release_gpu_b64(ptr addrspace(1) %addr, i64 %value) {
; CHECK-NEXT: ld.param.b64 %rd2, [test_st_async_release_gpu_b64_param_1];
; CHECK-NEXT: st.async.release.gpu.global.b64 [%rd1], %rd2;
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.scope.gpu.space.global.i64(ptr addrspace(1) %addr, i64 %value, /* isMultimem= */ i1 false)
+ call void @llvm.nvvm.st.async.gpu.i64(ptr addrspace(1) %addr, i64 %value, /* isMultimem= */ i1 false)
ret void
}
@@ -143,7 +143,7 @@ define void @test_st_async_mmio_release_sys_b8(ptr addrspace(1) %addr, i8 %value
; CHECK-NEXT: st.async.mmio.release.sys.global.b8 [%rd1], %st_async_val;
; CHECK-NEXT: }
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.mmio.scope.sys.space.global.i8(ptr addrspace(1) %addr, i8 %value)
+ call void @llvm.nvvm.st.async.mmio.sys.i8(ptr addrspace(1) %addr, i8 %value)
ret void
}
@@ -158,7 +158,7 @@ define void @test_st_async_mmio_release_sys_b16(ptr addrspace(1) %addr, i16 %val
; CHECK-NEXT: ld.param.b16 %rs1, [test_st_async_mmio_release_sys_b16_param_1];
; CHECK-NEXT: st.async.mmio.release.sys.global.b16 [%rd1], %rs1;
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.mmio.scope.sys.space.global.i16(ptr addrspace(1) %addr, i16 %value)
+ call void @llvm.nvvm.st.async.mmio.sys.i16(ptr addrspace(1) %addr, i16 %value)
ret void
}
@@ -173,7 +173,7 @@ define void @test_st_async_mmio_release_sys_b32(ptr addrspace(1) %addr, i32 %val
; CHECK-NEXT: ld.param.b32 %r1, [test_st_async_mmio_release_sys_b32_param_1];
; CHECK-NEXT: st.async.mmio.release.sys.global.b32 [%rd1], %r1;
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.mmio.scope.sys.space.global.i32(ptr addrspace(1) %addr, i32 %value)
+ call void @llvm.nvvm.st.async.mmio.sys.i32(ptr addrspace(1) %addr, i32 %value)
ret void
}
@@ -187,6 +187,6 @@ define void @test_st_async_mmio_release_sys_b64(ptr addrspace(1) %addr, i64 %val
; CHECK-NEXT: ld.param.b64 %rd2, [test_st_async_mmio_release_sys_b64_param_1];
; CHECK-NEXT: st.async.mmio.release.sys.global.b64 [%rd1], %rd2;
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.mmio.scope.sys.space.global.i64(ptr addrspace(1) %addr, i64 %value)
+ call void @llvm.nvvm.st.async.mmio.sys.i64(ptr addrspace(1) %addr, i64 %value)
ret void
}
diff --git a/llvm/test/CodeGen/NVPTX/st_async_release_multimem.ll b/llvm/test/CodeGen/NVPTX/st_async_release_multimem.ll
index 0631a08d3e15b..5ad6850bcd990 100644
--- a/llvm/test/CodeGen/NVPTX/st_async_release_multimem.ll
+++ b/llvm/test/CodeGen/NVPTX/st_async_release_multimem.ll
@@ -17,7 +17,7 @@ define void @test_multimem_st_async_release_sys_b8(ptr addrspace(1) %addr, i8 %v
; CHECK-NEXT: multimem.st.async.release.sys.global.b8 [%rd1], %st_async_val;
; CHECK-NEXT: }
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.scope.sys.space.global.i8(ptr addrspace(1) %addr, i8 %value, /* isMultimem= */ i1 true)
+ call void @llvm.nvvm.st.async.sys.i8(ptr addrspace(1) %addr, i8 %value, /* isMultimem= */ i1 true)
ret void
}
@@ -32,7 +32,7 @@ define void @test_multimem_st_async_release_sys_b16(ptr addrspace(1) %addr, i16
; CHECK-NEXT: ld.param.b16 %rs1, [test_multimem_st_async_release_sys_b16_param_1];
; CHECK-NEXT: multimem.st.async.release.sys.global.b16 [%rd1], %rs1;
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.scope.sys.space.global.i16(ptr addrspace(1) %addr, i16 %value, /* isMultimem= */ i1 true)
+ call void @llvm.nvvm.st.async.sys.i16(ptr addrspace(1) %addr, i16 %value, /* isMultimem= */ i1 true)
ret void
}
@@ -47,7 +47,7 @@ define void @test_multimem_st_async_release_sys_b32(ptr addrspace(1) %addr, i32
; CHECK-NEXT: ld.param.b32 %r1, [test_multimem_st_async_release_sys_b32_param_1];
; CHECK-NEXT: multimem.st.async.release.sys.global.b32 [%rd1], %r1;
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.scope.sys.space.global.i32(ptr addrspace(1) %addr, i32 %value, /* isMultimem= */ i1 true)
+ call void @llvm.nvvm.st.async.sys.i32(ptr addrspace(1) %addr, i32 %value, /* isMultimem= */ i1 true)
ret void
}
@@ -61,7 +61,7 @@ define void @test_multimem_st_async_release_sys_b64(ptr addrspace(1) %addr, i64
; CHECK-NEXT: ld.param.b64 %rd2, [test_multimem_st_async_release_sys_b64_param_1];
; CHECK-NEXT: multimem.st.async.release.sys.global.b64 [%rd1], %rd2;
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.scope.sys.space.global.i64(ptr addrspace(1) %addr, i64 %value, /* isMultimem= */ i1 true)
+ call void @llvm.nvvm.st.async.sys.i64(ptr addrspace(1) %addr, i64 %value, /* isMultimem= */ i1 true)
ret void
}
@@ -80,7 +80,7 @@ define void @test_multimem_st_async_release_gpu_b8(ptr addrspace(1) %addr, i8 %v
; CHECK-NEXT: multimem.st.async.release.gpu.global.b8 [%rd1], %st_async_val;
; CHECK-NEXT: }
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.scope.gpu.space.global.i8(ptr addrspace(1) %addr, i8 %value, /* isMultimem= */ i1 true)
+ call void @llvm.nvvm.st.async.gpu.i8(ptr addrspace(1) %addr, i8 %value, /* isMultimem= */ i1 true)
ret void
}
@@ -95,7 +95,7 @@ define void @test_multimem_st_async_release_gpu_b16(ptr addrspace(1) %addr, i16
; CHECK-NEXT: ld.param.b16 %rs1, [test_multimem_st_async_release_gpu_b16_param_1];
; CHECK-NEXT: multimem.st.async.release.gpu.global.b16 [%rd1], %rs1;
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.scope.gpu.space.global.i16(ptr addrspace(1) %addr, i16 %value, /* isMultimem= */ i1 true)
+ call void @llvm.nvvm.st.async.gpu.i16(ptr addrspace(1) %addr, i16 %value, /* isMultimem= */ i1 true)
ret void
}
@@ -110,7 +110,7 @@ define void @test_multimem_st_async_release_gpu_b32(ptr addrspace(1) %addr, i32
; CHECK-NEXT: ld.param.b32 %r1, [test_multimem_st_async_release_gpu_b32_param_1];
; CHECK-NEXT: multimem.st.async.release.gpu.global.b32 [%rd1], %r1;
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.scope.gpu.space.global.i32(ptr addrspace(1) %addr, i32 %value, /* isMultimem= */ i1 true)
+ call void @llvm.nvvm.st.async.gpu.i32(ptr addrspace(1) %addr, i32 %value, /* isMultimem= */ i1 true)
ret void
}
@@ -124,6 +124,6 @@ define void @test_multimem_st_async_release_gpu_b64(ptr addrspace(1) %addr, i64
; CHECK-NEXT: ld.param.b64 %rd2, [test_multimem_st_async_release_gpu_b64_param_1];
; CHECK-NEXT: multimem.st.async.release.gpu.global.b64 [%rd1], %rd2;
; CHECK-NEXT: ret;
- call void @llvm.nvvm.st.async.scope.gpu.space.global.i64(ptr addrspace(1) %addr, i64 %value, /* isMultimem= */ i1 true)
+ call void @llvm.nvvm.st.async.gpu.i64(ptr addrspace(1) %addr, i64 %value, /* isMultimem= */ i1 true)
ret void
}
>From 76d454703864d3942cfdcd465a5c9d542aab95f5 Mon Sep 17 00:00:00 2001
From: Srinivasa Ravi <srinivasar at nvidia.com>
Date: Tue, 9 Jun 2026 13:27:21 +0000
Subject: [PATCH 7/7] address comments
---
llvm/docs/NVPTXUsage.rst | 5 ++---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 4 ++--
2 files changed, 4 insertions(+), 5 deletions(-)
diff --git a/llvm/docs/NVPTXUsage.rst b/llvm/docs/NVPTXUsage.rst
index c1b2d7ccfbda8..c61883f007afc 100644
--- a/llvm/docs/NVPTXUsage.rst
+++ b/llvm/docs/NVPTXUsage.rst
@@ -3898,9 +3898,8 @@ is `0`, a regular ``st.async`` is emitted. When it is `1`, a
``multimem.st.async`` is emitted and the `%dest_addr` operand must be a
multimem address.
-The store operation is treated as a strong memory operation with ``.release``
-semantics. The effects of prior stores from the current thread are made visible
-to some operations in other threads in the scope.
+The store carries ``.release`` semantics — prior stores from the current thread
+are made visible to other threads in the scope.
The scope of this operation can be ``sys`` or ``gpu``.
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 33b2b132aca1e..77ae94336eb5e 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -1116,8 +1116,8 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM,
MVT::v64f32, MVT::v128f32},
Custom);
- // Custom lowering for tcgen05.st vector operands and the st.async mbarrier
- // i128 (.b128) operand. MVT::i8 is needed for the st.async.release b8
+ // Custom lowering for tcgen05.st vector operands and the st.async
+ // i128 (.b128) operand. MVT::i8 is needed for the st.async.{sys,gpu} b8
// variant.
setOperationAction(ISD::INTRINSIC_VOID,
{MVT::i8, MVT::v2i32, MVT::v4i32, MVT::v8i32, MVT::v16i32,
More information about the llvm-commits
mailing list