[llvm] [NVPTX] Add asynchronous store intrinsics (PR #200768)
Srinivasa Ravi via llvm-commits
llvm-commits at lists.llvm.org
Mon Jun 1 04:09:32 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/3] [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/3] 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/3] 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;
}
More information about the llvm-commits
mailing list