[llvm] [SPIRV] Add missing OpenCL atomic_fetch_min/max builtin mappings (PR #190443)
Paulius Velesko via llvm-commits
llvm-commits at lists.llvm.org
Fri Apr 3 21:24:17 PDT 2026
https://github.com/pvelesko created https://github.com/llvm/llvm-project/pull/190443
## Summary
The SPIR-V backend maps OpenCL `atomic_fetch_add`/`sub`/`or`/`xor`/`and` (and their `_explicit` variants) to SPIR-V atomic opcodes, but was missing support for `atomic_fetch_min`/`atomic_fetch_max`, their `_explicit` variants, and the legacy `atom_min`/`atom_max` builtins. This caused OpenCL programs using these atomics to emit unresolved function calls instead of the correct `OpAtomicSMin`/`OpAtomicSMax`/`OpAtomicUMin`/`OpAtomicUMax` instructions.
### Approach
Unlike add/sub/or/xor/and (which are sign-agnostic), min/max require distinct signed vs unsigned SPIR-V opcodes. Rather than inspecting the `OpTypeInt` signedness bit at runtime (which is always 0 in this backend), this patch uses the existing prefix-based builtin lookup mechanism in `lookupBuiltin`: the itanium demangler exposes the argument type (`int` vs `unsigned int`), and the lookup adds an `s_` or `u_` prefix accordingly, matching the correct `.td` entry.
### Changes
- **`SPIRVBuiltins.td`**: Add 12 prefixed `DemangledNativeBuiltin` entries (`s_`/`u_` variants for `atomic_fetch_min`, `atomic_fetch_max`, `atomic_fetch_min_explicit`, `atomic_fetch_max_explicit`, `atom_min`, `atom_max`)
- **`SPIRVBuiltins.cpp`**: Add `OpAtomicSMin`/`OpAtomicSMax`/`OpAtomicUMin`/`OpAtomicUMax` cases to the `generateAtomicInst` switch to route these opcodes to `buildAtomicRMWInst` (also fixes the pre-existing `__spirv_Atomic{S,U}{Min,Max}` builtins that were defined in `.td` but not handled in the switch)
- **New test**: `atomic_fetch_min_max.ll` covering all 6 builtins for both signed and unsigned `i32` types
>From adba5284ff882532753d3d126fa3df9a96c125bb Mon Sep 17 00:00:00 2001
From: Paulius Velesko <pvelesko at pglc.io>
Date: Tue, 17 Mar 2026 09:50:45 +0200
Subject: [PATCH] [SPIRV] Add missing OpenCL atomic_fetch_min/max builtin
mappings
The SPIR-V backend had mappings for atomic_fetch_add/sub/or/xor/and
and their _explicit variants, but was missing atomic_fetch_min/max,
atomic_fetch_min/max_explicit, and the legacy atom_min/max builtins.
This caused OpenCL programs using these atomics to emit unresolved
function calls instead of the correct OpAtomicSMin/OpAtomicSMax/
OpAtomicUMin/OpAtomicUMax instructions.
Use the existing prefix-based builtin lookup mechanism to select the
correct signed or unsigned opcode based on the demangled argument
type: signed int args get the "s_" prefix (OpAtomicSMin/SMax),
unsigned int args get the "u_" prefix (OpAtomicUMin/UMax).
Also adds the missing OpAtomicSMin/SMax/UMin/UMax cases to the
generateAtomicInst switch, which fixes the pre-existing __spirv_
Atomic{S,U}{Min,Max} builtins that were already defined in the .td
but not routed to buildAtomicRMWInst.
---
llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp | 4 +
llvm/lib/Target/SPIRV/SPIRVBuiltins.td | 12 +++
.../OpenCL/atomic_fetch_min_max.ll | 100 ++++++++++++++++++
3 files changed, 116 insertions(+)
create mode 100644 llvm/test/CodeGen/SPIRV/transcoding/OpenCL/atomic_fetch_min_max.ll
diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
index db5218ec73bb7..181dda97ba81a 100644
--- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
@@ -1800,6 +1800,10 @@ static bool generateAtomicInst(const SPIRV::IncomingCall *Call,
case SPIRV::OpAtomicXor:
case SPIRV::OpAtomicAnd:
case SPIRV::OpAtomicExchange:
+ case SPIRV::OpAtomicSMax:
+ case SPIRV::OpAtomicSMin:
+ case SPIRV::OpAtomicUMax:
+ case SPIRV::OpAtomicUMin:
return buildAtomicRMWInst(Call, Opcode, MIRBuilder, GR);
case SPIRV::OpMemoryBarrier:
return buildBarrierInst(Call, SPIRV::OpMemoryBarrier, MIRBuilder, GR);
diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.td b/llvm/lib/Target/SPIRV/SPIRVBuiltins.td
index f3daf7c22674e..97692b2497592 100644
--- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.td
+++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.td
@@ -633,6 +633,18 @@ defm : DemangledNativeBuiltin<"atomic_fetch_sub_explicit", OpenCL_std, Atomic, 3
defm : DemangledNativeBuiltin<"atomic_fetch_or_explicit", OpenCL_std, Atomic, 3, 4, OpAtomicOr>;
defm : DemangledNativeBuiltin<"atomic_fetch_xor_explicit", OpenCL_std, Atomic, 3, 4, OpAtomicXor>;
defm : DemangledNativeBuiltin<"atomic_fetch_and_explicit", OpenCL_std, Atomic, 3, 4, OpAtomicAnd>;
+defm : DemangledNativeBuiltin<"s_atomic_fetch_min", OpenCL_std, Atomic, 2, 4, OpAtomicSMin>;
+defm : DemangledNativeBuiltin<"s_atomic_fetch_max", OpenCL_std, Atomic, 2, 4, OpAtomicSMax>;
+defm : DemangledNativeBuiltin<"u_atomic_fetch_min", OpenCL_std, Atomic, 2, 4, OpAtomicUMin>;
+defm : DemangledNativeBuiltin<"u_atomic_fetch_max", OpenCL_std, Atomic, 2, 4, OpAtomicUMax>;
+defm : DemangledNativeBuiltin<"s_atomic_fetch_min_explicit", OpenCL_std, Atomic, 3, 4, OpAtomicSMin>;
+defm : DemangledNativeBuiltin<"s_atomic_fetch_max_explicit", OpenCL_std, Atomic, 3, 4, OpAtomicSMax>;
+defm : DemangledNativeBuiltin<"u_atomic_fetch_min_explicit", OpenCL_std, Atomic, 3, 4, OpAtomicUMin>;
+defm : DemangledNativeBuiltin<"u_atomic_fetch_max_explicit", OpenCL_std, Atomic, 3, 4, OpAtomicUMax>;
+defm : DemangledNativeBuiltin<"s_atom_min", OpenCL_std, Atomic, 2, 2, OpAtomicSMin>;
+defm : DemangledNativeBuiltin<"s_atom_max", OpenCL_std, Atomic, 2, 2, OpAtomicSMax>;
+defm : DemangledNativeBuiltin<"u_atom_min", OpenCL_std, Atomic, 2, 2, OpAtomicUMin>;
+defm : DemangledNativeBuiltin<"u_atom_max", OpenCL_std, Atomic, 2, 2, OpAtomicUMax>;
defm : DemangledNativeBuiltin<"atomic_flag_test_and_set", OpenCL_std, Atomic, 1, 1, OpAtomicFlagTestAndSet>;
defm : DemangledNativeBuiltin<"__spirv_AtomicFlagTestAndSet", OpenCL_std, Atomic, 3, 3, OpAtomicFlagTestAndSet>;
defm : DemangledNativeBuiltin<"atomic_flag_test_and_set_explicit", OpenCL_std, Atomic, 2, 3, OpAtomicFlagTestAndSet>;
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpenCL/atomic_fetch_min_max.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpenCL/atomic_fetch_min_max.ll
new file mode 100644
index 0000000000000..ea0b203d0b162
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpenCL/atomic_fetch_min_max.ll
@@ -0,0 +1,100 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+;; This test checks that the backend correctly translates OpenCL C
+;; atomic_fetch_min/max, atomic_fetch_min/max_explicit, and legacy
+;; atom_min/max built-in functions into the corresponding SPIR-V
+;; OpAtomicSMin/OpAtomicSMax/OpAtomicUMin/OpAtomicUMax instructions,
+;; selecting the signed or unsigned variant based on the argument type.
+
+;; __kernel void test_atomic_min_max_signed(__global int *p, int val) {
+;; atomic_fetch_min(p, val);
+;; atomic_fetch_max(p, val);
+;; atomic_fetch_min_explicit(p, val, memory_order_relaxed);
+;; atomic_fetch_max_explicit(p, val, memory_order_relaxed);
+;; atom_min(p, val);
+;; atom_max(p, val);
+;; }
+;;
+;; __kernel void test_atomic_min_max_unsigned(__global unsigned int *p, unsigned int val) {
+;; atomic_fetch_min(p, val);
+;; atomic_fetch_max(p, val);
+;; atomic_fetch_min_explicit(p, val, memory_order_relaxed);
+;; atomic_fetch_max_explicit(p, val, memory_order_relaxed);
+;; atom_min(p, val);
+;; atom_max(p, val);
+;; }
+
+; CHECK-SPIRV-DAG: %[[#UINT:]] = OpTypeInt 32 0
+; CHECK-SPIRV-DAG: %[[#UINT_PTR:]] = OpTypePointer CrossWorkgroup %[[#UINT]]
+
+;; 0x2 Workgroup
+; CHECK-SPIRV-DAG: %[[#WORKGROUP_SCOPE:]] = OpConstant %[[#UINT]] 2{{$}}
+
+;; 0x0 Relaxed
+; CHECK-SPIRV-DAG: %[[#RELAXED:]] = OpConstantNull %[[#UINT]]
+
+;; Signed variants: all should use OpAtomicSMin / OpAtomicSMax
+; CHECK-SPIRV: %[[#TEST_SIGNED:]] = OpFunction %[[#]]
+; CHECK-SPIRV: %[[#SPTR:]] = OpFunctionParameter %[[#UINT_PTR]]
+; CHECK-SPIRV: %[[#SVAL:]] = OpFunctionParameter %[[#UINT]]
+; CHECK-SPIRV: %[[#]] = OpAtomicSMin %[[#UINT]] %[[#SPTR]] %[[#WORKGROUP_SCOPE]] %[[#RELAXED]] %[[#SVAL]]
+; CHECK-SPIRV: %[[#]] = OpAtomicSMax %[[#UINT]] %[[#SPTR]] %[[#WORKGROUP_SCOPE]] %[[#RELAXED]] %[[#SVAL]]
+; CHECK-SPIRV: %[[#]] = OpAtomicSMin %[[#UINT]] %[[#SPTR]] %[[#WORKGROUP_SCOPE]] %[[#RELAXED]] %[[#SVAL]]
+; CHECK-SPIRV: %[[#]] = OpAtomicSMax %[[#UINT]] %[[#SPTR]] %[[#WORKGROUP_SCOPE]] %[[#RELAXED]] %[[#SVAL]]
+; CHECK-SPIRV: %[[#]] = OpAtomicSMin %[[#UINT]] %[[#SPTR]] %[[#WORKGROUP_SCOPE]] %[[#RELAXED]] %[[#SVAL]]
+; CHECK-SPIRV: %[[#]] = OpAtomicSMax %[[#UINT]] %[[#SPTR]] %[[#WORKGROUP_SCOPE]] %[[#RELAXED]] %[[#SVAL]]
+
+;; Unsigned variants: all should use OpAtomicUMin / OpAtomicUMax
+; CHECK-SPIRV: %[[#TEST_UNSIGNED:]] = OpFunction %[[#]]
+; CHECK-SPIRV: %[[#UPTR:]] = OpFunctionParameter %[[#UINT_PTR]]
+; CHECK-SPIRV: %[[#UVAL:]] = OpFunctionParameter %[[#UINT]]
+; CHECK-SPIRV: %[[#]] = OpAtomicUMin %[[#UINT]] %[[#UPTR]] %[[#WORKGROUP_SCOPE]] %[[#RELAXED]] %[[#UVAL]]
+; CHECK-SPIRV: %[[#]] = OpAtomicUMax %[[#UINT]] %[[#UPTR]] %[[#WORKGROUP_SCOPE]] %[[#RELAXED]] %[[#UVAL]]
+; CHECK-SPIRV: %[[#]] = OpAtomicUMin %[[#UINT]] %[[#UPTR]] %[[#WORKGROUP_SCOPE]] %[[#RELAXED]] %[[#UVAL]]
+; CHECK-SPIRV: %[[#]] = OpAtomicUMax %[[#UINT]] %[[#UPTR]] %[[#WORKGROUP_SCOPE]] %[[#RELAXED]] %[[#UVAL]]
+; CHECK-SPIRV: %[[#]] = OpAtomicUMin %[[#UINT]] %[[#UPTR]] %[[#WORKGROUP_SCOPE]] %[[#RELAXED]] %[[#UVAL]]
+; CHECK-SPIRV: %[[#]] = OpAtomicUMax %[[#UINT]] %[[#UPTR]] %[[#WORKGROUP_SCOPE]] %[[#RELAXED]] %[[#UVAL]]
+
+define dso_local spir_kernel void @test_atomic_min_max_signed(ptr addrspace(1) noundef %p, i32 noundef %val) local_unnamed_addr {
+entry:
+ %call0 = tail call spir_func i32 @_Z16atomic_fetch_minPU3AS1Vii(ptr addrspace(1) noundef %p, i32 noundef %val)
+ %call1 = tail call spir_func i32 @_Z16atomic_fetch_maxPU3AS1Vii(ptr addrspace(1) noundef %p, i32 noundef %val)
+ %call2 = tail call spir_func i32 @_Z25atomic_fetch_min_explicitPU3AS1Viii(ptr addrspace(1) noundef %p, i32 noundef %val, i32 noundef 0)
+ %call3 = tail call spir_func i32 @_Z25atomic_fetch_max_explicitPU3AS1Viii(ptr addrspace(1) noundef %p, i32 noundef %val, i32 noundef 0)
+ %call4 = tail call spir_func i32 @_Z8atom_minPU3AS1Vii(ptr addrspace(1) noundef %p, i32 noundef %val)
+ %call5 = tail call spir_func i32 @_Z8atom_maxPU3AS1Vii(ptr addrspace(1) noundef %p, i32 noundef %val)
+ ret void
+}
+
+define dso_local spir_kernel void @test_atomic_min_max_unsigned(ptr addrspace(1) noundef %p, i32 noundef %val) local_unnamed_addr {
+entry:
+ %call0 = tail call spir_func i32 @_Z16atomic_fetch_minPU3AS1Vjj(ptr addrspace(1) noundef %p, i32 noundef %val)
+ %call1 = tail call spir_func i32 @_Z16atomic_fetch_maxPU3AS1Vjj(ptr addrspace(1) noundef %p, i32 noundef %val)
+ %call2 = tail call spir_func i32 @_Z25atomic_fetch_min_explicitPU3AS1Vjji(ptr addrspace(1) noundef %p, i32 noundef %val, i32 noundef 0)
+ %call3 = tail call spir_func i32 @_Z25atomic_fetch_max_explicitPU3AS1Vjji(ptr addrspace(1) noundef %p, i32 noundef %val, i32 noundef 0)
+ %call4 = tail call spir_func i32 @_Z8atom_minPU3AS1Vjj(ptr addrspace(1) noundef %p, i32 noundef %val)
+ %call5 = tail call spir_func i32 @_Z8atom_maxPU3AS1Vjj(ptr addrspace(1) noundef %p, i32 noundef %val)
+ ret void
+}
+
+;; Signed builtins
+declare spir_func i32 @_Z16atomic_fetch_minPU3AS1Vii(ptr addrspace(1) noundef, i32 noundef) local_unnamed_addr
+declare spir_func i32 @_Z16atomic_fetch_maxPU3AS1Vii(ptr addrspace(1) noundef, i32 noundef) local_unnamed_addr
+declare spir_func i32 @_Z25atomic_fetch_min_explicitPU3AS1Viii(ptr addrspace(1) noundef, i32 noundef, i32 noundef) local_unnamed_addr
+declare spir_func i32 @_Z25atomic_fetch_max_explicitPU3AS1Viii(ptr addrspace(1) noundef, i32 noundef, i32 noundef) local_unnamed_addr
+declare spir_func i32 @_Z8atom_minPU3AS1Vii(ptr addrspace(1) noundef, i32 noundef) local_unnamed_addr
+declare spir_func i32 @_Z8atom_maxPU3AS1Vii(ptr addrspace(1) noundef, i32 noundef) local_unnamed_addr
+
+;; Unsigned builtins
+declare spir_func i32 @_Z16atomic_fetch_minPU3AS1Vjj(ptr addrspace(1) noundef, i32 noundef) local_unnamed_addr
+declare spir_func i32 @_Z16atomic_fetch_maxPU3AS1Vjj(ptr addrspace(1) noundef, i32 noundef) local_unnamed_addr
+declare spir_func i32 @_Z25atomic_fetch_min_explicitPU3AS1Vjji(ptr addrspace(1) noundef, i32 noundef, i32 noundef) local_unnamed_addr
+declare spir_func i32 @_Z25atomic_fetch_max_explicitPU3AS1Vjji(ptr addrspace(1) noundef, i32 noundef, i32 noundef) local_unnamed_addr
+declare spir_func i32 @_Z8atom_minPU3AS1Vjj(ptr addrspace(1) noundef, i32 noundef) local_unnamed_addr
+declare spir_func i32 @_Z8atom_maxPU3AS1Vjj(ptr addrspace(1) noundef, i32 noundef) local_unnamed_addr
+
+;; References:
+;; [1]: https://www.khronos.org/registry/OpenCL/specs/3.0-unified/html/OpenCL_C.html#atomic-functions
+;; [2]: https://www.khronos.org/registry/spir-v/specs/unified1/SPIRV.html#OpAtomicSMin
+;; [3]: https://www.khronos.org/registry/spir-v/specs/unified1/SPIRV.html#OpAtomicSMax
More information about the llvm-commits
mailing list