[llvm] [SPIRV] Add missing OpenCL atomic_fetch_min/max builtin mappings (PR #190443)
Paulius Velesko via llvm-commits
llvm-commits at lists.llvm.org
Sat Apr 4 01:25:40 PDT 2026
https://github.com/pvelesko updated https://github.com/llvm/llvm-project/pull/190443
>From ed077b2464d55eac616a369c9df7f6e414aa5b87 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 | 97 +++++++++++++++++++
3 files changed, 113 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..31c4e5cefd65d
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpenCL/atomic_fetch_min_max.ll
@@ -0,0 +1,97 @@
+; 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{{$}}
+
+;; 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]] %[[#]] %[[#SVAL]]
+; CHECK-SPIRV: %[[#]] = OpAtomicSMax %[[#UINT]] %[[#SPTR]] %[[#WORKGROUP_SCOPE]] %[[#]] %[[#SVAL]]
+; CHECK-SPIRV: %[[#]] = OpAtomicSMin %[[#UINT]] %[[#SPTR]] %[[#WORKGROUP_SCOPE]] %[[#]] %[[#SVAL]]
+; CHECK-SPIRV: %[[#]] = OpAtomicSMax %[[#UINT]] %[[#SPTR]] %[[#WORKGROUP_SCOPE]] %[[#]] %[[#SVAL]]
+; CHECK-SPIRV: %[[#]] = OpAtomicSMin %[[#UINT]] %[[#SPTR]] %[[#WORKGROUP_SCOPE]] %[[#]] %[[#SVAL]]
+; CHECK-SPIRV: %[[#]] = OpAtomicSMax %[[#UINT]] %[[#SPTR]] %[[#WORKGROUP_SCOPE]] %[[#]] %[[#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]] %[[#]] %[[#UVAL]]
+; CHECK-SPIRV: %[[#]] = OpAtomicUMax %[[#UINT]] %[[#UPTR]] %[[#WORKGROUP_SCOPE]] %[[#]] %[[#UVAL]]
+; CHECK-SPIRV: %[[#]] = OpAtomicUMin %[[#UINT]] %[[#UPTR]] %[[#WORKGROUP_SCOPE]] %[[#]] %[[#UVAL]]
+; CHECK-SPIRV: %[[#]] = OpAtomicUMax %[[#UINT]] %[[#UPTR]] %[[#WORKGROUP_SCOPE]] %[[#]] %[[#UVAL]]
+; CHECK-SPIRV: %[[#]] = OpAtomicUMin %[[#UINT]] %[[#UPTR]] %[[#WORKGROUP_SCOPE]] %[[#]] %[[#UVAL]]
+; CHECK-SPIRV: %[[#]] = OpAtomicUMax %[[#UINT]] %[[#UPTR]] %[[#WORKGROUP_SCOPE]] %[[#]] %[[#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