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