summaryrefslogtreecommitdiffstats
path: root/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0003-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/llvm10-0003-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch')
-rw-r--r--dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0003-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch982
1 files changed, 982 insertions, 0 deletions
diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0003-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0003-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch
new file mode 100644
index 00000000..49edd7e1
--- /dev/null
+++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0003-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch
@@ -0,0 +1,982 @@
1From fbc9996d6490a5d4720b85b47f38335e7fdc99d9 Mon Sep 17 00:00:00 2001
2From: haonanya <haonan.yang@intel.com>
3Date: Mon, 19 Jul 2021 10:14:20 +0800
4Subject: [PATCH 3/3] 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-100/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/OCL20ToSPIRV.cpp | 79 ++++++++++++++++--
12 lib/SPIRV/SPIRVToOCL.h | 3 +
13 lib/SPIRV/SPIRVToOCL12.cpp | 21 +++++
14 lib/SPIRV/SPIRVToOCL20.cpp | 28 ++++++-
15 lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h | 1 -
16 lib/SPIRV/libSPIRV/SPIRVOpCode.h | 8 +-
17 test/AtomicFAddEXTForOCL.ll | 64 +++++++++++++++
18 test/AtomicFAddExt.ll | 111 ++++++++-----------------
19 test/AtomicFMaxEXT.ll | 113 +++++++-------------------
20 test/AtomicFMaxEXTForOCL.ll | 64 +++++++++++++++
21 test/AtomicFMinEXT.ll | 113 +++++++-------------------
22 test/AtomicFMinEXTForOCL.ll | 64 +++++++++++++++
23 test/InvalidAtomicBuiltins.cl | 8 --
24 13 files changed, 417 insertions(+), 260 deletions(-)
25 create mode 100644 test/AtomicFAddEXTForOCL.ll
26 create mode 100644 test/AtomicFMaxEXTForOCL.ll
27 create mode 100644 test/AtomicFMinEXTForOCL.ll
28
29diff --git a/lib/SPIRV/OCL20ToSPIRV.cpp b/lib/SPIRV/OCL20ToSPIRV.cpp
30index e30aa5be..b676a009 100644
31--- a/lib/SPIRV/OCL20ToSPIRV.cpp
32+++ b/lib/SPIRV/OCL20ToSPIRV.cpp
33@@ -408,10 +408,63 @@ void OCL20ToSPIRV::visitCallInst(CallInst &CI) {
34 if (DemangledName.find(kOCLBuiltinName::AtomicPrefix) == 0 ||
35 DemangledName.find(kOCLBuiltinName::AtomPrefix) == 0) {
36
37- // Compute atomic builtins do not support floating types.
38- if (CI.getType()->isFloatingPointTy() &&
39- isComputeAtomicOCLBuiltin(DemangledName))
40- return;
41+ // Compute "atom" prefixed builtins do not support floating types.
42+ if (CI.getType()->isFloatingPointTy()) {
43+ if (DemangledName.find(kOCLBuiltinName::AtomPrefix) == 0)
44+ return;
45+ // handle functions which are "atomic_" prefixed.
46+ StringRef Stem = DemangledName;
47+ Stem = Stem.drop_front(strlen("atomic_"));
48+ // FP-typed atomic_{add, sub, inc, dec, exchange, min, max, or, and, xor,
49+ // fetch_or, fetch_xor, fetch_and, fetch_or_explicit, fetch_xor_explicit,
50+ // fetch_and_explicit} should be identified as function call
51+ bool IsFunctionCall = llvm::StringSwitch<bool>(Stem)
52+ .Case("add", true)
53+ .Case("sub", true)
54+ .Case("inc", true)
55+ .Case("dec", true)
56+ .Case("cmpxchg", true)
57+ .Case("min", true)
58+ .Case("max", true)
59+ .Case("or", true)
60+ .Case("xor", true)
61+ .Case("and", true)
62+ .Case("fetch_or", true)
63+ .Case("fetch_and", true)
64+ .Case("fetch_xor", true)
65+ .Case("fetch_or_explicit", true)
66+ .Case("fetch_xor_explicit", true)
67+ .Case("fetch_and_explicit", true)
68+ .Default(false);
69+ if (IsFunctionCall)
70+ return;
71+ if (F->arg_size() != 2) {
72+ IsFunctionCall = llvm::StringSwitch<bool>(Stem)
73+ .Case("exchange", true)
74+ .Case("fetch_add", true)
75+ .Case("fetch_sub", true)
76+ .Case("fetch_min", true)
77+ .Case("fetch_max", true)
78+ .Case("load", true)
79+ .Case("store", true)
80+ .Default(false);
81+ if (IsFunctionCall)
82+ return;
83+ }
84+ if (F->arg_size() != 3 && F->arg_size() != 4) {
85+ IsFunctionCall = llvm::StringSwitch<bool>(Stem)
86+ .Case("exchange_explicit", true)
87+ .Case("fetch_add_explicit", true)
88+ .Case("fetch_sub_explicit", true)
89+ .Case("fetch_min_explicit", true)
90+ .Case("fetch_max_explicit", true)
91+ .Case("load_explicit", true)
92+ .Case("store_explicit", true)
93+ .Default(false);
94+ if (IsFunctionCall)
95+ return;
96+ }
97+ }
98
99 auto PCI = &CI;
100 if (DemangledName == kOCLBuiltinName::AtomicInit) {
101@@ -819,7 +872,7 @@ void OCL20ToSPIRV::transAtomicBuiltin(CallInst *CI, OCLBuiltinTransInfo &Info) {
102 AttributeList Attrs = CI->getCalledFunction()->getAttributes();
103 mutateCallInstSPIRV(
104 M, CI,
105- [=](CallInst *CI, std::vector<Value *> &Args) {
106+ [=](CallInst *CI, std::vector<Value *> &Args) -> std::string {
107 Info.PostProc(Args);
108 // Order of args in OCL20:
109 // object, 0-2 other args, 1-2 order, scope
110@@ -864,7 +917,21 @@ void OCL20ToSPIRV::transAtomicBuiltin(CallInst *CI, OCLBuiltinTransInfo &Info) {
111 std::rotate(Args.begin() + 2, Args.begin() + OrderIdx,
112 Args.end() - Offset);
113 }
114- return getSPIRVFuncName(OCLSPIRVBuiltinMap::map(Info.UniqName));
115+ llvm::Type* AtomicBuiltinsReturnType =
116+ CI->getCalledFunction()->getReturnType();
117+ auto IsFPType = [](llvm::Type *ReturnType) {
118+ return ReturnType->isHalfTy() || ReturnType->isFloatTy() ||
119+ ReturnType->isDoubleTy();
120+ };
121+ auto SPIRVFunctionName =
122+ getSPIRVFuncName(OCLSPIRVBuiltinMap::map(Info.UniqName));
123+ if (!IsFPType(AtomicBuiltinsReturnType))
124+ return SPIRVFunctionName;
125+ // Translate FP-typed atomic builtins.
126+ return llvm::StringSwitch<std::string>(SPIRVFunctionName)
127+ .Case("__spirv_AtomicIAdd", "__spirv_AtomicFAddEXT")
128+ .Case("__spirv_AtomicSMax", "__spirv_AtomicFMaxEXT")
129+ .Case("__spirv_AtomicSMin", "__spirv_AtomicFMinEXT");
130 },
131 &Attrs);
132 }
133diff --git a/lib/SPIRV/SPIRVToOCL.h b/lib/SPIRV/SPIRVToOCL.h
134index ddeec0b6..006fb0b1 100644
135--- a/lib/SPIRV/SPIRVToOCL.h
136+++ b/lib/SPIRV/SPIRVToOCL.h
137@@ -178,6 +178,9 @@ public:
138 /// using separate maps for OpenCL 1.2 and OpenCL 2.0
139 virtual Instruction *mutateAtomicName(CallInst *CI, Op OC) = 0;
140
141+ // Transform FP atomic opcode to corresponding OpenCL function name
142+ virtual std::string mapFPAtomicName(Op OC) = 0;
143+
144 private:
145 /// Transform uniform group opcode to corresponding OpenCL function name,
146 /// example: GroupIAdd(Reduce) => group_iadd => work_group_reduce_add |
147diff --git a/lib/SPIRV/SPIRVToOCL12.cpp b/lib/SPIRV/SPIRVToOCL12.cpp
148index afddd596..d7f00de3 100644
149--- a/lib/SPIRV/SPIRVToOCL12.cpp
150+++ b/lib/SPIRV/SPIRVToOCL12.cpp
151@@ -104,6 +104,9 @@ public:
152 /// cl_khr_int64_base_atomics and cl_khr_int64_extended_atomics extensions.
153 std::string mapAtomicName(Op OC, Type *Ty);
154
155+ // Transform FP atomic opcode to corresponding OpenCL function name
156+ std::string mapFPAtomicName(Op OC) override;
157+
158 static char ID;
159 };
160
161@@ -338,6 +341,21 @@ Instruction *SPIRVToOCL12::visitCallSPIRVAtomicBuiltin(CallInst *CI, Op OC) {
162 return NewCI;
163 }
164
165+std::string SPIRVToOCL12::mapFPAtomicName(Op OC) {
166+ assert(isFPAtomicOpCode(OC) && "Not intended to handle other opcodes than "
167+ "AtomicF{Add/Min/Max}EXT!");
168+ switch (OC) {
169+ case OpAtomicFAddEXT:
170+ return "atomic_add";
171+ case OpAtomicFMinEXT:
172+ return "atomic_min";
173+ case OpAtomicFMaxEXT:
174+ return "atomic_max";
175+ default:
176+ llvm_unreachable("Unsupported opcode!");
177+ }
178+}
179+
180 Instruction *SPIRVToOCL12::mutateAtomicName(CallInst *CI, Op OC) {
181 AttributeList Attrs = CI->getCalledFunction()->getAttributes();
182 return mutateCallInstOCL(
183@@ -351,6 +369,9 @@ Instruction *SPIRVToOCL12::mutateAtomicName(CallInst *CI, Op OC) {
184 std::string SPIRVToOCL12::mapAtomicName(Op OC, Type *Ty) {
185 std::string Prefix = Ty->isIntegerTy(64) ? kOCLBuiltinName::AtomPrefix
186 : kOCLBuiltinName::AtomicPrefix;
187+ // Map fp atomic instructions to regular OpenCL built-ins.
188+ if (isFPAtomicOpCode(OC))
189+ return mapFPAtomicName(OC);
190 return Prefix += OCL12SPIRVBuiltinMap::rmap(OC);
191 }
192
193diff --git a/lib/SPIRV/SPIRVToOCL20.cpp b/lib/SPIRV/SPIRVToOCL20.cpp
194index d829ff42..01d088e9 100644
195--- a/lib/SPIRV/SPIRVToOCL20.cpp
196+++ b/lib/SPIRV/SPIRVToOCL20.cpp
197@@ -82,6 +82,9 @@ public:
198 /// compare_exchange_strong/weak_explicit
199 Instruction *visitCallSPIRVAtomicCmpExchg(CallInst *CI, Op OC) override;
200
201+ // Transform FP atomic opcode to corresponding OpenCL function name
202+ std::string mapFPAtomicName(Op OC) override;
203+
204 static char ID;
205 };
206
207@@ -144,11 +147,29 @@ void SPIRVToOCL20::visitCallSPIRVControlBarrier(CallInst *CI) {
208 &Attrs);
209 }
210
211+std::string SPIRVToOCL20::mapFPAtomicName(Op OC) {
212+ assert(isFPAtomicOpCode(OC) && "Not intended to handle other opcodes than "
213+ "AtomicF{Add/Min/Max}EXT!");
214+ switch (OC) {
215+ case OpAtomicFAddEXT:
216+ return "atomic_fetch_add_explicit";
217+ case OpAtomicFMinEXT:
218+ return "atomic_fetch_min_explicit";
219+ case OpAtomicFMaxEXT:
220+ return "atomic_fetch_max_explicit";
221+ default:
222+ llvm_unreachable("Unsupported opcode!");
223+ }
224+}
225+
226 Instruction *SPIRVToOCL20::mutateAtomicName(CallInst *CI, Op OC) {
227 AttributeList Attrs = CI->getCalledFunction()->getAttributes();
228 return mutateCallInstOCL(
229 M, CI,
230 [=](CallInst *, std::vector<Value *> &Args) {
231+ // Map fp atomic instructions to regular OpenCL built-ins.
232+ if (isFPAtomicOpCode(OC))
233+ return mapFPAtomicName(OC);
234 return OCLSPIRVBuiltinMap::rmap(OC);
235 },
236 &Attrs);
237@@ -215,7 +236,12 @@ CallInst *SPIRVToOCL20::mutateCommonAtomicArguments(CallInst *CI, Op OC) {
238 }
239 }
240 auto Ptr = findFirstPtr(Args);
241- auto Name = OCLSPIRVBuiltinMap::rmap(OC);
242+ std::string Name;
243+ // Map fp atomic instructions to regular OpenCL built-ins.
244+ if (isFPAtomicOpCode(OC))
245+ Name = mapFPAtomicName(OC);
246+ else
247+ Name = OCLSPIRVBuiltinMap::rmap(OC);
248 auto NumOrder = getSPIRVAtomicBuiltinNumMemoryOrderArgs(OC);
249 auto ScopeIdx = Ptr + 1;
250 auto OrderIdx = Ptr + 2;
251diff --git a/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h b/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h
252index 13f93fbe..7b707993 100644
253--- a/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h
254+++ b/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h
255@@ -521,7 +521,6 @@ template <> inline void SPIRVMap<Capability, std::string>::init() {
256 add(CapabilityAtomicFloat64AddEXT, "AtomicFloat64AddEXT");
257 add(CapabilityAtomicFloat32MinMaxEXT, "AtomicFloat32MinMaxEXT");
258 add(CapabilityAtomicFloat64MinMaxEXT, "AtomicFloat64MinMaxEXT");
259- add(CapabilityAtomicFloat16MinMaxEXT, "AtomicFloat16MinMaxEXT");
260 add(CapabilitySubgroupShuffleINTEL, "SubgroupShuffleINTEL");
261 add(CapabilitySubgroupBufferBlockIOINTEL, "SubgroupBufferBlockIOINTEL");
262 add(CapabilitySubgroupImageBlockIOINTEL, "SubgroupImageBlockIOINTEL");
263diff --git a/lib/SPIRV/libSPIRV/SPIRVOpCode.h b/lib/SPIRV/libSPIRV/SPIRVOpCode.h
264index feec70f6..8e595e83 100644
265--- a/lib/SPIRV/libSPIRV/SPIRVOpCode.h
266+++ b/lib/SPIRV/libSPIRV/SPIRVOpCode.h
267@@ -54,11 +54,17 @@ template <> inline void SPIRVMap<Op, std::string>::init() {
268 }
269 SPIRV_DEF_NAMEMAP(Op, OpCodeNameMap)
270
271+inline bool isFPAtomicOpCode(Op OpCode) {
272+ return OpCode == OpAtomicFAddEXT || OpCode == OpAtomicFMinEXT ||
273+ OpCode == OpAtomicFMaxEXT;
274+}
275+
276 inline bool isAtomicOpCode(Op OpCode) {
277 static_assert(OpAtomicLoad < OpAtomicXor, "");
278 return ((unsigned)OpCode >= OpAtomicLoad &&
279 (unsigned)OpCode <= OpAtomicXor) ||
280- OpCode == OpAtomicFlagTestAndSet || OpCode == OpAtomicFlagClear;
281+ OpCode == OpAtomicFlagTestAndSet || OpCode == OpAtomicFlagClear ||
282+ isFPAtomicOpCode(OpCode);
283 }
284 inline bool isBinaryOpCode(Op OpCode) {
285 return ((unsigned)OpCode >= OpIAdd && (unsigned)OpCode <= OpFMod) ||
286diff --git a/test/AtomicFAddEXTForOCL.ll b/test/AtomicFAddEXTForOCL.ll
287new file mode 100644
288index 00000000..fb146fb9
289--- /dev/null
290+++ b/test/AtomicFAddEXTForOCL.ll
291@@ -0,0 +1,64 @@
292+; RUN: llvm-as %s -o %t.bc
293+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_add -o %t.spv
294+; RUN: spirv-val %t.spv
295+; RUN: llvm-spirv -to-text %t.spv -o %t.spt
296+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
297+
298+; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc
299+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20
300+
301+; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc
302+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV
303+
304+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"
305+target triple = "spir-unknown-unknown"
306+
307+; CHECK-SPIRV: Capability AtomicFloat32AddEXT
308+; CHECK-SPIRV: Capability AtomicFloat64AddEXT
309+; CHECK-SPIRV: Extension "SPV_EXT_shader_atomic_float_add"
310+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32
311+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64
312+
313+
314+; Function Attrs: convergent norecurse nounwind
315+define dso_local spir_func void @test_atomic_float(float addrspace(1)* %a) local_unnamed_addr #0 {
316+entry:
317+ ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_32]]
318+ ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_add_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}})
319+ ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+fiif]]({{.*}})
320+ %call = tail call spir_func float @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2
321+ ret void
322+}
323+
324+; Function Attrs: convergent
325+declare spir_func float @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1
326+; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
327+
328+; Function Attrs: convergent norecurse nounwind
329+define dso_local spir_func void @test_atomic_double(double addrspace(1)* %a) local_unnamed_addr #0 {
330+entry:
331+ ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_64]]
332+ ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_add_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}})
333+ ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+diid]]({{.*}})
334+ %call = tail call spir_func double @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2
335+ ret void
336+}
337+; Function Attrs: convergent
338+declare spir_func double @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)*, double, i32) local_unnamed_addr #1
339+; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
340+
341+; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
342+; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
343+
344+attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
345+attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
346+attributes #2 = { convergent nounwind }
347+
348+!llvm.module.flags = !{!0}
349+!opencl.ocl.version = !{!1}
350+!opencl.spir.version = !{!1}
351+!llvm.ident = !{!2}
352+
353+!0 = !{i32 1, !"wchar_size", i32 4}
354+!1 = !{i32 2, i32 0}
355+!2 = !{!"clang version 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"}
356diff --git a/test/AtomicFAddExt.ll b/test/AtomicFAddExt.ll
357index 011dd8a7..42bdfeea 100644
358--- a/test/AtomicFAddExt.ll
359+++ b/test/AtomicFAddExt.ll
360@@ -4,20 +4,16 @@
361 ; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
362
363 ; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
364-; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefix=CHECK-LLVM
365+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL12
366
367-target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
368-target triple = "spir64-unknown-unknown-sycldevice"
369-
370-%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" }
371-%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" = type { [1 x i64] }
372-%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" }
373-
374-$_ZTSZZ3addIfEvvENKUlRN2cl4sycl7handlerEE19_14clES3_EUlNS1_4itemILi1ELb1EEEE23_37 = comdat any
375+; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc
376+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20
377
378-$_ZTSZZ3addIdEvvENKUlRN2cl4sycl7handlerEE19_14clES3_EUlNS1_4itemILi1ELb1EEEE23_37 = comdat any
379+; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc
380+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV
381
382-@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
383+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
384+target triple = "spir64-unknown-unknown-sycldevice"
385
386 ; CHECK-SPIRV: Capability AtomicFloat32AddEXT
387 ; CHECK-SPIRV: Capability AtomicFloat64AddEXT
388@@ -25,62 +21,43 @@ $_ZTSZZ3addIdEvvENKUlRN2cl4sycl7handlerEE19_14clES3_EUlNS1_4itemILi1ELb1EEEE23_3
389 ; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32
390 ; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64
391
392-; Function Attrs: convergent norecurse mustprogress
393-define weak_odr dso_local spir_kernel void @_ZTSZZ3addIfEvvENKUlRN2cl4sycl7handlerEE19_14clES3_EUlNS1_4itemILi1ELb1EEEE23_37(float addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3, float addrspace(1)* %_arg_4, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_6, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_8) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 {
394+; Function Attrs: convergent norecurse nounwind
395+define dso_local spir_func float @_Z14AtomicFloatIncRf(float addrspace(4)* align 4 dereferenceable(4) %Arg) local_unnamed_addr #0 {
396 entry:
397- %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
398- %1 = load i64, i64* %0, align 8
399- %add.ptr.i29 = getelementptr inbounds float, float addrspace(1)* %_arg_, i64 %1
400- %2 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_8, i64 0, i32 0, i32 0, i64 0
401- %3 = load i64, i64* %2, align 8
402- %add.ptr.i = getelementptr inbounds float, float addrspace(1)* %_arg_4, i64 %3
403- %4 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32, !noalias !5
404- %5 = extractelement <3 x i64> %4, i64 0
405+ %0 = addrspacecast float addrspace(4)* %Arg to float addrspace(1)*
406 ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_32]]
407- ; CHECK-LLVM: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+]]({{.*}})
408- %call3.i.i.i.i = tail call spir_func float @_Z21__spirv_AtomicFAddEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)* %add.ptr.i29, i32 1, i32 896, float 1.000000e+00) #2
409- %add.i.i = fadd float %call3.i.i.i.i, 1.000000e+00
410- %sext.i = shl i64 %5, 32
411- %conv5.i = ashr exact i64 %sext.i, 32
412- %ptridx.i.i = getelementptr inbounds float, float addrspace(1)* %add.ptr.i, i64 %conv5.i
413- %ptridx.ascast.i.i = addrspacecast float addrspace(1)* %ptridx.i.i to float addrspace(4)*
414- store float %add.i.i, float addrspace(4)* %ptridx.ascast.i.i, align 4, !tbaa !14
415- ret void
416+ ; CHECK-LLVM-CL12: call spir_func float @[[FLOAT_FUNC_NAME:_Z10atomic_add[[:alnum:]]+ff]]({{.*}})
417+ ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_add_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}})
418+ ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+fiif]]({{.*}})
419+ %call3.i.i = tail call spir_func float @_Z21__spirv_AtomicFAddEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)* %0, i32 1, i32 896, float 1.000000e+00) #2
420+ ret float %call3.i.i
421 }
422
423 ; Function Attrs: convergent
424-; CHECK-LLVM: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float addrspace(1)*, i32, i32, float)
425 declare dso_local spir_func float @_Z21__spirv_AtomicFAddEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)*, i32, i32, float) local_unnamed_addr #1
426+; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
427
428-; Function Attrs: convergent norecurse mustprogress
429-define weak_odr dso_local spir_kernel void @_ZTSZZ3addIdEvvENKUlRN2cl4sycl7handlerEE19_14clES3_EUlNS1_4itemILi1ELb1EEEE23_37(double addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3, double addrspace(1)* %_arg_4, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_6, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_8) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 {
430+; Function Attrs: convergent norecurse nounwind
431+define dso_local spir_func double @_Z15AtomicDoubleIncRd(double addrspace(4)* align 8 dereferenceable(8) %Arg) local_unnamed_addr #0 {
432 entry:
433- %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
434- %1 = load i64, i64* %0, align 8
435- %add.ptr.i29 = getelementptr inbounds double, double addrspace(1)* %_arg_, i64 %1
436- %2 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_8, i64 0, i32 0, i32 0, i64 0
437- %3 = load i64, i64* %2, align 8
438- %add.ptr.i = getelementptr inbounds double, double addrspace(1)* %_arg_4, i64 %3
439- %4 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32, !noalias !18
440- %5 = extractelement <3 x i64> %4, i64 0
441+ %0 = addrspacecast double addrspace(4)* %Arg to double addrspace(1)*
442 ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_64]]
443- ; CHECK-LLVM: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+]]({{.*}})
444- %call3.i.i.i.i = tail call spir_func double @_Z21__spirv_AtomicFAddEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)* %add.ptr.i29, i32 1, i32 896, double 1.000000e+00) #2
445- %add.i.i = fadd double %call3.i.i.i.i, 1.000000e+00
446- %sext.i = shl i64 %5, 32
447- %conv5.i = ashr exact i64 %sext.i, 32
448- %ptridx.i.i = getelementptr inbounds double, double addrspace(1)* %add.ptr.i, i64 %conv5.i
449- %ptridx.ascast.i.i = addrspacecast double addrspace(1)* %ptridx.i.i to double addrspace(4)*
450- store double %add.i.i, double addrspace(4)* %ptridx.ascast.i.i, align 8, !tbaa !27
451- ret void
452+ ; CHECK-LLVM-CL12: call spir_func double @[[DOUBLE_FUNC_NAME:_Z10atomic_add[[:alnum:]]+dd]]({{.*}})
453+ ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_add_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}})
454+ ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+diid]]({{.*}})
455+ %call3.i.i = tail call spir_func double @_Z21__spirv_AtomicFAddEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)* %0, i32 1, i32 896, double 1.000000e+00) #2
456+ ret double %call3.i.i
457 }
458
459 ; Function Attrs: convergent
460-; CHECK-LLVM: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double addrspace(1)*, i32, i32, double)
461 declare dso_local spir_func double @_Z21__spirv_AtomicFAddEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)*, i32, i32, double) local_unnamed_addr #1
462+; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
463
464-attributes #0 = { convergent norecurse }
465-attributes #1 = { convergent }
466+; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
467+; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
468+
469+attributes #0 = { convergent norecurse nounwind "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
470+attributes #1 = { convergent "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
471 attributes #2 = { convergent nounwind }
472
473 !llvm.module.flags = !{!0}
474@@ -91,29 +68,5 @@ attributes #2 = { convergent nounwind }
475 !0 = !{i32 1, !"wchar_size", i32 4}
476 !1 = !{i32 1, i32 2}
477 !2 = !{i32 4, i32 100000}
478-!3 = !{!"clang version 12.0.0"}
479-!4 = !{i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1}
480-!5 = !{!6, !8, !10, !12}
481-!6 = distinct !{!6, !7, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"}
482-!7 = distinct !{!7, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"}
483-!8 = distinct !{!8, !9, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"}
484-!9 = distinct !{!9, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"}
485-!10 = distinct !{!10, !11, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv: %agg.result"}
486-!11 = distinct !{!11, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv"}
487-!12 = distinct !{!12, !13, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE: %agg.result"}
488-!13 = distinct !{!13, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE"}
489-!14 = !{!15, !15, i64 0}
490-!15 = !{!"float", !16, i64 0}
491-!16 = !{!"omnipotent char", !17, i64 0}
492-!17 = !{!"Simple C++ TBAA"}
493-!18 = !{!19, !21, !23, !25}
494-!19 = distinct !{!19, !20, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"}
495-!20 = distinct !{!20, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"}
496-!21 = distinct !{!21, !22, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"}
497-!22 = distinct !{!22, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"}
498-!23 = distinct !{!23, !24, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv: %agg.result"}
499-!24 = distinct !{!24, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv"}
500-!25 = distinct !{!25, !26, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE: %agg.result"}
501-!26 = distinct !{!26, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE"}
502-!27 = !{!28, !28, i64 0}
503-!28 = !{!"double", !16, i64 0}
504+!3 = !{!"clang version 13.0.0"}
505+
506diff --git a/test/AtomicFMaxEXT.ll b/test/AtomicFMaxEXT.ll
507index 1b81e53b..1c2eec93 100644
508--- a/test/AtomicFMaxEXT.ll
509+++ b/test/AtomicFMaxEXT.ll
510@@ -4,20 +4,16 @@
511 ; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
512
513 ; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
514-; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefix=CHECK-LLVM
515+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL12
516
517-target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
518-target triple = "spir64-unknown-unknown-sycldevice"
519-
520-%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" }
521-%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" = type { [1 x i64] }
522-%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" }
523-
524-$_ZTSZZ8max_testIfEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37 = comdat any
525+; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc
526+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20
527
528-$_ZTSZZ8max_testIdEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37 = comdat any
529+; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc
530+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV
531
532-@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
533+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
534+target triple = "spir64-unknown-unknown-sycldevice"
535
536 ; CHECK-SPIRV: Capability AtomicFloat32MinMaxEXT
537 ; CHECK-SPIRV: Capability AtomicFloat64MinMaxEXT
538@@ -25,65 +21,42 @@ $_ZTSZZ8max_testIdEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4item
539 ; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32
540 ; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64
541
542-; Function Attrs: convergent norecurse
543-define weak_odr dso_local spir_kernel void @_ZTSZZ8max_testIfEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37(float addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3, float addrspace(1)* %_arg_4, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_6, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_8) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 {
544+; Function Attrs: convergent norecurse nounwind
545+define dso_local spir_func float @_Z14AtomicFloatMaxRf(float addrspace(4)* align 4 dereferenceable(4) %Arg) local_unnamed_addr #0 {
546 entry:
547- %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
548- %1 = load i64, i64* %0, align 8
549- %add.ptr.i29 = getelementptr inbounds float, float addrspace(1)* %_arg_, i64 %1
550- %2 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_8, i64 0, i32 0, i32 0, i64 0
551- %3 = load i64, i64* %2, align 8
552- %add.ptr.i = getelementptr inbounds float, float addrspace(1)* %_arg_4, i64 %3
553- %4 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32, !noalias !5
554- %5 = extractelement <3 x i64> %4, i64 0
555- %conv.i = trunc i64 %5 to i32
556- %conv3.i = sitofp i32 %conv.i to float
557- %add.i = fadd float %conv3.i, 1.000000e+00
558+ %0 = addrspacecast float addrspace(4)* %Arg to float addrspace(1)*
559 ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_32]]
560- ; CHECK-LLVM: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+]]({{.*}})
561- %call3.i.i.i = tail call spir_func float @_Z21__spirv_AtomicFMaxEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)* %add.ptr.i29, i32 1, i32 896, float %add.i) #2
562- %sext.i = shl i64 %5, 32
563- %conv6.i = ashr exact i64 %sext.i, 32
564- %ptridx.i.i = getelementptr inbounds float, float addrspace(1)* %add.ptr.i, i64 %conv6.i
565- %ptridx.ascast.i.i = addrspacecast float addrspace(1)* %ptridx.i.i to float addrspace(4)*
566- store float %call3.i.i.i, float addrspace(4)* %ptridx.ascast.i.i, align 4, !tbaa !14
567- ret void
568+ ; CHECK-LLVM-CL12: call spir_func float @[[FLOAT_FUNC_NAME:_Z10atomic_max[[:alnum:]]+ff]]({{.*}})
569+ ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_max_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}})
570+ ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+fiif]]({{.*}})
571+ %call.i.i.i = tail call spir_func float @_Z21__spirv_AtomicFMaxEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)* %0, i32 1, i32 896, float 1.000000e+00) #2
572+ ret float %call.i.i.i
573 }
574
575 ; Function Attrs: convergent
576-; CHECK-LLVM: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float addrspace(1)*, i32, i32, float)
577 declare dso_local spir_func float @_Z21__spirv_AtomicFMaxEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)*, i32, i32, float) local_unnamed_addr #1
578+; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
579
580-; Function Attrs: convergent norecurse
581-define weak_odr dso_local spir_kernel void @_ZTSZZ8max_testIdEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37(double addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3, double addrspace(1)* %_arg_4, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_6, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_8) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 {
582+; Function Attrs: convergent norecurse nounwind
583+define dso_local spir_func double @_Z15AtomicDoubleMaxRd(double addrspace(4)* align 8 dereferenceable(8) %Arg) local_unnamed_addr #0 {
584 entry:
585- %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
586- %1 = load i64, i64* %0, align 8
587- %add.ptr.i29 = getelementptr inbounds double, double addrspace(1)* %_arg_, i64 %1
588- %2 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_8, i64 0, i32 0, i32 0, i64 0
589- %3 = load i64, i64* %2, align 8
590- %add.ptr.i = getelementptr inbounds double, double addrspace(1)* %_arg_4, i64 %3
591- %4 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32, !noalias !18
592- %5 = extractelement <3 x i64> %4, i64 0
593- %conv.i = trunc i64 %5 to i32
594- %conv3.i = sitofp i32 %conv.i to double
595- %add.i = fadd double %conv3.i, 1.000000e+00
596+ %0 = addrspacecast double addrspace(4)* %Arg to double addrspace(1)*
597 ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_64]]
598- ; CHECK-LLVM: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+]]({{.*}})
599- %call3.i.i.i = tail call spir_func double @_Z21__spirv_AtomicFMaxEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)* %add.ptr.i29, i32 1, i32 896, double %add.i) #2
600- %sext.i = shl i64 %5, 32
601- %conv6.i = ashr exact i64 %sext.i, 32
602- %ptridx.i.i = getelementptr inbounds double, double addrspace(1)* %add.ptr.i, i64 %conv6.i
603- %ptridx.ascast.i.i = addrspacecast double addrspace(1)* %ptridx.i.i to double addrspace(4)*
604- store double %call3.i.i.i, double addrspace(4)* %ptridx.ascast.i.i, align 8, !tbaa !27
605- ret void
606+ ; CHECK-LLVM-CL12: call spir_func double @[[DOUBLE_FUNC_NAME:_Z10atomic_max[[:alnum:]]+dd]]({{.*}})
607+ ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_max_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}})
608+ ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+diid]]({{.*}})
609+ %call.i.i.i = tail call spir_func double @_Z21__spirv_AtomicFMaxEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)* %0, i32 1, i32 896, double 1.000000e+00) #2
610+ ret double %call.i.i.i
611 }
612
613 ; Function Attrs: convergent
614-; CHECK-LLVM: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double addrspace(1)*, i32, i32, double)
615 declare dso_local spir_func double @_Z21__spirv_AtomicFMaxEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)*, i32, i32, double) local_unnamed_addr #1
616+; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
617
618-attributes #0 = { convergent norecurse "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
619+; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
620+; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
621+
622+attributes #0 = { convergent norecurse nounwind "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
623 attributes #1 = { convergent "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
624 attributes #2 = { convergent nounwind }
625
626@@ -95,29 +68,5 @@ attributes #2 = { convergent nounwind }
627 !0 = !{i32 1, !"wchar_size", i32 4}
628 !1 = !{i32 1, i32 2}
629 !2 = !{i32 4, i32 100000}
630-!3 = !{!"clang version 12.0.0"}
631-!4 = !{i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1}
632-!5 = !{!6, !8, !10, !12}
633-!6 = distinct !{!6, !7, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"}
634-!7 = distinct !{!7, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"}
635-!8 = distinct !{!8, !9, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"}
636-!9 = distinct !{!9, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"}
637-!10 = distinct !{!10, !11, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv: %agg.result"}
638-!11 = distinct !{!11, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv"}
639-!12 = distinct !{!12, !13, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE: %agg.result"}
640-!13 = distinct !{!13, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE"}
641-!14 = !{!15, !15, i64 0}
642-!15 = !{!"float", !16, i64 0}
643-!16 = !{!"omnipotent char", !17, i64 0}
644-!17 = !{!"Simple C++ TBAA"}
645-!18 = !{!19, !21, !23, !25}
646-!19 = distinct !{!19, !20, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"}
647-!20 = distinct !{!20, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"}
648-!21 = distinct !{!21, !22, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"}
649-!22 = distinct !{!22, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"}
650-!23 = distinct !{!23, !24, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv: %agg.result"}
651-!24 = distinct !{!24, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv"}
652-!25 = distinct !{!25, !26, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE: %agg.result"}
653-!26 = distinct !{!26, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE"}
654-!27 = !{!28, !28, i64 0}
655-!28 = !{!"double", !16, i64 0}
656+!3 = !{!"clang version 13.0.0"}
657+
658diff --git a/test/AtomicFMaxEXTForOCL.ll b/test/AtomicFMaxEXTForOCL.ll
659new file mode 100644
660index 00000000..1f2530d9
661--- /dev/null
662+++ b/test/AtomicFMaxEXTForOCL.ll
663@@ -0,0 +1,64 @@
664+; RUN: llvm-as %s -o %t.bc
665+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_min_max -o %t.spv
666+; RUN: spirv-val %t.spv
667+; RUN: llvm-spirv -to-text %t.spv -o %t.spt
668+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
669+
670+; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc
671+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20
672+
673+; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc
674+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV
675+
676+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"
677+target triple = "spir-unknown-unknown"
678+
679+; CHECK-SPIRV: Capability AtomicFloat32MinMaxEXT
680+; CHECK-SPIRV: Capability AtomicFloat64MinMaxEXT
681+; CHECK-SPIRV: Extension "SPV_EXT_shader_atomic_float_min_max"
682+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32
683+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64
684+
685+; Function Attrs: convergent norecurse nounwind
686+define dso_local spir_func void @test_float(float addrspace(1)* %a) local_unnamed_addr #0 {
687+entry:
688+ ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_32]]
689+ ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_max_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}})
690+ ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+fiif]]({{.*}})
691+ %call = tail call spir_func float @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2
692+ ret void
693+}
694+
695+; Function Attrs: convergent
696+declare spir_func float @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1
697+; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
698+
699+; Function Attrs: convergent norecurse nounwind
700+define dso_local spir_func void @test_double(double addrspace(1)* %a) local_unnamed_addr #0 {
701+entry:
702+ ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_64]]
703+ ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_max_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}})
704+ ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+diid]]({{.*}})
705+ %call = tail call spir_func double @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2
706+ ret void
707+}
708+
709+; Function Attrs: convergent
710+declare spir_func double @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)*, double, i32) local_unnamed_addr #1
711+; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
712+
713+; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
714+; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
715+
716+attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
717+attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
718+attributes #2 = { convergent nounwind }
719+
720+!llvm.module.flags = !{!0}
721+!opencl.ocl.version = !{!1}
722+!opencl.spir.version = !{!1}
723+!llvm.ident = !{!2}
724+
725+!0 = !{i32 1, !"wchar_size", i32 4}
726+!1 = !{i32 2, i32 0}
727+!2 = !{!"clang version 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"}
728diff --git a/test/AtomicFMinEXT.ll b/test/AtomicFMinEXT.ll
729index 98c98b8e..9e40a669 100644
730--- a/test/AtomicFMinEXT.ll
731+++ b/test/AtomicFMinEXT.ll
732@@ -4,20 +4,16 @@
733 ; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
734
735 ; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
736-; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefix=CHECK-LLVM
737+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL12
738
739-target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
740-target triple = "spir64-unknown-unknown-sycldevice"
741-
742-%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" }
743-%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" = type { [1 x i64] }
744-%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" }
745-
746-$_ZTSZZ8min_testIfEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37 = comdat any
747+; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc
748+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20
749
750-$_ZTSZZ8min_testIdEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37 = comdat any
751+; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc
752+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV
753
754-@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
755+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
756+target triple = "spir64-unknown-unknown-sycldevice"
757
758 ; CHECK-SPIRV: Capability AtomicFloat32MinMaxEXT
759 ; CHECK-SPIRV: Capability AtomicFloat64MinMaxEXT
760@@ -25,65 +21,42 @@ $_ZTSZZ8min_testIdEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4item
761 ; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32
762 ; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64
763
764-; Function Attrs: convergent norecurse
765-define weak_odr dso_local spir_kernel void @_ZTSZZ8min_testIfEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37(float addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3, float addrspace(1)* %_arg_4, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_6, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_8) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 {
766+; Function Attrs: convergent norecurse nounwind
767+define dso_local spir_func float @_Z14AtomicFloatMinRf(float addrspace(4)* align 4 dereferenceable(4) %Arg) local_unnamed_addr #0 {
768 entry:
769- %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
770- %1 = load i64, i64* %0, align 8
771- %add.ptr.i29 = getelementptr inbounds float, float addrspace(1)* %_arg_, i64 %1
772- %2 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_8, i64 0, i32 0, i32 0, i64 0
773- %3 = load i64, i64* %2, align 8
774- %add.ptr.i = getelementptr inbounds float, float addrspace(1)* %_arg_4, i64 %3
775- %4 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32, !noalias !5
776- %5 = extractelement <3 x i64> %4, i64 0
777- %conv.i = trunc i64 %5 to i32
778- %conv3.i = sitofp i32 %conv.i to float
779- %add.i = fadd float %conv3.i, 1.000000e+00
780+ %0 = addrspacecast float addrspace(4)* %Arg to float addrspace(1)*
781 ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_32]]
782- ; CHECK-LLVM: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+]]({{.*}})
783- %call3.i.i.i = tail call spir_func float @_Z21__spirv_AtomicFMinEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)* %add.ptr.i29, i32 1, i32 896, float %add.i) #2
784- %sext.i = shl i64 %5, 32
785- %conv6.i = ashr exact i64 %sext.i, 32
786- %ptridx.i.i = getelementptr inbounds float, float addrspace(1)* %add.ptr.i, i64 %conv6.i
787- %ptridx.ascast.i.i = addrspacecast float addrspace(1)* %ptridx.i.i to float addrspace(4)*
788- store float %call3.i.i.i, float addrspace(4)* %ptridx.ascast.i.i, align 4, !tbaa !14
789- ret void
790+ ; CHECK-LLVM-CL12: call spir_func float @[[FLOAT_FUNC_NAME:_Z10atomic_min[[:alnum:]]+ff]]({{.*}})
791+ ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_min_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}})
792+ ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+fiif]]({{.*}})
793+ %call.i.i.i = tail call spir_func float @_Z21__spirv_AtomicFMinEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)* %0, i32 1, i32 896, float 1.000000e+00) #2
794+ ret float %call.i.i.i
795 }
796
797 ; Function Attrs: convergent
798-; CHECK-LLVM: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float addrspace(1)*, i32, i32, float)
799 declare dso_local spir_func float @_Z21__spirv_AtomicFMinEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)*, i32, i32, float) local_unnamed_addr #1
800+; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
801
802-; Function Attrs: convergent norecurse
803-define weak_odr dso_local spir_kernel void @_ZTSZZ8min_testIdEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37(double addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3, double addrspace(1)* %_arg_4, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_6, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_8) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 {
804+; Function Attrs: convergent norecurse nounwind
805+define dso_local spir_func double @_Z15AtomicDoubleMinRd(double addrspace(4)* align 8 dereferenceable(8) %Arg) local_unnamed_addr #0 {
806 entry:
807- %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
808- %1 = load i64, i64* %0, align 8
809- %add.ptr.i29 = getelementptr inbounds double, double addrspace(1)* %_arg_, i64 %1
810- %2 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_8, i64 0, i32 0, i32 0, i64 0
811- %3 = load i64, i64* %2, align 8
812- %add.ptr.i = getelementptr inbounds double, double addrspace(1)* %_arg_4, i64 %3
813- %4 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32, !noalias !18
814- %5 = extractelement <3 x i64> %4, i64 0
815- %conv.i = trunc i64 %5 to i32
816- %conv3.i = sitofp i32 %conv.i to double
817- %add.i = fadd double %conv3.i, 1.000000e+00
818+ %0 = addrspacecast double addrspace(4)* %Arg to double addrspace(1)*
819 ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_64]]
820- ; CHECK-LLVM: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+]]({{.*}})
821- %call3.i.i.i = tail call spir_func double @_Z21__spirv_AtomicFMinEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)* %add.ptr.i29, i32 1, i32 896, double %add.i) #2
822- %sext.i = shl i64 %5, 32
823- %conv6.i = ashr exact i64 %sext.i, 32
824- %ptridx.i.i = getelementptr inbounds double, double addrspace(1)* %add.ptr.i, i64 %conv6.i
825- %ptridx.ascast.i.i = addrspacecast double addrspace(1)* %ptridx.i.i to double addrspace(4)*
826- store double %call3.i.i.i, double addrspace(4)* %ptridx.ascast.i.i, align 8, !tbaa !27
827- ret void
828+ ; CHECK-LLVM-CL12: call spir_func double @[[DOUBLE_FUNC_NAME:_Z10atomic_min[[:alnum:]]+dd]]({{.*}})
829+ ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_min_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}})
830+ ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+diid]]({{.*}})
831+ %call.i.i.i = tail call spir_func double @_Z21__spirv_AtomicFMinEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)* %0, i32 1, i32 896, double 1.000000e+00) #2
832+ ret double %call.i.i.i
833 }
834
835 ; Function Attrs: convergent
836-; CHECK-LLVM: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double addrspace(1)*, i32, i32, double)
837 declare dso_local spir_func double @_Z21__spirv_AtomicFMinEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)*, i32, i32, double) local_unnamed_addr #1
838+; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
839
840-attributes #0 = { convergent norecurse "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
841+; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
842+; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
843+
844+attributes #0 = { convergent norecurse nounwind "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
845 attributes #1 = { convergent "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
846 attributes #2 = { convergent nounwind }
847
848@@ -95,29 +68,5 @@ attributes #2 = { convergent nounwind }
849 !0 = !{i32 1, !"wchar_size", i32 4}
850 !1 = !{i32 1, i32 2}
851 !2 = !{i32 4, i32 100000}
852-!3 = !{!"clang version 12.0.0 (https://github.com/otcshare/llvm.git 67add71766d55d6a8d8d894822f583d6365a3b7d)"}
853-!4 = !{i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1}
854-!5 = !{!6, !8, !10, !12}
855-!6 = distinct !{!6, !7, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"}
856-!7 = distinct !{!7, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"}
857-!8 = distinct !{!8, !9, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"}
858-!9 = distinct !{!9, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"}
859-!10 = distinct !{!10, !11, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv: %agg.result"}
860-!11 = distinct !{!11, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv"}
861-!12 = distinct !{!12, !13, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE: %agg.result"}
862-!13 = distinct !{!13, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE"}
863-!14 = !{!15, !15, i64 0}
864-!15 = !{!"float", !16, i64 0}
865-!16 = !{!"omnipotent char", !17, i64 0}
866-!17 = !{!"Simple C++ TBAA"}
867-!18 = !{!19, !21, !23, !25}
868-!19 = distinct !{!19, !20, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"}
869-!20 = distinct !{!20, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"}
870-!21 = distinct !{!21, !22, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"}
871-!22 = distinct !{!22, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"}
872-!23 = distinct !{!23, !24, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv: %agg.result"}
873-!24 = distinct !{!24, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv"}
874-!25 = distinct !{!25, !26, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE: %agg.result"}
875-!26 = distinct !{!26, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE"}
876-!27 = !{!28, !28, i64 0}
877-!28 = !{!"double", !16, i64 0}
878+!3 = !{!"clang version 13.0.0"}
879+
880diff --git a/test/AtomicFMinEXTForOCL.ll b/test/AtomicFMinEXTForOCL.ll
881new file mode 100644
882index 00000000..6196b0f8
883--- /dev/null
884+++ b/test/AtomicFMinEXTForOCL.ll
885@@ -0,0 +1,64 @@
886+; RUN: llvm-as %s -o %t.bc
887+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_min_max -o %t.spv
888+; RUN: spirv-val %t.spv
889+; RUN: llvm-spirv -to-text %t.spv -o %t.spt
890+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
891+
892+; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc
893+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20
894+
895+; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc
896+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV
897+
898+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"
899+target triple = "spir-unknown-unknown"
900+
901+; CHECK-SPIRV: Capability AtomicFloat32MinMaxEXT
902+; CHECK-SPIRV: Capability AtomicFloat64MinMaxEXT
903+; CHECK-SPIRV: Extension "SPV_EXT_shader_atomic_float_min_max"
904+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32
905+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64
906+
907+; Function Attrs: convergent norecurse nounwind
908+define dso_local spir_func void @test_float(float addrspace(1)* %a) local_unnamed_addr #0 {
909+entry:
910+ ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_32]]
911+ ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_min_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}})
912+ ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+fiif]]({{.*}})
913+ %call = tail call spir_func float @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2
914+ ret void
915+}
916+
917+; Function Attrs: convergent
918+declare spir_func float @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1
919+; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
920+
921+; Function Attrs: convergent norecurse nounwind
922+define dso_local spir_func void @test_double(double addrspace(1)* %a) local_unnamed_addr #0 {
923+entry:
924+ ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_64]]
925+ ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_min_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}})
926+ ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+diid]]({{.*}})
927+ %call = tail call spir_func double @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2
928+ ret void
929+}
930+
931+; Function Attrs: convergent
932+declare spir_func double @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)*, double, i32) local_unnamed_addr #1
933+; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
934+
935+; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
936+; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
937+
938+attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
939+attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
940+attributes #2 = { convergent nounwind }
941+
942+!llvm.module.flags = !{!0}
943+!opencl.ocl.version = !{!1}
944+!opencl.spir.version = !{!1}
945+!llvm.ident = !{!2}
946+
947+!0 = !{i32 1, !"wchar_size", i32 4}
948+!1 = !{i32 2, i32 0}
949+!2 = !{!"clang version 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"}
950diff --git a/test/InvalidAtomicBuiltins.cl b/test/InvalidAtomicBuiltins.cl
951index b8ec5b89..2182f070 100644
952--- a/test/InvalidAtomicBuiltins.cl
953+++ b/test/InvalidAtomicBuiltins.cl
954@@ -41,13 +41,9 @@ float __attribute__((overloadable)) atomic_fetch_xor(volatile generic atomic_flo
955 double __attribute__((overloadable)) atomic_fetch_and(volatile generic atomic_double *object, double operand, memory_order order);
956 double __attribute__((overloadable)) atomic_fetch_max(volatile generic atomic_double *object, double operand, memory_order order);
957 double __attribute__((overloadable)) atomic_fetch_min(volatile generic atomic_double *object, double operand, memory_order order);
958-float __attribute__((overloadable)) atomic_fetch_add_explicit(volatile generic atomic_float *object, float operand, memory_order order);
959-float __attribute__((overloadable)) atomic_fetch_sub_explicit(volatile generic atomic_float *object, float operand, memory_order order);
960 float __attribute__((overloadable)) atomic_fetch_or_explicit(volatile generic atomic_float *object, float operand, memory_order order);
961 float __attribute__((overloadable)) atomic_fetch_xor_explicit(volatile generic atomic_float *object, float operand, memory_order order);
962 double __attribute__((overloadable)) atomic_fetch_and_explicit(volatile generic atomic_double *object, double operand, memory_order order);
963-double __attribute__((overloadable)) atomic_fetch_max_explicit(volatile generic atomic_double *object, double operand, memory_order order);
964-double __attribute__((overloadable)) atomic_fetch_min_explicit(volatile generic atomic_double *object, double operand, memory_order order);
965
966 __kernel void test_atomic_fn(volatile __global float *p,
967 volatile __global double *pp,
968@@ -86,11 +82,7 @@ __kernel void test_atomic_fn(volatile __global float *p,
969 d = atomic_fetch_and(pp, val, order);
970 d = atomic_fetch_min(pp, val, order);
971 d = atomic_fetch_max(pp, val, order);
972- f = atomic_fetch_add_explicit(p, val, order);
973- f = atomic_fetch_sub_explicit(p, val, order);
974 f = atomic_fetch_or_explicit(p, val, order);
975 f = atomic_fetch_xor_explicit(p, val, order);
976 d = atomic_fetch_and_explicit(pp, val, order);
977- d = atomic_fetch_min_explicit(pp, val, order);
978- d = atomic_fetch_max_explicit(pp, val, order);
979 }
980--
9812.17.1
982
generated by cgit v1.2.3-54-g00ecf (git 2.39.0) at 2024-05-03 18:45:50 +0000