https://github.com/intel/opencl-clang/tree/ocl-open-110/patches
Signed-off-by: Naveen Saini <[email protected]> --- ...h => llvm11-0001-OpenCL-3.0-support.patch} | 6 +- ...0001-llvm-spirv-skip-building-tests.patch} | 12 +- ...-cl_ext_float_atomics-in-SPIRVWriter.patch | 433 ++++++++++++++++++ ...y-leak-fix-for-Managed-Static-Mutex.patch} | 6 +- ...11-0003-Remove-repo-name-in-LLVM-IR.patch} | 9 +- ...UPPORT__-macro-for-SPIR-since-SPIR-d.patch | 51 +++ ...rseCommandLineOptions-in-BackendUtil.patch | 52 +++ ...-OpenCL-support-cl_ext_float_atomics.patch | 353 ++++++++++++++ .../clang/llvm-project-source.bbappend | 14 +- 9 files changed, 913 insertions(+), 23 deletions(-) rename dynamic-layers/clang-layer/recipes-devtools/clang/files/{llvm11-OpenCL-3.0-support.patch => llvm11-0001-OpenCL-3.0-support.patch} (99%) rename dynamic-layers/clang-layer/recipes-devtools/clang/files/{llvm11-skip-building-tests.patch => llvm11-0001-llvm-spirv-skip-building-tests.patch} (81%) create mode 100644 dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0002-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch rename dynamic-layers/clang-layer/recipes-devtools/clang/files/{0001-Memory-leak-fix-for-Managed-Static-Mutex.patch => llvm11-0002-Memory-leak-fix-for-Managed-Static-Mutex.patch} (89%) rename dynamic-layers/clang-layer/recipes-devtools/clang/files/{llvm11-Remove-repo-name-in-LLVM-IR.patch => llvm11-0003-Remove-repo-name-in-LLVM-IR.patch} (91%) create mode 100644 dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0004-Remove-__IMAGE_SUPPORT__-macro-for-SPIR-since-SPIR-d.patch create mode 100644 dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0005-Avoid-calling-ParseCommandLineOptions-in-BackendUtil.patch create mode 100644 dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0006-OpenCL-support-cl_ext_float_atomics.patch diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-OpenCL-3.0-support.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0001-OpenCL-3.0-support.patch similarity index 99% rename from dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-OpenCL-3.0-support.patch rename to dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0001-OpenCL-3.0-support.patch index 98545db0..af433e14 100644 --- a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-OpenCL-3.0-support.patch +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0001-OpenCL-3.0-support.patch @@ -1,13 +1,11 @@ -From d767afa79d1c8153081eac1ef33e348cadbea5bb Mon Sep 17 00:00:00 2001 +From 36d87f69fee9c3d3f399f8e4027ab707ad050e80 Mon Sep 17 00:00:00 2001 From: Anton Zabaznov <[email protected]> Date: Tue, 22 Sep 2020 19:03:50 +0300 -Subject: [PATCH] OpenCL 3.0 support +Subject: [PATCH 1/6] OpenCL 3.0 support Upstream-Status: Backport [Taken from opencl-clang patches, https://github.com/intel/opencl-clang/blob/ocl-open-110/patches/clang/0001-OpenCL-3.0-support.patch] Signed-off-by: Anton Zabaznov <[email protected]> Signed-off-by: Naveen Saini <[email protected]> - - --- clang/include/clang/Basic/Builtins.def | 65 +- clang/include/clang/Basic/Builtins.h | 13 +- diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-skip-building-tests.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0001-llvm-spirv-skip-building-tests.patch similarity index 81% rename from dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-skip-building-tests.patch rename to dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0001-llvm-spirv-skip-building-tests.patch index 011c09ee..237dec51 100644 --- a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-skip-building-tests.patch +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0001-llvm-spirv-skip-building-tests.patch @@ -1,7 +1,7 @@ -From d362652617c5e840089273df0c6623a9745c92a2 Mon Sep 17 00:00:00 2001 +From 6690d77f9007ce82984dc1b6ae12585cb3e04785 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/2] llvm-spirv: skip building tests Some of these need clang to be built and since we're building this in-tree, that leads to problems when compiling libcxx, compiler-rt which aren't built @@ -19,10 +19,10 @@ Signed-off-by: Naveen Saini <[email protected]> 1 file changed, 10 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt -index ecebb4cb..578ca602 100644 +index ec61fb95..d723c0a5 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt -@@ -25,13 +25,6 @@ if(LLVM_SPIRV_BUILD_EXTERNAL) +@@ -26,13 +26,6 @@ if(LLVM_SPIRV_BUILD_EXTERNAL) set(CMAKE_CXX_STANDARD 14) set(CMAKE_CXX_STANDARD_REQUIRED ON) @@ -36,7 +36,7 @@ index ecebb4cb..578ca602 100644 find_package(LLVM ${BASE_LLVM_VERSION} REQUIRED COMPONENTS Analysis -@@ -62,9 +55,6 @@ set(LLVM_SPIRV_INCLUDE_DIRS ${CMAKE_CURRENT_SOURCE_DIR}/include) +@@ -65,9 +58,6 @@ set(LLVM_SPIRV_INCLUDE_DIRS ${CMAKE_CURRENT_SOURCE_DIR}/include) add_subdirectory(lib/SPIRV) add_subdirectory(tools/llvm-spirv) @@ -47,5 +47,5 @@ index ecebb4cb..578ca602 100644 install( FILES -- -2.26.2 +2.17.1 diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0002-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0002-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch new file mode 100644 index 00000000..14e370f7 --- /dev/null +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0002-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch @@ -0,0 +1,433 @@ +From 8e12d8fb3cdbdafca73fe8ed4f0cde773b1788b4 Mon Sep 17 00:00:00 2001 +From: haonanya <[email protected]> +Date: Wed, 28 Jul 2021 11:43:20 +0800 +Subject: [PATCH 2/2] Add support for cl_ext_float_atomics in SPIRVWriter + +Upstream-Status: Backport [Taken from opencl-clang patches, https://github.com/intel/opencl-clang/blob/ocl-open-110/patches/spirv/0001-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch] + +Signed-off-by: haonanya <[email protected]> +Signed-off-by: Naveen Saini <[email protected]> +--- + lib/SPIRV/OCLToSPIRV.cpp | 80 +++++++++++++++++++++++-- + lib/SPIRV/OCLUtil.cpp | 26 -------- + lib/SPIRV/OCLUtil.h | 4 -- + test/negative/InvalidAtomicBuiltins.cl | 12 +--- + test/transcoding/AtomicFAddEXTForOCL.ll | 64 ++++++++++++++++++++ + test/transcoding/AtomicFMaxEXTForOCL.ll | 64 ++++++++++++++++++++ + test/transcoding/AtomicFMinEXTForOCL.ll | 64 ++++++++++++++++++++ + 7 files changed, 269 insertions(+), 45 deletions(-) + create mode 100644 test/transcoding/AtomicFAddEXTForOCL.ll + create mode 100644 test/transcoding/AtomicFMaxEXTForOCL.ll + create mode 100644 test/transcoding/AtomicFMinEXTForOCL.ll + +diff --git a/lib/SPIRV/OCLToSPIRV.cpp b/lib/SPIRV/OCLToSPIRV.cpp +index 04d51586..f00f5f7b 100644 +--- a/lib/SPIRV/OCLToSPIRV.cpp ++++ b/lib/SPIRV/OCLToSPIRV.cpp +@@ -421,10 +421,63 @@ void OCLToSPIRVBase::visitCallInst(CallInst &CI) { + if (DemangledName.find(kOCLBuiltinName::AtomicPrefix) == 0 || + DemangledName.find(kOCLBuiltinName::AtomPrefix) == 0) { + +- // Compute atomic builtins do not support floating types. +- if (CI.getType()->isFloatingPointTy() && +- isComputeAtomicOCLBuiltin(DemangledName)) +- return; ++ // Compute "atom" prefixed builtins do not support floating types. ++ if (CI.getType()->isFloatingPointTy()) { ++ if (DemangledName.find(kOCLBuiltinName::AtomPrefix) == 0) ++ return; ++ // handle functions which are "atomic_" prefixed. ++ StringRef Stem = DemangledName; ++ Stem = Stem.drop_front(strlen("atomic_")); ++ // FP-typed atomic_{add, sub, inc, dec, exchange, min, max, or, and, xor, ++ // fetch_or, fetch_xor, fetch_and, fetch_or_explicit, fetch_xor_explicit, ++ // fetch_and_explicit} should be identified as function call ++ bool IsFunctionCall = llvm::StringSwitch<bool>(Stem) ++ .Case("add", true) ++ .Case("sub", true) ++ .Case("inc", true) ++ .Case("dec", true) ++ .Case("cmpxchg", true) ++ .Case("min", true) ++ .Case("max", true) ++ .Case("or", true) ++ .Case("xor", true) ++ .Case("and", true) ++ .Case("fetch_or", true) ++ .Case("fetch_and", true) ++ .Case("fetch_xor", true) ++ .Case("fetch_or_explicit", true) ++ .Case("fetch_xor_explicit", true) ++ .Case("fetch_and_explicit", true) ++ .Default(false); ++ if (IsFunctionCall) ++ return; ++ if (F->arg_size() != 2) { ++ IsFunctionCall = llvm::StringSwitch<bool>(Stem) ++ .Case("exchange", true) ++ .Case("fetch_add", true) ++ .Case("fetch_sub", true) ++ .Case("fetch_min", true) ++ .Case("fetch_max", true) ++ .Case("load", true) ++ .Case("store", true) ++ .Default(false); ++ if (IsFunctionCall) ++ return; ++ } ++ if (F->arg_size() != 3 && F->arg_size() != 4) { ++ IsFunctionCall = llvm::StringSwitch<bool>(Stem) ++ .Case("exchange_explicit", true) ++ .Case("fetch_add_explicit", true) ++ .Case("fetch_sub_explicit", true) ++ .Case("fetch_min_explicit", true) ++ .Case("fetch_max_explicit", true) ++ .Case("load_explicit", true) ++ .Case("store_explicit", true) ++ .Default(false); ++ if (IsFunctionCall) ++ return; ++ } ++ } + + auto PCI = &CI; + if (DemangledName == kOCLBuiltinName::AtomicInit) { +@@ -839,7 +892,7 @@ void OCLToSPIRVBase::transAtomicBuiltin(CallInst *CI, + AttributeList Attrs = CI->getCalledFunction()->getAttributes(); + mutateCallInstSPIRV( + M, CI, +- [=](CallInst *CI, std::vector<Value *> &Args) { ++ [=](CallInst *CI, std::vector<Value *> &Args) -> std::string { + Info.PostProc(Args); + // Order of args in OCL20: + // object, 0-2 other args, 1-2 order, scope +@@ -868,7 +921,22 @@ void OCLToSPIRVBase::transAtomicBuiltin(CallInst *CI, + std::rotate(Args.begin() + 2, Args.begin() + OrderIdx, + Args.end() - Offset); + } +- return getSPIRVFuncName(OCLSPIRVBuiltinMap::map(Info.UniqName)); ++ ++ llvm::Type* AtomicBuiltinsReturnType = ++ CI->getCalledFunction()->getReturnType(); ++ auto IsFPType = [](llvm::Type *ReturnType) { ++ return ReturnType->isHalfTy() || ReturnType->isFloatTy() || ++ ReturnType->isDoubleTy(); ++ }; ++ auto SPIRVFunctionName = ++ getSPIRVFuncName(OCLSPIRVBuiltinMap::map(Info.UniqName)); ++ if (!IsFPType(AtomicBuiltinsReturnType)) ++ return SPIRVFunctionName; ++ // Translate FP-typed atomic builtins. ++ return llvm::StringSwitch<std::string>(SPIRVFunctionName) ++ .Case("__spirv_AtomicIAdd", "__spirv_AtomicFAddEXT") ++ .Case("__spirv_AtomicSMax", "__spirv_AtomicFMaxEXT") ++ .Case("__spirv_AtomicSMin", "__spirv_AtomicFMinEXT"); + }, + &Attrs); + } +diff --git a/lib/SPIRV/OCLUtil.cpp b/lib/SPIRV/OCLUtil.cpp +index 2de3f152..85155e39 100644 +--- a/lib/SPIRV/OCLUtil.cpp ++++ b/lib/SPIRV/OCLUtil.cpp +@@ -662,32 +662,6 @@ size_t getSPIRVAtomicBuiltinNumMemoryOrderArgs(Op OC) { + return 1; + } + +-bool isComputeAtomicOCLBuiltin(StringRef DemangledName) { +- if (!DemangledName.startswith(kOCLBuiltinName::AtomicPrefix) && +- !DemangledName.startswith(kOCLBuiltinName::AtomPrefix)) +- return false; +- +- return llvm::StringSwitch<bool>(DemangledName) +- .EndsWith("add", true) +- .EndsWith("sub", true) +- .EndsWith("inc", true) +- .EndsWith("dec", true) +- .EndsWith("cmpxchg", true) +- .EndsWith("min", true) +- .EndsWith("max", true) +- .EndsWith("and", true) +- .EndsWith("or", true) +- .EndsWith("xor", true) +- .EndsWith("add_explicit", true) +- .EndsWith("sub_explicit", true) +- .EndsWith("or_explicit", true) +- .EndsWith("xor_explicit", true) +- .EndsWith("and_explicit", true) +- .EndsWith("min_explicit", true) +- .EndsWith("max_explicit", true) +- .Default(false); +-} +- + BarrierLiterals getBarrierLiterals(CallInst *CI) { + auto N = CI->getNumArgOperands(); + assert(N == 1 || N == 2); +diff --git a/lib/SPIRV/OCLUtil.h b/lib/SPIRV/OCLUtil.h +index 4c05c672..c8577e9b 100644 +--- a/lib/SPIRV/OCLUtil.h ++++ b/lib/SPIRV/OCLUtil.h +@@ -394,10 +394,6 @@ size_t getAtomicBuiltinNumMemoryOrderArgs(StringRef Name); + /// Get number of memory order arguments for spirv atomic builtin function. + size_t getSPIRVAtomicBuiltinNumMemoryOrderArgs(Op OC); + +-/// Return true for OpenCL builtins which do compute operations +-/// (like add, sub, min, max, inc, dec, ...) atomically +-bool isComputeAtomicOCLBuiltin(StringRef DemangledName); +- + /// Get OCL version from metadata opencl.ocl.version. + /// \param AllowMulti Allows multiple operands if true. + /// \return OCL version encoded as Major*10^5+Minor*10^3+Rev, +diff --git a/test/negative/InvalidAtomicBuiltins.cl b/test/negative/InvalidAtomicBuiltins.cl +index b8ec5b89..23dcc4e3 100644 +--- a/test/negative/InvalidAtomicBuiltins.cl ++++ b/test/negative/InvalidAtomicBuiltins.cl +@@ -1,7 +1,9 @@ + // Check that translator doesn't generate atomic instructions for atomic builtins + // which are not defined in the spec. + +-// RUN: %clang_cc1 -triple spir -O1 -cl-std=cl2.0 -fdeclare-opencl-builtins -finclude-default-header %s -emit-llvm-bc -o %t.bc ++// To drop `fdeclare-opencl-builtins` option, since FP-typed atomic function ++// TableGen definitions have not been introduced. ++// RUN: %clang_cc1 -triple spir -O1 -cl-std=cl2.0 -finclude-default-header %s -emit-llvm-bc -o %t.bc + // RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s + // RUN: llvm-spirv %t.bc -o %t.spv + // RUN: spirv-val %t.spv +@@ -41,13 +43,9 @@ float __attribute__((overloadable)) atomic_fetch_xor(volatile generic atomic_flo + double __attribute__((overloadable)) atomic_fetch_and(volatile generic atomic_double *object, double operand, memory_order order); + double __attribute__((overloadable)) atomic_fetch_max(volatile generic atomic_double *object, double operand, memory_order order); + double __attribute__((overloadable)) atomic_fetch_min(volatile generic atomic_double *object, double operand, memory_order order); +-float __attribute__((overloadable)) atomic_fetch_add_explicit(volatile generic atomic_float *object, float operand, memory_order order); +-float __attribute__((overloadable)) atomic_fetch_sub_explicit(volatile generic atomic_float *object, float operand, memory_order order); + float __attribute__((overloadable)) atomic_fetch_or_explicit(volatile generic atomic_float *object, float operand, memory_order order); + float __attribute__((overloadable)) atomic_fetch_xor_explicit(volatile generic atomic_float *object, float operand, memory_order order); + double __attribute__((overloadable)) atomic_fetch_and_explicit(volatile generic atomic_double *object, double operand, memory_order order); +-double __attribute__((overloadable)) atomic_fetch_max_explicit(volatile generic atomic_double *object, double operand, memory_order order); +-double __attribute__((overloadable)) atomic_fetch_min_explicit(volatile generic atomic_double *object, double operand, memory_order order); + + __kernel void test_atomic_fn(volatile __global float *p, + volatile __global double *pp, +@@ -86,11 +84,7 @@ __kernel void test_atomic_fn(volatile __global float *p, + d = atomic_fetch_and(pp, val, order); + d = atomic_fetch_min(pp, val, order); + d = atomic_fetch_max(pp, val, order); +- f = atomic_fetch_add_explicit(p, val, order); +- f = atomic_fetch_sub_explicit(p, val, order); + f = atomic_fetch_or_explicit(p, val, order); + f = atomic_fetch_xor_explicit(p, val, order); + d = atomic_fetch_and_explicit(pp, val, order); +- d = atomic_fetch_min_explicit(pp, val, order); +- d = atomic_fetch_max_explicit(pp, val, order); + } +diff --git a/test/transcoding/AtomicFAddEXTForOCL.ll b/test/transcoding/AtomicFAddEXTForOCL.ll +new file mode 100644 +index 00000000..fb146fb9 +--- /dev/null ++++ b/test/transcoding/AtomicFAddEXTForOCL.ll +@@ -0,0 +1,64 @@ ++; RUN: llvm-as %s -o %t.bc ++; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_add -o %t.spv ++; RUN: spirv-val %t.spv ++; RUN: llvm-spirv -to-text %t.spv -o %t.spt ++; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV ++ ++; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc ++; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20 ++ ++; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc ++; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV ++ ++target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" ++target triple = "spir-unknown-unknown" ++ ++; CHECK-SPIRV: Capability AtomicFloat32AddEXT ++; CHECK-SPIRV: Capability AtomicFloat64AddEXT ++; CHECK-SPIRV: Extension "SPV_EXT_shader_atomic_float_add" ++; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32 ++; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64 ++ ++ ++; Function Attrs: convergent norecurse nounwind ++define dso_local spir_func void @test_atomic_float(float addrspace(1)* %a) local_unnamed_addr #0 { ++entry: ++ ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_32]] ++ ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_add_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}}) ++ ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+fiif]]({{.*}}) ++ %call = tail call spir_func float @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2 ++ ret void ++} ++ ++; Function Attrs: convergent ++declare spir_func float @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1 ++; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float ++ ++; Function Attrs: convergent norecurse nounwind ++define dso_local spir_func void @test_atomic_double(double addrspace(1)* %a) local_unnamed_addr #0 { ++entry: ++ ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_64]] ++ ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_add_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}}) ++ ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+diid]]({{.*}}) ++ %call = tail call spir_func double @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2 ++ ret void ++} ++; Function Attrs: convergent ++declare spir_func double @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)*, double, i32) local_unnamed_addr #1 ++; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double ++ ++; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float ++; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double ++ ++attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } ++attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } ++attributes #2 = { convergent nounwind } ++ ++!llvm.module.flags = !{!0} ++!opencl.ocl.version = !{!1} ++!opencl.spir.version = !{!1} ++!llvm.ident = !{!2} ++ ++!0 = !{i32 1, !"wchar_size", i32 4} ++!1 = !{i32 2, i32 0} ++!2 = !{!"clang version 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"} +diff --git a/test/transcoding/AtomicFMaxEXTForOCL.ll b/test/transcoding/AtomicFMaxEXTForOCL.ll +new file mode 100644 +index 00000000..1f2530d9 +--- /dev/null ++++ b/test/transcoding/AtomicFMaxEXTForOCL.ll +@@ -0,0 +1,64 @@ ++; RUN: llvm-as %s -o %t.bc ++; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_min_max -o %t.spv ++; RUN: spirv-val %t.spv ++; RUN: llvm-spirv -to-text %t.spv -o %t.spt ++; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV ++ ++; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc ++; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20 ++ ++; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc ++; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV ++ ++target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" ++target triple = "spir-unknown-unknown" ++ ++; CHECK-SPIRV: Capability AtomicFloat32MinMaxEXT ++; CHECK-SPIRV: Capability AtomicFloat64MinMaxEXT ++; CHECK-SPIRV: Extension "SPV_EXT_shader_atomic_float_min_max" ++; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32 ++; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64 ++ ++; Function Attrs: convergent norecurse nounwind ++define dso_local spir_func void @test_float(float addrspace(1)* %a) local_unnamed_addr #0 { ++entry: ++ ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_32]] ++ ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_max_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}}) ++ ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+fiif]]({{.*}}) ++ %call = tail call spir_func float @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2 ++ ret void ++} ++ ++; Function Attrs: convergent ++declare spir_func float @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1 ++; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float ++ ++; Function Attrs: convergent norecurse nounwind ++define dso_local spir_func void @test_double(double addrspace(1)* %a) local_unnamed_addr #0 { ++entry: ++ ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_64]] ++ ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_max_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}}) ++ ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+diid]]({{.*}}) ++ %call = tail call spir_func double @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2 ++ ret void ++} ++ ++; Function Attrs: convergent ++declare spir_func double @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)*, double, i32) local_unnamed_addr #1 ++; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double ++ ++; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float ++; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double ++ ++attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } ++attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } ++attributes #2 = { convergent nounwind } ++ ++!llvm.module.flags = !{!0} ++!opencl.ocl.version = !{!1} ++!opencl.spir.version = !{!1} ++!llvm.ident = !{!2} ++ ++!0 = !{i32 1, !"wchar_size", i32 4} ++!1 = !{i32 2, i32 0} ++!2 = !{!"clang version 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"} +diff --git a/test/transcoding/AtomicFMinEXTForOCL.ll b/test/transcoding/AtomicFMinEXTForOCL.ll +new file mode 100644 +index 00000000..6196b0f8 +--- /dev/null ++++ b/test/transcoding/AtomicFMinEXTForOCL.ll +@@ -0,0 +1,64 @@ ++; RUN: llvm-as %s -o %t.bc ++; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_min_max -o %t.spv ++; RUN: spirv-val %t.spv ++; RUN: llvm-spirv -to-text %t.spv -o %t.spt ++; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV ++ ++; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc ++; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20 ++ ++; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc ++; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV ++ ++target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" ++target triple = "spir-unknown-unknown" ++ ++; CHECK-SPIRV: Capability AtomicFloat32MinMaxEXT ++; CHECK-SPIRV: Capability AtomicFloat64MinMaxEXT ++; CHECK-SPIRV: Extension "SPV_EXT_shader_atomic_float_min_max" ++; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32 ++; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64 ++ ++; Function Attrs: convergent norecurse nounwind ++define dso_local spir_func void @test_float(float addrspace(1)* %a) local_unnamed_addr #0 { ++entry: ++ ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_32]] ++ ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_min_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}}) ++ ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+fiif]]({{.*}}) ++ %call = tail call spir_func float @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2 ++ ret void ++} ++ ++; Function Attrs: convergent ++declare spir_func float @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1 ++; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float ++ ++; Function Attrs: convergent norecurse nounwind ++define dso_local spir_func void @test_double(double addrspace(1)* %a) local_unnamed_addr #0 { ++entry: ++ ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_64]] ++ ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_min_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}}) ++ ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+diid]]({{.*}}) ++ %call = tail call spir_func double @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2 ++ ret void ++} ++ ++; Function Attrs: convergent ++declare spir_func double @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)*, double, i32) local_unnamed_addr #1 ++; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double ++ ++; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float ++; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double ++ ++attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } ++attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } ++attributes #2 = { convergent nounwind } ++ ++!llvm.module.flags = !{!0} ++!opencl.ocl.version = !{!1} ++!opencl.spir.version = !{!1} ++!llvm.ident = !{!2} ++ ++!0 = !{i32 1, !"wchar_size", i32 4} ++!1 = !{i32 2, i32 0} ++!2 = !{!"clang version 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"} +-- +2.17.1 + diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/0001-Memory-leak-fix-for-Managed-Static-Mutex.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0002-Memory-leak-fix-for-Managed-Static-Mutex.patch similarity index 89% rename from dynamic-layers/clang-layer/recipes-devtools/clang/files/0001-Memory-leak-fix-for-Managed-Static-Mutex.patch rename to dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0002-Memory-leak-fix-for-Managed-Static-Mutex.patch index 103dad5e..5b1f207e 100644 --- a/dynamic-layers/clang-layer/recipes-devtools/clang/files/0001-Memory-leak-fix-for-Managed-Static-Mutex.patch +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0002-Memory-leak-fix-for-Managed-Static-Mutex.patch @@ -1,7 +1,7 @@ -From c86c43b70e029b102543e8a85d269cbeb5c00279 Mon Sep 17 00:00:00 2001 +From ef27f1f99ad661c9604b7ff10efb1122466c508b Mon Sep 17 00:00:00 2001 From: juanrod2 <> Date: Tue, 22 Dec 2020 08:33:08 +0800 -Subject: [PATCH] Memory leak fix for Managed Static Mutex +Subject: [PATCH 2/6] Memory leak fix for Managed Static Mutex Upstream-Status: Backport [Taken from opencl-clang patches; https://github.com/intel/opencl-clang/blob/ocl-open-100/patches/llvm/0001-Memory-leak-fix-for-Managed-Static-Mutex.patch] @@ -31,5 +31,5 @@ index 053493f72fb5..6571580ccecf 100644 + ManagedStaticMutex = nullptr; } -- -2.29.2 +2.17.1 diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-Remove-repo-name-in-LLVM-IR.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0003-Remove-repo-name-in-LLVM-IR.patch similarity index 91% rename from dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-Remove-repo-name-in-LLVM-IR.patch rename to dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0003-Remove-repo-name-in-LLVM-IR.patch index 09089432..15c4f9e2 100644 --- a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-Remove-repo-name-in-LLVM-IR.patch +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0003-Remove-repo-name-in-LLVM-IR.patch @@ -1,18 +1,17 @@ -From ff0a6da84b94c16c4519c649f1f7bed3cdf89bbb Mon Sep 17 00:00:00 2001 +From a71ab6fb04b918b856f1dd802cfdb4a7ccd53290 Mon Sep 17 00:00:00 2001 From: Feng Zou <[email protected]> Date: Tue, 20 Oct 2020 11:29:04 +0800 -Subject: [PATCH] Remove repo name in LLVM IR +Subject: [PATCH 3/6] Remove repo name in LLVM IR Upstream-Status: Backport [Taken from opencl-clang patches, https://github.com/intel/opencl-clang/blob/ocl-open-110/patches/llvm/0002-Remove-repo-name-in-LLVM-IR.patch] Signed-off-by: Feng Zou <[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 18edbeabe3e..2d965263478 100644 +index 18edbeabe3e4..2d9652634787 100644 --- a/llvm/cmake/modules/VersionFromVCS.cmake +++ b/llvm/cmake/modules/VersionFromVCS.cmake @@ -33,17 +33,18 @@ function(get_source_info path revision repository) @@ -46,5 +45,5 @@ index 18edbeabe3e..2d965263478 100644 else() message(WARNING "Git not found. Version cannot be determined.") -- -2.18.1 +2.17.1 diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0004-Remove-__IMAGE_SUPPORT__-macro-for-SPIR-since-SPIR-d.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0004-Remove-__IMAGE_SUPPORT__-macro-for-SPIR-since-SPIR-d.patch new file mode 100644 index 00000000..25d88367 --- /dev/null +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0004-Remove-__IMAGE_SUPPORT__-macro-for-SPIR-since-SPIR-d.patch @@ -0,0 +1,51 @@ +From 546d9089fe5e21cccc671a0a89555cd4d5f8c817 Mon Sep 17 00:00:00 2001 +From: Naveen Saini <[email protected]> +Date: Thu, 19 Aug 2021 15:52:24 +0800 +Subject: [PATCH 4/6] Remove __IMAGE_SUPPORT__ macro for SPIR since SPIR + doesn't require image support + +Upstream-Status: Backport [Taken from opencl-clang patches; https://github.com/intel/opencl-clang/blob/ocl-open-110/patches/clang/0002-Remove-__IMAGE_SUPPORT__-macro-for-SPIR.patch] + +Signed-off-by: haonanya <[email protected]> +Signed-off-by: Naveen Saini <[email protected]> +--- + clang/lib/Frontend/InitPreprocessor.cpp | 3 --- + clang/test/Preprocessor/predefined-macros.c | 2 -- + 2 files changed, 5 deletions(-) + +diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp +index 5bb489c11909..cf3b48cb65d2 100644 +--- a/clang/lib/Frontend/InitPreprocessor.cpp ++++ b/clang/lib/Frontend/InitPreprocessor.cpp +@@ -1115,9 +1115,6 @@ static void InitializePredefinedMacros(const TargetInfo &TI, + if (TI.getSupportedOpenCLOpts().isSupported(#Ext)) \ + Builder.defineMacro(#Ext); + #include "clang/Basic/OpenCLExtensions.def" +- +- if (TI.getTriple().isSPIR()) +- Builder.defineMacro("__IMAGE_SUPPORT__"); + } + + if (TI.hasInt128Type() && LangOpts.CPlusPlus && LangOpts.GNUMode) { +diff --git a/clang/test/Preprocessor/predefined-macros.c b/clang/test/Preprocessor/predefined-macros.c +index 6c80517ec4d4..b5e5d7e2d546 100644 +--- a/clang/test/Preprocessor/predefined-macros.c ++++ b/clang/test/Preprocessor/predefined-macros.c +@@ -186,14 +186,12 @@ + + // RUN: %clang_cc1 %s -E -dM -o - -x cl -triple spir-unknown-unknown \ + // RUN: | FileCheck -match-full-lines %s --check-prefix=CHECK-SPIR +-// CHECK-SPIR-DAG: #define __IMAGE_SUPPORT__ 1 + // CHECK-SPIR-DAG: #define __SPIR__ 1 + // CHECK-SPIR-DAG: #define __SPIR32__ 1 + // CHECK-SPIR-NOT: #define __SPIR64__ 1 + + // RUN: %clang_cc1 %s -E -dM -o - -x cl -triple spir64-unknown-unknown \ + // RUN: | FileCheck -match-full-lines %s --check-prefix=CHECK-SPIR64 +-// CHECK-SPIR64-DAG: #define __IMAGE_SUPPORT__ 1 + // CHECK-SPIR64-DAG: #define __SPIR__ 1 + // CHECK-SPIR64-DAG: #define __SPIR64__ 1 + // CHECK-SPIR64-NOT: #define __SPIR32__ 1 +-- +2.17.1 + diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0005-Avoid-calling-ParseCommandLineOptions-in-BackendUtil.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0005-Avoid-calling-ParseCommandLineOptions-in-BackendUtil.patch new file mode 100644 index 00000000..2b86532c --- /dev/null +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0005-Avoid-calling-ParseCommandLineOptions-in-BackendUtil.patch @@ -0,0 +1,52 @@ +From 747e48959e18ac8b586078a82472a0799d12925c Mon Sep 17 00:00:00 2001 +From: Raphael Isemann <[email protected]> +Date: Thu, 1 Apr 2021 18:41:44 +0200 +Subject: [PATCH 5/6] Avoid calling ParseCommandLineOptions in BackendUtil if + possible + +Calling `ParseCommandLineOptions` should only be called from `main` as the +CommandLine setup code isn't thread-safe. As BackendUtil is part of the +generic Clang FrontendAction logic, a process which has several threads executing +Clang FrontendActions will randomly crash in the unsafe setup code. + +This patch avoids calling the function unless either the debug-pass option or +limit-float-precision option is set. Without these two options set the +`ParseCommandLineOptions` call doesn't do anything beside parsing +the command line `clang` which doesn't set any options. + +See also D99652 where LLDB received a workaround for this crash. + +Reviewed By: JDevlieghere + +Differential Revision: https://reviews.llvm.org/D99740 + +Upstream-Status: Backport [Taken from opencl-clang patches; https://github.com/intel/opencl-clang/blob/ocl-open-110/patches/clang/0003-Avoid-calling-ParseCommandLineOptions-in-BackendUtil.patch] + +Signed-off-by: Naveen Saini <[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 dce0940670a2..ab478090ed1c 100644 +--- a/clang/lib/CodeGen/BackendUtil.cpp ++++ b/clang/lib/CodeGen/BackendUtil.cpp +@@ -797,7 +797,15 @@ static void setCommandLineOpts(const CodeGenOptions &CodeGenOpts) { + BackendArgs.push_back("-limit-float-precision"); + BackendArgs.push_back(CodeGenOpts.LimitFloatPrecision.c_str()); + } ++ // Check for the default "clang" invocation that won't set any cl::opt values. ++ // Skip trying to parse the command line invocation to avoid the issues ++ // described below. ++ if (BackendArgs.size() == 1) ++ return; + BackendArgs.push_back(nullptr); ++ // FIXME: The command line parser below is not thread-safe and shares a global ++ // state, so this call might crash or overwrite the options of another Clang ++ // instance in the same process. + llvm::cl::ParseCommandLineOptions(BackendArgs.size() - 1, + BackendArgs.data()); + } +-- +2.17.1 + diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0006-OpenCL-support-cl_ext_float_atomics.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0006-OpenCL-support-cl_ext_float_atomics.patch new file mode 100644 index 00000000..0178fd43 --- /dev/null +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm11-0006-OpenCL-support-cl_ext_float_atomics.patch @@ -0,0 +1,353 @@ +From a1b924d76cdacfa3f9dbb79a9e3edddcd75f61ca Mon Sep 17 00:00:00 2001 +From: Naveen Saini <[email protected]> +Date: Thu, 19 Aug 2021 16:06:33 +0800 +Subject: [PATCH 6/6] [OpenCL] support cl_ext_float_atomics + +Upstream-Status: Backport [Taken from opencl-clang patches; https://github.com/intel/opencl-clang/blob/ocl-open-110/patches/clang/0004-OpenCL-support-cl_ext_float_atomics.patch] + +Signed-off-by: haonanya <[email protected]> +Signed-off-by: Naveen Saini <[email protected]> +--- + clang/lib/Headers/opencl-c-base.h | 25 ++++ + clang/lib/Headers/opencl-c.h | 195 ++++++++++++++++++++++++++ + clang/test/Headers/opencl-c-header.cl | 85 +++++++++++ + 3 files changed, 305 insertions(+) + +diff --git a/clang/lib/Headers/opencl-c-base.h b/clang/lib/Headers/opencl-c-base.h +index afa900ab24d9..9a3ee8529acf 100644 +--- a/clang/lib/Headers/opencl-c-base.h ++++ b/clang/lib/Headers/opencl-c-base.h +@@ -62,6 +62,31 @@ + #endif + #endif // defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ == CL_VERSION_2_0) + ++#if (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200) ++// For SPIR all extensions are supported. ++#if defined(__SPIR__) ++#define cl_ext_float_atomics ++#ifdef cl_khr_fp16 ++#define __opencl_c_ext_fp16_global_atomic_load_store 1 ++#define __opencl_c_ext_fp16_local_atomic_load_store 1 ++#define __opencl_c_ext_fp16_global_atomic_add 1 ++#define __opencl_c_ext_fp16_local_atomic_add 1 ++#define __opencl_c_ext_fp16_global_atomic_min_max 1 ++#define __opencl_c_ext_fp16_local_atomic_min_max 1 ++#endif ++#ifdef __opencl_c_fp64 ++#define __opencl_c_ext_fp64_global_atomic_add 1 ++#define __opencl_c_ext_fp64_local_atomic_add 1 ++#define __opencl_c_ext_fp64_global_atomic_min_max 1 ++#define __opencl_c_ext_fp64_local_atomic_min_max 1 ++#endif ++#define __opencl_c_ext_fp32_global_atomic_add 1 ++#define __opencl_c_ext_fp32_local_atomic_add 1 ++#define __opencl_c_ext_fp32_global_atomic_min_max 1 ++#define __opencl_c_ext_fp32_local_atomic_min_max 1 ++#endif // defined(__SPIR__) ++#endif // (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200) ++ + // built-in scalar data types: + + /** +diff --git a/clang/lib/Headers/opencl-c.h b/clang/lib/Headers/opencl-c.h +index 67d900eb1c3d..bda0f5c6df80 100644 +--- a/clang/lib/Headers/opencl-c.h ++++ b/clang/lib/Headers/opencl-c.h +@@ -14354,6 +14354,201 @@ intptr_t __ovld atomic_fetch_max_explicit( + // defined(cl_khr_int64_extended_atomics) + #endif // (__OPENCL_C_VERSION__ >= CL_VERSION_3_0) + ++#if defined(cl_ext_float_atomics) ++ ++#if defined(__opencl_c_ext_fp32_global_atomic_min_max) ++float __ovld atomic_fetch_min(volatile __global atomic_float *object, ++ float operand); ++float __ovld atomic_fetch_max(volatile __global atomic_float *object, ++ float operand); ++float __ovld atomic_fetch_min_explicit(volatile __global atomic_float *object, ++ float operand, memory_order order); ++float __ovld atomic_fetch_max_explicit(volatile __global atomic_float *object, ++ float operand, memory_order order); ++float __ovld atomic_fetch_min_explicit(volatile __global atomic_float *object, ++ float operand, memory_order order, ++ memory_scope scope); ++float __ovld atomic_fetch_max_explicit(volatile __global atomic_float *object, ++ float operand, memory_order order, ++ memory_scope scope); ++#endif ++#if defined(__opencl_c_ext_fp32_local_atomic_min_max) ++float __ovld atomic_fetch_min(volatile __local atomic_float *object, ++ float operand); ++float __ovld atomic_fetch_max(volatile __local atomic_float *object, ++ float operand); ++float __ovld atomic_fetch_min_explicit(volatile __local atomic_float *object, ++ float operand, memory_order order); ++float __ovld atomic_fetch_max_explicit(volatile __local atomic_float *object, ++ float operand, memory_order order); ++float __ovld atomic_fetch_min_explicit(volatile __local atomic_float *object, ++ float operand, memory_order order, ++ memory_scope scope); ++float __ovld atomic_fetch_max_explicit(volatile __local atomic_float *object, ++ float operand, memory_order order, ++ memory_scope scope); ++#endif ++#if defined(__opencl_c_ext_fp32_global_atomic_min_max) || \ ++ defined(__opencl_c_ext_fp32_local_atomic_min_max) ++float __ovld atomic_fetch_min(volatile atomic_float *object, float operand); ++float __ovld atomic_fetch_max(volatile atomic_float *object, float operand); ++float __ovld atomic_fetch_min_explicit(volatile atomic_float *object, ++ float operand, memory_order order); ++float __ovld atomic_fetch_max_explicit(volatile atomic_float *object, ++ float operand, memory_order order); ++float __ovld atomic_fetch_min_explicit(volatile atomic_float *object, ++ float operand, memory_order order, ++ memory_scope scope); ++float __ovld atomic_fetch_max_explicit(volatile atomic_float *object, ++ float operand, memory_order order, ++ memory_scope scope); ++#endif ++#if defined(__opencl_c_ext_fp64_global_atomic_min_max) ++double __ovld atomic_fetch_min(volatile __global atomic_double *object, ++ double operand); ++double __ovld atomic_fetch_max(volatile __global atomic_double *object, ++ double operand); ++double __ovld atomic_fetch_min_explicit(volatile __global atomic_double *object, ++ double operand, memory_order order); ++double __ovld atomic_fetch_max_explicit(volatile __global atomic_double *object, ++ double operand, memory_order order); ++double __ovld atomic_fetch_min_explicit(volatile __global atomic_double *object, ++ double operand, memory_order order, ++ memory_scope scope); ++double __ovld atomic_fetch_max_explicit(volatile __global atomic_double *object, ++ double operand, memory_order order, ++ memory_scope scope); ++#endif ++#if defined(__opencl_c_ext_fp64_local_atomic_min_max) ++double __ovld atomic_fetch_min(volatile __local atomic_double *object, ++ double operand); ++double __ovld atomic_fetch_max(volatile __local atomic_double *object, ++ double operand); ++double __ovld atomic_fetch_min_explicit(volatile __local atomic_double *object, ++ double operand, memory_order order); ++double __ovld atomic_fetch_max_explicit(volatile __local atomic_double *object, ++ double operand, memory_order order); ++double __ovld atomic_fetch_min_explicit(volatile __local atomic_double *object, ++ double operand, memory_order order, ++ memory_scope scope); ++double __ovld atomic_fetch_max_explicit(volatile __local atomic_double *object, ++ double operand, memory_order order, ++ memory_scope scope); ++#endif ++#if defined(__opencl_c_ext_fp64_global_atomic_min_max) || \ ++ defined(__opencl_c_ext_fp64_local_atomic_min_max) ++double __ovld atomic_fetch_min(volatile atomic_double *object, double operand); ++double __ovld atomic_fetch_max(volatile atomic_double *object, double operand); ++double __ovld atomic_fetch_min_explicit(volatile atomic_double *object, ++ double operand, memory_order order); ++double __ovld atomic_fetch_max_explicit(volatile atomic_double *object, ++ double operand, memory_order order); ++double __ovld atomic_fetch_min_explicit(volatile atomic_double *object, ++ double operand, memory_order order, ++ memory_scope scope); ++double __ovld atomic_fetch_max_explicit(volatile atomic_double *object, ++ double operand, memory_order order, ++ memory_scope scope); ++#endif ++ ++#if defined(__opencl_c_ext_fp32_global_atomic_add) ++float __ovld atomic_fetch_add(volatile __global atomic_float *object, ++ float operand); ++float __ovld atomic_fetch_sub(volatile __global atomic_float *object, ++ float operand); ++float __ovld atomic_fetch_add_explicit(volatile __global atomic_float *object, ++ float operand, memory_order order); ++float __ovld atomic_fetch_sub_explicit(volatile __global atomic_float *object, ++ float operand, memory_order order); ++float __ovld atomic_fetch_add_explicit(volatile __global atomic_float *object, ++ float operand, memory_order order, ++ memory_scope scope); ++float __ovld atomic_fetch_sub_explicit(volatile __global atomic_float *object, ++ float operand, memory_order order, ++ memory_scope scope); ++#endif ++#if defined(__opencl_c_ext_fp32_local_atomic_add) ++float __ovld atomic_fetch_add(volatile __local atomic_float *object, ++ float operand); ++float __ovld atomic_fetch_sub(volatile __local atomic_float *object, ++ float operand); ++float __ovld atomic_fetch_add_explicit(volatile __local atomic_float *object, ++ float operand, memory_order order); ++float __ovld atomic_fetch_sub_explicit(volatile __local atomic_float *object, ++ float operand, memory_order order); ++float __ovld atomic_fetch_add_explicit(volatile __local atomic_float *object, ++ float operand, memory_order order, ++ memory_scope scope); ++float __ovld atomic_fetch_sub_explicit(volatile __local atomic_float *object, ++ float operand, memory_order order, ++ memory_scope scope); ++#endif ++#if defined(__opencl_c_ext_fp32_global_atomic_add) || \ ++ defined(__opencl_c_ext_fp32_local_atomic_add) ++float __ovld atomic_fetch_add(volatile atomic_float *object, float operand); ++float __ovld atomic_fetch_sub(volatile atomic_float *object, float operand); ++float __ovld atomic_fetch_add_explicit(volatile atomic_float *object, ++ float operand, memory_order order); ++float __ovld atomic_fetch_sub_explicit(volatile atomic_float *object, ++ float operand, memory_order order); ++float __ovld atomic_fetch_add_explicit(volatile atomic_float *object, ++ float operand, memory_order order, ++ memory_scope scope); ++float __ovld atomic_fetch_sub_explicit(volatile atomic_float *object, ++ float operand, memory_order order, ++ memory_scope scope); ++#endif ++ ++#if defined(__opencl_c_ext_fp64_global_atomic_add) ++double __ovld atomic_fetch_add(volatile __global atomic_double *object, ++ double operand); ++double __ovld atomic_fetch_sub(volatile __global atomic_double *object, ++ double operand); ++double __ovld atomic_fetch_add_explicit(volatile __global atomic_double *object, ++ double operand, memory_order order); ++double __ovld atomic_fetch_sub_explicit(volatile __global atomic_double *object, ++ double operand, memory_order order); ++double __ovld atomic_fetch_add_explicit(volatile __global atomic_double *object, ++ double operand, memory_order order, ++ memory_scope scope); ++double __ovld atomic_fetch_sub_explicit(volatile __global atomic_double *object, ++ double operand, memory_order order, ++ memory_scope scope); ++#endif ++#if defined(__opencl_c_ext_fp64_local_atomic_add) ++double __ovld atomic_fetch_add(volatile __local atomic_double *object, ++ double operand); ++double __ovld atomic_fetch_sub(volatile __local atomic_double *object, ++ double operand); ++double __ovld atomic_fetch_add_explicit(volatile __local atomic_double *object, ++ double operand, memory_order order); ++double __ovld atomic_fetch_sub_explicit(volatile __local atomic_double *object, ++ double operand, memory_order order); ++double __ovld atomic_fetch_add_explicit(volatile __local atomic_double *object, ++ double operand, memory_order order, ++ memory_scope scope); ++double __ovld atomic_fetch_sub_explicit(volatile __local atomic_double *object, ++ double operand, memory_order order, ++ memory_scope scope); ++#endif ++#if defined(__opencl_c_ext_fp64_global_atomic_add) || \ ++ defined(__opencl_c_ext_fp64_local_atomic_add) ++double __ovld atomic_fetch_add(volatile atomic_double *object, double operand); ++double __ovld atomic_fetch_sub(volatile atomic_double *object, double operand); ++double __ovld atomic_fetch_add_explicit(volatile atomic_double *object, ++ double operand, memory_order order); ++double __ovld atomic_fetch_sub_explicit(volatile atomic_double *object, ++ double operand, memory_order order); ++double __ovld atomic_fetch_add_explicit(volatile atomic_double *object, ++ double operand, memory_order order, ++ memory_scope scope); ++double __ovld atomic_fetch_sub_explicit(volatile atomic_double *object, ++ double operand, memory_order order, ++ memory_scope scope); ++#endif ++ ++#endif // cl_ext_float_atomics ++ + // atomic_store() + + #if defined(__opencl_c_atomic_scope_device) && \ +diff --git a/clang/test/Headers/opencl-c-header.cl b/clang/test/Headers/opencl-c-header.cl +index 2716076acdcf..6b3eca84e8b9 100644 +--- a/clang/test/Headers/opencl-c-header.cl ++++ b/clang/test/Headers/opencl-c-header.cl +@@ -98,3 +98,88 @@ global atomic_int z = ATOMIC_VAR_INIT(99); + #pragma OPENCL EXTENSION cl_intel_planar_yuv : enable + + // CHECK-MOD: Reading modules ++ ++// For SPIR all extensions are supported. ++#if defined(__SPIR__) ++ ++#if (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200) ++ ++#if __opencl_c_ext_fp16_global_atomic_load_store != 1 ++#error "Incorrectly defined __opencl_c_ext_fp16_global_atomic_load_store" ++#endif ++#if __opencl_c_ext_fp16_local_atomic_load_store != 1 ++#error "Incorrectly defined __opencl_c_ext_fp16_local_atomic_load_store" ++#endif ++#if __opencl_c_ext_fp16_global_atomic_add != 1 ++#error "Incorrectly defined __opencl_c_ext_fp16_global_atomic_add" ++#endif ++#if __opencl_c_ext_fp32_global_atomic_add != 1 ++#error "Incorrectly defined __opencl_c_ext_fp32_global_atomic_add" ++#endif ++#if __opencl_c_ext_fp16_local_atomic_add != 1 ++#error "Incorrectly defined __opencl_c_ext_fp16_local_atomic_add" ++#endif ++#if __opencl_c_ext_fp32_local_atomic_add != 1 ++#error "Incorrectly defined __opencl_c_ext_fp32_local_atomic_add" ++#endif ++#if __opencl_c_ext_fp16_global_atomic_min_max != 1 ++#error "Incorrectly defined __opencl_c_ext_fp16_global_atomic_min_max" ++#endif ++#if __opencl_c_ext_fp32_global_atomic_min_max != 1 ++#error "Incorrectly defined __opencl_c_ext_fp32_global_atomic_min_max" ++#endif ++#if __opencl_c_ext_fp16_local_atomic_min_max != 1 ++#error "Incorrectly defined __opencl_c_ext_fp16_local_atomic_min_max" ++#endif ++#if __opencl_c_ext_fp32_local_atomic_min_max != 1 ++#error "Incorrectly defined __opencl_c_ext_fp32_local_atomic_min_max" ++#endif ++ ++#else ++#ifdef __opencl_c_ext_fp16_global_atomic_load_store ++#error "Incorrectly __opencl_c_ext_fp16_global_atomic_load_store defined" ++#endif ++#ifdef __opencl_c_ext_fp16_local_atomic_load_store ++#error "Incorrectly __opencl_c_ext_fp16_local_atomic_load_store defined" ++#endif ++#ifdef __opencl_c_ext_fp16_global_atomic_add ++#error "Incorrectly __opencl_c_ext_fp16_global_atomic_add defined" ++#endif ++#ifdef __opencl_c_ext_fp32_global_atomic_add ++#error "Incorrectly __opencl_c_ext_fp32_global_atomic_add defined" ++#endif ++#ifdef __opencl_c_ext_fp64_global_atomic_add ++#error "Incorrectly __opencl_c_ext_fp64_global_atomic_add defined" ++#endif ++#ifdef __opencl_c_ext_fp16_local_atomic_add ++#error "Incorrectly __opencl_c_ext_fp16_local_atomic_add defined" ++#endif ++#ifdef __opencl_c_ext_fp32_local_atomic_add ++#error "Incorrectly __opencl_c_ext_fp32_local_atomic_add defined" ++#endif ++#ifdef __opencl_c_ext_fp64_local_atomic_add ++#error "Incorrectly __opencl_c_ext_fp64_local_atomic_add defined" ++#endif ++#ifdef __opencl_c_ext_fp16_global_atomic_min_max ++#error "Incorrectly __opencl_c_ext_fp16_global_atomic_min_max defined" ++#endif ++#ifdef __opencl_c_ext_fp32_global_atomic_min_max ++#error "Incorrectly __opencl_c_ext_fp32_global_atomic_min_max defined" ++#endif ++#ifdef __opencl_c_ext_fp64_global_atomic_min_max ++#error "Incorrectly __opencl_c_ext_fp64_global_atomic_min_max defined" ++#endif ++#ifdef __opencl_c_ext_fp16_local_atomic_min_max ++#error "Incorrectly __opencl_c_ext_fp16_local_atomic_min_max defined" ++#endif ++#ifdef __opencl_c_ext_fp32_local_atomic_min_max ++#error "Incorrectly __opencl_c_ext_fp32_local_atomic_min_max defined" ++#endif ++#ifdef __opencl_c_ext_fp64_local_atomic_min_max ++#error "Incorrectly __opencl_c_ext_fp64_local_atomic_min_max defined" ++#endif ++ ++#endif //(defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200) ++ ++#endif // defined(__SPIR__) ++ +-- +2.17.1 + diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/llvm-project-source.bbappend b/dynamic-layers/clang-layer/recipes-devtools/clang/llvm-project-source.bbappend index ac34321c..84192714 100644 --- a/dynamic-layers/clang-layer/recipes-devtools/clang/llvm-project-source.bbappend +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/llvm-project-source.bbappend @@ -1,7 +1,7 @@ FILESEXTRAPATHS:prepend:intel-x86-common := "${THISDIR}/files:" SPIRV10_SRCREV = "fe4d6b767363a1995ccbfca27f79efb10dcfe110" -SPIRV11_SRCREV = "2a8c1e6c9778deaa720a23e08c293006dc5d56fd" +SPIRV11_SRCREV = "ca3a50e6e3193e399d26446d4f74a90e2a531f3a" SPIRV_SRCREV = "${@bb.utils.contains('LLVMVERSION', '10.0.1', '${SPIRV10_SRCREV}', '${SPIRV11_SRCREV}', d)}" @@ -21,10 +21,14 @@ SRC_URI_LLVM10_PATCHES = " \ " SRC_URI_LLVM11_PATCHES = " \ - file://llvm11-skip-building-tests.patch;patchdir=llvm/projects/llvm-spirv \ - file://llvm11-OpenCL-3.0-support.patch \ - file://0001-Memory-leak-fix-for-Managed-Static-Mutex.patch \ - file://llvm11-Remove-repo-name-in-LLVM-IR.patch \ + file://llvm11-0001-llvm-spirv-skip-building-tests.patch;patchdir=llvm/projects/llvm-spirv \ + file://llvm11-0002-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch;patchdir=llvm/projects/llvm-spirv \ + file://llvm11-0001-OpenCL-3.0-support.patch \ + file://llvm11-0002-Memory-leak-fix-for-Managed-Static-Mutex.patch \ + file://llvm11-0003-Remove-repo-name-in-LLVM-IR.patch \ + file://llvm11-0004-Remove-__IMAGE_SUPPORT__-macro-for-SPIR-since-SPIR-d.patch \ + file://llvm11-0005-Avoid-calling-ParseCommandLineOptions-in-BackendUtil.patch \ + file://llvm11-0006-OpenCL-support-cl_ext_float_atomics.patch \ " SRC_URI_LLVM12_PATCHES = " \ file://0001-Remove-__IMAGE_SUPPORT__-macro-for-SPIR.patch \ -- 2.17.1
-=-=-=-=-=-=-=-=-=-=-=- Links: You receive all messages sent to this group. View/Reply Online (#7204): https://lists.yoctoproject.org/g/meta-intel/message/7204 Mute This Topic: https://lists.yoctoproject.org/mt/85011614/21656 Group Owner: [email protected] Unsubscribe: https://lists.yoctoproject.org/g/meta-intel/unsub [[email protected]] -=-=-=-=-=-=-=-=-=-=-=-
