[llvm] fbbf5d2 - [NVPTX] Preserve volatile on atomic local loads and stores (#224719)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Sep 21 13:51:49 PDT 2026
Author: Alex MacLean
Date: 2026-09-21T13:51:43-07:00
New Revision: fbbf5d2daec93118bb5e478b7d6089fbbb857f94
URL: https://github.com/llvm/llvm-project/commit/fbbf5d2daec93118bb5e478b7d6089fbbb857f94
DIFF: https://github.com/llvm/llvm-project/commit/fbbf5d2daec93118bb5e478b7d6089fbbb857f94.diff
LOG: [NVPTX] Preserve volatile on atomic local loads and stores (#224719)
Local loads and stores discard atomic ordering during instruction
selection, but volatile accesses with acquire, release, or seq_cst
ordering also lose the volatile qualifier. Preserve volatile for all
valid load/store orderings when local volatile instructions are
supported. Follow up to #217764.
Added:
Modified:
llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
llvm/lib/Target/NVPTX/NVPTXSubtarget.h
llvm/test/CodeGen/NVPTX/volatile-local.ll
Removed:
################################################################################
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
index 81badbc1c3b10..7411aa11824ff 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
@@ -706,12 +706,6 @@ getOperationOrderings(MemSDNode *N, const NVPTXSubtarget *Subtarget) {
bool HasMemoryOrdering = Subtarget->hasMemoryOrdering();
bool HasRelaxedMMIO = Subtarget->hasRelaxedMMIO();
- bool IsSupportedLocalVolatile = CodeAddrSpace == NVPTX::AddressSpace::Local &&
- Subtarget->hasFeature(NVPTX::PTX91) &&
- N->isVolatile() &&
- (Ordering == AtomicOrdering::NotAtomic ||
- Ordering == AtomicOrdering::Unordered ||
- Ordering == AtomicOrdering::Monotonic);
// clang-format off
@@ -736,8 +730,8 @@ getOperationOrderings(MemSDNode *N, const NVPTXSubtarget *Subtarget) {
// | Relaxed | Yes | Generic,Shared [0] | .volatile | .volatile |
// | Relaxed | Yes | Global [0] | .volatile | .mmio.relaxed.sys (PTX 8.2+) |
// | | | | | or .volatile (PTX 8.1-) |
- // | Relaxed | Yes | Local (PTX 9.0-) | plain [1] | .weak [1] |
- // | Relaxed | Yes | Local (PTX 9.1+) | .volatile | .volatile |
+ // | Yes | Yes | Local (PTX 9.0-) | plain [1] | .weak [1] |
+ // | Yes | Yes | Local (PTX 9.1+) | .volatile | .volatile |
// | Relaxed | Yes | Const,Param | plain [1] | .weak [1] |
// | Other | Yes | Generic, Shared, | Error [2] | <atomic sem> [3] |
// | | | / Global [0] | | |
@@ -800,9 +794,15 @@ getOperationOrderings(MemSDNode *N, const NVPTXSubtarget *Subtarget) {
// preserve the side-effect using the weak memory instruction and
// another instruction, such as a dead dummy volatile load.
- if ((CodeAddrSpace == NVPTX::AddressSpace::Local &&
- !IsSupportedLocalVolatile) ||
- CodeAddrSpace == NVPTX::AddressSpace::Const ||
+ if (CodeAddrSpace == NVPTX::AddressSpace::Local) {
+ // Local memory is private to a thread. Drop atomic ordering but preserve
+ // volatile accesses where supported.
+ return Subtarget->hasLocalVolatile() && N->isVolatile()
+ ? NVPTX::Ordering::Volatile
+ : NVPTX::Ordering::NotAtomic;
+ }
+
+ if (CodeAddrSpace == NVPTX::AddressSpace::Const ||
CodeAddrSpace == NVPTX::AddressSpace::EntryParam ||
CodeAddrSpace == NVPTX::AddressSpace::DeviceParam) {
return NVPTX::Ordering::NotAtomic;
@@ -824,11 +824,10 @@ getOperationOrderings(MemSDNode *N, const NVPTXSubtarget *Subtarget) {
// [3]: TODO: these should eventually use .mmio<.atomic sem>; for now we drop
// the volatile semantics and preserve the atomic ones.
- // PTX atomics are not available outside generic, global, or shared memory.
- // PTX volatile operations additionally support local memory in PTX 9.1+.
+ // Apart from local volatile accesses handled above, PTX atomics and volatile
+ // operations are only available in generic, global, or shared memory.
bool AddrSupportsVolatileOrAtomic =
- (IsSupportedLocalVolatile ||
- CodeAddrSpace == NVPTX::AddressSpace::Generic ||
+ (CodeAddrSpace == NVPTX::AddressSpace::Generic ||
CodeAddrSpace == NVPTX::AddressSpace::Global ||
CodeAddrSpace == NVPTX::AddressSpace::Shared ||
CodeAddrSpace == NVPTX::AddressSpace::SharedCluster);
diff --git a/llvm/lib/Target/NVPTX/NVPTXSubtarget.h b/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
index ade564f027e4b..ee8adec2da060 100644
--- a/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
+++ b/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
@@ -119,6 +119,7 @@ class NVPTXSubtarget : public NVPTXGenSubtargetInfo {
bool hasRelaxedMMIO() const {
return hasFeature(NVPTX::SM70) && hasFeature(NVPTX::PTX82);
}
+ bool hasLocalVolatile() const { return hasFeature(NVPTX::PTX91); }
bool hasDotInstructions() const { return hasFeature(NVPTX::SM61); }
bool hasCLMAD() const {
return hasFeature(NVPTX::SM80) && hasFeature(NVPTX::PTX93);
diff --git a/llvm/test/CodeGen/NVPTX/volatile-local.ll b/llvm/test/CodeGen/NVPTX/volatile-local.ll
index 225d363904828..5dffb0e3232c8 100644
--- a/llvm/test/CodeGen/NVPTX/volatile-local.ll
+++ b/llvm/test/CodeGen/NVPTX/volatile-local.ll
@@ -54,3 +54,120 @@ define void @store_local_volatile(ptr addrspace(5) %p, i32 %v) {
store volatile i32 %v, ptr addrspace(5) %p, align 4
ret void
}
+
+; Local accesses discard atomic ordering but preserve volatile in PTX 9.1.
+define i32 @local_volatile_unordered(ptr addrspace(5) %p) {
+; PTX90-LABEL: local_volatile_unordered(
+; PTX90: {
+; PTX90-NEXT: .reg .b32 %r<2>;
+; PTX90-NEXT: .reg .b64 %rd<2>;
+; PTX90-EMPTY:
+; PTX90-NEXT: // %bb.0:
+; PTX90-NEXT: ld.param::func.b64 %rd1, [local_volatile_unordered_param_0];
+; PTX90-NEXT: ld.local.b32 %r1, [%rd1];
+; PTX90-NEXT: st.local.b32 [%rd1], %r1;
+; PTX90-NEXT: st.param::func.b32 [func_retval0], %r1;
+; PTX90-NEXT: ret;
+;
+; PTX91-LABEL: local_volatile_unordered(
+; PTX91: {
+; PTX91-NEXT: .reg .b32 %r<2>;
+; PTX91-NEXT: .reg .b64 %rd<2>;
+; PTX91-EMPTY:
+; PTX91-NEXT: // %bb.0:
+; PTX91-NEXT: ld.param::func.b64 %rd1, [local_volatile_unordered_param_0];
+; PTX91-NEXT: ld.volatile.local.b32 %r1, [%rd1];
+; PTX91-NEXT: st.volatile.local.b32 [%rd1], %r1;
+; PTX91-NEXT: st.param::func.b32 [func_retval0], %r1;
+; PTX91-NEXT: ret;
+ %v = load atomic volatile i32, ptr addrspace(5) %p unordered, align 4
+ store atomic volatile i32 %v, ptr addrspace(5) %p unordered, align 4
+ ret i32 %v
+}
+
+define i32 @local_volatile_monotonic(ptr addrspace(5) %p) {
+; PTX90-LABEL: local_volatile_monotonic(
+; PTX90: {
+; PTX90-NEXT: .reg .b32 %r<2>;
+; PTX90-NEXT: .reg .b64 %rd<2>;
+; PTX90-EMPTY:
+; PTX90-NEXT: // %bb.0:
+; PTX90-NEXT: ld.param::func.b64 %rd1, [local_volatile_monotonic_param_0];
+; PTX90-NEXT: ld.local.b32 %r1, [%rd1];
+; PTX90-NEXT: st.local.b32 [%rd1], %r1;
+; PTX90-NEXT: st.param::func.b32 [func_retval0], %r1;
+; PTX90-NEXT: ret;
+;
+; PTX91-LABEL: local_volatile_monotonic(
+; PTX91: {
+; PTX91-NEXT: .reg .b32 %r<2>;
+; PTX91-NEXT: .reg .b64 %rd<2>;
+; PTX91-EMPTY:
+; PTX91-NEXT: // %bb.0:
+; PTX91-NEXT: ld.param::func.b64 %rd1, [local_volatile_monotonic_param_0];
+; PTX91-NEXT: ld.volatile.local.b32 %r1, [%rd1];
+; PTX91-NEXT: st.volatile.local.b32 [%rd1], %r1;
+; PTX91-NEXT: st.param::func.b32 [func_retval0], %r1;
+; PTX91-NEXT: ret;
+ %v = load atomic volatile i32, ptr addrspace(5) %p monotonic, align 4
+ store atomic volatile i32 %v, ptr addrspace(5) %p monotonic, align 4
+ ret i32 %v
+}
+
+define i32 @local_volatile_acquire_release(ptr addrspace(5) %p) {
+; PTX90-LABEL: local_volatile_acquire_release(
+; PTX90: {
+; PTX90-NEXT: .reg .b32 %r<2>;
+; PTX90-NEXT: .reg .b64 %rd<2>;
+; PTX90-EMPTY:
+; PTX90-NEXT: // %bb.0:
+; PTX90-NEXT: ld.param::func.b64 %rd1, [local_volatile_acquire_release_param_0];
+; PTX90-NEXT: ld.local.b32 %r1, [%rd1];
+; PTX90-NEXT: st.local.b32 [%rd1], %r1;
+; PTX90-NEXT: st.param::func.b32 [func_retval0], %r1;
+; PTX90-NEXT: ret;
+;
+; PTX91-LABEL: local_volatile_acquire_release(
+; PTX91: {
+; PTX91-NEXT: .reg .b32 %r<2>;
+; PTX91-NEXT: .reg .b64 %rd<2>;
+; PTX91-EMPTY:
+; PTX91-NEXT: // %bb.0:
+; PTX91-NEXT: ld.param::func.b64 %rd1, [local_volatile_acquire_release_param_0];
+; PTX91-NEXT: ld.volatile.local.b32 %r1, [%rd1];
+; PTX91-NEXT: st.volatile.local.b32 [%rd1], %r1;
+; PTX91-NEXT: st.param::func.b32 [func_retval0], %r1;
+; PTX91-NEXT: ret;
+ %v = load atomic volatile i32, ptr addrspace(5) %p acquire, align 4
+ store atomic volatile i32 %v, ptr addrspace(5) %p release, align 4
+ ret i32 %v
+}
+
+define i32 @local_volatile_seq_cst(ptr addrspace(5) %p) {
+; PTX90-LABEL: local_volatile_seq_cst(
+; PTX90: {
+; PTX90-NEXT: .reg .b32 %r<2>;
+; PTX90-NEXT: .reg .b64 %rd<2>;
+; PTX90-EMPTY:
+; PTX90-NEXT: // %bb.0:
+; PTX90-NEXT: ld.param::func.b64 %rd1, [local_volatile_seq_cst_param_0];
+; PTX90-NEXT: ld.local.b32 %r1, [%rd1];
+; PTX90-NEXT: st.local.b32 [%rd1], %r1;
+; PTX90-NEXT: st.param::func.b32 [func_retval0], %r1;
+; PTX90-NEXT: ret;
+;
+; PTX91-LABEL: local_volatile_seq_cst(
+; PTX91: {
+; PTX91-NEXT: .reg .b32 %r<2>;
+; PTX91-NEXT: .reg .b64 %rd<2>;
+; PTX91-EMPTY:
+; PTX91-NEXT: // %bb.0:
+; PTX91-NEXT: ld.param::func.b64 %rd1, [local_volatile_seq_cst_param_0];
+; PTX91-NEXT: ld.volatile.local.b32 %r1, [%rd1];
+; PTX91-NEXT: st.volatile.local.b32 [%rd1], %r1;
+; PTX91-NEXT: st.param::func.b32 [func_retval0], %r1;
+; PTX91-NEXT: ret;
+ %v = load atomic volatile i32, ptr addrspace(5) %p seq_cst, align 4
+ store atomic volatile i32 %v, ptr addrspace(5) %p seq_cst, align 4
+ ret i32 %v
+}
More information about the llvm-commits
mailing list