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>
2025-07-19 21:09:03 +02:00 · 2021-08-20 09:45:24 +08:00 · 2021-08-20 09:45:24 +08:00 · 5448b52ae6
commit 5448b52ae6
parent 109fe96793
9 changed files with 913 additions and 23 deletions
--- a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0001-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 @@
-From d767afa79d1c8153081eac1ef33e348cadbea5bb Mon Sep 17 00:00:00 2001
+From 36d87f69fee9c3d3f399f8e4027ab707ad050e80 Mon Sep 17 00:00:00 2001
 From: Anton Zabaznov <anton.zabaznov@intel.com>
 Date: Tue, 22 Sep 2020 19:03:50 +0300
-Subject: [PATCH] OpenCL 3.0 support
+Subject: [PATCH 1/6] OpenCL 3.0 support

 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]
 Signed-off-by: Anton Zabaznov <anton.zabaznov@intel.com>
 Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com>
-
-
 ---
 clang/include/clang/Basic/Builtins.def        |   65 +-
 clang/include/clang/Basic/Builtins.h          |   13 +-
--- a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0001-llvm-spirv-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 @@
-From d362652617c5e840089273df0c6623a9745c92a2 Mon Sep 17 00:00:00 2001
+From 6690d77f9007ce82984dc1b6ae12585cb3e04785 Mon Sep 17 00:00:00 2001
 From: Naveen Saini <naveen.kumar.saini@intel.com>
 Date: Wed, 21 Aug 2019 14:35:31 +0800
-Subject: [PATCH] llvm-spirv: skip building tests
+Subject: [PATCH 1/2] llvm-spirv: skip building tests

 Some of these need clang to be built and since we're building this in-tree,
 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>
 1 file changed, 10 deletions(-)

 diff --git a/CMakeLists.txt b/CMakeLists.txt
-index ecebb4cb..578ca602 100644
+index ec61fb95..d723c0a5 100644
 --- a/CMakeLists.txt
 +++ b/CMakeLists.txt
-@@ -25,13 +25,6 @@ if(LLVM_SPIRV_BUILD_EXTERNAL)
+@@ -26,13 +26,6 @@ if(LLVM_SPIRV_BUILD_EXTERNAL)
   set(CMAKE_CXX_STANDARD 14)
   set(CMAKE_CXX_STANDARD_REQUIRED ON)
 
@ -36,7 +36,7 @@ index ecebb4cb..578ca602 100644
   find_package(LLVM ${BASE_LLVM_VERSION} REQUIRED
     COMPONENTS
       Analysis
-@@ -62,9 +55,6 @@ set(LLVM_SPIRV_INCLUDE_DIRS ${CMAKE_CURRENT_SOURCE_DIR}/include)
+@@ -65,9 +58,6 @@ set(LLVM_SPIRV_INCLUDE_DIRS ${CMAKE_CURRENT_SOURCE_DIR}/include)
 
 add_subdirectory(lib/SPIRV)
 add_subdirectory(tools/llvm-spirv)
@ -47,5 +47,5 @@ index ecebb4cb..578ca602 100644
 install(
   FILES
 -- 
-2.26.2
+2.17.1

--- 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
@ -0,0 +1,433 @@
+From 8e12d8fb3cdbdafca73fe8ed4f0cde773b1788b4 Mon Sep 17 00:00:00 2001
+From: haonanya <haonan.yang@intel.com>
+Date: Wed, 28 Jul 2021 11:43:20 +0800
+Subject: [PATCH 2/2] Add support for cl_ext_float_atomics in SPIRVWriter
+
+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]
+
+Signed-off-by: haonanya <haonan.yang@intel.com>
+Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com>
+---
+ lib/SPIRV/OCLToSPIRV.cpp                | 80 +++++++++++++++++++++++--
+ lib/SPIRV/OCLUtil.cpp                   | 26 --------
+ lib/SPIRV/OCLUtil.h                     |  4 --
+ test/negative/InvalidAtomicBuiltins.cl  | 12 +---
+ test/transcoding/AtomicFAddEXTForOCL.ll | 64 ++++++++++++++++++++
+ test/transcoding/AtomicFMaxEXTForOCL.ll | 64 ++++++++++++++++++++
+ test/transcoding/AtomicFMinEXTForOCL.ll | 64 ++++++++++++++++++++
+ 7 files changed, 269 insertions(+), 45 deletions(-)
+ create mode 100644 test/transcoding/AtomicFAddEXTForOCL.ll
+ create mode 100644 test/transcoding/AtomicFMaxEXTForOCL.ll
+ create mode 100644 test/transcoding/AtomicFMinEXTForOCL.ll
+
+diff --git a/lib/SPIRV/OCLToSPIRV.cpp b/lib/SPIRV/OCLToSPIRV.cpp
+index 04d51586..f00f5f7b 100644
+--- a/lib/SPIRV/OCLToSPIRV.cpp
+++ b/lib/SPIRV/OCLToSPIRV.cpp
+@@ -421,10 +421,63 @@ void OCLToSPIRVBase::visitCallInst(CallInst &CI) {
+   if (DemangledName.find(kOCLBuiltinName::AtomicPrefix) == 0 ||
+       DemangledName.find(kOCLBuiltinName::AtomPrefix) == 0) {
+ 
+-    // Compute atomic builtins do not support floating types.
+-    if (CI.getType()->isFloatingPointTy() &&
+-        isComputeAtomicOCLBuiltin(DemangledName))
+-      return;
+    // Compute "atom" prefixed builtins do not support floating types.
+    if (CI.getType()->isFloatingPointTy()) {
+      if (DemangledName.find(kOCLBuiltinName::AtomPrefix) == 0)
+        return;
+      // handle functions which are "atomic_" prefixed.
+      StringRef Stem = DemangledName;
+      Stem = Stem.drop_front(strlen("atomic_"));
+      // FP-typed atomic_{add, sub, inc, dec, exchange, min, max, or, and, xor,
+      // fetch_or, fetch_xor, fetch_and, fetch_or_explicit, fetch_xor_explicit,
+      // fetch_and_explicit} should be identified as function call
+      bool IsFunctionCall = llvm::StringSwitch<bool>(Stem)
+                                .Case("add", true)
+                                .Case("sub", true)
+                                .Case("inc", true)
+                                .Case("dec", true)
+                                .Case("cmpxchg", true)
+                                .Case("min", true)
+                                .Case("max", true)
+                                .Case("or", true)
+                                .Case("xor", true)
+                                .Case("and", true)
+                                .Case("fetch_or", true)
+                                .Case("fetch_and", true)
+                                .Case("fetch_xor", true)
+                                .Case("fetch_or_explicit", true)
+                                .Case("fetch_xor_explicit", true)
+                                .Case("fetch_and_explicit", true)
+                                .Default(false);
+      if (IsFunctionCall)
+        return;
+      if (F->arg_size() != 2) {
+        IsFunctionCall = llvm::StringSwitch<bool>(Stem)
+                             .Case("exchange", true)
+                             .Case("fetch_add", true)
+                             .Case("fetch_sub", true)
+                             .Case("fetch_min", true)
+                             .Case("fetch_max", true)
+                             .Case("load", true)
+                             .Case("store", true)
+                             .Default(false);
+        if (IsFunctionCall)
+          return;
+      }
+      if (F->arg_size() != 3 && F->arg_size() != 4) {
+        IsFunctionCall = llvm::StringSwitch<bool>(Stem)
+                             .Case("exchange_explicit", true)
+                             .Case("fetch_add_explicit", true)
+                             .Case("fetch_sub_explicit", true)
+                             .Case("fetch_min_explicit", true)
+                             .Case("fetch_max_explicit", true)
+                             .Case("load_explicit", true)
+                             .Case("store_explicit", true)
+                             .Default(false);
+        if (IsFunctionCall)
+          return;
+      }
+    }
+ 
+     auto PCI = &CI;
+     if (DemangledName == kOCLBuiltinName::AtomicInit) {
+@@ -839,7 +892,7 @@ void OCLToSPIRVBase::transAtomicBuiltin(CallInst *CI,
+   AttributeList Attrs = CI->getCalledFunction()->getAttributes();
+   mutateCallInstSPIRV(
+       M, CI,
+-      [=](CallInst *CI, std::vector<Value *> &Args) {
+      [=](CallInst *CI, std::vector<Value *> &Args) -> std::string {
+         Info.PostProc(Args);
+         // Order of args in OCL20:
+         // object, 0-2 other args, 1-2 order, scope
+@@ -868,7 +921,22 @@ void OCLToSPIRVBase::transAtomicBuiltin(CallInst *CI,
+           std::rotate(Args.begin() + 2, Args.begin() + OrderIdx,
+                       Args.end() - Offset);
+         }
+-        return getSPIRVFuncName(OCLSPIRVBuiltinMap::map(Info.UniqName));
+
+        llvm::Type* AtomicBuiltinsReturnType =
+            CI->getCalledFunction()->getReturnType();
+        auto IsFPType = [](llvm::Type *ReturnType) {
+          return ReturnType->isHalfTy() || ReturnType->isFloatTy() ||
+                 ReturnType->isDoubleTy();
+        };
+        auto SPIRVFunctionName =
+            getSPIRVFuncName(OCLSPIRVBuiltinMap::map(Info.UniqName));
+        if (!IsFPType(AtomicBuiltinsReturnType))
+          return SPIRVFunctionName;
+        // Translate FP-typed atomic builtins.
+        return llvm::StringSwitch<std::string>(SPIRVFunctionName)
+            .Case("__spirv_AtomicIAdd", "__spirv_AtomicFAddEXT")
+            .Case("__spirv_AtomicSMax", "__spirv_AtomicFMaxEXT")
+            .Case("__spirv_AtomicSMin", "__spirv_AtomicFMinEXT");
+       },
+       &Attrs);
+ }
+diff --git a/lib/SPIRV/OCLUtil.cpp b/lib/SPIRV/OCLUtil.cpp
+index 2de3f152..85155e39 100644
+--- a/lib/SPIRV/OCLUtil.cpp
+++ b/lib/SPIRV/OCLUtil.cpp
+@@ -662,32 +662,6 @@ size_t getSPIRVAtomicBuiltinNumMemoryOrderArgs(Op OC) {
+   return 1;
+ }
+ 
+-bool isComputeAtomicOCLBuiltin(StringRef DemangledName) {
+-  if (!DemangledName.startswith(kOCLBuiltinName::AtomicPrefix) &&
+-      !DemangledName.startswith(kOCLBuiltinName::AtomPrefix))
+-    return false;
+-
+-  return llvm::StringSwitch<bool>(DemangledName)
+-      .EndsWith("add", true)
+-      .EndsWith("sub", true)
+-      .EndsWith("inc", true)
+-      .EndsWith("dec", true)
+-      .EndsWith("cmpxchg", true)
+-      .EndsWith("min", true)
+-      .EndsWith("max", true)
+-      .EndsWith("and", true)
+-      .EndsWith("or", true)
+-      .EndsWith("xor", true)
+-      .EndsWith("add_explicit", true)
+-      .EndsWith("sub_explicit", true)
+-      .EndsWith("or_explicit", true)
+-      .EndsWith("xor_explicit", true)
+-      .EndsWith("and_explicit", true)
+-      .EndsWith("min_explicit", true)
+-      .EndsWith("max_explicit", true)
+-      .Default(false);
+-}
+-
+ BarrierLiterals getBarrierLiterals(CallInst *CI) {
+   auto N = CI->getNumArgOperands();
+   assert(N == 1 || N == 2);
+diff --git a/lib/SPIRV/OCLUtil.h b/lib/SPIRV/OCLUtil.h
+index 4c05c672..c8577e9b 100644
+--- a/lib/SPIRV/OCLUtil.h
+++ b/lib/SPIRV/OCLUtil.h
+@@ -394,10 +394,6 @@ size_t getAtomicBuiltinNumMemoryOrderArgs(StringRef Name);
+ /// Get number of memory order arguments for spirv atomic builtin function.
+ size_t getSPIRVAtomicBuiltinNumMemoryOrderArgs(Op OC);
+ 
+-/// Return true for OpenCL builtins which do compute operations
+-/// (like add, sub, min, max, inc, dec, ...) atomically
+-bool isComputeAtomicOCLBuiltin(StringRef DemangledName);
+-
+ /// Get OCL version from metadata opencl.ocl.version.
+ /// \param AllowMulti Allows multiple operands if true.
+ /// \return OCL version encoded as Major*10^5+Minor*10^3+Rev,
+diff --git a/test/negative/InvalidAtomicBuiltins.cl b/test/negative/InvalidAtomicBuiltins.cl
+index b8ec5b89..23dcc4e3 100644
+--- a/test/negative/InvalidAtomicBuiltins.cl
+++ b/test/negative/InvalidAtomicBuiltins.cl
+@@ -1,7 +1,9 @@
+ // Check that translator doesn't generate atomic instructions for atomic builtins
+ // which are not defined in the spec.
+ 
+-// RUN: %clang_cc1 -triple spir -O1 -cl-std=cl2.0 -fdeclare-opencl-builtins -finclude-default-header %s -emit-llvm-bc -o %t.bc
+// To drop `fdeclare-opencl-builtins` option, since FP-typed atomic function
+// TableGen definitions have not been introduced.
+// RUN: %clang_cc1 -triple spir -O1 -cl-std=cl2.0 -finclude-default-header %s -emit-llvm-bc -o %t.bc
+ // RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s
+ // RUN: llvm-spirv %t.bc -o %t.spv
+ // RUN: spirv-val %t.spv
+@@ -41,13 +43,9 @@ float __attribute__((overloadable)) atomic_fetch_xor(volatile generic atomic_flo
+ double __attribute__((overloadable)) atomic_fetch_and(volatile generic atomic_double *object, double operand, memory_order order);
+ double __attribute__((overloadable)) atomic_fetch_max(volatile generic atomic_double *object, double operand, memory_order order);
+ double __attribute__((overloadable)) atomic_fetch_min(volatile generic atomic_double *object, double operand, memory_order order);
+-float __attribute__((overloadable)) atomic_fetch_add_explicit(volatile generic atomic_float *object, float operand, memory_order order);
+-float __attribute__((overloadable)) atomic_fetch_sub_explicit(volatile generic atomic_float *object, float operand, memory_order order);
+ float __attribute__((overloadable)) atomic_fetch_or_explicit(volatile generic atomic_float *object, float operand, memory_order order);
+ float __attribute__((overloadable)) atomic_fetch_xor_explicit(volatile generic atomic_float *object, float operand, memory_order order);
+ double __attribute__((overloadable)) atomic_fetch_and_explicit(volatile generic atomic_double *object, double operand, memory_order order);
+-double __attribute__((overloadable)) atomic_fetch_max_explicit(volatile generic atomic_double *object, double operand, memory_order order);
+-double __attribute__((overloadable)) atomic_fetch_min_explicit(volatile generic atomic_double *object, double operand, memory_order order);
+ 
+ __kernel void test_atomic_fn(volatile __global float *p,
+                              volatile __global double *pp,
+@@ -86,11 +84,7 @@ __kernel void test_atomic_fn(volatile __global float *p,
+     d = atomic_fetch_and(pp, val, order);
+     d = atomic_fetch_min(pp, val, order);
+     d = atomic_fetch_max(pp, val, order);
+-    f = atomic_fetch_add_explicit(p, val, order);
+-    f = atomic_fetch_sub_explicit(p, val, order);
+     f = atomic_fetch_or_explicit(p, val, order);
+     f = atomic_fetch_xor_explicit(p, val, order);
+     d = atomic_fetch_and_explicit(pp, val, order);
+-    d = atomic_fetch_min_explicit(pp, val, order);
+-    d = atomic_fetch_max_explicit(pp, val, order);
+ }
+diff --git a/test/transcoding/AtomicFAddEXTForOCL.ll b/test/transcoding/AtomicFAddEXTForOCL.ll
+new file mode 100644
+index 00000000..fb146fb9
+--- /dev/null
+++ b/test/transcoding/AtomicFAddEXTForOCL.ll
+@@ -0,0 +1,64 @@
+; RUN: llvm-as %s -o %t.bc
+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_add -o %t.spv
+; RUN: spirv-val %t.spv
+; RUN: llvm-spirv -to-text %t.spv -o %t.spt
+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
+
+; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc
+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20
+
+; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc
+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV
+
+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"
+target triple = "spir-unknown-unknown"
+
+; CHECK-SPIRV: Capability AtomicFloat32AddEXT
+; CHECK-SPIRV: Capability AtomicFloat64AddEXT
+; CHECK-SPIRV: Extension "SPV_EXT_shader_atomic_float_add"
+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32
+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64
+
+
+; Function Attrs: convergent norecurse nounwind
+define dso_local spir_func void @test_atomic_float(float addrspace(1)* %a) local_unnamed_addr #0 {
+entry:
+  ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_32]]
+  ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_add_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}})
+  ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+fiif]]({{.*}})
+  %call = tail call spir_func float @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2
+  ret void
+}
+
+; Function Attrs: convergent
+declare spir_func float @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1
+; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
+
+; Function Attrs: convergent norecurse nounwind
+define dso_local spir_func void @test_atomic_double(double addrspace(1)* %a) local_unnamed_addr #0 {
+entry:
+  ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_64]]
+  ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_add_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}})
+  ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+diid]]({{.*}})
+  %call = tail call spir_func double @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2
+  ret void
+}
+; Function Attrs: convergent
+declare spir_func double @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)*, double, i32) local_unnamed_addr #1
+; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
+
+; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
+; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
+
+attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #2 = { convergent nounwind }
+
+!llvm.module.flags = !{!0}
+!opencl.ocl.version = !{!1}
+!opencl.spir.version = !{!1}
+!llvm.ident = !{!2}
+
+!0 = !{i32 1, !"wchar_size", i32 4}
+!1 = !{i32 2, i32 0}
+!2 = !{!"clang version 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"}
+diff --git a/test/transcoding/AtomicFMaxEXTForOCL.ll b/test/transcoding/AtomicFMaxEXTForOCL.ll
+new file mode 100644
+index 00000000..1f2530d9
+--- /dev/null
+++ b/test/transcoding/AtomicFMaxEXTForOCL.ll
+@@ -0,0 +1,64 @@
+; RUN: llvm-as %s -o %t.bc
+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_min_max -o %t.spv
+; RUN: spirv-val %t.spv
+; RUN: llvm-spirv -to-text %t.spv -o %t.spt
+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
+
+; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc
+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20
+
+; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc
+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV
+
+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"
+target triple = "spir-unknown-unknown"
+
+; CHECK-SPIRV: Capability AtomicFloat32MinMaxEXT
+; CHECK-SPIRV: Capability AtomicFloat64MinMaxEXT
+; CHECK-SPIRV: Extension "SPV_EXT_shader_atomic_float_min_max"
+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32
+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64
+
+; Function Attrs: convergent norecurse nounwind
+define dso_local spir_func void @test_float(float addrspace(1)* %a) local_unnamed_addr #0 {
+entry:
+  ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_32]]
+  ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_max_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}})
+  ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+fiif]]({{.*}})
+  %call = tail call spir_func float @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2
+  ret void
+}
+
+; Function Attrs: convergent
+declare spir_func float @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1
+; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
+
+; Function Attrs: convergent norecurse nounwind
+define dso_local spir_func void @test_double(double addrspace(1)* %a) local_unnamed_addr #0 {
+entry:
+  ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_64]]
+  ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_max_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}})
+  ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+diid]]({{.*}})
+  %call = tail call spir_func double @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2
+  ret void
+}
+
+; Function Attrs: convergent
+declare spir_func double @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)*, double, i32) local_unnamed_addr #1
+; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
+
+; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
+; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
+
+attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #2 = { convergent nounwind }
+
+!llvm.module.flags = !{!0}
+!opencl.ocl.version = !{!1}
+!opencl.spir.version = !{!1}
+!llvm.ident = !{!2}
+
+!0 = !{i32 1, !"wchar_size", i32 4}
+!1 = !{i32 2, i32 0}
+!2 = !{!"clang version 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"}
+diff --git a/test/transcoding/AtomicFMinEXTForOCL.ll b/test/transcoding/AtomicFMinEXTForOCL.ll
+new file mode 100644
+index 00000000..6196b0f8
+--- /dev/null
+++ b/test/transcoding/AtomicFMinEXTForOCL.ll
+@@ -0,0 +1,64 @@
+; RUN: llvm-as %s -o %t.bc
+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_min_max -o %t.spv
+; RUN: spirv-val %t.spv
+; RUN: llvm-spirv -to-text %t.spv -o %t.spt
+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
+
+; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc
+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20
+
+; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc
+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV
+
+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"
+target triple = "spir-unknown-unknown"
+
+; CHECK-SPIRV: Capability AtomicFloat32MinMaxEXT
+; CHECK-SPIRV: Capability AtomicFloat64MinMaxEXT
+; CHECK-SPIRV: Extension "SPV_EXT_shader_atomic_float_min_max"
+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32
+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64
+
+; Function Attrs: convergent norecurse nounwind
+define dso_local spir_func void @test_float(float addrspace(1)* %a) local_unnamed_addr #0 {
+entry:
+  ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_32]]
+  ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_min_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}})
+  ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+fiif]]({{.*}})
+  %call = tail call spir_func float @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2
+  ret void
+}
+
+; Function Attrs: convergent
+declare spir_func float @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1
+; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
+
+; Function Attrs: convergent norecurse nounwind
+define dso_local spir_func void @test_double(double addrspace(1)* %a) local_unnamed_addr #0 {
+entry:
+  ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_64]]
+  ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_min_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}})
+  ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+diid]]({{.*}})
+  %call = tail call spir_func double @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2
+  ret void
+}
+
+; Function Attrs: convergent
+declare spir_func double @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)*, double, i32) local_unnamed_addr #1
+; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
+
+; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
+; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
+
+attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #2 = { convergent nounwind }
+
+!llvm.module.flags = !{!0}
+!opencl.ocl.version = !{!1}
+!opencl.spir.version = !{!1}
+!llvm.ident = !{!2}
+
+!0 = !{i32 1, !"wchar_size", i32 4}
+!1 = !{i32 2, i32 0}
+!2 = !{!"clang version 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"}
+-- 
+2.17.1
+
--- a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0002-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 @@
-From c86c43b70e029b102543e8a85d269cbeb5c00279 Mon Sep 17 00:00:00 2001
+From ef27f1f99ad661c9604b7ff10efb1122466c508b Mon Sep 17 00:00:00 2001
 From: juanrod2 <>
 Date: Tue, 22 Dec 2020 08:33:08 +0800
-Subject: [PATCH] Memory leak fix for Managed Static Mutex
+Subject: [PATCH 2/6] Memory leak fix for Managed Static Mutex

 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]

@ -31,5 +31,5 @@ index 053493f72fb5..6571580ccecf 100644
 +  ManagedStaticMutex = nullptr;
 }
 -- 
-2.29.2
+2.17.1

--- a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0003-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 @@
-From ff0a6da84b94c16c4519c649f1f7bed3cdf89bbb Mon Sep 17 00:00:00 2001
+From a71ab6fb04b918b856f1dd802cfdb4a7ccd53290 Mon Sep 17 00:00:00 2001
 From: Feng Zou <feng.zou@intel.com>
 Date: Tue, 20 Oct 2020 11:29:04 +0800
-Subject: [PATCH] Remove repo name in LLVM IR
+Subject: [PATCH 3/6] Remove repo name in LLVM IR

 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]
 Signed-off-by: Feng Zou <feng.zou@intel.com>
 Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com>
-
 ---
 llvm/cmake/modules/VersionFromVCS.cmake | 23 ++++++++++++-----------
 1 file changed, 12 insertions(+), 11 deletions(-)

 diff --git a/llvm/cmake/modules/VersionFromVCS.cmake b/llvm/cmake/modules/VersionFromVCS.cmake
-index 18edbeabe3e..2d965263478 100644
+index 18edbeabe3e4..2d9652634787 100644
 --- a/llvm/cmake/modules/VersionFromVCS.cmake
 +++ b/llvm/cmake/modules/VersionFromVCS.cmake
@@ -33,17 +33,18 @@ function(get_source_info path revision repository)
@ -46,5 +45,5 @@ index 18edbeabe3e..2d965263478 100644
   else()
     message(WARNING "Git not found. Version cannot be determined.")
 -- 
-2.18.1
+2.17.1

--- 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
@ -0,0 +1,51 @@
+From 546d9089fe5e21cccc671a0a89555cd4d5f8c817 Mon Sep 17 00:00:00 2001
+From: Naveen Saini <naveen.kumar.saini@intel.com>
+Date: Thu, 19 Aug 2021 15:52:24 +0800
+Subject: [PATCH 4/6] Remove __IMAGE_SUPPORT__ macro for SPIR since SPIR
+ doesn't require image support
+
+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]
+
+Signed-off-by: haonanya <haonan.yang@intel.com>
+Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com>
+---
+ clang/lib/Frontend/InitPreprocessor.cpp     | 3 ---
+ clang/test/Preprocessor/predefined-macros.c | 2 --
+ 2 files changed, 5 deletions(-)
+
+diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp
+index 5bb489c11909..cf3b48cb65d2 100644
+--- a/clang/lib/Frontend/InitPreprocessor.cpp
+++ b/clang/lib/Frontend/InitPreprocessor.cpp
+@@ -1115,9 +1115,6 @@ static void InitializePredefinedMacros(const TargetInfo &TI,
+   if (TI.getSupportedOpenCLOpts().isSupported(#Ext))                           \
+     Builder.defineMacro(#Ext);
+ #include "clang/Basic/OpenCLExtensions.def"
+-
+-    if (TI.getTriple().isSPIR())
+-      Builder.defineMacro("__IMAGE_SUPPORT__");
+   }
+ 
+   if (TI.hasInt128Type() && LangOpts.CPlusPlus && LangOpts.GNUMode) {
+diff --git a/clang/test/Preprocessor/predefined-macros.c b/clang/test/Preprocessor/predefined-macros.c
+index 6c80517ec4d4..b5e5d7e2d546 100644
+--- a/clang/test/Preprocessor/predefined-macros.c
+++ b/clang/test/Preprocessor/predefined-macros.c
+@@ -186,14 +186,12 @@
+ 
+ // RUN: %clang_cc1 %s -E -dM -o - -x cl -triple spir-unknown-unknown \
+ // RUN:   | FileCheck -match-full-lines %s --check-prefix=CHECK-SPIR
+-// CHECK-SPIR-DAG: #define __IMAGE_SUPPORT__ 1
+ // CHECK-SPIR-DAG: #define __SPIR__ 1
+ // CHECK-SPIR-DAG: #define __SPIR32__ 1
+ // CHECK-SPIR-NOT: #define __SPIR64__ 1
+ 
+ // RUN: %clang_cc1 %s -E -dM -o - -x cl -triple spir64-unknown-unknown \
+ // RUN:   | FileCheck -match-full-lines %s --check-prefix=CHECK-SPIR64
+-// CHECK-SPIR64-DAG: #define __IMAGE_SUPPORT__ 1
+ // CHECK-SPIR64-DAG: #define __SPIR__ 1
+ // CHECK-SPIR64-DAG: #define __SPIR64__ 1
+ // CHECK-SPIR64-NOT: #define __SPIR32__ 1
+-- 
+2.17.1
+
--- 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
@ -0,0 +1,52 @@
+From 747e48959e18ac8b586078a82472a0799d12925c Mon Sep 17 00:00:00 2001
+From: Raphael Isemann <teemperor@gmail.com>
+Date: Thu, 1 Apr 2021 18:41:44 +0200
+Subject: [PATCH 5/6] Avoid calling ParseCommandLineOptions in BackendUtil if
+ possible
+
+Calling `ParseCommandLineOptions` should only be called from `main` as the
+CommandLine setup code isn't thread-safe. As BackendUtil is part of the
+generic Clang FrontendAction logic, a process which has several threads executing
+Clang FrontendActions will randomly crash in the unsafe setup code.
+
+This patch avoids calling the function unless either the debug-pass option or
+limit-float-precision option is set. Without these two options set the
+`ParseCommandLineOptions` call doesn't do anything beside parsing
+the command line `clang` which doesn't set any options.
+
+See also D99652 where LLDB received a workaround for this crash.
+
+Reviewed By: JDevlieghere
+
+Differential Revision: https://reviews.llvm.org/D99740
+
+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]
+
+Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com>
+---
+ clang/lib/CodeGen/BackendUtil.cpp | 8 ++++++++
+ 1 file changed, 8 insertions(+)
+
+diff --git a/clang/lib/CodeGen/BackendUtil.cpp b/clang/lib/CodeGen/BackendUtil.cpp
+index dce0940670a2..ab478090ed1c 100644
+--- a/clang/lib/CodeGen/BackendUtil.cpp
+++ b/clang/lib/CodeGen/BackendUtil.cpp
+@@ -797,7 +797,15 @@ static void setCommandLineOpts(const CodeGenOptions &CodeGenOpts) {
+     BackendArgs.push_back("-limit-float-precision");
+     BackendArgs.push_back(CodeGenOpts.LimitFloatPrecision.c_str());
+   }
+  // Check for the default "clang" invocation that won't set any cl::opt values.
+  // Skip trying to parse the command line invocation to avoid the issues
+  // described below.
+  if (BackendArgs.size() == 1)
+    return;
+   BackendArgs.push_back(nullptr);
+  // FIXME: The command line parser below is not thread-safe and shares a global
+  // state, so this call might crash or overwrite the options of another Clang
+  // instance in the same process.
+   llvm::cl::ParseCommandLineOptions(BackendArgs.size() - 1,
+                                     BackendArgs.data());
+ }
+-- 
+2.17.1
+
--- 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
@ -0,0 +1,353 @@
+From a1b924d76cdacfa3f9dbb79a9e3edddcd75f61ca Mon Sep 17 00:00:00 2001
+From: Naveen Saini <naveen.kumar.saini@intel.com>
+Date: Thu, 19 Aug 2021 16:06:33 +0800
+Subject: [PATCH 6/6] [OpenCL] support cl_ext_float_atomics
+
+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]
+
+Signed-off-by: haonanya <haonan.yang@intel.com>
+Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com>
+---
+ clang/lib/Headers/opencl-c-base.h     |  25 ++++
+ clang/lib/Headers/opencl-c.h          | 195 ++++++++++++++++++++++++++
+ clang/test/Headers/opencl-c-header.cl |  85 +++++++++++
+ 3 files changed, 305 insertions(+)
+
+diff --git a/clang/lib/Headers/opencl-c-base.h b/clang/lib/Headers/opencl-c-base.h
+index afa900ab24d9..9a3ee8529acf 100644
+--- a/clang/lib/Headers/opencl-c-base.h
+++ b/clang/lib/Headers/opencl-c-base.h
+@@ -62,6 +62,31 @@
+ #endif
+ #endif // defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ == CL_VERSION_2_0)
+ 
+#if (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200)
+// For SPIR all extensions are supported.
+#if defined(__SPIR__)
+#define cl_ext_float_atomics
+#ifdef cl_khr_fp16
+#define __opencl_c_ext_fp16_global_atomic_load_store 1
+#define __opencl_c_ext_fp16_local_atomic_load_store 1
+#define __opencl_c_ext_fp16_global_atomic_add 1
+#define __opencl_c_ext_fp16_local_atomic_add 1
+#define __opencl_c_ext_fp16_global_atomic_min_max 1
+#define __opencl_c_ext_fp16_local_atomic_min_max 1
+#endif
+#ifdef __opencl_c_fp64
+#define __opencl_c_ext_fp64_global_atomic_add 1
+#define __opencl_c_ext_fp64_local_atomic_add 1
+#define __opencl_c_ext_fp64_global_atomic_min_max 1
+#define __opencl_c_ext_fp64_local_atomic_min_max 1
+#endif
+#define __opencl_c_ext_fp32_global_atomic_add 1
+#define __opencl_c_ext_fp32_local_atomic_add 1
+#define __opencl_c_ext_fp32_global_atomic_min_max 1
+#define __opencl_c_ext_fp32_local_atomic_min_max 1
+#endif // defined(__SPIR__)
+#endif // (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200)
+
+ // built-in scalar data types:
+ 
+ /**
+diff --git a/clang/lib/Headers/opencl-c.h b/clang/lib/Headers/opencl-c.h
+index 67d900eb1c3d..bda0f5c6df80 100644
+--- a/clang/lib/Headers/opencl-c.h
+++ b/clang/lib/Headers/opencl-c.h
+@@ -14354,6 +14354,201 @@ intptr_t __ovld atomic_fetch_max_explicit(
+        // defined(cl_khr_int64_extended_atomics)
+ #endif // (__OPENCL_C_VERSION__ >= CL_VERSION_3_0)
+ 
+#if defined(cl_ext_float_atomics)
+
+#if defined(__opencl_c_ext_fp32_global_atomic_min_max)
+float __ovld atomic_fetch_min(volatile __global atomic_float *object,
+                              float operand);
+float __ovld atomic_fetch_max(volatile __global atomic_float *object,
+                              float operand);
+float __ovld atomic_fetch_min_explicit(volatile __global atomic_float *object,
+                                       float operand, memory_order order);
+float __ovld atomic_fetch_max_explicit(volatile __global atomic_float *object,
+                                       float operand, memory_order order);
+float __ovld atomic_fetch_min_explicit(volatile __global atomic_float *object,
+                                       float operand, memory_order order,
+                                       memory_scope scope);
+float __ovld atomic_fetch_max_explicit(volatile __global atomic_float *object,
+                                       float operand, memory_order order,
+                                       memory_scope scope);
+#endif
+#if defined(__opencl_c_ext_fp32_local_atomic_min_max)
+float __ovld atomic_fetch_min(volatile __local atomic_float *object,
+                              float operand);
+float __ovld atomic_fetch_max(volatile __local atomic_float *object,
+                              float operand);
+float __ovld atomic_fetch_min_explicit(volatile __local atomic_float *object,
+                                       float operand, memory_order order);
+float __ovld atomic_fetch_max_explicit(volatile __local atomic_float *object,
+                                       float operand, memory_order order);
+float __ovld atomic_fetch_min_explicit(volatile __local atomic_float *object,
+                                       float operand, memory_order order,
+                                       memory_scope scope);
+float __ovld atomic_fetch_max_explicit(volatile __local atomic_float *object,
+                                       float operand, memory_order order,
+                                       memory_scope scope);
+#endif
+#if defined(__opencl_c_ext_fp32_global_atomic_min_max) ||                      \
+    defined(__opencl_c_ext_fp32_local_atomic_min_max)
+float __ovld atomic_fetch_min(volatile atomic_float *object, float operand);
+float __ovld atomic_fetch_max(volatile atomic_float *object, float operand);
+float __ovld atomic_fetch_min_explicit(volatile atomic_float *object,
+                                       float operand, memory_order order);
+float __ovld atomic_fetch_max_explicit(volatile atomic_float *object,
+                                       float operand, memory_order order);
+float __ovld atomic_fetch_min_explicit(volatile atomic_float *object,
+                                       float operand, memory_order order,
+                                       memory_scope scope);
+float __ovld atomic_fetch_max_explicit(volatile atomic_float *object,
+                                       float operand, memory_order order,
+                                       memory_scope scope);
+#endif
+#if defined(__opencl_c_ext_fp64_global_atomic_min_max)
+double __ovld atomic_fetch_min(volatile __global atomic_double *object,
+                               double operand);
+double __ovld atomic_fetch_max(volatile __global atomic_double *object,
+                               double operand);
+double __ovld atomic_fetch_min_explicit(volatile __global atomic_double *object,
+                                        double operand, memory_order order);
+double __ovld atomic_fetch_max_explicit(volatile __global atomic_double *object,
+                                        double operand, memory_order order);
+double __ovld atomic_fetch_min_explicit(volatile __global atomic_double *object,
+                                        double operand, memory_order order,
+                                        memory_scope scope);
+double __ovld atomic_fetch_max_explicit(volatile __global atomic_double *object,
+                                        double operand, memory_order order,
+                                        memory_scope scope);
+#endif
+#if defined(__opencl_c_ext_fp64_local_atomic_min_max)
+double __ovld atomic_fetch_min(volatile __local atomic_double *object,
+                               double operand);
+double __ovld atomic_fetch_max(volatile __local atomic_double *object,
+                               double operand);
+double __ovld atomic_fetch_min_explicit(volatile __local atomic_double *object,
+                                        double operand, memory_order order);
+double __ovld atomic_fetch_max_explicit(volatile __local atomic_double *object,
+                                        double operand, memory_order order);
+double __ovld atomic_fetch_min_explicit(volatile __local atomic_double *object,
+                                        double operand, memory_order order,
+                                        memory_scope scope);
+double __ovld atomic_fetch_max_explicit(volatile __local atomic_double *object,
+                                        double operand, memory_order order,
+                                        memory_scope scope);
+#endif
+#if defined(__opencl_c_ext_fp64_global_atomic_min_max) ||                      \
+    defined(__opencl_c_ext_fp64_local_atomic_min_max)
+double __ovld atomic_fetch_min(volatile atomic_double *object, double operand);
+double __ovld atomic_fetch_max(volatile atomic_double *object, double operand);
+double __ovld atomic_fetch_min_explicit(volatile atomic_double *object,
+                                        double operand, memory_order order);
+double __ovld atomic_fetch_max_explicit(volatile atomic_double *object,
+                                        double operand, memory_order order);
+double __ovld atomic_fetch_min_explicit(volatile atomic_double *object,
+                                        double operand, memory_order order,
+                                        memory_scope scope);
+double __ovld atomic_fetch_max_explicit(volatile atomic_double *object,
+                                        double operand, memory_order order,
+                                        memory_scope scope);
+#endif
+
+#if defined(__opencl_c_ext_fp32_global_atomic_add)
+float __ovld atomic_fetch_add(volatile __global atomic_float *object,
+                              float operand);
+float __ovld atomic_fetch_sub(volatile __global atomic_float *object,
+                              float operand);
+float __ovld atomic_fetch_add_explicit(volatile __global atomic_float *object,
+                                       float operand, memory_order order);
+float __ovld atomic_fetch_sub_explicit(volatile __global atomic_float *object,
+                                       float operand, memory_order order);
+float __ovld atomic_fetch_add_explicit(volatile __global atomic_float *object,
+                                       float operand, memory_order order,
+                                       memory_scope scope);
+float __ovld atomic_fetch_sub_explicit(volatile __global atomic_float *object,
+                                       float operand, memory_order order,
+                                       memory_scope scope);
+#endif
+#if defined(__opencl_c_ext_fp32_local_atomic_add)
+float __ovld atomic_fetch_add(volatile __local atomic_float *object,
+                              float operand);
+float __ovld atomic_fetch_sub(volatile __local atomic_float *object,
+                              float operand);
+float __ovld atomic_fetch_add_explicit(volatile __local atomic_float *object,
+                                       float operand, memory_order order);
+float __ovld atomic_fetch_sub_explicit(volatile __local atomic_float *object,
+                                       float operand, memory_order order);
+float __ovld atomic_fetch_add_explicit(volatile __local atomic_float *object,
+                                       float operand, memory_order order,
+                                       memory_scope scope);
+float __ovld atomic_fetch_sub_explicit(volatile __local atomic_float *object,
+                                       float operand, memory_order order,
+                                       memory_scope scope);
+#endif
+#if defined(__opencl_c_ext_fp32_global_atomic_add) ||                          \
+    defined(__opencl_c_ext_fp32_local_atomic_add)
+float __ovld atomic_fetch_add(volatile atomic_float *object, float operand);
+float __ovld atomic_fetch_sub(volatile atomic_float *object, float operand);
+float __ovld atomic_fetch_add_explicit(volatile atomic_float *object,
+                                       float operand, memory_order order);
+float __ovld atomic_fetch_sub_explicit(volatile atomic_float *object,
+                                       float operand, memory_order order);
+float __ovld atomic_fetch_add_explicit(volatile atomic_float *object,
+                                       float operand, memory_order order,
+                                       memory_scope scope);
+float __ovld atomic_fetch_sub_explicit(volatile atomic_float *object,
+                                       float operand, memory_order order,
+                                       memory_scope scope);
+#endif
+
+#if defined(__opencl_c_ext_fp64_global_atomic_add)
+double __ovld atomic_fetch_add(volatile __global atomic_double *object,
+                               double operand);
+double __ovld atomic_fetch_sub(volatile __global atomic_double *object,
+                               double operand);
+double __ovld atomic_fetch_add_explicit(volatile __global atomic_double *object,
+                                        double operand, memory_order order);
+double __ovld atomic_fetch_sub_explicit(volatile __global atomic_double *object,
+                                        double operand, memory_order order);
+double __ovld atomic_fetch_add_explicit(volatile __global atomic_double *object,
+                                        double operand, memory_order order,
+                                        memory_scope scope);
+double __ovld atomic_fetch_sub_explicit(volatile __global atomic_double *object,
+                                        double operand, memory_order order,
+                                        memory_scope scope);
+#endif
+#if defined(__opencl_c_ext_fp64_local_atomic_add)
+double __ovld atomic_fetch_add(volatile __local atomic_double *object,
+                               double operand);
+double __ovld atomic_fetch_sub(volatile __local atomic_double *object,
+                               double operand);
+double __ovld atomic_fetch_add_explicit(volatile __local atomic_double *object,
+                                        double operand, memory_order order);
+double __ovld atomic_fetch_sub_explicit(volatile __local atomic_double *object,
+                                        double operand, memory_order order);
+double __ovld atomic_fetch_add_explicit(volatile __local atomic_double *object,
+                                        double operand, memory_order order,
+                                        memory_scope scope);
+double __ovld atomic_fetch_sub_explicit(volatile __local atomic_double *object,
+                                        double operand, memory_order order,
+                                        memory_scope scope);
+#endif
+#if defined(__opencl_c_ext_fp64_global_atomic_add) ||                          \
+    defined(__opencl_c_ext_fp64_local_atomic_add)
+double __ovld atomic_fetch_add(volatile atomic_double *object, double operand);
+double __ovld atomic_fetch_sub(volatile atomic_double *object, double operand);
+double __ovld atomic_fetch_add_explicit(volatile atomic_double *object,
+                                        double operand, memory_order order);
+double __ovld atomic_fetch_sub_explicit(volatile atomic_double *object,
+                                        double operand, memory_order order);
+double __ovld atomic_fetch_add_explicit(volatile atomic_double *object,
+                                        double operand, memory_order order,
+                                        memory_scope scope);
+double __ovld atomic_fetch_sub_explicit(volatile atomic_double *object,
+                                        double operand, memory_order order,
+                                        memory_scope scope);
+#endif
+
+#endif // cl_ext_float_atomics
+
+ // atomic_store()
+ 
+ #if defined(__opencl_c_atomic_scope_device) &&                                 \
+diff --git a/clang/test/Headers/opencl-c-header.cl b/clang/test/Headers/opencl-c-header.cl
+index 2716076acdcf..6b3eca84e8b9 100644
+--- a/clang/test/Headers/opencl-c-header.cl
+++ b/clang/test/Headers/opencl-c-header.cl
+@@ -98,3 +98,88 @@ global atomic_int z = ATOMIC_VAR_INIT(99);
+ #pragma OPENCL EXTENSION cl_intel_planar_yuv : enable
+ 
+ // CHECK-MOD: Reading modules
+
+// For SPIR all extensions are supported.
+#if defined(__SPIR__)
+
+#if (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200)
+
+#if __opencl_c_ext_fp16_global_atomic_load_store != 1
+#error "Incorrectly defined __opencl_c_ext_fp16_global_atomic_load_store"
+#endif
+#if __opencl_c_ext_fp16_local_atomic_load_store != 1
+#error "Incorrectly defined __opencl_c_ext_fp16_local_atomic_load_store"
+#endif
+#if __opencl_c_ext_fp16_global_atomic_add != 1
+#error "Incorrectly defined __opencl_c_ext_fp16_global_atomic_add"
+#endif
+#if __opencl_c_ext_fp32_global_atomic_add != 1
+#error "Incorrectly defined __opencl_c_ext_fp32_global_atomic_add"
+#endif
+#if __opencl_c_ext_fp16_local_atomic_add != 1
+#error "Incorrectly defined __opencl_c_ext_fp16_local_atomic_add"
+#endif
+#if __opencl_c_ext_fp32_local_atomic_add != 1
+#error "Incorrectly defined __opencl_c_ext_fp32_local_atomic_add"
+#endif
+#if __opencl_c_ext_fp16_global_atomic_min_max != 1
+#error "Incorrectly defined __opencl_c_ext_fp16_global_atomic_min_max"
+#endif
+#if __opencl_c_ext_fp32_global_atomic_min_max != 1
+#error "Incorrectly defined __opencl_c_ext_fp32_global_atomic_min_max"
+#endif
+#if __opencl_c_ext_fp16_local_atomic_min_max != 1
+#error "Incorrectly defined __opencl_c_ext_fp16_local_atomic_min_max"
+#endif
+#if __opencl_c_ext_fp32_local_atomic_min_max != 1
+#error "Incorrectly defined __opencl_c_ext_fp32_local_atomic_min_max"
+#endif
+
+#else
+#ifdef __opencl_c_ext_fp16_global_atomic_load_store
+#error "Incorrectly __opencl_c_ext_fp16_global_atomic_load_store defined"
+#endif
+#ifdef __opencl_c_ext_fp16_local_atomic_load_store
+#error "Incorrectly __opencl_c_ext_fp16_local_atomic_load_store defined"
+#endif
+#ifdef __opencl_c_ext_fp16_global_atomic_add
+#error "Incorrectly __opencl_c_ext_fp16_global_atomic_add defined"
+#endif
+#ifdef __opencl_c_ext_fp32_global_atomic_add
+#error "Incorrectly __opencl_c_ext_fp32_global_atomic_add defined"
+#endif
+#ifdef __opencl_c_ext_fp64_global_atomic_add
+#error "Incorrectly __opencl_c_ext_fp64_global_atomic_add defined"
+#endif
+#ifdef __opencl_c_ext_fp16_local_atomic_add
+#error "Incorrectly __opencl_c_ext_fp16_local_atomic_add defined"
+#endif
+#ifdef __opencl_c_ext_fp32_local_atomic_add
+#error "Incorrectly __opencl_c_ext_fp32_local_atomic_add defined"
+#endif
+#ifdef __opencl_c_ext_fp64_local_atomic_add
+#error "Incorrectly __opencl_c_ext_fp64_local_atomic_add defined"
+#endif
+#ifdef __opencl_c_ext_fp16_global_atomic_min_max
+#error "Incorrectly __opencl_c_ext_fp16_global_atomic_min_max defined"
+#endif
+#ifdef __opencl_c_ext_fp32_global_atomic_min_max
+#error "Incorrectly __opencl_c_ext_fp32_global_atomic_min_max defined"
+#endif
+#ifdef __opencl_c_ext_fp64_global_atomic_min_max
+#error "Incorrectly __opencl_c_ext_fp64_global_atomic_min_max defined"
+#endif
+#ifdef __opencl_c_ext_fp16_local_atomic_min_max
+#error "Incorrectly __opencl_c_ext_fp16_local_atomic_min_max defined"
+#endif
+#ifdef __opencl_c_ext_fp32_local_atomic_min_max
+#error "Incorrectly __opencl_c_ext_fp32_local_atomic_min_max defined"
+#endif
+#ifdef __opencl_c_ext_fp64_local_atomic_min_max
+#error "Incorrectly __opencl_c_ext_fp64_local_atomic_min_max defined"
+#endif
+
+#endif //(defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200)
+
+#endif // defined(__SPIR__)
+
+-- 
+2.17.1
+
--- 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 @@
 FILESEXTRAPATHS:prepend:intel-x86-common := "${THISDIR}/files:"

 SPIRV10_SRCREV = "fe4d6b767363a1995ccbfca27f79efb10dcfe110"
-SPIRV11_SRCREV = "2a8c1e6c9778deaa720a23e08c293006dc5d56fd"
+SPIRV11_SRCREV = "ca3a50e6e3193e399d26446d4f74a90e2a531f3a"

 SPIRV_SRCREV = "${@bb.utils.contains('LLVMVERSION', '10.0.1', '${SPIRV10_SRCREV}', '${SPIRV11_SRCREV}', d)}"

@ -21,10 +21,14 @@ SRC_URI_LLVM10_PATCHES = " \
                   "

 SRC_URI_LLVM11_PATCHES = " \
-                   file://llvm11-skip-building-tests.patch;patchdir=llvm/projects/llvm-spirv \
-                   file://llvm11-OpenCL-3.0-support.patch \
-                   file://0001-Memory-leak-fix-for-Managed-Static-Mutex.patch \
-                   file://llvm11-Remove-repo-name-in-LLVM-IR.patch \
+                   file://llvm11-0001-llvm-spirv-skip-building-tests.patch;patchdir=llvm/projects/llvm-spirv \
+                   file://llvm11-0002-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch;patchdir=llvm/projects/llvm-spirv \
+                   file://llvm11-0001-OpenCL-3.0-support.patch \
+                   file://llvm11-0002-Memory-leak-fix-for-Managed-Static-Mutex.patch \
+                   file://llvm11-0003-Remove-repo-name-in-LLVM-IR.patch \
+                   file://llvm11-0004-Remove-__IMAGE_SUPPORT__-macro-for-SPIR-since-SPIR-d.patch \
+                   file://llvm11-0005-Avoid-calling-ParseCommandLineOptions-in-BackendUtil.patch \
+                   file://llvm11-0006-OpenCL-support-cl_ext_float_atomics.patch \
                   "
 SRC_URI_LLVM12_PATCHES = " \
                   file://0001-Remove-__IMAGE_SUPPORT__-macro-for-SPIR.patch \