https://github.com/intel/opencl-clang/tree/ocl-open-100/patches
Signed-off-by: Naveen Saini <[email protected]> --- ...of-work-item-builtin-translation-745.patch | 119 --- ...h => llvm10-0001-OpenCL-3.0-support.patch} | 4 +- ...0001-llvm-spirv-skip-building-tests.patch} | 12 +- ...cl_khr_extended_subgroup-extensions.patch} | 4 +- ...ree-with-cmake-DLLVM_LINK_LLVM_DYLI.patch} | 7 +- ...-cl_ext_float_atomics-in-SPIRVWriter.patch | 982 ++++++++++++++++++ ...ry-leak-fix-for-Managed-Static-Mutex.patch | 35 + ...10-0004-Remove-repo-name-in-LLVM-IR.patch} | 9 +- ...UPPORT__-macro-for-SPIR-since-SPIR-d.patch | 47 + ...rseCommandLineOptions-in-BackendUtil.patch | 53 + ...10-0007-support-cl_ext_float_atomics.patch | 377 +++++++ .../clang/llvm-project-source.bbappend | 19 +- 12 files changed, 1524 insertions(+), 144 deletions(-) delete mode 100644 dynamic-layers/clang-layer/recipes-devtools/clang/files/0001-Fix-debug-info-of-work-item-builtin-translation-745.patch rename dynamic-layers/clang-layer/recipes-devtools/clang/files/{llvm10-OpenCL-3.0-support.patch => llvm10-0001-OpenCL-3.0-support.patch} (99%) rename dynamic-layers/clang-layer/recipes-devtools/clang/files/{llvm10-skip-building-tests.patch => llvm10-0001-llvm-spirv-skip-building-tests.patch} (81%) rename dynamic-layers/clang-layer/recipes-devtools/clang/files/{0002-Add-cl_khr_extended_subgroup-extensions.patch => llvm10-0002-Add-cl_khr_extended_subgroup-extensions.patch} (99%) rename dynamic-layers/clang-layer/recipes-devtools/clang/files/{fix-shared-libs.patch => llvm10-0002-Fix-building-in-tree-with-cmake-DLLVM_LINK_LLVM_DYLI.patch} (85%) create mode 100644 dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0003-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch create mode 100644 dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0003-Memory-leak-fix-for-Managed-Static-Mutex.patch rename dynamic-layers/clang-layer/recipes-devtools/clang/files/{llvm10-Remove-repo-name-in-LLVM-IR.patch => llvm10-0004-Remove-repo-name-in-LLVM-IR.patch} (91%) create mode 100644 dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0005-Remove-__IMAGE_SUPPORT__-macro-for-SPIR-since-SPIR-d.patch create mode 100644 dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0006-Avoid-calling-ParseCommandLineOptions-in-BackendUtil.patch create mode 100644 dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0007-support-cl_ext_float_atomics.patch diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/0001-Fix-debug-info-of-work-item-builtin-translation-745.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/0001-Fix-debug-info-of-work-item-builtin-translation-745.patch deleted file mode 100644 index 923b871f..00000000 --- a/dynamic-layers/clang-layer/recipes-devtools/clang/files/0001-Fix-debug-info-of-work-item-builtin-translation-745.patch +++ /dev/null @@ -1,119 +0,0 @@ -From 200c200eb19602ffd7c8f29d0b2df9df1fd311bf Mon Sep 17 00:00:00 2001 -From: Naveen Saini <[email protected]> -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 <[email protected]> -Signed-off-by: Naveen Saini <[email protected]> ---- - 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 - diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-OpenCL-3.0-support.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0001-OpenCL-3.0-support.patch similarity index 99% rename from dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-OpenCL-3.0-support.patch rename to dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0001-OpenCL-3.0-support.patch index 53395ea0..1ab00df0 100644 --- a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-OpenCL-3.0-support.patch +++ b/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 <[email protected]> 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 <[email protected]> diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-skip-building-tests.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0001-llvm-spirv-skip-building-tests.patch similarity index 81% rename from dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-skip-building-tests.patch rename to dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0001-llvm-spirv-skip-building-tests.patch index 8e58ec25..84a4ba19 100644 --- a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-skip-building-tests.patch +++ b/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 <[email protected]> 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 <[email protected]> 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 diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/0002-Add-cl_khr_extended_subgroup-extensions.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0002-Add-cl_khr_extended_subgroup-extensions.patch similarity index 99% rename from dynamic-layers/clang-layer/recipes-devtools/clang/files/0002-Add-cl_khr_extended_subgroup-extensions.patch rename to dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0002-Add-cl_khr_extended_subgroup-extensions.patch index cbe492c4..3f1b24e7 100644 --- a/dynamic-layers/clang-layer/recipes-devtools/clang/files/0002-Add-cl_khr_extended_subgroup-extensions.patch +++ b/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 <[email protected]> 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. diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/fix-shared-libs.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0002-Fix-building-in-tree-with-cmake-DLLVM_LINK_LLVM_DYLI.patch similarity index 85% rename from dynamic-layers/clang-layer/recipes-devtools/clang/files/fix-shared-libs.patch rename to dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0002-Fix-building-in-tree-with-cmake-DLLVM_LINK_LLVM_DYLI.patch index d69d2a97..1aff65e7 100644 --- a/dynamic-layers/clang-layer/recipes-devtools/clang/files/fix-shared-libs.patch +++ b/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 <[email protected]> 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 + diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0003-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0003-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch new file mode 100644 index 00000000..49edd7e1 --- /dev/null +++ b/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 <[email protected]> +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 <[email protected]> +Signed-off-by: Naveen Saini <[email protected]> +--- + 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 + diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0003-Memory-leak-fix-for-Managed-Static-Mutex.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0003-Memory-leak-fix-for-Managed-Static-Mutex.patch new file mode 100644 index 00000000..3b035f47 --- /dev/null +++ b/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 <[email protected]> + +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 + diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-Remove-repo-name-in-LLVM-IR.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0004-Remove-repo-name-in-LLVM-IR.patch similarity index 91% rename from dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-Remove-repo-name-in-LLVM-IR.patch rename to dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0004-Remove-repo-name-in-LLVM-IR.patch index 232ae063..f8dec996 100644 --- a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-Remove-repo-name-in-LLVM-IR.patch +++ b/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 <[email protected]> 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 <[email protected]> Signed-off-by: Naveen Saini <[email protected]> - --- 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 diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0005-Remove-__IMAGE_SUPPORT__-macro-for-SPIR-since-SPIR-d.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0005-Remove-__IMAGE_SUPPORT__-macro-for-SPIR-since-SPIR-d.patch new file mode 100644 index 00000000..f8f177e5 --- /dev/null +++ b/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 <[email protected]> +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 <[email protected]> +Signed-off-by: Naveen Saini <[email protected]> +--- + 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 + diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0006-Avoid-calling-ParseCommandLineOptions-in-BackendUtil.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0006-Avoid-calling-ParseCommandLineOptions-in-BackendUtil.patch new file mode 100644 index 00000000..0b4ee8c7 --- /dev/null +++ b/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 <[email protected]> +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 <[email protected]> +Signed-off-by: Naveen Saini <[email protected]> +--- + 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 + diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0007-support-cl_ext_float_atomics.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0007-support-cl_ext_float_atomics.patch new file mode 100644 index 00000000..f7d191ff --- /dev/null +++ b/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 <[email protected]> +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 <[email protected]> +Signed-off-by: Naveen Saini <[email protected]> +--- + 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 + diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/llvm-project-source.bbappend b/dynamic-layers/clang-layer/recipes-devtools/clang/llvm-project-source.bbappend index a09343b3..ac34321c 100644 --- a/dynamic-layers/clang-layer/recipes-devtools/clang/llvm-project-source.bbappend +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/llvm-project-source.bbappend @@ -1,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 = " \ -- 2.17.1
-=-=-=-=-=-=-=-=-=-=-=- Links: You receive all messages sent to this group. View/Reply Online (#7203): https://lists.yoctoproject.org/g/meta-intel/message/7203 Mute This Topic: https://lists.yoctoproject.org/mt/85011613/21656 Group Owner: [email protected] Unsubscribe: https://lists.yoctoproject.org/g/meta-intel/unsub [[email protected]] -=-=-=-=-=-=-=-=-=-=-=-
