summaryrefslogtreecommitdiffstats
path: root/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0002-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch
diff options
context:
space:
mode:
Diffstat (limited to 'dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0002-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch')
-rw-r--r--dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0002-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch433
1 files changed, 433 insertions, 0 deletions
diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0002-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0002-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch
new file mode 100644
index 00000000..14e370f7
--- /dev/null
+++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0002-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch
@@ -0,0 +1,433 @@
1From 8e12d8fb3cdbdafca73fe8ed4f0cde773b1788b4 Mon Sep 17 00:00:00 2001
2From: haonanya <haonan.yang@intel.com>
3Date: Wed, 28 Jul 2021 11:43:20 +0800
4Subject: [PATCH 2/2] Add support for cl_ext_float_atomics in SPIRVWriter
5
6Upstream-Status: Backport [Taken from opencl-clang patches, https://github.com/intel/opencl-clang/blob/ocl-open-110/patches/spirv/0001-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch]
7
8Signed-off-by: haonanya <haonan.yang@intel.com>
9Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com>
10---
11 lib/SPIRV/OCLToSPIRV.cpp | 80 +++++++++++++++++++++++--
12 lib/SPIRV/OCLUtil.cpp | 26 --------
13 lib/SPIRV/OCLUtil.h | 4 --
14 test/negative/InvalidAtomicBuiltins.cl | 12 +---
15 test/transcoding/AtomicFAddEXTForOCL.ll | 64 ++++++++++++++++++++
16 test/transcoding/AtomicFMaxEXTForOCL.ll | 64 ++++++++++++++++++++
17 test/transcoding/AtomicFMinEXTForOCL.ll | 64 ++++++++++++++++++++
18 7 files changed, 269 insertions(+), 45 deletions(-)
19 create mode 100644 test/transcoding/AtomicFAddEXTForOCL.ll
20 create mode 100644 test/transcoding/AtomicFMaxEXTForOCL.ll
21 create mode 100644 test/transcoding/AtomicFMinEXTForOCL.ll
22
23diff --git a/lib/SPIRV/OCLToSPIRV.cpp b/lib/SPIRV/OCLToSPIRV.cpp
24index 04d51586..f00f5f7b 100644
25--- a/lib/SPIRV/OCLToSPIRV.cpp
26+++ b/lib/SPIRV/OCLToSPIRV.cpp
27@@ -421,10 +421,63 @@ void OCLToSPIRVBase::visitCallInst(CallInst &CI) {
28 if (DemangledName.find(kOCLBuiltinName::AtomicPrefix) == 0 ||
29 DemangledName.find(kOCLBuiltinName::AtomPrefix) == 0) {
30
31- // Compute atomic builtins do not support floating types.
32- if (CI.getType()->isFloatingPointTy() &&
33- isComputeAtomicOCLBuiltin(DemangledName))
34- return;
35+ // Compute "atom" prefixed builtins do not support floating types.
36+ if (CI.getType()->isFloatingPointTy()) {
37+ if (DemangledName.find(kOCLBuiltinName::AtomPrefix) == 0)
38+ return;
39+ // handle functions which are "atomic_" prefixed.
40+ StringRef Stem = DemangledName;
41+ Stem = Stem.drop_front(strlen("atomic_"));
42+ // FP-typed atomic_{add, sub, inc, dec, exchange, min, max, or, and, xor,
43+ // fetch_or, fetch_xor, fetch_and, fetch_or_explicit, fetch_xor_explicit,
44+ // fetch_and_explicit} should be identified as function call
45+ bool IsFunctionCall = llvm::StringSwitch<bool>(Stem)
46+ .Case("add", true)
47+ .Case("sub", true)
48+ .Case("inc", true)
49+ .Case("dec", true)
50+ .Case("cmpxchg", true)
51+ .Case("min", true)
52+ .Case("max", true)
53+ .Case("or", true)
54+ .Case("xor", true)
55+ .Case("and", true)
56+ .Case("fetch_or", true)
57+ .Case("fetch_and", true)
58+ .Case("fetch_xor", true)
59+ .Case("fetch_or_explicit", true)
60+ .Case("fetch_xor_explicit", true)
61+ .Case("fetch_and_explicit", true)
62+ .Default(false);
63+ if (IsFunctionCall)
64+ return;
65+ if (F->arg_size() != 2) {
66+ IsFunctionCall = llvm::StringSwitch<bool>(Stem)
67+ .Case("exchange", true)
68+ .Case("fetch_add", true)
69+ .Case("fetch_sub", true)
70+ .Case("fetch_min", true)
71+ .Case("fetch_max", true)
72+ .Case("load", true)
73+ .Case("store", true)
74+ .Default(false);
75+ if (IsFunctionCall)
76+ return;
77+ }
78+ if (F->arg_size() != 3 && F->arg_size() != 4) {
79+ IsFunctionCall = llvm::StringSwitch<bool>(Stem)
80+ .Case("exchange_explicit", true)
81+ .Case("fetch_add_explicit", true)
82+ .Case("fetch_sub_explicit", true)
83+ .Case("fetch_min_explicit", true)
84+ .Case("fetch_max_explicit", true)
85+ .Case("load_explicit", true)
86+ .Case("store_explicit", true)
87+ .Default(false);
88+ if (IsFunctionCall)
89+ return;
90+ }
91+ }
92
93 auto PCI = &CI;
94 if (DemangledName == kOCLBuiltinName::AtomicInit) {
95@@ -839,7 +892,7 @@ void OCLToSPIRVBase::transAtomicBuiltin(CallInst *CI,
96 AttributeList Attrs = CI->getCalledFunction()->getAttributes();
97 mutateCallInstSPIRV(
98 M, CI,
99- [=](CallInst *CI, std::vector<Value *> &Args) {
100+ [=](CallInst *CI, std::vector<Value *> &Args) -> std::string {
101 Info.PostProc(Args);
102 // Order of args in OCL20:
103 // object, 0-2 other args, 1-2 order, scope
104@@ -868,7 +921,22 @@ void OCLToSPIRVBase::transAtomicBuiltin(CallInst *CI,
105 std::rotate(Args.begin() + 2, Args.begin() + OrderIdx,
106 Args.end() - Offset);
107 }
108- return getSPIRVFuncName(OCLSPIRVBuiltinMap::map(Info.UniqName));
109+
110+ llvm::Type* AtomicBuiltinsReturnType =
111+ CI->getCalledFunction()->getReturnType();
112+ auto IsFPType = [](llvm::Type *ReturnType) {
113+ return ReturnType->isHalfTy() || ReturnType->isFloatTy() ||
114+ ReturnType->isDoubleTy();
115+ };
116+ auto SPIRVFunctionName =
117+ getSPIRVFuncName(OCLSPIRVBuiltinMap::map(Info.UniqName));
118+ if (!IsFPType(AtomicBuiltinsReturnType))
119+ return SPIRVFunctionName;
120+ // Translate FP-typed atomic builtins.
121+ return llvm::StringSwitch<std::string>(SPIRVFunctionName)
122+ .Case("__spirv_AtomicIAdd", "__spirv_AtomicFAddEXT")
123+ .Case("__spirv_AtomicSMax", "__spirv_AtomicFMaxEXT")
124+ .Case("__spirv_AtomicSMin", "__spirv_AtomicFMinEXT");
125 },
126 &Attrs);
127 }
128diff --git a/lib/SPIRV/OCLUtil.cpp b/lib/SPIRV/OCLUtil.cpp
129index 2de3f152..85155e39 100644
130--- a/lib/SPIRV/OCLUtil.cpp
131+++ b/lib/SPIRV/OCLUtil.cpp
132@@ -662,32 +662,6 @@ size_t getSPIRVAtomicBuiltinNumMemoryOrderArgs(Op OC) {
133 return 1;
134 }
135
136-bool isComputeAtomicOCLBuiltin(StringRef DemangledName) {
137- if (!DemangledName.startswith(kOCLBuiltinName::AtomicPrefix) &&
138- !DemangledName.startswith(kOCLBuiltinName::AtomPrefix))
139- return false;
140-
141- return llvm::StringSwitch<bool>(DemangledName)
142- .EndsWith("add", true)
143- .EndsWith("sub", true)
144- .EndsWith("inc", true)
145- .EndsWith("dec", true)
146- .EndsWith("cmpxchg", true)
147- .EndsWith("min", true)
148- .EndsWith("max", true)
149- .EndsWith("and", true)
150- .EndsWith("or", true)
151- .EndsWith("xor", true)
152- .EndsWith("add_explicit", true)
153- .EndsWith("sub_explicit", true)
154- .EndsWith("or_explicit", true)
155- .EndsWith("xor_explicit", true)
156- .EndsWith("and_explicit", true)
157- .EndsWith("min_explicit", true)
158- .EndsWith("max_explicit", true)
159- .Default(false);
160-}
161-
162 BarrierLiterals getBarrierLiterals(CallInst *CI) {
163 auto N = CI->getNumArgOperands();
164 assert(N == 1 || N == 2);
165diff --git a/lib/SPIRV/OCLUtil.h b/lib/SPIRV/OCLUtil.h
166index 4c05c672..c8577e9b 100644
167--- a/lib/SPIRV/OCLUtil.h
168+++ b/lib/SPIRV/OCLUtil.h
169@@ -394,10 +394,6 @@ size_t getAtomicBuiltinNumMemoryOrderArgs(StringRef Name);
170 /// Get number of memory order arguments for spirv atomic builtin function.
171 size_t getSPIRVAtomicBuiltinNumMemoryOrderArgs(Op OC);
172
173-/// Return true for OpenCL builtins which do compute operations
174-/// (like add, sub, min, max, inc, dec, ...) atomically
175-bool isComputeAtomicOCLBuiltin(StringRef DemangledName);
176-
177 /// Get OCL version from metadata opencl.ocl.version.
178 /// \param AllowMulti Allows multiple operands if true.
179 /// \return OCL version encoded as Major*10^5+Minor*10^3+Rev,
180diff --git a/test/negative/InvalidAtomicBuiltins.cl b/test/negative/InvalidAtomicBuiltins.cl
181index b8ec5b89..23dcc4e3 100644
182--- a/test/negative/InvalidAtomicBuiltins.cl
183+++ b/test/negative/InvalidAtomicBuiltins.cl
184@@ -1,7 +1,9 @@
185 // Check that translator doesn't generate atomic instructions for atomic builtins
186 // which are not defined in the spec.
187
188-// RUN: %clang_cc1 -triple spir -O1 -cl-std=cl2.0 -fdeclare-opencl-builtins -finclude-default-header %s -emit-llvm-bc -o %t.bc
189+// To drop `fdeclare-opencl-builtins` option, since FP-typed atomic function
190+// TableGen definitions have not been introduced.
191+// RUN: %clang_cc1 -triple spir -O1 -cl-std=cl2.0 -finclude-default-header %s -emit-llvm-bc -o %t.bc
192 // RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s
193 // RUN: llvm-spirv %t.bc -o %t.spv
194 // RUN: spirv-val %t.spv
195@@ -41,13 +43,9 @@ float __attribute__((overloadable)) atomic_fetch_xor(volatile generic atomic_flo
196 double __attribute__((overloadable)) atomic_fetch_and(volatile generic atomic_double *object, double operand, memory_order order);
197 double __attribute__((overloadable)) atomic_fetch_max(volatile generic atomic_double *object, double operand, memory_order order);
198 double __attribute__((overloadable)) atomic_fetch_min(volatile generic atomic_double *object, double operand, memory_order order);
199-float __attribute__((overloadable)) atomic_fetch_add_explicit(volatile generic atomic_float *object, float operand, memory_order order);
200-float __attribute__((overloadable)) atomic_fetch_sub_explicit(volatile generic atomic_float *object, float operand, memory_order order);
201 float __attribute__((overloadable)) atomic_fetch_or_explicit(volatile generic atomic_float *object, float operand, memory_order order);
202 float __attribute__((overloadable)) atomic_fetch_xor_explicit(volatile generic atomic_float *object, float operand, memory_order order);
203 double __attribute__((overloadable)) atomic_fetch_and_explicit(volatile generic atomic_double *object, double operand, memory_order order);
204-double __attribute__((overloadable)) atomic_fetch_max_explicit(volatile generic atomic_double *object, double operand, memory_order order);
205-double __attribute__((overloadable)) atomic_fetch_min_explicit(volatile generic atomic_double *object, double operand, memory_order order);
206
207 __kernel void test_atomic_fn(volatile __global float *p,
208 volatile __global double *pp,
209@@ -86,11 +84,7 @@ __kernel void test_atomic_fn(volatile __global float *p,
210 d = atomic_fetch_and(pp, val, order);
211 d = atomic_fetch_min(pp, val, order);
212 d = atomic_fetch_max(pp, val, order);
213- f = atomic_fetch_add_explicit(p, val, order);
214- f = atomic_fetch_sub_explicit(p, val, order);
215 f = atomic_fetch_or_explicit(p, val, order);
216 f = atomic_fetch_xor_explicit(p, val, order);
217 d = atomic_fetch_and_explicit(pp, val, order);
218- d = atomic_fetch_min_explicit(pp, val, order);
219- d = atomic_fetch_max_explicit(pp, val, order);
220 }
221diff --git a/test/transcoding/AtomicFAddEXTForOCL.ll b/test/transcoding/AtomicFAddEXTForOCL.ll
222new file mode 100644
223index 00000000..fb146fb9
224--- /dev/null
225+++ b/test/transcoding/AtomicFAddEXTForOCL.ll
226@@ -0,0 +1,64 @@
227+; RUN: llvm-as %s -o %t.bc
228+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_add -o %t.spv
229+; RUN: spirv-val %t.spv
230+; RUN: llvm-spirv -to-text %t.spv -o %t.spt
231+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
232+
233+; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc
234+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20
235+
236+; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc
237+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV
238+
239+target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
240+target triple = "spir-unknown-unknown"
241+
242+; CHECK-SPIRV: Capability AtomicFloat32AddEXT
243+; CHECK-SPIRV: Capability AtomicFloat64AddEXT
244+; CHECK-SPIRV: Extension "SPV_EXT_shader_atomic_float_add"
245+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32
246+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64
247+
248+
249+; Function Attrs: convergent norecurse nounwind
250+define dso_local spir_func void @test_atomic_float(float addrspace(1)* %a) local_unnamed_addr #0 {
251+entry:
252+ ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_32]]
253+ ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_add_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}})
254+ ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+fiif]]({{.*}})
255+ %call = tail call spir_func float @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2
256+ ret void
257+}
258+
259+; Function Attrs: convergent
260+declare spir_func float @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1
261+; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
262+
263+; Function Attrs: convergent norecurse nounwind
264+define dso_local spir_func void @test_atomic_double(double addrspace(1)* %a) local_unnamed_addr #0 {
265+entry:
266+ ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_64]]
267+ ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_add_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}})
268+ ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+diid]]({{.*}})
269+ %call = tail call spir_func double @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2
270+ ret void
271+}
272+; Function Attrs: convergent
273+declare spir_func double @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)*, double, i32) local_unnamed_addr #1
274+; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
275+
276+; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
277+; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
278+
279+attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
280+attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
281+attributes #2 = { convergent nounwind }
282+
283+!llvm.module.flags = !{!0}
284+!opencl.ocl.version = !{!1}
285+!opencl.spir.version = !{!1}
286+!llvm.ident = !{!2}
287+
288+!0 = !{i32 1, !"wchar_size", i32 4}
289+!1 = !{i32 2, i32 0}
290+!2 = !{!"clang version 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"}
291diff --git a/test/transcoding/AtomicFMaxEXTForOCL.ll b/test/transcoding/AtomicFMaxEXTForOCL.ll
292new file mode 100644
293index 00000000..1f2530d9
294--- /dev/null
295+++ b/test/transcoding/AtomicFMaxEXTForOCL.ll
296@@ -0,0 +1,64 @@
297+; RUN: llvm-as %s -o %t.bc
298+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_min_max -o %t.spv
299+; RUN: spirv-val %t.spv
300+; RUN: llvm-spirv -to-text %t.spv -o %t.spt
301+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
302+
303+; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc
304+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20
305+
306+; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc
307+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV
308+
309+target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
310+target triple = "spir-unknown-unknown"
311+
312+; CHECK-SPIRV: Capability AtomicFloat32MinMaxEXT
313+; CHECK-SPIRV: Capability AtomicFloat64MinMaxEXT
314+; CHECK-SPIRV: Extension "SPV_EXT_shader_atomic_float_min_max"
315+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32
316+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64
317+
318+; Function Attrs: convergent norecurse nounwind
319+define dso_local spir_func void @test_float(float addrspace(1)* %a) local_unnamed_addr #0 {
320+entry:
321+ ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_32]]
322+ ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_max_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}})
323+ ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+fiif]]({{.*}})
324+ %call = tail call spir_func float @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2
325+ ret void
326+}
327+
328+; Function Attrs: convergent
329+declare spir_func float @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1
330+; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
331+
332+; Function Attrs: convergent norecurse nounwind
333+define dso_local spir_func void @test_double(double addrspace(1)* %a) local_unnamed_addr #0 {
334+entry:
335+ ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_64]]
336+ ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_max_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}})
337+ ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+diid]]({{.*}})
338+ %call = tail call spir_func double @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2
339+ ret void
340+}
341+
342+; Function Attrs: convergent
343+declare spir_func double @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)*, double, i32) local_unnamed_addr #1
344+; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
345+
346+; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
347+; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
348+
349+attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
350+attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
351+attributes #2 = { convergent nounwind }
352+
353+!llvm.module.flags = !{!0}
354+!opencl.ocl.version = !{!1}
355+!opencl.spir.version = !{!1}
356+!llvm.ident = !{!2}
357+
358+!0 = !{i32 1, !"wchar_size", i32 4}
359+!1 = !{i32 2, i32 0}
360+!2 = !{!"clang version 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"}
361diff --git a/test/transcoding/AtomicFMinEXTForOCL.ll b/test/transcoding/AtomicFMinEXTForOCL.ll
362new file mode 100644
363index 00000000..6196b0f8
364--- /dev/null
365+++ b/test/transcoding/AtomicFMinEXTForOCL.ll
366@@ -0,0 +1,64 @@
367+; RUN: llvm-as %s -o %t.bc
368+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_min_max -o %t.spv
369+; RUN: spirv-val %t.spv
370+; RUN: llvm-spirv -to-text %t.spv -o %t.spt
371+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
372+
373+; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc
374+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20
375+
376+; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc
377+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV
378+
379+target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
380+target triple = "spir-unknown-unknown"
381+
382+; CHECK-SPIRV: Capability AtomicFloat32MinMaxEXT
383+; CHECK-SPIRV: Capability AtomicFloat64MinMaxEXT
384+; CHECK-SPIRV: Extension "SPV_EXT_shader_atomic_float_min_max"
385+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32
386+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64
387+
388+; Function Attrs: convergent norecurse nounwind
389+define dso_local spir_func void @test_float(float addrspace(1)* %a) local_unnamed_addr #0 {
390+entry:
391+ ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_32]]
392+ ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_min_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}})
393+ ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+fiif]]({{.*}})
394+ %call = tail call spir_func float @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2
395+ ret void
396+}
397+
398+; Function Attrs: convergent
399+declare spir_func float @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1
400+; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
401+
402+; Function Attrs: convergent norecurse nounwind
403+define dso_local spir_func void @test_double(double addrspace(1)* %a) local_unnamed_addr #0 {
404+entry:
405+ ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_64]]
406+ ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_min_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}})
407+ ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+diid]]({{.*}})
408+ %call = tail call spir_func double @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2
409+ ret void
410+}
411+
412+; Function Attrs: convergent
413+declare spir_func double @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)*, double, i32) local_unnamed_addr #1
414+; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
415+
416+; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
417+; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
418+
419+attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
420+attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
421+attributes #2 = { convergent nounwind }
422+
423+!llvm.module.flags = !{!0}
424+!opencl.ocl.version = !{!1}
425+!opencl.spir.version = !{!1}
426+!llvm.ident = !{!2}
427+
428+!0 = !{i32 1, !"wchar_size", i32 4}
429+!1 = !{i32 2, i32 0}
430+!2 = !{!"clang version 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"}
431--
4322.17.1
433
generated by cgit v1.2.3-54-g00ecf (git 2.39.0) at 2024-05-28 19:32:30 +0000