diff options
author | Naveen Saini <naveen.kumar.saini@intel.com> | 2021-08-20 09:45:24 +0800 |
---|---|---|
committer | Anuj Mittal <anuj.mittal@intel.com> | 2021-08-24 10:41:24 +0800 |
commit | 5448b52ae684c250bf79df1cd40c3b16efcc86dc (patch) | |
tree | ac00846d84838b62b83fb6baf08ae4b449ef8212 /dynamic-layers | |
parent | 109fe9679337315fe80c1f97491fe4059fdf05cb (diff) | |
download | meta-intel-5448b52ae684c250bf79df1cd40c3b16efcc86dc.tar.gz |
llvm/11.0.0: apply opencl-clang recommend patches
https://github.com/intel/opencl-clang/tree/ocl-open-110/patches
Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com>
Signed-off-by: Anuj Mittal <anuj.mittal@intel.com>
Diffstat (limited to 'dynamic-layers')
9 files changed, 913 insertions, 23 deletions
diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-OpenCL-3.0-support.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0001-OpenCL-3.0-support.patch index 98545db0..af433e14 100644 --- a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-OpenCL-3.0-support.patch +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0001-OpenCL-3.0-support.patch | |||
@@ -1,13 +1,11 @@ | |||
1 | From d767afa79d1c8153081eac1ef33e348cadbea5bb Mon Sep 17 00:00:00 2001 | 1 | From 36d87f69fee9c3d3f399f8e4027ab707ad050e80 Mon Sep 17 00:00:00 2001 |
2 | From: Anton Zabaznov <anton.zabaznov@intel.com> | 2 | From: Anton Zabaznov <anton.zabaznov@intel.com> |
3 | Date: Tue, 22 Sep 2020 19:03:50 +0300 | 3 | Date: Tue, 22 Sep 2020 19:03:50 +0300 |
4 | Subject: [PATCH] OpenCL 3.0 support | 4 | Subject: [PATCH 1/6] OpenCL 3.0 support |
5 | 5 | ||
6 | Upstream-Status: Backport [Taken from opencl-clang patches, https://github.com/intel/opencl-clang/blob/ocl-open-110/patches/clang/0001-OpenCL-3.0-support.patch] | 6 | Upstream-Status: Backport [Taken from opencl-clang patches, https://github.com/intel/opencl-clang/blob/ocl-open-110/patches/clang/0001-OpenCL-3.0-support.patch] |
7 | Signed-off-by: Anton Zabaznov <anton.zabaznov@intel.com> | 7 | Signed-off-by: Anton Zabaznov <anton.zabaznov@intel.com> |
8 | Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com> | 8 | Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com> |
9 | |||
10 | |||
11 | --- | 9 | --- |
12 | clang/include/clang/Basic/Builtins.def | 65 +- | 10 | clang/include/clang/Basic/Builtins.def | 65 +- |
13 | clang/include/clang/Basic/Builtins.h | 13 +- | 11 | clang/include/clang/Basic/Builtins.h | 13 +- |
diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-skip-building-tests.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0001-llvm-spirv-skip-building-tests.patch index 011c09ee..237dec51 100644 --- a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-skip-building-tests.patch +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0001-llvm-spirv-skip-building-tests.patch | |||
@@ -1,7 +1,7 @@ | |||
1 | From d362652617c5e840089273df0c6623a9745c92a2 Mon Sep 17 00:00:00 2001 | 1 | From 6690d77f9007ce82984dc1b6ae12585cb3e04785 Mon Sep 17 00:00:00 2001 |
2 | From: Naveen Saini <naveen.kumar.saini@intel.com> | 2 | From: Naveen Saini <naveen.kumar.saini@intel.com> |
3 | Date: Wed, 21 Aug 2019 14:35:31 +0800 | 3 | Date: Wed, 21 Aug 2019 14:35:31 +0800 |
4 | Subject: [PATCH] llvm-spirv: skip building tests | 4 | Subject: [PATCH 1/2] llvm-spirv: skip building tests |
5 | 5 | ||
6 | Some of these need clang to be built and since we're building this in-tree, | 6 | Some of these need clang to be built and since we're building this in-tree, |
7 | that leads to problems when compiling libcxx, compiler-rt which aren't built | 7 | that leads to problems when compiling libcxx, compiler-rt which aren't built |
@@ -19,10 +19,10 @@ Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com> | |||
19 | 1 file changed, 10 deletions(-) | 19 | 1 file changed, 10 deletions(-) |
20 | 20 | ||
21 | diff --git a/CMakeLists.txt b/CMakeLists.txt | 21 | diff --git a/CMakeLists.txt b/CMakeLists.txt |
22 | index ecebb4cb..578ca602 100644 | 22 | index ec61fb95..d723c0a5 100644 |
23 | --- a/CMakeLists.txt | 23 | --- a/CMakeLists.txt |
24 | +++ b/CMakeLists.txt | 24 | +++ b/CMakeLists.txt |
25 | @@ -25,13 +25,6 @@ if(LLVM_SPIRV_BUILD_EXTERNAL) | 25 | @@ -26,13 +26,6 @@ if(LLVM_SPIRV_BUILD_EXTERNAL) |
26 | set(CMAKE_CXX_STANDARD 14) | 26 | set(CMAKE_CXX_STANDARD 14) |
27 | set(CMAKE_CXX_STANDARD_REQUIRED ON) | 27 | set(CMAKE_CXX_STANDARD_REQUIRED ON) |
28 | 28 | ||
@@ -36,7 +36,7 @@ index ecebb4cb..578ca602 100644 | |||
36 | find_package(LLVM ${BASE_LLVM_VERSION} REQUIRED | 36 | find_package(LLVM ${BASE_LLVM_VERSION} REQUIRED |
37 | COMPONENTS | 37 | COMPONENTS |
38 | Analysis | 38 | Analysis |
39 | @@ -62,9 +55,6 @@ set(LLVM_SPIRV_INCLUDE_DIRS ${CMAKE_CURRENT_SOURCE_DIR}/include) | 39 | @@ -65,9 +58,6 @@ set(LLVM_SPIRV_INCLUDE_DIRS ${CMAKE_CURRENT_SOURCE_DIR}/include) |
40 | 40 | ||
41 | add_subdirectory(lib/SPIRV) | 41 | add_subdirectory(lib/SPIRV) |
42 | add_subdirectory(tools/llvm-spirv) | 42 | add_subdirectory(tools/llvm-spirv) |
@@ -47,5 +47,5 @@ index ecebb4cb..578ca602 100644 | |||
47 | install( | 47 | install( |
48 | FILES | 48 | FILES |
49 | -- | 49 | -- |
50 | 2.26.2 | 50 | 2.17.1 |
51 | 51 | ||
diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0002-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0002-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch new file mode 100644 index 00000000..14e370f7 --- /dev/null +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0002-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch | |||
@@ -0,0 +1,433 @@ | |||
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 | |||
diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/0001-Memory-leak-fix-for-Managed-Static-Mutex.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0002-Memory-leak-fix-for-Managed-Static-Mutex.patch index 103dad5e..5b1f207e 100644 --- a/dynamic-layers/clang-layer/recipes-devtools/clang/files/0001-Memory-leak-fix-for-Managed-Static-Mutex.patch +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0002-Memory-leak-fix-for-Managed-Static-Mutex.patch | |||
@@ -1,7 +1,7 @@ | |||
1 | From c86c43b70e029b102543e8a85d269cbeb5c00279 Mon Sep 17 00:00:00 2001 | 1 | From ef27f1f99ad661c9604b7ff10efb1122466c508b Mon Sep 17 00:00:00 2001 |
2 | From: juanrod2 <> | 2 | From: juanrod2 <> |
3 | Date: Tue, 22 Dec 2020 08:33:08 +0800 | 3 | Date: Tue, 22 Dec 2020 08:33:08 +0800 |
4 | Subject: [PATCH] Memory leak fix for Managed Static Mutex | 4 | Subject: [PATCH 2/6] Memory leak fix for Managed Static Mutex |
5 | 5 | ||
6 | Upstream-Status: Backport [Taken from opencl-clang patches; https://github.com/intel/opencl-clang/blob/ocl-open-100/patches/llvm/0001-Memory-leak-fix-for-Managed-Static-Mutex.patch] | 6 | Upstream-Status: Backport [Taken from opencl-clang patches; https://github.com/intel/opencl-clang/blob/ocl-open-100/patches/llvm/0001-Memory-leak-fix-for-Managed-Static-Mutex.patch] |
7 | 7 | ||
@@ -31,5 +31,5 @@ index 053493f72fb5..6571580ccecf 100644 | |||
31 | + ManagedStaticMutex = nullptr; | 31 | + ManagedStaticMutex = nullptr; |
32 | } | 32 | } |
33 | -- | 33 | -- |
34 | 2.29.2 | 34 | 2.17.1 |
35 | 35 | ||
diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-Remove-repo-name-in-LLVM-IR.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0003-Remove-repo-name-in-LLVM-IR.patch index 09089432..15c4f9e2 100644 --- a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-Remove-repo-name-in-LLVM-IR.patch +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0003-Remove-repo-name-in-LLVM-IR.patch | |||
@@ -1,18 +1,17 @@ | |||
1 | From ff0a6da84b94c16c4519c649f1f7bed3cdf89bbb Mon Sep 17 00:00:00 2001 | 1 | From a71ab6fb04b918b856f1dd802cfdb4a7ccd53290 Mon Sep 17 00:00:00 2001 |
2 | From: Feng Zou <feng.zou@intel.com> | 2 | From: Feng Zou <feng.zou@intel.com> |
3 | Date: Tue, 20 Oct 2020 11:29:04 +0800 | 3 | Date: Tue, 20 Oct 2020 11:29:04 +0800 |
4 | Subject: [PATCH] Remove repo name in LLVM IR | 4 | Subject: [PATCH 3/6] Remove repo name in LLVM IR |
5 | 5 | ||
6 | Upstream-Status: Backport [Taken from opencl-clang patches, https://github.com/intel/opencl-clang/blob/ocl-open-110/patches/llvm/0002-Remove-repo-name-in-LLVM-IR.patch] | 6 | Upstream-Status: Backport [Taken from opencl-clang patches, https://github.com/intel/opencl-clang/blob/ocl-open-110/patches/llvm/0002-Remove-repo-name-in-LLVM-IR.patch] |
7 | Signed-off-by: Feng Zou <feng.zou@intel.com> | 7 | Signed-off-by: Feng Zou <feng.zou@intel.com> |
8 | Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com> | 8 | Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com> |
9 | |||
10 | --- | 9 | --- |
11 | llvm/cmake/modules/VersionFromVCS.cmake | 23 ++++++++++++----------- | 10 | llvm/cmake/modules/VersionFromVCS.cmake | 23 ++++++++++++----------- |
12 | 1 file changed, 12 insertions(+), 11 deletions(-) | 11 | 1 file changed, 12 insertions(+), 11 deletions(-) |
13 | 12 | ||
14 | diff --git a/llvm/cmake/modules/VersionFromVCS.cmake b/llvm/cmake/modules/VersionFromVCS.cmake | 13 | diff --git a/llvm/cmake/modules/VersionFromVCS.cmake b/llvm/cmake/modules/VersionFromVCS.cmake |
15 | index 18edbeabe3e..2d965263478 100644 | 14 | index 18edbeabe3e4..2d9652634787 100644 |
16 | --- a/llvm/cmake/modules/VersionFromVCS.cmake | 15 | --- a/llvm/cmake/modules/VersionFromVCS.cmake |
17 | +++ b/llvm/cmake/modules/VersionFromVCS.cmake | 16 | +++ b/llvm/cmake/modules/VersionFromVCS.cmake |
18 | @@ -33,17 +33,18 @@ function(get_source_info path revision repository) | 17 | @@ -33,17 +33,18 @@ function(get_source_info path revision repository) |
@@ -46,5 +45,5 @@ index 18edbeabe3e..2d965263478 100644 | |||
46 | else() | 45 | else() |
47 | message(WARNING "Git not found. Version cannot be determined.") | 46 | message(WARNING "Git not found. Version cannot be determined.") |
48 | -- | 47 | -- |
49 | 2.18.1 | 48 | 2.17.1 |
50 | 49 | ||
diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0004-Remove-__IMAGE_SUPPORT__-macro-for-SPIR-since-SPIR-d.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0004-Remove-__IMAGE_SUPPORT__-macro-for-SPIR-since-SPIR-d.patch new file mode 100644 index 00000000..25d88367 --- /dev/null +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0004-Remove-__IMAGE_SUPPORT__-macro-for-SPIR-since-SPIR-d.patch | |||
@@ -0,0 +1,51 @@ | |||
1 | From 546d9089fe5e21cccc671a0a89555cd4d5f8c817 Mon Sep 17 00:00:00 2001 | ||
2 | From: Naveen Saini <naveen.kumar.saini@intel.com> | ||
3 | Date: Thu, 19 Aug 2021 15:52:24 +0800 | ||
4 | Subject: [PATCH 4/6] Remove __IMAGE_SUPPORT__ macro for SPIR since SPIR | ||
5 | doesn't require image support | ||
6 | |||
7 | Upstream-Status: Backport [Taken from opencl-clang patches; https://github.com/intel/opencl-clang/blob/ocl-open-110/patches/clang/0002-Remove-__IMAGE_SUPPORT__-macro-for-SPIR.patch] | ||
8 | |||
9 | Signed-off-by: haonanya <haonan.yang@intel.com> | ||
10 | Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com> | ||
11 | --- | ||
12 | clang/lib/Frontend/InitPreprocessor.cpp | 3 --- | ||
13 | clang/test/Preprocessor/predefined-macros.c | 2 -- | ||
14 | 2 files changed, 5 deletions(-) | ||
15 | |||
16 | diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp | ||
17 | index 5bb489c11909..cf3b48cb65d2 100644 | ||
18 | --- a/clang/lib/Frontend/InitPreprocessor.cpp | ||
19 | +++ b/clang/lib/Frontend/InitPreprocessor.cpp | ||
20 | @@ -1115,9 +1115,6 @@ static void InitializePredefinedMacros(const TargetInfo &TI, | ||
21 | if (TI.getSupportedOpenCLOpts().isSupported(#Ext)) \ | ||
22 | Builder.defineMacro(#Ext); | ||
23 | #include "clang/Basic/OpenCLExtensions.def" | ||
24 | - | ||
25 | - if (TI.getTriple().isSPIR()) | ||
26 | - Builder.defineMacro("__IMAGE_SUPPORT__"); | ||
27 | } | ||
28 | |||
29 | if (TI.hasInt128Type() && LangOpts.CPlusPlus && LangOpts.GNUMode) { | ||
30 | diff --git a/clang/test/Preprocessor/predefined-macros.c b/clang/test/Preprocessor/predefined-macros.c | ||
31 | index 6c80517ec4d4..b5e5d7e2d546 100644 | ||
32 | --- a/clang/test/Preprocessor/predefined-macros.c | ||
33 | +++ b/clang/test/Preprocessor/predefined-macros.c | ||
34 | @@ -186,14 +186,12 @@ | ||
35 | |||
36 | // RUN: %clang_cc1 %s -E -dM -o - -x cl -triple spir-unknown-unknown \ | ||
37 | // RUN: | FileCheck -match-full-lines %s --check-prefix=CHECK-SPIR | ||
38 | -// CHECK-SPIR-DAG: #define __IMAGE_SUPPORT__ 1 | ||
39 | // CHECK-SPIR-DAG: #define __SPIR__ 1 | ||
40 | // CHECK-SPIR-DAG: #define __SPIR32__ 1 | ||
41 | // CHECK-SPIR-NOT: #define __SPIR64__ 1 | ||
42 | |||
43 | // RUN: %clang_cc1 %s -E -dM -o - -x cl -triple spir64-unknown-unknown \ | ||
44 | // RUN: | FileCheck -match-full-lines %s --check-prefix=CHECK-SPIR64 | ||
45 | -// CHECK-SPIR64-DAG: #define __IMAGE_SUPPORT__ 1 | ||
46 | // CHECK-SPIR64-DAG: #define __SPIR__ 1 | ||
47 | // CHECK-SPIR64-DAG: #define __SPIR64__ 1 | ||
48 | // CHECK-SPIR64-NOT: #define __SPIR32__ 1 | ||
49 | -- | ||
50 | 2.17.1 | ||
51 | |||
diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0005-Avoid-calling-ParseCommandLineOptions-in-BackendUtil.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0005-Avoid-calling-ParseCommandLineOptions-in-BackendUtil.patch new file mode 100644 index 00000000..2b86532c --- /dev/null +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0005-Avoid-calling-ParseCommandLineOptions-in-BackendUtil.patch | |||
@@ -0,0 +1,52 @@ | |||
1 | From 747e48959e18ac8b586078a82472a0799d12925c Mon Sep 17 00:00:00 2001 | ||
2 | From: Raphael Isemann <teemperor@gmail.com> | ||
3 | Date: Thu, 1 Apr 2021 18:41:44 +0200 | ||
4 | Subject: [PATCH 5/6] Avoid calling ParseCommandLineOptions in BackendUtil if | ||
5 | possible | ||
6 | |||
7 | Calling `ParseCommandLineOptions` should only be called from `main` as the | ||
8 | CommandLine setup code isn't thread-safe. As BackendUtil is part of the | ||
9 | generic Clang FrontendAction logic, a process which has several threads executing | ||
10 | Clang FrontendActions will randomly crash in the unsafe setup code. | ||
11 | |||
12 | This patch avoids calling the function unless either the debug-pass option or | ||
13 | limit-float-precision option is set. Without these two options set the | ||
14 | `ParseCommandLineOptions` call doesn't do anything beside parsing | ||
15 | the command line `clang` which doesn't set any options. | ||
16 | |||
17 | See also D99652 where LLDB received a workaround for this crash. | ||
18 | |||
19 | Reviewed By: JDevlieghere | ||
20 | |||
21 | Differential Revision: https://reviews.llvm.org/D99740 | ||
22 | |||
23 | Upstream-Status: Backport [Taken from opencl-clang patches; https://github.com/intel/opencl-clang/blob/ocl-open-110/patches/clang/0003-Avoid-calling-ParseCommandLineOptions-in-BackendUtil.patch] | ||
24 | |||
25 | Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com> | ||
26 | --- | ||
27 | clang/lib/CodeGen/BackendUtil.cpp | 8 ++++++++ | ||
28 | 1 file changed, 8 insertions(+) | ||
29 | |||
30 | diff --git a/clang/lib/CodeGen/BackendUtil.cpp b/clang/lib/CodeGen/BackendUtil.cpp | ||
31 | index dce0940670a2..ab478090ed1c 100644 | ||
32 | --- a/clang/lib/CodeGen/BackendUtil.cpp | ||
33 | +++ b/clang/lib/CodeGen/BackendUtil.cpp | ||
34 | @@ -797,7 +797,15 @@ static void setCommandLineOpts(const CodeGenOptions &CodeGenOpts) { | ||
35 | BackendArgs.push_back("-limit-float-precision"); | ||
36 | BackendArgs.push_back(CodeGenOpts.LimitFloatPrecision.c_str()); | ||
37 | } | ||
38 | + // Check for the default "clang" invocation that won't set any cl::opt values. | ||
39 | + // Skip trying to parse the command line invocation to avoid the issues | ||
40 | + // described below. | ||
41 | + if (BackendArgs.size() == 1) | ||
42 | + return; | ||
43 | BackendArgs.push_back(nullptr); | ||
44 | + // FIXME: The command line parser below is not thread-safe and shares a global | ||
45 | + // state, so this call might crash or overwrite the options of another Clang | ||
46 | + // instance in the same process. | ||
47 | llvm::cl::ParseCommandLineOptions(BackendArgs.size() - 1, | ||
48 | BackendArgs.data()); | ||
49 | } | ||
50 | -- | ||
51 | 2.17.1 | ||
52 | |||
diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0006-OpenCL-support-cl_ext_float_atomics.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0006-OpenCL-support-cl_ext_float_atomics.patch new file mode 100644 index 00000000..0178fd43 --- /dev/null +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0006-OpenCL-support-cl_ext_float_atomics.patch | |||
@@ -0,0 +1,353 @@ | |||
1 | From a1b924d76cdacfa3f9dbb79a9e3edddcd75f61ca Mon Sep 17 00:00:00 2001 | ||
2 | From: Naveen Saini <naveen.kumar.saini@intel.com> | ||
3 | Date: Thu, 19 Aug 2021 16:06:33 +0800 | ||
4 | Subject: [PATCH 6/6] [OpenCL] support cl_ext_float_atomics | ||
5 | |||
6 | Upstream-Status: Backport [Taken from opencl-clang patches; https://github.com/intel/opencl-clang/blob/ocl-open-110/patches/clang/0004-OpenCL-support-cl_ext_float_atomics.patch] | ||
7 | |||
8 | Signed-off-by: haonanya <haonan.yang@intel.com> | ||
9 | Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com> | ||
10 | --- | ||
11 | clang/lib/Headers/opencl-c-base.h | 25 ++++ | ||
12 | clang/lib/Headers/opencl-c.h | 195 ++++++++++++++++++++++++++ | ||
13 | clang/test/Headers/opencl-c-header.cl | 85 +++++++++++ | ||
14 | 3 files changed, 305 insertions(+) | ||
15 | |||
16 | diff --git a/clang/lib/Headers/opencl-c-base.h b/clang/lib/Headers/opencl-c-base.h | ||
17 | index afa900ab24d9..9a3ee8529acf 100644 | ||
18 | --- a/clang/lib/Headers/opencl-c-base.h | ||
19 | +++ b/clang/lib/Headers/opencl-c-base.h | ||
20 | @@ -62,6 +62,31 @@ | ||
21 | #endif | ||
22 | #endif // defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ == CL_VERSION_2_0) | ||
23 | |||
24 | +#if (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200) | ||
25 | +// For SPIR all extensions are supported. | ||
26 | +#if defined(__SPIR__) | ||
27 | +#define cl_ext_float_atomics | ||
28 | +#ifdef cl_khr_fp16 | ||
29 | +#define __opencl_c_ext_fp16_global_atomic_load_store 1 | ||
30 | +#define __opencl_c_ext_fp16_local_atomic_load_store 1 | ||
31 | +#define __opencl_c_ext_fp16_global_atomic_add 1 | ||
32 | +#define __opencl_c_ext_fp16_local_atomic_add 1 | ||
33 | +#define __opencl_c_ext_fp16_global_atomic_min_max 1 | ||
34 | +#define __opencl_c_ext_fp16_local_atomic_min_max 1 | ||
35 | +#endif | ||
36 | +#ifdef __opencl_c_fp64 | ||
37 | +#define __opencl_c_ext_fp64_global_atomic_add 1 | ||
38 | +#define __opencl_c_ext_fp64_local_atomic_add 1 | ||
39 | +#define __opencl_c_ext_fp64_global_atomic_min_max 1 | ||
40 | +#define __opencl_c_ext_fp64_local_atomic_min_max 1 | ||
41 | +#endif | ||
42 | +#define __opencl_c_ext_fp32_global_atomic_add 1 | ||
43 | +#define __opencl_c_ext_fp32_local_atomic_add 1 | ||
44 | +#define __opencl_c_ext_fp32_global_atomic_min_max 1 | ||
45 | +#define __opencl_c_ext_fp32_local_atomic_min_max 1 | ||
46 | +#endif // defined(__SPIR__) | ||
47 | +#endif // (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200) | ||
48 | + | ||
49 | // built-in scalar data types: | ||
50 | |||
51 | /** | ||
52 | diff --git a/clang/lib/Headers/opencl-c.h b/clang/lib/Headers/opencl-c.h | ||
53 | index 67d900eb1c3d..bda0f5c6df80 100644 | ||
54 | --- a/clang/lib/Headers/opencl-c.h | ||
55 | +++ b/clang/lib/Headers/opencl-c.h | ||
56 | @@ -14354,6 +14354,201 @@ intptr_t __ovld atomic_fetch_max_explicit( | ||
57 | // defined(cl_khr_int64_extended_atomics) | ||
58 | #endif // (__OPENCL_C_VERSION__ >= CL_VERSION_3_0) | ||
59 | |||
60 | +#if defined(cl_ext_float_atomics) | ||
61 | + | ||
62 | +#if defined(__opencl_c_ext_fp32_global_atomic_min_max) | ||
63 | +float __ovld atomic_fetch_min(volatile __global atomic_float *object, | ||
64 | + float operand); | ||
65 | +float __ovld atomic_fetch_max(volatile __global atomic_float *object, | ||
66 | + float operand); | ||
67 | +float __ovld atomic_fetch_min_explicit(volatile __global atomic_float *object, | ||
68 | + float operand, memory_order order); | ||
69 | +float __ovld atomic_fetch_max_explicit(volatile __global atomic_float *object, | ||
70 | + float operand, memory_order order); | ||
71 | +float __ovld atomic_fetch_min_explicit(volatile __global atomic_float *object, | ||
72 | + float operand, memory_order order, | ||
73 | + memory_scope scope); | ||
74 | +float __ovld atomic_fetch_max_explicit(volatile __global atomic_float *object, | ||
75 | + float operand, memory_order order, | ||
76 | + memory_scope scope); | ||
77 | +#endif | ||
78 | +#if defined(__opencl_c_ext_fp32_local_atomic_min_max) | ||
79 | +float __ovld atomic_fetch_min(volatile __local atomic_float *object, | ||
80 | + float operand); | ||
81 | +float __ovld atomic_fetch_max(volatile __local atomic_float *object, | ||
82 | + float operand); | ||
83 | +float __ovld atomic_fetch_min_explicit(volatile __local atomic_float *object, | ||
84 | + float operand, memory_order order); | ||
85 | +float __ovld atomic_fetch_max_explicit(volatile __local atomic_float *object, | ||
86 | + float operand, memory_order order); | ||
87 | +float __ovld atomic_fetch_min_explicit(volatile __local atomic_float *object, | ||
88 | + float operand, memory_order order, | ||
89 | + memory_scope scope); | ||
90 | +float __ovld atomic_fetch_max_explicit(volatile __local atomic_float *object, | ||
91 | + float operand, memory_order order, | ||
92 | + memory_scope scope); | ||
93 | +#endif | ||
94 | +#if defined(__opencl_c_ext_fp32_global_atomic_min_max) || \ | ||
95 | + defined(__opencl_c_ext_fp32_local_atomic_min_max) | ||
96 | +float __ovld atomic_fetch_min(volatile atomic_float *object, float operand); | ||
97 | +float __ovld atomic_fetch_max(volatile atomic_float *object, float operand); | ||
98 | +float __ovld atomic_fetch_min_explicit(volatile atomic_float *object, | ||
99 | + float operand, memory_order order); | ||
100 | +float __ovld atomic_fetch_max_explicit(volatile atomic_float *object, | ||
101 | + float operand, memory_order order); | ||
102 | +float __ovld atomic_fetch_min_explicit(volatile atomic_float *object, | ||
103 | + float operand, memory_order order, | ||
104 | + memory_scope scope); | ||
105 | +float __ovld atomic_fetch_max_explicit(volatile atomic_float *object, | ||
106 | + float operand, memory_order order, | ||
107 | + memory_scope scope); | ||
108 | +#endif | ||
109 | +#if defined(__opencl_c_ext_fp64_global_atomic_min_max) | ||
110 | +double __ovld atomic_fetch_min(volatile __global atomic_double *object, | ||
111 | + double operand); | ||
112 | +double __ovld atomic_fetch_max(volatile __global atomic_double *object, | ||
113 | + double operand); | ||
114 | +double __ovld atomic_fetch_min_explicit(volatile __global atomic_double *object, | ||
115 | + double operand, memory_order order); | ||
116 | +double __ovld atomic_fetch_max_explicit(volatile __global atomic_double *object, | ||
117 | + double operand, memory_order order); | ||
118 | +double __ovld atomic_fetch_min_explicit(volatile __global atomic_double *object, | ||
119 | + double operand, memory_order order, | ||
120 | + memory_scope scope); | ||
121 | +double __ovld atomic_fetch_max_explicit(volatile __global atomic_double *object, | ||
122 | + double operand, memory_order order, | ||
123 | + memory_scope scope); | ||
124 | +#endif | ||
125 | +#if defined(__opencl_c_ext_fp64_local_atomic_min_max) | ||
126 | +double __ovld atomic_fetch_min(volatile __local atomic_double *object, | ||
127 | + double operand); | ||
128 | +double __ovld atomic_fetch_max(volatile __local atomic_double *object, | ||
129 | + double operand); | ||
130 | +double __ovld atomic_fetch_min_explicit(volatile __local atomic_double *object, | ||
131 | + double operand, memory_order order); | ||
132 | +double __ovld atomic_fetch_max_explicit(volatile __local atomic_double *object, | ||
133 | + double operand, memory_order order); | ||
134 | +double __ovld atomic_fetch_min_explicit(volatile __local atomic_double *object, | ||
135 | + double operand, memory_order order, | ||
136 | + memory_scope scope); | ||
137 | +double __ovld atomic_fetch_max_explicit(volatile __local atomic_double *object, | ||
138 | + double operand, memory_order order, | ||
139 | + memory_scope scope); | ||
140 | +#endif | ||
141 | +#if defined(__opencl_c_ext_fp64_global_atomic_min_max) || \ | ||
142 | + defined(__opencl_c_ext_fp64_local_atomic_min_max) | ||
143 | +double __ovld atomic_fetch_min(volatile atomic_double *object, double operand); | ||
144 | +double __ovld atomic_fetch_max(volatile atomic_double *object, double operand); | ||
145 | +double __ovld atomic_fetch_min_explicit(volatile atomic_double *object, | ||
146 | + double operand, memory_order order); | ||
147 | +double __ovld atomic_fetch_max_explicit(volatile atomic_double *object, | ||
148 | + double operand, memory_order order); | ||
149 | +double __ovld atomic_fetch_min_explicit(volatile atomic_double *object, | ||
150 | + double operand, memory_order order, | ||
151 | + memory_scope scope); | ||
152 | +double __ovld atomic_fetch_max_explicit(volatile atomic_double *object, | ||
153 | + double operand, memory_order order, | ||
154 | + memory_scope scope); | ||
155 | +#endif | ||
156 | + | ||
157 | +#if defined(__opencl_c_ext_fp32_global_atomic_add) | ||
158 | +float __ovld atomic_fetch_add(volatile __global atomic_float *object, | ||
159 | + float operand); | ||
160 | +float __ovld atomic_fetch_sub(volatile __global atomic_float *object, | ||
161 | + float operand); | ||
162 | +float __ovld atomic_fetch_add_explicit(volatile __global atomic_float *object, | ||
163 | + float operand, memory_order order); | ||
164 | +float __ovld atomic_fetch_sub_explicit(volatile __global atomic_float *object, | ||
165 | + float operand, memory_order order); | ||
166 | +float __ovld atomic_fetch_add_explicit(volatile __global atomic_float *object, | ||
167 | + float operand, memory_order order, | ||
168 | + memory_scope scope); | ||
169 | +float __ovld atomic_fetch_sub_explicit(volatile __global atomic_float *object, | ||
170 | + float operand, memory_order order, | ||
171 | + memory_scope scope); | ||
172 | +#endif | ||
173 | +#if defined(__opencl_c_ext_fp32_local_atomic_add) | ||
174 | +float __ovld atomic_fetch_add(volatile __local atomic_float *object, | ||
175 | + float operand); | ||
176 | +float __ovld atomic_fetch_sub(volatile __local atomic_float *object, | ||
177 | + float operand); | ||
178 | +float __ovld atomic_fetch_add_explicit(volatile __local atomic_float *object, | ||
179 | + float operand, memory_order order); | ||
180 | +float __ovld atomic_fetch_sub_explicit(volatile __local atomic_float *object, | ||
181 | + float operand, memory_order order); | ||
182 | +float __ovld atomic_fetch_add_explicit(volatile __local atomic_float *object, | ||
183 | + float operand, memory_order order, | ||
184 | + memory_scope scope); | ||
185 | +float __ovld atomic_fetch_sub_explicit(volatile __local atomic_float *object, | ||
186 | + float operand, memory_order order, | ||
187 | + memory_scope scope); | ||
188 | +#endif | ||
189 | +#if defined(__opencl_c_ext_fp32_global_atomic_add) || \ | ||
190 | + defined(__opencl_c_ext_fp32_local_atomic_add) | ||
191 | +float __ovld atomic_fetch_add(volatile atomic_float *object, float operand); | ||
192 | +float __ovld atomic_fetch_sub(volatile atomic_float *object, float operand); | ||
193 | +float __ovld atomic_fetch_add_explicit(volatile atomic_float *object, | ||
194 | + float operand, memory_order order); | ||
195 | +float __ovld atomic_fetch_sub_explicit(volatile atomic_float *object, | ||
196 | + float operand, memory_order order); | ||
197 | +float __ovld atomic_fetch_add_explicit(volatile atomic_float *object, | ||
198 | + float operand, memory_order order, | ||
199 | + memory_scope scope); | ||
200 | +float __ovld atomic_fetch_sub_explicit(volatile atomic_float *object, | ||
201 | + float operand, memory_order order, | ||
202 | + memory_scope scope); | ||
203 | +#endif | ||
204 | + | ||
205 | +#if defined(__opencl_c_ext_fp64_global_atomic_add) | ||
206 | +double __ovld atomic_fetch_add(volatile __global atomic_double *object, | ||
207 | + double operand); | ||
208 | +double __ovld atomic_fetch_sub(volatile __global atomic_double *object, | ||
209 | + double operand); | ||
210 | +double __ovld atomic_fetch_add_explicit(volatile __global atomic_double *object, | ||
211 | + double operand, memory_order order); | ||
212 | +double __ovld atomic_fetch_sub_explicit(volatile __global atomic_double *object, | ||
213 | + double operand, memory_order order); | ||
214 | +double __ovld atomic_fetch_add_explicit(volatile __global atomic_double *object, | ||
215 | + double operand, memory_order order, | ||
216 | + memory_scope scope); | ||
217 | +double __ovld atomic_fetch_sub_explicit(volatile __global atomic_double *object, | ||
218 | + double operand, memory_order order, | ||
219 | + memory_scope scope); | ||
220 | +#endif | ||
221 | +#if defined(__opencl_c_ext_fp64_local_atomic_add) | ||
222 | +double __ovld atomic_fetch_add(volatile __local atomic_double *object, | ||
223 | + double operand); | ||
224 | +double __ovld atomic_fetch_sub(volatile __local atomic_double *object, | ||
225 | + double operand); | ||
226 | +double __ovld atomic_fetch_add_explicit(volatile __local atomic_double *object, | ||
227 | + double operand, memory_order order); | ||
228 | +double __ovld atomic_fetch_sub_explicit(volatile __local atomic_double *object, | ||
229 | + double operand, memory_order order); | ||
230 | +double __ovld atomic_fetch_add_explicit(volatile __local atomic_double *object, | ||
231 | + double operand, memory_order order, | ||
232 | + memory_scope scope); | ||
233 | +double __ovld atomic_fetch_sub_explicit(volatile __local atomic_double *object, | ||
234 | + double operand, memory_order order, | ||
235 | + memory_scope scope); | ||
236 | +#endif | ||
237 | +#if defined(__opencl_c_ext_fp64_global_atomic_add) || \ | ||
238 | + defined(__opencl_c_ext_fp64_local_atomic_add) | ||
239 | +double __ovld atomic_fetch_add(volatile atomic_double *object, double operand); | ||
240 | +double __ovld atomic_fetch_sub(volatile atomic_double *object, double operand); | ||
241 | +double __ovld atomic_fetch_add_explicit(volatile atomic_double *object, | ||
242 | + double operand, memory_order order); | ||
243 | +double __ovld atomic_fetch_sub_explicit(volatile atomic_double *object, | ||
244 | + double operand, memory_order order); | ||
245 | +double __ovld atomic_fetch_add_explicit(volatile atomic_double *object, | ||
246 | + double operand, memory_order order, | ||
247 | + memory_scope scope); | ||
248 | +double __ovld atomic_fetch_sub_explicit(volatile atomic_double *object, | ||
249 | + double operand, memory_order order, | ||
250 | + memory_scope scope); | ||
251 | +#endif | ||
252 | + | ||
253 | +#endif // cl_ext_float_atomics | ||
254 | + | ||
255 | // atomic_store() | ||
256 | |||
257 | #if defined(__opencl_c_atomic_scope_device) && \ | ||
258 | diff --git a/clang/test/Headers/opencl-c-header.cl b/clang/test/Headers/opencl-c-header.cl | ||
259 | index 2716076acdcf..6b3eca84e8b9 100644 | ||
260 | --- a/clang/test/Headers/opencl-c-header.cl | ||
261 | +++ b/clang/test/Headers/opencl-c-header.cl | ||
262 | @@ -98,3 +98,88 @@ global atomic_int z = ATOMIC_VAR_INIT(99); | ||
263 | #pragma OPENCL EXTENSION cl_intel_planar_yuv : enable | ||
264 | |||
265 | // CHECK-MOD: Reading modules | ||
266 | + | ||
267 | +// For SPIR all extensions are supported. | ||
268 | +#if defined(__SPIR__) | ||
269 | + | ||
270 | +#if (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200) | ||
271 | + | ||
272 | +#if __opencl_c_ext_fp16_global_atomic_load_store != 1 | ||
273 | +#error "Incorrectly defined __opencl_c_ext_fp16_global_atomic_load_store" | ||
274 | +#endif | ||
275 | +#if __opencl_c_ext_fp16_local_atomic_load_store != 1 | ||
276 | +#error "Incorrectly defined __opencl_c_ext_fp16_local_atomic_load_store" | ||
277 | +#endif | ||
278 | +#if __opencl_c_ext_fp16_global_atomic_add != 1 | ||
279 | +#error "Incorrectly defined __opencl_c_ext_fp16_global_atomic_add" | ||
280 | +#endif | ||
281 | +#if __opencl_c_ext_fp32_global_atomic_add != 1 | ||
282 | +#error "Incorrectly defined __opencl_c_ext_fp32_global_atomic_add" | ||
283 | +#endif | ||
284 | +#if __opencl_c_ext_fp16_local_atomic_add != 1 | ||
285 | +#error "Incorrectly defined __opencl_c_ext_fp16_local_atomic_add" | ||
286 | +#endif | ||
287 | +#if __opencl_c_ext_fp32_local_atomic_add != 1 | ||
288 | +#error "Incorrectly defined __opencl_c_ext_fp32_local_atomic_add" | ||
289 | +#endif | ||
290 | +#if __opencl_c_ext_fp16_global_atomic_min_max != 1 | ||
291 | +#error "Incorrectly defined __opencl_c_ext_fp16_global_atomic_min_max" | ||
292 | +#endif | ||
293 | +#if __opencl_c_ext_fp32_global_atomic_min_max != 1 | ||
294 | +#error "Incorrectly defined __opencl_c_ext_fp32_global_atomic_min_max" | ||
295 | +#endif | ||
296 | +#if __opencl_c_ext_fp16_local_atomic_min_max != 1 | ||
297 | +#error "Incorrectly defined __opencl_c_ext_fp16_local_atomic_min_max" | ||
298 | +#endif | ||
299 | +#if __opencl_c_ext_fp32_local_atomic_min_max != 1 | ||
300 | +#error "Incorrectly defined __opencl_c_ext_fp32_local_atomic_min_max" | ||
301 | +#endif | ||
302 | + | ||
303 | +#else | ||
304 | +#ifdef __opencl_c_ext_fp16_global_atomic_load_store | ||
305 | +#error "Incorrectly __opencl_c_ext_fp16_global_atomic_load_store defined" | ||
306 | +#endif | ||
307 | +#ifdef __opencl_c_ext_fp16_local_atomic_load_store | ||
308 | +#error "Incorrectly __opencl_c_ext_fp16_local_atomic_load_store defined" | ||
309 | +#endif | ||
310 | +#ifdef __opencl_c_ext_fp16_global_atomic_add | ||
311 | +#error "Incorrectly __opencl_c_ext_fp16_global_atomic_add defined" | ||
312 | +#endif | ||
313 | +#ifdef __opencl_c_ext_fp32_global_atomic_add | ||
314 | +#error "Incorrectly __opencl_c_ext_fp32_global_atomic_add defined" | ||
315 | +#endif | ||
316 | +#ifdef __opencl_c_ext_fp64_global_atomic_add | ||
317 | +#error "Incorrectly __opencl_c_ext_fp64_global_atomic_add defined" | ||
318 | +#endif | ||
319 | +#ifdef __opencl_c_ext_fp16_local_atomic_add | ||
320 | +#error "Incorrectly __opencl_c_ext_fp16_local_atomic_add defined" | ||
321 | +#endif | ||
322 | +#ifdef __opencl_c_ext_fp32_local_atomic_add | ||
323 | +#error "Incorrectly __opencl_c_ext_fp32_local_atomic_add defined" | ||
324 | +#endif | ||
325 | +#ifdef __opencl_c_ext_fp64_local_atomic_add | ||
326 | +#error "Incorrectly __opencl_c_ext_fp64_local_atomic_add defined" | ||
327 | +#endif | ||
328 | +#ifdef __opencl_c_ext_fp16_global_atomic_min_max | ||
329 | +#error "Incorrectly __opencl_c_ext_fp16_global_atomic_min_max defined" | ||
330 | +#endif | ||
331 | +#ifdef __opencl_c_ext_fp32_global_atomic_min_max | ||
332 | +#error "Incorrectly __opencl_c_ext_fp32_global_atomic_min_max defined" | ||
333 | +#endif | ||
334 | +#ifdef __opencl_c_ext_fp64_global_atomic_min_max | ||
335 | +#error "Incorrectly __opencl_c_ext_fp64_global_atomic_min_max defined" | ||
336 | +#endif | ||
337 | +#ifdef __opencl_c_ext_fp16_local_atomic_min_max | ||
338 | +#error "Incorrectly __opencl_c_ext_fp16_local_atomic_min_max defined" | ||
339 | +#endif | ||
340 | +#ifdef __opencl_c_ext_fp32_local_atomic_min_max | ||
341 | +#error "Incorrectly __opencl_c_ext_fp32_local_atomic_min_max defined" | ||
342 | +#endif | ||
343 | +#ifdef __opencl_c_ext_fp64_local_atomic_min_max | ||
344 | +#error "Incorrectly __opencl_c_ext_fp64_local_atomic_min_max defined" | ||
345 | +#endif | ||
346 | + | ||
347 | +#endif //(defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200) | ||
348 | + | ||
349 | +#endif // defined(__SPIR__) | ||
350 | + | ||
351 | -- | ||
352 | 2.17.1 | ||
353 | |||
diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/llvm-project-source.bbappend b/dynamic-layers/clang-layer/recipes-devtools/clang/llvm-project-source.bbappend index ac34321c..84192714 100644 --- a/dynamic-layers/clang-layer/recipes-devtools/clang/llvm-project-source.bbappend +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/llvm-project-source.bbappend | |||
@@ -1,7 +1,7 @@ | |||
1 | FILESEXTRAPATHS:prepend:intel-x86-common := "${THISDIR}/files:" | 1 | FILESEXTRAPATHS:prepend:intel-x86-common := "${THISDIR}/files:" |
2 | 2 | ||
3 | SPIRV10_SRCREV = "fe4d6b767363a1995ccbfca27f79efb10dcfe110" | 3 | SPIRV10_SRCREV = "fe4d6b767363a1995ccbfca27f79efb10dcfe110" |
4 | SPIRV11_SRCREV = "2a8c1e6c9778deaa720a23e08c293006dc5d56fd" | 4 | SPIRV11_SRCREV = "ca3a50e6e3193e399d26446d4f74a90e2a531f3a" |
5 | 5 | ||
6 | SPIRV_SRCREV = "${@bb.utils.contains('LLVMVERSION', '10.0.1', '${SPIRV10_SRCREV}', '${SPIRV11_SRCREV}', d)}" | 6 | SPIRV_SRCREV = "${@bb.utils.contains('LLVMVERSION', '10.0.1', '${SPIRV10_SRCREV}', '${SPIRV11_SRCREV}', d)}" |
7 | 7 | ||
@@ -21,10 +21,14 @@ SRC_URI_LLVM10_PATCHES = " \ | |||
21 | " | 21 | " |
22 | 22 | ||
23 | SRC_URI_LLVM11_PATCHES = " \ | 23 | SRC_URI_LLVM11_PATCHES = " \ |
24 | file://llvm11-skip-building-tests.patch;patchdir=llvm/projects/llvm-spirv \ | 24 | file://llvm11-0001-llvm-spirv-skip-building-tests.patch;patchdir=llvm/projects/llvm-spirv \ |
25 | file://llvm11-OpenCL-3.0-support.patch \ | 25 | file://llvm11-0002-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch;patchdir=llvm/projects/llvm-spirv \ |
26 | file://0001-Memory-leak-fix-for-Managed-Static-Mutex.patch \ | 26 | file://llvm11-0001-OpenCL-3.0-support.patch \ |
27 | file://llvm11-Remove-repo-name-in-LLVM-IR.patch \ | 27 | file://llvm11-0002-Memory-leak-fix-for-Managed-Static-Mutex.patch \ |
28 | file://llvm11-0003-Remove-repo-name-in-LLVM-IR.patch \ | ||
29 | file://llvm11-0004-Remove-__IMAGE_SUPPORT__-macro-for-SPIR-since-SPIR-d.patch \ | ||
30 | file://llvm11-0005-Avoid-calling-ParseCommandLineOptions-in-BackendUtil.patch \ | ||
31 | file://llvm11-0006-OpenCL-support-cl_ext_float_atomics.patch \ | ||
28 | " | 32 | " |
29 | SRC_URI_LLVM12_PATCHES = " \ | 33 | SRC_URI_LLVM12_PATCHES = " \ |
30 | file://0001-Remove-__IMAGE_SUPPORT__-macro-for-SPIR.patch \ | 34 | file://0001-Remove-__IMAGE_SUPPORT__-macro-for-SPIR.patch \ |