llvm/10.0.0: apply opencl-clang recommend patches

https://github.com/intel/opencl-clang/tree/ocl-open-100/patches

Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com>
Signed-off-by: Anuj Mittal <anuj.mittal@intel.com>

dynamic-layers/clang-layer/recipes-devtools/clang/files/0001-Fix-debug-info-of-work-item-builtin-translation-745.patch

@@ -1,119 +0,0 @@
-From 200c200eb19602ffd7c8f29d0b2df9df1fd311bf Mon Sep 17 00:00:00 2001
-From: Naveen Saini <naveen.kumar.saini@intel.com>
-Date: Wed, 7 Apr 2021 17:44:20 +0800
-Subject: [PATCH] Fix debug info of work-item builtin translation (#745)
-
-debug info of work-item builtins are lost in both llvm IR -> spirv and
-spirv -> llvm IR translations. See #744
-
-Upstream-Status: Backport [https://github.com/KhronosGroup/SPIRV-LLVM-Translator/commit/c734c5c8bbd3012a09c610e4be68e90cc603c580]
-Signed-off-by: Wenju He <wenju.he@intel.com>
-Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com>
----
- lib/SPIRV/OCL20ToSPIRV.cpp              |  5 ++-
- lib/SPIRV/SPIRVReader.cpp               |  1 +
- test/DebugInfo/builtin-get-global-id.ll | 60 +++++++++++++++++++++++++
- 3 files changed, 65 insertions(+), 1 deletion(-)
- create mode 100644 test/DebugInfo/builtin-get-global-id.ll
-
-diff --git a/lib/SPIRV/OCL20ToSPIRV.cpp b/lib/SPIRV/OCL20ToSPIRV.cpp
-index 1262c48c..a742c8cf 100644
---- a/lib/SPIRV/OCL20ToSPIRV.cpp
-+++ b/lib/SPIRV/OCL20ToSPIRV.cpp
-@@ -1297,11 +1297,14 @@ void OCL20ToSPIRV::transWorkItemBuiltinsToVariables() {
-     for (auto UI = I.user_begin(), UE = I.user_end(); UI != UE; ++UI) {
-       auto CI = dyn_cast<CallInst>(*UI);
-       assert(CI && "invalid instruction");
--      Value *NewValue = new LoadInst(BV, "", CI);
-+      const DebugLoc &DLoc = CI->getDebugLoc();
-+      Instruction *NewValue = new LoadInst(BV, "", CI);
-+      NewValue->setDebugLoc(DLoc);
-       LLVM_DEBUG(dbgs() << "Transform: " << *CI << " => " << *NewValue << '\n');
-       if (IsVec) {
-         NewValue =
-             ExtractElementInst::Create(NewValue, CI->getArgOperand(0), "", CI);
-+        NewValue->setDebugLoc(DLoc);
-         LLVM_DEBUG(dbgs() << *NewValue << '\n');
-       }
-       NewValue->takeName(CI);
-diff --git a/lib/SPIRV/SPIRVReader.cpp b/lib/SPIRV/SPIRVReader.cpp
-index 16a3dd38..528f6663 100644
---- a/lib/SPIRV/SPIRVReader.cpp
-+++ b/lib/SPIRV/SPIRVReader.cpp
-@@ -307,6 +307,7 @@ bool SPIRVToLLVM::transOCLBuiltinFromVariable(GlobalVariable *GV,
-   auto Replace = [&](std::vector<Value *> Arg, Instruction *I) {
-     auto Call = CallInst::Create(Func, Arg, "", I);
-     Call->takeName(I);
-+    Call->setDebugLoc(I->getDebugLoc());
-     setAttrByCalledFunc(Call);
-     SPIRVDBG(dbgs() << "[transOCLBuiltinFromVariable] " << *I << " -> " << *Call
-                     << '\n';)
-diff --git a/test/DebugInfo/builtin-get-global-id.ll b/test/DebugInfo/builtin-get-global-id.ll
-new file mode 100644
-index 00000000..a4a00e63
---- /dev/null
-+++ b/test/DebugInfo/builtin-get-global-id.ll
-@@ -0,0 +1,60 @@
-+; Check debug info of builtin get_global_id is preserved from LLVM IR to spirv
-+; and spirv to LLVM IR translation.
-+
-+; Original .cl source:
-+; kernel void test() {
-+;   size_t gid = get_global_id(0);
-+; }
-+
-+; Command line:
-+; ./clang -cc1 1.cl -triple spir64 -cl-std=cl2.0 -emit-llvm -finclude-default-header -debug-info-kind=line-tables-only -O0
-+
-+; RUN: llvm-as %s -o %t.bc
-+; RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s --check-prefix CHECK-SPIRV
-+; RUN: llvm-spirv %t.bc -o %t.spv
-+; RUN: llvm-spirv -r %t.spv -o - | llvm-dis -o - | FileCheck %s
-+
-+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
-+target triple = "spir64"
-+
-+; CHECK-SPIRV: ExtInst {{.*}} DebugScope
-+; CHECK-SPIRV-NEXT: Line {{[0-9]+}} 2 16
-+; CHECK-SPIRV-NEXT: Load {{[0-9]+}} [[LoadRes:[0-9]+]]
-+; CHECK-SPIRV-NEXT: CompositeExtract {{[0-9]+}} {{[0-9]+}} [[LoadRes]] 0
-+
-+; Function Attrs: convergent noinline norecurse nounwind optnone
-+define spir_kernel void @test() #0 !dbg !7 !kernel_arg_addr_space !2 !kernel_arg_access_qual !2 !kernel_arg_type !2 !kernel_arg_base_type !2 !kernel_arg_type_qual !2 {
-+entry:
-+  %gid = alloca i64, align 8
-+  %call = call spir_func i64 @_Z13get_global_idj(i32 0) #2, !dbg !10
-+; CHECK: %call = call spir_func i64 @_Z13get_global_idj(i32 0) #1, !dbg [[DBG:![0-9]+]]
-+  store i64 %call, i64* %gid, align 8, !dbg !11
-+  ret void, !dbg !12
-+}
-+
-+; Function Attrs: convergent nounwind readnone
-+declare spir_func i64 @_Z13get_global_idj(i32) #1
-+
-+attributes #0 = { convergent noinline norecurse nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="false" "unsafe-fp-math"="false" "use-soft-float"="false" }
-+attributes #1 = { convergent nounwind readnone "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
-+attributes #2 = { convergent nounwind readnone }
-+
-+!llvm.dbg.cu = !{!0}
-+!llvm.module.flags = !{!3, !4}
-+!opencl.ocl.version = !{!5}
-+!opencl.spir.version = !{!5}
-+!llvm.ident = !{!6}
-+
-+!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang version 12.0.0 (https://github.com/llvm/llvm-project.git b5bc56da8aa23dc57db9d286b0591dbcf9b1bdd3)", isOptimized: false, runtimeVersion: 0, emissionKind: LineTablesOnly, enums: !2, nameTableKind: None)
-+!1 = !DIFile(filename: "<stdin>", directory: "")
-+!2 = !{}
-+!3 = !{i32 2, !"Debug Info Version", i32 3}
-+!4 = !{i32 1, !"wchar_size", i32 4}
-+!5 = !{i32 2, i32 0}
-+!6 = !{!"clang version 12.0.0 (https://github.com/llvm/llvm-project.git b5bc56da8aa23dc57db9d286b0591dbcf9b1bdd3)"}
-+!7 = distinct !DISubprogram(name: "test", scope: !8, file: !8, line: 1, type: !9, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0, retainedNodes: !2)
-+!8 = !DIFile(filename: "1.cl", directory: "")
-+!9 = !DISubroutineType(types: !2)
-+!10 = !DILocation(line: 2, column: 16, scope: !7)
-+!11 = !DILocation(line: 2, column: 10, scope: !7)
-+!12 = !DILocation(line: 3, column: 1, scope: !7)
--- 
-2.17.1
-

dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0001-OpenCL-3.0-support.patch

@@ -1,7 +1,7 @@
-From 31ec702cb365f4d02dd2146fb4329d642b8fc30b Mon Sep 17 00:00:00 2001
+From 8dbdb2f26674a938ff43b5bfe5b3bf3d1117f9e4 Mon Sep 17 00:00:00 2001
 From: Naveen Saini <naveen.kumar.saini@intel.com>
 Date: Wed, 7 Apr 2021 16:36:10 +0800
-Subject: [PATCH 1/2] OpenCL 3.0 support
+Subject: [PATCH 1/7] OpenCL 3.0 support
 
 Upstream-Status: Backport [Taken from opencl-clang patches, https://github.com/intel/opencl-clang/blob/ocl-open-100/patches/clang/0001-OpenCL-3.0-support.patch]
 Signed-off-by: Anton Zabaznov <anton.zabaznov@intel.com>

dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0001-llvm-spirv-skip-building-tests.patch

@@ -1,7 +1,7 @@
-From 455ce9c25df5313f4a6649cc27075bdfbe25af18 Mon Sep 17 00:00:00 2001
+From 661021749a168c423d69d0ba7cdfa16fed860836 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/3] 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 b718c00..9805140 100644
+index 92c50370..80999c98 100644
 --- a/CMakeLists.txt
 +++ b/CMakeLists.txt
-@@ -24,13 +24,6 @@ if(LLVM_SPIRV_BUILD_EXTERNAL)
+@@ -25,13 +25,6 @@ if(LLVM_SPIRV_BUILD_EXTERNAL)
    set(CMAKE_CXX_STANDARD 14)
    set(CMAKE_CXX_STANDARD_REQUIRED ON)
  
@@ -36,7 +36,7 @@ index b718c00..9805140 100644
    find_package(LLVM 10.0.0 REQUIRED
      COMPONENTS
        Analysis
-@@ -61,9 +54,6 @@ set(LLVM_SPIRV_INCLUDE_DIRS ${CMAKE_CURRENT_SOURCE_DIR}/include)
+@@ -63,9 +56,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 b718c00..9805140 100644
  install(
    FILES
 -- 
-2.7.4
+2.17.1
 

dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0002-Add-cl_khr_extended_subgroup-extensions.patch

@@ -1,7 +1,7 @@
-From 27d47f1a17c8921b07acc8cdc26e38cc609de4a9 Mon Sep 17 00:00:00 2001
+From 3f544cfe44ee5f113a3fb554aca2cf5d64996062 Mon Sep 17 00:00:00 2001
 From: Naveen Saini <naveen.kumar.saini@intel.com>
 Date: Wed, 7 Apr 2021 16:38:38 +0800
-Subject: [PATCH 2/2] Add cl_khr_extended_subgroup extensions.
+Subject: [PATCH 2/7] Add cl_khr_extended_subgroup extensions.
 
 Added extensions and their function declarations into
 the standard header.

dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0002-Fix-building-in-tree-with-cmake-DLLVM_LINK_LLVM_DYLI.patch

@@ -1,7 +1,7 @@
-From a6d4ccf082858e63e139ca06c02a071c343d2657 Mon Sep 17 00:00:00 2001
+From 331e323ae2633a8999a660314022491d670c442c Mon Sep 17 00:00:00 2001
 From: Andrea Bocci <andrea.bocci@cern.ch>
 Date: Sun, 15 Mar 2020 17:35:44 +0100
-Subject: [PATCH] Fix building in-tree with cmake -DLLVM_LINK_LLVM_DYLIB=ON
+Subject: [PATCH 2/3] Fix building in-tree with cmake -DLLVM_LINK_LLVM_DYLIB=ON
 
 Building in-tree with LLVM 11.0 master with the LLVM_LINK_LLVM_DYLIB
 cmake flag fails to link with the LLVMSPIRVLib library.
@@ -28,3 +28,6 @@ index 9aa96d9c..501c0daf 100644
    target_link_libraries(llvm-spirv PRIVATE LLVMSPIRVLib)
  endif()
  
+-- 
+2.17.1
+

dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0003-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch

@@ -0,0 +1,982 @@
+From fbc9996d6490a5d4720b85b47f38335e7fdc99d9 Mon Sep 17 00:00:00 2001
+From: haonanya <haonan.yang@intel.com>
+Date: Mon, 19 Jul 2021 10:14:20 +0800
+Subject: [PATCH 3/3] 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-100/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/OCL20ToSPIRV.cpp            |  79 ++++++++++++++++--
+ lib/SPIRV/SPIRVToOCL.h                |   3 +
+ lib/SPIRV/SPIRVToOCL12.cpp            |  21 +++++
+ lib/SPIRV/SPIRVToOCL20.cpp            |  28 ++++++-
+ lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h |   1 -
+ lib/SPIRV/libSPIRV/SPIRVOpCode.h      |   8 +-
+ test/AtomicFAddEXTForOCL.ll           |  64 +++++++++++++++
+ test/AtomicFAddExt.ll                 | 111 ++++++++-----------------
+ test/AtomicFMaxEXT.ll                 | 113 +++++++-------------------
+ test/AtomicFMaxEXTForOCL.ll           |  64 +++++++++++++++
+ test/AtomicFMinEXT.ll                 | 113 +++++++-------------------
+ test/AtomicFMinEXTForOCL.ll           |  64 +++++++++++++++
+ test/InvalidAtomicBuiltins.cl         |   8 --
+ 13 files changed, 417 insertions(+), 260 deletions(-)
+ create mode 100644 test/AtomicFAddEXTForOCL.ll
+ create mode 100644 test/AtomicFMaxEXTForOCL.ll
+ create mode 100644 test/AtomicFMinEXTForOCL.ll
+
+diff --git a/lib/SPIRV/OCL20ToSPIRV.cpp b/lib/SPIRV/OCL20ToSPIRV.cpp
+index e30aa5be..b676a009 100644
+--- a/lib/SPIRV/OCL20ToSPIRV.cpp
++++ b/lib/SPIRV/OCL20ToSPIRV.cpp
+@@ -408,10 +408,63 @@ void OCL20ToSPIRV::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) {
+@@ -819,7 +872,7 @@ void OCL20ToSPIRV::transAtomicBuiltin(CallInst *CI, OCLBuiltinTransInfo &Info) {
+   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
+@@ -864,7 +917,21 @@ void OCL20ToSPIRV::transAtomicBuiltin(CallInst *CI, OCLBuiltinTransInfo &Info) {
+           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/SPIRVToOCL.h b/lib/SPIRV/SPIRVToOCL.h
+index ddeec0b6..006fb0b1 100644
+--- a/lib/SPIRV/SPIRVToOCL.h
++++ b/lib/SPIRV/SPIRVToOCL.h
+@@ -178,6 +178,9 @@ public:
+   /// using separate maps for OpenCL 1.2 and OpenCL 2.0
+   virtual Instruction *mutateAtomicName(CallInst *CI, Op OC) = 0;
+ 
++  // Transform FP atomic opcode to corresponding OpenCL function name
++  virtual std::string mapFPAtomicName(Op OC) = 0;
++
+ private:
+   /// Transform uniform group opcode to corresponding OpenCL function name,
+   /// example: GroupIAdd(Reduce) => group_iadd => work_group_reduce_add |
+diff --git a/lib/SPIRV/SPIRVToOCL12.cpp b/lib/SPIRV/SPIRVToOCL12.cpp
+index afddd596..d7f00de3 100644
+--- a/lib/SPIRV/SPIRVToOCL12.cpp
++++ b/lib/SPIRV/SPIRVToOCL12.cpp
+@@ -104,6 +104,9 @@ public:
+   /// cl_khr_int64_base_atomics and cl_khr_int64_extended_atomics extensions.
+   std::string mapAtomicName(Op OC, Type *Ty);
+ 
++  // Transform FP atomic opcode to corresponding OpenCL function name
++  std::string mapFPAtomicName(Op OC) override;
++
+   static char ID;
+ };
+ 
+@@ -338,6 +341,21 @@ Instruction *SPIRVToOCL12::visitCallSPIRVAtomicBuiltin(CallInst *CI, Op OC) {
+   return NewCI;
+ }
+ 
++std::string SPIRVToOCL12::mapFPAtomicName(Op OC) {
++  assert(isFPAtomicOpCode(OC) && "Not intended to handle other opcodes than "
++                                 "AtomicF{Add/Min/Max}EXT!");
++  switch (OC) {
++  case OpAtomicFAddEXT:
++    return "atomic_add";
++  case OpAtomicFMinEXT:
++    return "atomic_min";
++  case OpAtomicFMaxEXT:
++    return "atomic_max";
++  default:
++    llvm_unreachable("Unsupported opcode!");
++  }
++}
++
+ Instruction *SPIRVToOCL12::mutateAtomicName(CallInst *CI, Op OC) {
+   AttributeList Attrs = CI->getCalledFunction()->getAttributes();
+   return mutateCallInstOCL(
+@@ -351,6 +369,9 @@ Instruction *SPIRVToOCL12::mutateAtomicName(CallInst *CI, Op OC) {
+ std::string SPIRVToOCL12::mapAtomicName(Op OC, Type *Ty) {
+   std::string Prefix = Ty->isIntegerTy(64) ? kOCLBuiltinName::AtomPrefix
+                                            : kOCLBuiltinName::AtomicPrefix;
++  // Map fp atomic instructions to regular OpenCL built-ins.
++  if (isFPAtomicOpCode(OC))
++    return mapFPAtomicName(OC);
+   return Prefix += OCL12SPIRVBuiltinMap::rmap(OC);
+ }
+ 
+diff --git a/lib/SPIRV/SPIRVToOCL20.cpp b/lib/SPIRV/SPIRVToOCL20.cpp
+index d829ff42..01d088e9 100644
+--- a/lib/SPIRV/SPIRVToOCL20.cpp
++++ b/lib/SPIRV/SPIRVToOCL20.cpp
+@@ -82,6 +82,9 @@ public:
+   /// compare_exchange_strong/weak_explicit
+   Instruction *visitCallSPIRVAtomicCmpExchg(CallInst *CI, Op OC) override;
+ 
++  // Transform FP atomic opcode to corresponding OpenCL function name
++  std::string mapFPAtomicName(Op OC) override;
++
+   static char ID;
+ };
+ 
+@@ -144,11 +147,29 @@ void SPIRVToOCL20::visitCallSPIRVControlBarrier(CallInst *CI) {
+       &Attrs);
+ }
+ 
++std::string SPIRVToOCL20::mapFPAtomicName(Op OC) {
++  assert(isFPAtomicOpCode(OC) && "Not intended to handle other opcodes than "
++                                 "AtomicF{Add/Min/Max}EXT!");
++  switch (OC) {
++  case OpAtomicFAddEXT:
++    return "atomic_fetch_add_explicit";
++  case OpAtomicFMinEXT:
++    return "atomic_fetch_min_explicit";
++  case OpAtomicFMaxEXT:
++    return "atomic_fetch_max_explicit";
++  default:
++    llvm_unreachable("Unsupported opcode!");
++  }
++}
++
+ Instruction *SPIRVToOCL20::mutateAtomicName(CallInst *CI, Op OC) {
+   AttributeList Attrs = CI->getCalledFunction()->getAttributes();
+   return mutateCallInstOCL(
+       M, CI,
+       [=](CallInst *, std::vector<Value *> &Args) {
++        // Map fp atomic instructions to regular OpenCL built-ins.
++        if (isFPAtomicOpCode(OC))
++          return mapFPAtomicName(OC);
+         return OCLSPIRVBuiltinMap::rmap(OC);
+       },
+       &Attrs);
+@@ -215,7 +236,12 @@ CallInst *SPIRVToOCL20::mutateCommonAtomicArguments(CallInst *CI, Op OC) {
+           }
+         }
+         auto Ptr = findFirstPtr(Args);
+-        auto Name = OCLSPIRVBuiltinMap::rmap(OC);
++	std::string Name;
++	// Map fp atomic instructions to regular OpenCL built-ins.
++        if (isFPAtomicOpCode(OC))
++          Name = mapFPAtomicName(OC);
++        else
++          Name = OCLSPIRVBuiltinMap::rmap(OC);
+         auto NumOrder = getSPIRVAtomicBuiltinNumMemoryOrderArgs(OC);
+         auto ScopeIdx = Ptr + 1;
+         auto OrderIdx = Ptr + 2;
+diff --git a/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h b/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h
+index 13f93fbe..7b707993 100644
+--- a/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h
++++ b/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h
+@@ -521,7 +521,6 @@ template <> inline void SPIRVMap<Capability, std::string>::init() {
+   add(CapabilityAtomicFloat64AddEXT, "AtomicFloat64AddEXT");
+   add(CapabilityAtomicFloat32MinMaxEXT, "AtomicFloat32MinMaxEXT");
+   add(CapabilityAtomicFloat64MinMaxEXT, "AtomicFloat64MinMaxEXT");
+-  add(CapabilityAtomicFloat16MinMaxEXT, "AtomicFloat16MinMaxEXT");
+   add(CapabilitySubgroupShuffleINTEL, "SubgroupShuffleINTEL");
+   add(CapabilitySubgroupBufferBlockIOINTEL, "SubgroupBufferBlockIOINTEL");
+   add(CapabilitySubgroupImageBlockIOINTEL, "SubgroupImageBlockIOINTEL");
+diff --git a/lib/SPIRV/libSPIRV/SPIRVOpCode.h b/lib/SPIRV/libSPIRV/SPIRVOpCode.h
+index feec70f6..8e595e83 100644
+--- a/lib/SPIRV/libSPIRV/SPIRVOpCode.h
++++ b/lib/SPIRV/libSPIRV/SPIRVOpCode.h
+@@ -54,11 +54,17 @@ template <> inline void SPIRVMap<Op, std::string>::init() {
+ }
+ SPIRV_DEF_NAMEMAP(Op, OpCodeNameMap)
+ 
++inline bool isFPAtomicOpCode(Op OpCode) {
++  return OpCode == OpAtomicFAddEXT || OpCode == OpAtomicFMinEXT ||
++         OpCode == OpAtomicFMaxEXT;
++}
++
+ inline bool isAtomicOpCode(Op OpCode) {
+   static_assert(OpAtomicLoad < OpAtomicXor, "");
+   return ((unsigned)OpCode >= OpAtomicLoad &&
+           (unsigned)OpCode <= OpAtomicXor) ||
+-         OpCode == OpAtomicFlagTestAndSet || OpCode == OpAtomicFlagClear;
++         OpCode == OpAtomicFlagTestAndSet || OpCode == OpAtomicFlagClear ||
++	 isFPAtomicOpCode(OpCode);
+ }
+ inline bool isBinaryOpCode(Op OpCode) {
+   return ((unsigned)OpCode >= OpIAdd && (unsigned)OpCode <= OpFMod) ||
+diff --git a/test/AtomicFAddEXTForOCL.ll b/test/AtomicFAddEXTForOCL.ll
+new file mode 100644
+index 00000000..fb146fb9
+--- /dev/null
++++ b/test/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/AtomicFAddExt.ll b/test/AtomicFAddExt.ll
+index 011dd8a7..42bdfeea 100644
+--- a/test/AtomicFAddExt.ll
++++ b/test/AtomicFAddExt.ll
+@@ -4,20 +4,16 @@
+ ; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
+ 
+ ; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
+-; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefix=CHECK-LLVM
++; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL12
+ 
+-target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
+-target triple = "spir64-unknown-unknown-sycldevice"
+-
+-%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" }
+-%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" = type { [1 x i64] }
+-%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" }
+-
+-$_ZTSZZ3addIfEvvENKUlRN2cl4sycl7handlerEE19_14clES3_EUlNS1_4itemILi1ELb1EEEE23_37 = comdat any
++; 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
+ 
+-$_ZTSZZ3addIdEvvENKUlRN2cl4sycl7handlerEE19_14clES3_EUlNS1_4itemILi1ELb1EEEE23_37 = comdat any
++; 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
+ 
+-@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
++target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
++target triple = "spir64-unknown-unknown-sycldevice"
+ 
+ ; CHECK-SPIRV: Capability AtomicFloat32AddEXT
+ ; CHECK-SPIRV: Capability AtomicFloat64AddEXT
+@@ -25,62 +21,43 @@ $_ZTSZZ3addIdEvvENKUlRN2cl4sycl7handlerEE19_14clES3_EUlNS1_4itemILi1ELb1EEEE23_3
+ ; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32
+ ; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64
+ 
+-; Function Attrs: convergent norecurse mustprogress
+-define weak_odr dso_local spir_kernel void @_ZTSZZ3addIfEvvENKUlRN2cl4sycl7handlerEE19_14clES3_EUlNS1_4itemILi1ELb1EEEE23_37(float addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3, float addrspace(1)* %_arg_4, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_6, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_8) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 {
++; Function Attrs: convergent norecurse nounwind
++define dso_local spir_func float @_Z14AtomicFloatIncRf(float addrspace(4)* align 4 dereferenceable(4) %Arg) local_unnamed_addr #0 {
+ entry:
+-  %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
+-  %1 = load i64, i64* %0, align 8
+-  %add.ptr.i29 = getelementptr inbounds float, float addrspace(1)* %_arg_, i64 %1
+-  %2 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_8, i64 0, i32 0, i32 0, i64 0
+-  %3 = load i64, i64* %2, align 8
+-  %add.ptr.i = getelementptr inbounds float, float addrspace(1)* %_arg_4, i64 %3
+-  %4 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32, !noalias !5
+-  %5 = extractelement <3 x i64> %4, i64 0
++  %0 = addrspacecast float addrspace(4)* %Arg to float addrspace(1)*
+   ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_32]]
+-  ; CHECK-LLVM: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+]]({{.*}})
+-  %call3.i.i.i.i = tail call spir_func float @_Z21__spirv_AtomicFAddEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)* %add.ptr.i29, i32 1, i32 896, float 1.000000e+00) #2
+-  %add.i.i = fadd float %call3.i.i.i.i, 1.000000e+00
+-  %sext.i = shl i64 %5, 32
+-  %conv5.i = ashr exact i64 %sext.i, 32
+-  %ptridx.i.i = getelementptr inbounds float, float addrspace(1)* %add.ptr.i, i64 %conv5.i
+-  %ptridx.ascast.i.i = addrspacecast float addrspace(1)* %ptridx.i.i to float addrspace(4)*
+-  store float %add.i.i, float addrspace(4)* %ptridx.ascast.i.i, align 4, !tbaa !14
+-  ret void
++  ; CHECK-LLVM-CL12: call spir_func float @[[FLOAT_FUNC_NAME:_Z10atomic_add[[:alnum:]]+ff]]({{.*}})
++  ; 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]]({{.*}})
++  %call3.i.i = tail call spir_func float @_Z21__spirv_AtomicFAddEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)* %0, i32 1, i32 896, float 1.000000e+00) #2
++  ret float %call3.i.i
+ }
+ 
+ ; Function Attrs: convergent
+-; CHECK-LLVM: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float addrspace(1)*, i32, i32, float)
+ declare dso_local spir_func float @_Z21__spirv_AtomicFAddEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)*, i32, i32, float) local_unnamed_addr #1
++; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
+ 
+-; Function Attrs: convergent norecurse mustprogress
+-define weak_odr dso_local spir_kernel void @_ZTSZZ3addIdEvvENKUlRN2cl4sycl7handlerEE19_14clES3_EUlNS1_4itemILi1ELb1EEEE23_37(double addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3, double addrspace(1)* %_arg_4, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_6, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_8) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 {
++; Function Attrs: convergent norecurse nounwind
++define dso_local spir_func double @_Z15AtomicDoubleIncRd(double addrspace(4)* align 8 dereferenceable(8) %Arg) local_unnamed_addr #0 {
+ entry:
+-  %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
+-  %1 = load i64, i64* %0, align 8
+-  %add.ptr.i29 = getelementptr inbounds double, double addrspace(1)* %_arg_, i64 %1
+-  %2 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_8, i64 0, i32 0, i32 0, i64 0
+-  %3 = load i64, i64* %2, align 8
+-  %add.ptr.i = getelementptr inbounds double, double addrspace(1)* %_arg_4, i64 %3
+-  %4 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32, !noalias !18
+-  %5 = extractelement <3 x i64> %4, i64 0
++  %0 = addrspacecast double addrspace(4)* %Arg to double addrspace(1)*
+   ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_64]]
+-  ; CHECK-LLVM: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+]]({{.*}})
+-  %call3.i.i.i.i = tail call spir_func double @_Z21__spirv_AtomicFAddEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)* %add.ptr.i29, i32 1, i32 896, double 1.000000e+00) #2
+-  %add.i.i = fadd double %call3.i.i.i.i, 1.000000e+00
+-  %sext.i = shl i64 %5, 32
+-  %conv5.i = ashr exact i64 %sext.i, 32
+-  %ptridx.i.i = getelementptr inbounds double, double addrspace(1)* %add.ptr.i, i64 %conv5.i
+-  %ptridx.ascast.i.i = addrspacecast double addrspace(1)* %ptridx.i.i to double addrspace(4)*
+-  store double %add.i.i, double addrspace(4)* %ptridx.ascast.i.i, align 8, !tbaa !27
+-  ret void
++  ; CHECK-LLVM-CL12: call spir_func double @[[DOUBLE_FUNC_NAME:_Z10atomic_add[[:alnum:]]+dd]]({{.*}})
++  ; 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]]({{.*}})
++  %call3.i.i = tail call spir_func double @_Z21__spirv_AtomicFAddEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)* %0, i32 1, i32 896, double 1.000000e+00) #2
++  ret double %call3.i.i
+ }
+ 
+ ; Function Attrs: convergent
+-; CHECK-LLVM: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double addrspace(1)*, i32, i32, double)
+ declare dso_local spir_func double @_Z21__spirv_AtomicFAddEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)*, i32, i32, double) local_unnamed_addr #1
++; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
+ 
+-attributes #0 = { convergent norecurse }
+-attributes #1 = { convergent }
++; 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 "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
++attributes #1 = { convergent "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
+ attributes #2 = { convergent nounwind }
+ 
+ !llvm.module.flags = !{!0}
+@@ -91,29 +68,5 @@ attributes #2 = { convergent nounwind }
+ !0 = !{i32 1, !"wchar_size", i32 4}
+ !1 = !{i32 1, i32 2}
+ !2 = !{i32 4, i32 100000}
+-!3 = !{!"clang version 12.0.0"}
+-!4 = !{i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1}
+-!5 = !{!6, !8, !10, !12}
+-!6 = distinct !{!6, !7, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"}
+-!7 = distinct !{!7, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"}
+-!8 = distinct !{!8, !9, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"}
+-!9 = distinct !{!9, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"}
+-!10 = distinct !{!10, !11, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv: %agg.result"}
+-!11 = distinct !{!11, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv"}
+-!12 = distinct !{!12, !13, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE: %agg.result"}
+-!13 = distinct !{!13, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE"}
+-!14 = !{!15, !15, i64 0}
+-!15 = !{!"float", !16, i64 0}
+-!16 = !{!"omnipotent char", !17, i64 0}
+-!17 = !{!"Simple C++ TBAA"}
+-!18 = !{!19, !21, !23, !25}
+-!19 = distinct !{!19, !20, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"}
+-!20 = distinct !{!20, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"}
+-!21 = distinct !{!21, !22, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"}
+-!22 = distinct !{!22, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"}
+-!23 = distinct !{!23, !24, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv: %agg.result"}
+-!24 = distinct !{!24, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv"}
+-!25 = distinct !{!25, !26, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE: %agg.result"}
+-!26 = distinct !{!26, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE"}
+-!27 = !{!28, !28, i64 0}
+-!28 = !{!"double", !16, i64 0}
++!3 = !{!"clang version 13.0.0"}
++
+diff --git a/test/AtomicFMaxEXT.ll b/test/AtomicFMaxEXT.ll
+index 1b81e53b..1c2eec93 100644
+--- a/test/AtomicFMaxEXT.ll
++++ b/test/AtomicFMaxEXT.ll
+@@ -4,20 +4,16 @@
+ ; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
+ 
+ ; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
+-; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefix=CHECK-LLVM
++; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL12
+ 
+-target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
+-target triple = "spir64-unknown-unknown-sycldevice"
+-
+-%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" }
+-%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" = type { [1 x i64] }
+-%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" }
+-
+-$_ZTSZZ8max_testIfEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37 = comdat any
++; 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
+ 
+-$_ZTSZZ8max_testIdEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37 = comdat any
++; 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
+ 
+-@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
++target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
++target triple = "spir64-unknown-unknown-sycldevice"
+ 
+ ; CHECK-SPIRV: Capability AtomicFloat32MinMaxEXT
+ ; CHECK-SPIRV: Capability AtomicFloat64MinMaxEXT
+@@ -25,65 +21,42 @@ $_ZTSZZ8max_testIdEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4item
+ ; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32
+ ; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64
+ 
+-; Function Attrs: convergent norecurse
+-define weak_odr dso_local spir_kernel void @_ZTSZZ8max_testIfEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37(float addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3, float addrspace(1)* %_arg_4, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_6, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_8) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 {
++; Function Attrs: convergent norecurse nounwind
++define dso_local spir_func float @_Z14AtomicFloatMaxRf(float addrspace(4)* align 4 dereferenceable(4) %Arg) local_unnamed_addr #0 {
+ entry:
+-  %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
+-  %1 = load i64, i64* %0, align 8
+-  %add.ptr.i29 = getelementptr inbounds float, float addrspace(1)* %_arg_, i64 %1
+-  %2 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_8, i64 0, i32 0, i32 0, i64 0
+-  %3 = load i64, i64* %2, align 8
+-  %add.ptr.i = getelementptr inbounds float, float addrspace(1)* %_arg_4, i64 %3
+-  %4 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32, !noalias !5
+-  %5 = extractelement <3 x i64> %4, i64 0
+-  %conv.i = trunc i64 %5 to i32
+-  %conv3.i = sitofp i32 %conv.i to float
+-  %add.i = fadd float %conv3.i, 1.000000e+00
++  %0 = addrspacecast float addrspace(4)* %Arg to float addrspace(1)*
+   ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_32]]
+-  ; CHECK-LLVM: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+]]({{.*}})
+-  %call3.i.i.i = tail call spir_func float @_Z21__spirv_AtomicFMaxEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)* %add.ptr.i29, i32 1, i32 896, float %add.i) #2
+-  %sext.i = shl i64 %5, 32
+-  %conv6.i = ashr exact i64 %sext.i, 32
+-  %ptridx.i.i = getelementptr inbounds float, float addrspace(1)* %add.ptr.i, i64 %conv6.i
+-  %ptridx.ascast.i.i = addrspacecast float addrspace(1)* %ptridx.i.i to float addrspace(4)*
+-  store float %call3.i.i.i, float addrspace(4)* %ptridx.ascast.i.i, align 4, !tbaa !14
+-  ret void
++  ; CHECK-LLVM-CL12: call spir_func float @[[FLOAT_FUNC_NAME:_Z10atomic_max[[:alnum:]]+ff]]({{.*}})
++  ; 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.i.i.i = tail call spir_func float @_Z21__spirv_AtomicFMaxEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)* %0, i32 1, i32 896, float 1.000000e+00) #2
++  ret float %call.i.i.i
+ }
+ 
+ ; Function Attrs: convergent
+-; CHECK-LLVM: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float addrspace(1)*, i32, i32, float)
+ declare dso_local spir_func float @_Z21__spirv_AtomicFMaxEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)*, i32, i32, float) local_unnamed_addr #1
++; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
+ 
+-; Function Attrs: convergent norecurse
+-define weak_odr dso_local spir_kernel void @_ZTSZZ8max_testIdEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37(double addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3, double addrspace(1)* %_arg_4, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_6, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_8) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 {
++; Function Attrs: convergent norecurse nounwind
++define dso_local spir_func double @_Z15AtomicDoubleMaxRd(double addrspace(4)* align 8 dereferenceable(8) %Arg) local_unnamed_addr #0 {
+ entry:
+-  %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
+-  %1 = load i64, i64* %0, align 8
+-  %add.ptr.i29 = getelementptr inbounds double, double addrspace(1)* %_arg_, i64 %1
+-  %2 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_8, i64 0, i32 0, i32 0, i64 0
+-  %3 = load i64, i64* %2, align 8
+-  %add.ptr.i = getelementptr inbounds double, double addrspace(1)* %_arg_4, i64 %3
+-  %4 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32, !noalias !18
+-  %5 = extractelement <3 x i64> %4, i64 0
+-  %conv.i = trunc i64 %5 to i32
+-  %conv3.i = sitofp i32 %conv.i to double
+-  %add.i = fadd double %conv3.i, 1.000000e+00
++  %0 = addrspacecast double addrspace(4)* %Arg to double addrspace(1)*
+   ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_64]]
+-  ; CHECK-LLVM: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+]]({{.*}})
+-  %call3.i.i.i = tail call spir_func double @_Z21__spirv_AtomicFMaxEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)* %add.ptr.i29, i32 1, i32 896, double %add.i) #2
+-  %sext.i = shl i64 %5, 32
+-  %conv6.i = ashr exact i64 %sext.i, 32
+-  %ptridx.i.i = getelementptr inbounds double, double addrspace(1)* %add.ptr.i, i64 %conv6.i
+-  %ptridx.ascast.i.i = addrspacecast double addrspace(1)* %ptridx.i.i to double addrspace(4)*
+-  store double %call3.i.i.i, double addrspace(4)* %ptridx.ascast.i.i, align 8, !tbaa !27
+-  ret void
++  ; CHECK-LLVM-CL12: call spir_func double @[[DOUBLE_FUNC_NAME:_Z10atomic_max[[:alnum:]]+dd]]({{.*}})
++  ; 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.i.i.i = tail call spir_func double @_Z21__spirv_AtomicFMaxEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)* %0, i32 1, i32 896, double 1.000000e+00) #2
++  ret double %call.i.i.i
+ }
+ 
+ ; Function Attrs: convergent
+-; CHECK-LLVM: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double addrspace(1)*, i32, i32, double)
+ declare dso_local spir_func double @_Z21__spirv_AtomicFMaxEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)*, i32, i32, double) local_unnamed_addr #1
++; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
+ 
+-attributes #0 = { convergent norecurse "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
++; 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 "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
+ attributes #1 = { convergent "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
+ attributes #2 = { convergent nounwind }
+ 
+@@ -95,29 +68,5 @@ attributes #2 = { convergent nounwind }
+ !0 = !{i32 1, !"wchar_size", i32 4}
+ !1 = !{i32 1, i32 2}
+ !2 = !{i32 4, i32 100000}
+-!3 = !{!"clang version 12.0.0"}
+-!4 = !{i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1}
+-!5 = !{!6, !8, !10, !12}
+-!6 = distinct !{!6, !7, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"}
+-!7 = distinct !{!7, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"}
+-!8 = distinct !{!8, !9, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"}
+-!9 = distinct !{!9, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"}
+-!10 = distinct !{!10, !11, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv: %agg.result"}
+-!11 = distinct !{!11, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv"}
+-!12 = distinct !{!12, !13, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE: %agg.result"}
+-!13 = distinct !{!13, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE"}
+-!14 = !{!15, !15, i64 0}
+-!15 = !{!"float", !16, i64 0}
+-!16 = !{!"omnipotent char", !17, i64 0}
+-!17 = !{!"Simple C++ TBAA"}
+-!18 = !{!19, !21, !23, !25}
+-!19 = distinct !{!19, !20, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"}
+-!20 = distinct !{!20, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"}
+-!21 = distinct !{!21, !22, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"}
+-!22 = distinct !{!22, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"}
+-!23 = distinct !{!23, !24, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv: %agg.result"}
+-!24 = distinct !{!24, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv"}
+-!25 = distinct !{!25, !26, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE: %agg.result"}
+-!26 = distinct !{!26, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE"}
+-!27 = !{!28, !28, i64 0}
+-!28 = !{!"double", !16, i64 0}
++!3 = !{!"clang version 13.0.0"}
++
+diff --git a/test/AtomicFMaxEXTForOCL.ll b/test/AtomicFMaxEXTForOCL.ll
+new file mode 100644
+index 00000000..1f2530d9
+--- /dev/null
++++ b/test/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/AtomicFMinEXT.ll b/test/AtomicFMinEXT.ll
+index 98c98b8e..9e40a669 100644
+--- a/test/AtomicFMinEXT.ll
++++ b/test/AtomicFMinEXT.ll
+@@ -4,20 +4,16 @@
+ ; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
+ 
+ ; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
+-; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefix=CHECK-LLVM
++; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL12
+ 
+-target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
+-target triple = "spir64-unknown-unknown-sycldevice"
+-
+-%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" }
+-%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" = type { [1 x i64] }
+-%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" }
+-
+-$_ZTSZZ8min_testIfEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37 = comdat any
++; 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
+ 
+-$_ZTSZZ8min_testIdEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37 = comdat any
++; 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
+ 
+-@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
++target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
++target triple = "spir64-unknown-unknown-sycldevice"
+ 
+ ; CHECK-SPIRV: Capability AtomicFloat32MinMaxEXT
+ ; CHECK-SPIRV: Capability AtomicFloat64MinMaxEXT
+@@ -25,65 +21,42 @@ $_ZTSZZ8min_testIdEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4item
+ ; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32
+ ; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64
+ 
+-; Function Attrs: convergent norecurse
+-define weak_odr dso_local spir_kernel void @_ZTSZZ8min_testIfEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37(float addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3, float addrspace(1)* %_arg_4, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_6, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_8) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 {
++; Function Attrs: convergent norecurse nounwind
++define dso_local spir_func float @_Z14AtomicFloatMinRf(float addrspace(4)* align 4 dereferenceable(4) %Arg) local_unnamed_addr #0 {
+ entry:
+-  %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
+-  %1 = load i64, i64* %0, align 8
+-  %add.ptr.i29 = getelementptr inbounds float, float addrspace(1)* %_arg_, i64 %1
+-  %2 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_8, i64 0, i32 0, i32 0, i64 0
+-  %3 = load i64, i64* %2, align 8
+-  %add.ptr.i = getelementptr inbounds float, float addrspace(1)* %_arg_4, i64 %3
+-  %4 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32, !noalias !5
+-  %5 = extractelement <3 x i64> %4, i64 0
+-  %conv.i = trunc i64 %5 to i32
+-  %conv3.i = sitofp i32 %conv.i to float
+-  %add.i = fadd float %conv3.i, 1.000000e+00
++  %0 = addrspacecast float addrspace(4)* %Arg to float addrspace(1)*
+   ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_32]]
+-  ; CHECK-LLVM: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+]]({{.*}})
+-  %call3.i.i.i = tail call spir_func float @_Z21__spirv_AtomicFMinEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)* %add.ptr.i29, i32 1, i32 896, float %add.i) #2
+-  %sext.i = shl i64 %5, 32
+-  %conv6.i = ashr exact i64 %sext.i, 32
+-  %ptridx.i.i = getelementptr inbounds float, float addrspace(1)* %add.ptr.i, i64 %conv6.i
+-  %ptridx.ascast.i.i = addrspacecast float addrspace(1)* %ptridx.i.i to float addrspace(4)*
+-  store float %call3.i.i.i, float addrspace(4)* %ptridx.ascast.i.i, align 4, !tbaa !14
+-  ret void
++  ; CHECK-LLVM-CL12: call spir_func float @[[FLOAT_FUNC_NAME:_Z10atomic_min[[:alnum:]]+ff]]({{.*}})
++  ; 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.i.i.i = tail call spir_func float @_Z21__spirv_AtomicFMinEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)* %0, i32 1, i32 896, float 1.000000e+00) #2
++  ret float %call.i.i.i
+ }
+ 
+ ; Function Attrs: convergent
+-; CHECK-LLVM: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float addrspace(1)*, i32, i32, float)
+ declare dso_local spir_func float @_Z21__spirv_AtomicFMinEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)*, i32, i32, float) local_unnamed_addr #1
++; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
+ 
+-; Function Attrs: convergent norecurse
+-define weak_odr dso_local spir_kernel void @_ZTSZZ8min_testIdEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37(double addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3, double addrspace(1)* %_arg_4, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_6, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_8) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 {
++; Function Attrs: convergent norecurse nounwind
++define dso_local spir_func double @_Z15AtomicDoubleMinRd(double addrspace(4)* align 8 dereferenceable(8) %Arg) local_unnamed_addr #0 {
+ entry:
+-  %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
+-  %1 = load i64, i64* %0, align 8
+-  %add.ptr.i29 = getelementptr inbounds double, double addrspace(1)* %_arg_, i64 %1
+-  %2 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_8, i64 0, i32 0, i32 0, i64 0
+-  %3 = load i64, i64* %2, align 8
+-  %add.ptr.i = getelementptr inbounds double, double addrspace(1)* %_arg_4, i64 %3
+-  %4 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32, !noalias !18
+-  %5 = extractelement <3 x i64> %4, i64 0
+-  %conv.i = trunc i64 %5 to i32
+-  %conv3.i = sitofp i32 %conv.i to double
+-  %add.i = fadd double %conv3.i, 1.000000e+00
++  %0 = addrspacecast double addrspace(4)* %Arg to double addrspace(1)*
+   ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_64]]
+-  ; CHECK-LLVM: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+]]({{.*}})
+-  %call3.i.i.i = tail call spir_func double @_Z21__spirv_AtomicFMinEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)* %add.ptr.i29, i32 1, i32 896, double %add.i) #2
+-  %sext.i = shl i64 %5, 32
+-  %conv6.i = ashr exact i64 %sext.i, 32
+-  %ptridx.i.i = getelementptr inbounds double, double addrspace(1)* %add.ptr.i, i64 %conv6.i
+-  %ptridx.ascast.i.i = addrspacecast double addrspace(1)* %ptridx.i.i to double addrspace(4)*
+-  store double %call3.i.i.i, double addrspace(4)* %ptridx.ascast.i.i, align 8, !tbaa !27
+-  ret void
++  ; CHECK-LLVM-CL12: call spir_func double @[[DOUBLE_FUNC_NAME:_Z10atomic_min[[:alnum:]]+dd]]({{.*}})
++  ; 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.i.i.i = tail call spir_func double @_Z21__spirv_AtomicFMinEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)* %0, i32 1, i32 896, double 1.000000e+00) #2
++  ret double %call.i.i.i
+ }
+ 
+ ; Function Attrs: convergent
+-; CHECK-LLVM: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double addrspace(1)*, i32, i32, double)
+ declare dso_local spir_func double @_Z21__spirv_AtomicFMinEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)*, i32, i32, double) local_unnamed_addr #1
++; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
+ 
+-attributes #0 = { convergent norecurse "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
++; 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 "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
+ attributes #1 = { convergent "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
+ attributes #2 = { convergent nounwind }
+ 
+@@ -95,29 +68,5 @@ attributes #2 = { convergent nounwind }
+ !0 = !{i32 1, !"wchar_size", i32 4}
+ !1 = !{i32 1, i32 2}
+ !2 = !{i32 4, i32 100000}
+-!3 = !{!"clang version 12.0.0 (https://github.com/otcshare/llvm.git 67add71766d55d6a8d8d894822f583d6365a3b7d)"}
+-!4 = !{i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1}
+-!5 = !{!6, !8, !10, !12}
+-!6 = distinct !{!6, !7, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"}
+-!7 = distinct !{!7, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"}
+-!8 = distinct !{!8, !9, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"}
+-!9 = distinct !{!9, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"}
+-!10 = distinct !{!10, !11, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv: %agg.result"}
+-!11 = distinct !{!11, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv"}
+-!12 = distinct !{!12, !13, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE: %agg.result"}
+-!13 = distinct !{!13, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE"}
+-!14 = !{!15, !15, i64 0}
+-!15 = !{!"float", !16, i64 0}
+-!16 = !{!"omnipotent char", !17, i64 0}
+-!17 = !{!"Simple C++ TBAA"}
+-!18 = !{!19, !21, !23, !25}
+-!19 = distinct !{!19, !20, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"}
+-!20 = distinct !{!20, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"}
+-!21 = distinct !{!21, !22, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"}
+-!22 = distinct !{!22, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"}
+-!23 = distinct !{!23, !24, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv: %agg.result"}
+-!24 = distinct !{!24, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv"}
+-!25 = distinct !{!25, !26, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE: %agg.result"}
+-!26 = distinct !{!26, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE"}
+-!27 = !{!28, !28, i64 0}
+-!28 = !{!"double", !16, i64 0}
++!3 = !{!"clang version 13.0.0"}
++
+diff --git a/test/AtomicFMinEXTForOCL.ll b/test/AtomicFMinEXTForOCL.ll
+new file mode 100644
+index 00000000..6196b0f8
+--- /dev/null
++++ b/test/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)"}
+diff --git a/test/InvalidAtomicBuiltins.cl b/test/InvalidAtomicBuiltins.cl
+index b8ec5b89..2182f070 100644
+--- a/test/InvalidAtomicBuiltins.cl
++++ b/test/InvalidAtomicBuiltins.cl
+@@ -41,13 +41,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 +82,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);
+ }
+-- 
+2.17.1
+

dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0003-Memory-leak-fix-for-Managed-Static-Mutex.patch

@@ -0,0 +1,35 @@
+From cfb18b75e8a353bc7486f337541476a36994b063 Mon Sep 17 00:00:00 2001
+From: juanrod2 <>
+Date: Tue, 22 Dec 2020 08:33:08 +0800
+Subject: [PATCH 3/7] 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]
+
+Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com>
+
+Cleaning a mutex inside ManagedStatic llvm class.
+---
+ llvm/lib/Support/ManagedStatic.cpp | 6 +++++-
+ 1 file changed, 5 insertions(+), 1 deletion(-)
+
+diff --git a/llvm/lib/Support/ManagedStatic.cpp b/llvm/lib/Support/ManagedStatic.cpp
+index 053493f72fb5..6571580ccecf 100644
+--- a/llvm/lib/Support/ManagedStatic.cpp
++++ b/llvm/lib/Support/ManagedStatic.cpp
+@@ -76,8 +76,12 @@ void ManagedStaticBase::destroy() const {
+ 
+ /// llvm_shutdown - Deallocate and destroy all ManagedStatic variables.
+ void llvm::llvm_shutdown() {
+-  std::lock_guard<std::recursive_mutex> Lock(*getManagedStaticMutex());
++  getManagedStaticMutex()->lock();
+ 
+   while (StaticList)
+     StaticList->destroy();
++
++  getManagedStaticMutex()->unlock();
++  delete ManagedStaticMutex;
++  ManagedStaticMutex = nullptr;
+ }
+-- 
+2.17.1
+

dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0004-Remove-repo-name-in-LLVM-IR.patch

@@ -1,18 +1,17 @@
-From b53fd86ffdeacb9b13624bdb110fd25e8c35cb92 Mon Sep 17 00:00:00 2001
+From b794037bf1f90a93efa4c542855ad569cb13b4c5 Mon Sep 17 00:00:00 2001
 From: Feng Zou <feng.zou@intel.com>
 Date: Mon, 19 Oct 2020 14:43:38 +0800
-Subject: [PATCH] Remove repo name in LLVM IR
+Subject: [PATCH 4/7] Remove repo name in LLVM IR
 
 Upstream-Status: Backport [Taken from opencl-clang patches, https://github.com/intel/opencl-clang/blob/ocl-open-100/patches/llvm/0003-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 1b6519b4b7c..8fd6b23bb34 100644
+index 1b6519b4b7c4..8fd6b23bb345 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 1b6519b4b7c..8fd6b23bb34 100644
    endif()
  endfunction()
 -- 
-2.18.1
+2.17.1
 

dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0005-Remove-__IMAGE_SUPPORT__-macro-for-SPIR-since-SPIR-d.patch

@@ -0,0 +1,47 @@
+From 3dd4766499d25e5978a5d90001f18e657e875da0 Mon Sep 17 00:00:00 2001
+From: haonanya <haonan.yang@intel.com>
+Date: Thu, 12 Aug 2021 15:48:34 +0800
+Subject: [PATCH 5/7] 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-100/patches/clang/0003-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 | 4 ----
+ 2 files changed, 7 deletions(-)
+
+diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp
+index aefd208e6cd3..b4a84636673a 100644
+--- a/clang/lib/Frontend/InitPreprocessor.cpp
++++ b/clang/lib/Frontend/InitPreprocessor.cpp
+@@ -1108,9 +1108,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 b088a37ba665..39a222d02faf 100644
+--- a/clang/test/Preprocessor/predefined-macros.c
++++ b/clang/test/Preprocessor/predefined-macros.c
+@@ -184,10 +184,6 @@
+ // MSCOPE:#define __OPENCL_MEMORY_SCOPE_WORK_GROUP 1
+ // MSCOPE:#define __OPENCL_MEMORY_SCOPE_WORK_ITEM 0
+ 
+-// 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: #define __IMAGE_SUPPORT__ 1
+-
+ // RUN: %clang_cc1 %s -E -dM -o - -x hip -triple amdgcn-amd-amdhsa \
+ // RUN:   | FileCheck -match-full-lines %s --check-prefix=CHECK-HIP
+ // CHECK-HIP-NOT: #define __CUDA_ARCH__
+-- 
+2.17.1
+

dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0006-Avoid-calling-ParseCommandLineOptions-in-BackendUtil.patch

@@ -0,0 +1,53 @@
+From 2c53abd0008bbecfcfe871c6060f4bbf1c94c74a Mon Sep 17 00:00:00 2001
+From: Raphael Isemann <teemperor@gmail.com>
+Date: Thu, 1 Apr 2021 18:41:44 +0200
+Subject: [PATCH 6/7] 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-100/patches/clang/0004-Avoid-calling-ParseCommandLineOptions-in-BackendUtil.patch]
+
+Signed-off-by: Raphael Isemann <teemperor@gmail.com>
+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 0bfcab88a3a9..db8fd4166d7a 100644
+--- a/clang/lib/CodeGen/BackendUtil.cpp
++++ b/clang/lib/CodeGen/BackendUtil.cpp
+@@ -743,7 +743,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
+

dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0007-support-cl_ext_float_atomics.patch

@@ -0,0 +1,377 @@
+From a685de6fc45afcdbe4a7120e9d5b33e175dd71cd Mon Sep 17 00:00:00 2001
+From: haonanya <haonan.yang@intel.com>
+Date: Fri, 13 Aug 2021 10:00:02 +0800
+Subject: [PATCH 7/7] support cl_ext_float_atomics
+
+Upstream-Status: Backport [Taken from opencl-clang patches; https://github.com/intel/opencl-clang/blob/ocl-open-100/patches/clang/0005-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          | 208 ++++++++++++++++++++++++++
+ clang/test/Headers/opencl-c-header.cl |  96 ++++++++++++
+ 3 files changed, 329 insertions(+)
+
+diff --git a/clang/lib/Headers/opencl-c-base.h b/clang/lib/Headers/opencl-c-base.h
+index 2cc688ccc3da..86bbee12fdf8 100644
+--- a/clang/lib/Headers/opencl-c-base.h
++++ b/clang/lib/Headers/opencl-c-base.h
+@@ -14,6 +14,31 @@
+   #define CL_VERSION_3_0 300
+ #endif
+ 
++#if (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200)
++// For SPIR all extensions are supported.
++#if defined(__SPIR__)
++#define cl_ext_float_atomics 1
++#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)
++
+ // Define features for 2.0 for header backward compatibility
+ #ifndef __opencl_c_int64
+   #define __opencl_c_int64 1
+diff --git a/clang/lib/Headers/opencl-c.h b/clang/lib/Headers/opencl-c.h
+index 67d900eb1c3d..b463e702d95e 100644
+--- a/clang/lib/Headers/opencl-c.h
++++ b/clang/lib/Headers/opencl-c.h
+@@ -14354,6 +14354,214 @@ 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 // defined(__opencl_c_ext_fp32_global_atomic_min_max)
++
++#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 // defined(__opencl_c_ext_fp32_local_atomic_min_max)
++
++#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 // defined(__opencl_c_ext_fp32_global_atomic_min_max) ||                      \
++    defined(__opencl_c_ext_fp32_local_atomic_min_max)
++
++#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 // defined(__opencl_c_ext_fp64_global_atomic_min_max)
++
++#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 // defined(__opencl_c_ext_fp64_local_atomic_min_max)
++
++#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 // defined(__opencl_c_ext_fp64_global_atomic_min_max) ||                      \
++    defined(__opencl_c_ext_fp64_local_atomic_min_max)
++
++#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 // defined(__opencl_c_ext_fp32_global_atomic_add)
++
++#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 // defined(__opencl_c_ext_fp32_local_atomic_add)
++
++#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 // defined(__opencl_c_ext_fp32_global_atomic_add) ||                          \
++    defined(__opencl_c_ext_fp32_local_atomic_add)
++
++#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 // defined(__opencl_c_ext_fp64_global_atomic_add)
++
++#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 // defined(__opencl_c_ext_fp64_local_atomic_add)
++
++#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 // defined(__opencl_c_ext_fp64_global_atomic_add) ||                          \
++    defined(__opencl_c_ext_fp64_local_atomic_add)
++
++#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..7f720cf28142 100644
+--- a/clang/test/Headers/opencl-c-header.cl
++++ b/clang/test/Headers/opencl-c-header.cl
+@@ -98,3 +98,99 @@ 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_fp64_global_atomic_add != 1
++#error "Incorrectly defined __opencl_c_ext_fp64_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_fp64_local_atomic_add != 1
++#error "Incorrectly defined __opencl_c_ext_fp64_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_fp64_global_atomic_min_max != 1
++#error "Incorrectly defined __opencl_c_ext_fp64_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
++#if __opencl_c_ext_fp64_local_atomic_min_max != 1
++#error "Incorrectly defined __opencl_c_ext_fp64_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
+

dynamic-layers/clang-layer/recipes-devtools/clang/llvm-project-source.bbappend

@@ -1,20 +1,23 @@
 FILESEXTRAPATHS:prepend:intel-x86-common := "${THISDIR}/files:"
 
-SPIRV10_SRCREV = "576abae62cecd171992017a4a786e3831221ab8d"
+SPIRV10_SRCREV = "fe4d6b767363a1995ccbfca27f79efb10dcfe110"
 SPIRV11_SRCREV = "2a8c1e6c9778deaa720a23e08c293006dc5d56fd"
 
 SPIRV_SRCREV = "${@bb.utils.contains('LLVMVERSION', '10.0.1', '${SPIRV10_SRCREV}', '${SPIRV11_SRCREV}', d)}"
 
 SRC_URI_LLVM10_PATCHES = " \
-                   file://llvm10-skip-building-tests.patch;patchdir=llvm/projects/llvm-spirv \
-                   file://fix-shared-libs.patch;patchdir=llvm/projects/llvm-spirv \
+                   file://llvm10-0001-llvm-spirv-skip-building-tests.patch;patchdir=llvm/projects/llvm-spirv \
+                   file://llvm10-0002-Fix-building-in-tree-with-cmake-DLLVM_LINK_LLVM_DYLI.patch;patchdir=llvm/projects/llvm-spirv \
+                   file://llvm10-0003-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch;patchdir=llvm/projects/llvm-spirv \
                    file://BasicBlockUtils-Add-metadata-fixing-in-SplitBlockPre.patch;patchdir=llvm \
                    file://IndVarSimplify-Do-not-use-SCEV-expander-for-IVCount-.patch;patchdir=llvm \
-                   file://llvm10-OpenCL-3.0-support.patch \
-                   file://0002-Add-cl_khr_extended_subgroup-extensions.patch \
-                   file://0001-Memory-leak-fix-for-Managed-Static-Mutex.patch \
-                   file://llvm10-Remove-repo-name-in-LLVM-IR.patch \
-                   file://0001-Fix-debug-info-of-work-item-builtin-translation-745.patch;patchdir=llvm/projects/llvm-spirv \
+                   file://llvm10-0001-OpenCL-3.0-support.patch \
+                   file://llvm10-0002-Add-cl_khr_extended_subgroup-extensions.patch \
+                   file://llvm10-0003-Memory-leak-fix-for-Managed-Static-Mutex.patch \
+                   file://llvm10-0004-Remove-repo-name-in-LLVM-IR.patch \
+                   file://llvm10-0005-Remove-__IMAGE_SUPPORT__-macro-for-SPIR-since-SPIR-d.patch \
+                   file://llvm10-0006-Avoid-calling-ParseCommandLineOptions-in-BackendUtil.patch \
+                   file://llvm10-0007-support-cl_ext_float_atomics.patch \
                    "
 
 SRC_URI_LLVM11_PATCHES = " \