diff options
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.patch | 433 |
1 files changed, 0 insertions, 433 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 deleted file mode 100644 index 14e370f7..00000000 --- a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0002-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch +++ /dev/null | |||
@@ -1,433 +0,0 @@ | |||
1 | From 8e12d8fb3cdbdafca73fe8ed4f0cde773b1788b4 Mon Sep 17 00:00:00 2001 | ||
2 | From: haonanya <haonan.yang@intel.com> | ||
3 | Date: Wed, 28 Jul 2021 11:43:20 +0800 | ||
4 | Subject: [PATCH 2/2] Add support for cl_ext_float_atomics in SPIRVWriter | ||
5 | |||
6 | Upstream-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 | |||
8 | Signed-off-by: haonanya <haonan.yang@intel.com> | ||
9 | Signed-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 | |||
23 | diff --git a/lib/SPIRV/OCLToSPIRV.cpp b/lib/SPIRV/OCLToSPIRV.cpp | ||
24 | index 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 | } | ||
128 | diff --git a/lib/SPIRV/OCLUtil.cpp b/lib/SPIRV/OCLUtil.cpp | ||
129 | index 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); | ||
165 | diff --git a/lib/SPIRV/OCLUtil.h b/lib/SPIRV/OCLUtil.h | ||
166 | index 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, | ||
180 | diff --git a/test/negative/InvalidAtomicBuiltins.cl b/test/negative/InvalidAtomicBuiltins.cl | ||
181 | index 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 | } | ||
221 | diff --git a/test/transcoding/AtomicFAddEXTForOCL.ll b/test/transcoding/AtomicFAddEXTForOCL.ll | ||
222 | new file mode 100644 | ||
223 | index 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)"} | ||
291 | diff --git a/test/transcoding/AtomicFMaxEXTForOCL.ll b/test/transcoding/AtomicFMaxEXTForOCL.ll | ||
292 | new file mode 100644 | ||
293 | index 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)"} | ||
361 | diff --git a/test/transcoding/AtomicFMinEXTForOCL.ll b/test/transcoding/AtomicFMinEXTForOCL.ll | ||
362 | new file mode 100644 | ||
363 | index 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 | -- | ||
432 | 2.17.1 | ||
433 | |||