mirror of
git://git.yoctoproject.org/meta-intel.git
synced 2025-07-19 12:59:03 +02:00
Remove support for gatesgarth
Building with oe-core gatesgarth is no longer supported. Remove from LAYERSERIES_COMPAT and remove the LLVM 11 patches as well. Signed-off-by: Anuj Mittal <anuj.mittal@intel.com>
This commit is contained in:
Anuj Mittal
parent
467d15d57a
commit
a809b8c531
|
|
@ -18,7 +18,7 @@ LAYERRECOMMENDS_intel = "dpdk"
|
|||
# This should only be incremented on significant changes that will
|
||||
# cause compatibility issues with other layers
|
||||
LAYERVERSION_intel = "5"
|
||||
LAYERSERIES_COMPAT_intel = "dunfell gatesgarth hardknott honister"
|
||||
LAYERSERIES_COMPAT_intel = "dunfell hardknott honister"
|
||||
|
||||
BBFILES_DYNAMIC += " \
|
||||
clang-layer:${LAYERDIR}/dynamic-layers/clang-layer/*/*/*.bb \
|
||||
|
|
|
|||
|
|
@ -27,11 +27,9 @@ PREFERRED_PROVIDER_libva-utils-native = "libva-intel-utils-native"
|
|||
PREFERRED_PROVIDER_nativesdk-libva-utils = "nativesdk-libva-intel-utils"
|
||||
|
||||
PREFERRED_VERSION_opencl-clang ?= "${@bb.utils.contains('LLVMVERSION', '10.0.1', '10.0.0', \
|
||||
bb.utils.contains('LLVMVERSION', '11.1.0', '11.0.0', \
|
||||
bb.utils.contains('LLVMVERSION', '12.0.0', '12.0.0', '13.0.0', d), d), d)}"
|
||||
bb.utils.contains('LLVMVERSION', '12.0.0', '12.0.0', '13.0.0', d), d)}"
|
||||
PREFERRED_VERSION_opencl-clang-native ?= "${@bb.utils.contains('LLVMVERSION', '10.0.1', '10.0.0', \
|
||||
bb.utils.contains('LLVMVERSION', '11.1.0', '11.0.0', \
|
||||
bb.utils.contains('LLVMVERSION', '12.0.0', '12.0.0', '13.0.0', d), d), d)}"
|
||||
bb.utils.contains('LLVMVERSION', '12.0.0', '12.0.0', '13.0.0', d), d)}"
|
||||
|
||||
XSERVER_X86_ASPEED_AST = "xf86-video-ast \
|
||||
"
|
||||
|
|
|
|||
File diff suppressed because it is too large
Load Diff
|
|
@ -1,51 +0,0 @@
|
|||
From 6690d77f9007ce82984dc1b6ae12585cb3e04785 Mon Sep 17 00:00:00 2001
|
||||
From: Naveen Saini <naveen.kumar.saini@intel.com>
|
||||
Date: Wed, 21 Aug 2019 14:35:31 +0800
|
||||
Subject: [PATCH 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
|
||||
in-tree.
|
||||
|
||||
Instead of using SPIRV_SKIP_CLANG_BUILD to skip clang build and adding this to
|
||||
all components, disable the building of tests altogether.
|
||||
|
||||
Upstream-Status: Inappropriate
|
||||
|
||||
Signed-off-by: Anuj Mittal <anuj.mittal@intel.com>
|
||||
Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com>
|
||||
---
|
||||
CMakeLists.txt | 10 ----------
|
||||
1 file changed, 10 deletions(-)
|
||||
|
||||
diff --git a/CMakeLists.txt b/CMakeLists.txt
|
||||
index ec61fb95..d723c0a5 100644
|
||||
--- a/CMakeLists.txt
|
||||
+++ b/CMakeLists.txt
|
||||
@@ -26,13 +26,6 @@ if(LLVM_SPIRV_BUILD_EXTERNAL)
|
||||
set(CMAKE_CXX_STANDARD 14)
|
||||
set(CMAKE_CXX_STANDARD_REQUIRED ON)
|
||||
|
||||
- if(LLVM_SPIRV_INCLUDE_TESTS)
|
||||
- set(LLVM_TEST_COMPONENTS
|
||||
- llvm-as
|
||||
- llvm-dis
|
||||
- )
|
||||
- endif(LLVM_SPIRV_INCLUDE_TESTS)
|
||||
-
|
||||
find_package(LLVM ${BASE_LLVM_VERSION} REQUIRED
|
||||
COMPONENTS
|
||||
Analysis
|
||||
@@ -65,9 +58,6 @@ set(LLVM_SPIRV_INCLUDE_DIRS ${CMAKE_CURRENT_SOURCE_DIR}/include)
|
||||
|
||||
add_subdirectory(lib/SPIRV)
|
||||
add_subdirectory(tools/llvm-spirv)
|
||||
-if(LLVM_SPIRV_INCLUDE_TESTS)
|
||||
- add_subdirectory(test)
|
||||
-endif(LLVM_SPIRV_INCLUDE_TESTS)
|
||||
|
||||
install(
|
||||
FILES
|
||||
--
|
||||
2.17.1
|
||||
|
||||
|
|
@ -1,433 +0,0 @@
|
|||
From 8e12d8fb3cdbdafca73fe8ed4f0cde773b1788b4 Mon Sep 17 00:00:00 2001
|
||||
From: haonanya <haonan.yang@intel.com>
|
||||
Date: Wed, 28 Jul 2021 11:43:20 +0800
|
||||
Subject: [PATCH 2/2] Add support for cl_ext_float_atomics in SPIRVWriter
|
||||
|
||||
Upstream-Status: Backport [Taken from opencl-clang patches, https://github.com/intel/opencl-clang/blob/ocl-open-110/patches/spirv/0001-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch]
|
||||
|
||||
Signed-off-by: haonanya <haonan.yang@intel.com>
|
||||
Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com>
|
||||
---
|
||||
lib/SPIRV/OCLToSPIRV.cpp | 80 +++++++++++++++++++++++--
|
||||
lib/SPIRV/OCLUtil.cpp | 26 --------
|
||||
lib/SPIRV/OCLUtil.h | 4 --
|
||||
test/negative/InvalidAtomicBuiltins.cl | 12 +---
|
||||
test/transcoding/AtomicFAddEXTForOCL.ll | 64 ++++++++++++++++++++
|
||||
test/transcoding/AtomicFMaxEXTForOCL.ll | 64 ++++++++++++++++++++
|
||||
test/transcoding/AtomicFMinEXTForOCL.ll | 64 ++++++++++++++++++++
|
||||
7 files changed, 269 insertions(+), 45 deletions(-)
|
||||
create mode 100644 test/transcoding/AtomicFAddEXTForOCL.ll
|
||||
create mode 100644 test/transcoding/AtomicFMaxEXTForOCL.ll
|
||||
create mode 100644 test/transcoding/AtomicFMinEXTForOCL.ll
|
||||
|
||||
diff --git a/lib/SPIRV/OCLToSPIRV.cpp b/lib/SPIRV/OCLToSPIRV.cpp
|
||||
index 04d51586..f00f5f7b 100644
|
||||
--- a/lib/SPIRV/OCLToSPIRV.cpp
|
||||
+++ b/lib/SPIRV/OCLToSPIRV.cpp
|
||||
@@ -421,10 +421,63 @@ void OCLToSPIRVBase::visitCallInst(CallInst &CI) {
|
||||
if (DemangledName.find(kOCLBuiltinName::AtomicPrefix) == 0 ||
|
||||
DemangledName.find(kOCLBuiltinName::AtomPrefix) == 0) {
|
||||
|
||||
- // Compute atomic builtins do not support floating types.
|
||||
- if (CI.getType()->isFloatingPointTy() &&
|
||||
- isComputeAtomicOCLBuiltin(DemangledName))
|
||||
- return;
|
||||
+ // Compute "atom" prefixed builtins do not support floating types.
|
||||
+ if (CI.getType()->isFloatingPointTy()) {
|
||||
+ if (DemangledName.find(kOCLBuiltinName::AtomPrefix) == 0)
|
||||
+ return;
|
||||
+ // handle functions which are "atomic_" prefixed.
|
||||
+ StringRef Stem = DemangledName;
|
||||
+ Stem = Stem.drop_front(strlen("atomic_"));
|
||||
+ // FP-typed atomic_{add, sub, inc, dec, exchange, min, max, or, and, xor,
|
||||
+ // fetch_or, fetch_xor, fetch_and, fetch_or_explicit, fetch_xor_explicit,
|
||||
+ // fetch_and_explicit} should be identified as function call
|
||||
+ bool IsFunctionCall = llvm::StringSwitch<bool>(Stem)
|
||||
+ .Case("add", true)
|
||||
+ .Case("sub", true)
|
||||
+ .Case("inc", true)
|
||||
+ .Case("dec", true)
|
||||
+ .Case("cmpxchg", true)
|
||||
+ .Case("min", true)
|
||||
+ .Case("max", true)
|
||||
+ .Case("or", true)
|
||||
+ .Case("xor", true)
|
||||
+ .Case("and", true)
|
||||
+ .Case("fetch_or", true)
|
||||
+ .Case("fetch_and", true)
|
||||
+ .Case("fetch_xor", true)
|
||||
+ .Case("fetch_or_explicit", true)
|
||||
+ .Case("fetch_xor_explicit", true)
|
||||
+ .Case("fetch_and_explicit", true)
|
||||
+ .Default(false);
|
||||
+ if (IsFunctionCall)
|
||||
+ return;
|
||||
+ if (F->arg_size() != 2) {
|
||||
+ IsFunctionCall = llvm::StringSwitch<bool>(Stem)
|
||||
+ .Case("exchange", true)
|
||||
+ .Case("fetch_add", true)
|
||||
+ .Case("fetch_sub", true)
|
||||
+ .Case("fetch_min", true)
|
||||
+ .Case("fetch_max", true)
|
||||
+ .Case("load", true)
|
||||
+ .Case("store", true)
|
||||
+ .Default(false);
|
||||
+ if (IsFunctionCall)
|
||||
+ return;
|
||||
+ }
|
||||
+ if (F->arg_size() != 3 && F->arg_size() != 4) {
|
||||
+ IsFunctionCall = llvm::StringSwitch<bool>(Stem)
|
||||
+ .Case("exchange_explicit", true)
|
||||
+ .Case("fetch_add_explicit", true)
|
||||
+ .Case("fetch_sub_explicit", true)
|
||||
+ .Case("fetch_min_explicit", true)
|
||||
+ .Case("fetch_max_explicit", true)
|
||||
+ .Case("load_explicit", true)
|
||||
+ .Case("store_explicit", true)
|
||||
+ .Default(false);
|
||||
+ if (IsFunctionCall)
|
||||
+ return;
|
||||
+ }
|
||||
+ }
|
||||
|
||||
auto PCI = &CI;
|
||||
if (DemangledName == kOCLBuiltinName::AtomicInit) {
|
||||
@@ -839,7 +892,7 @@ void OCLToSPIRVBase::transAtomicBuiltin(CallInst *CI,
|
||||
AttributeList Attrs = CI->getCalledFunction()->getAttributes();
|
||||
mutateCallInstSPIRV(
|
||||
M, CI,
|
||||
- [=](CallInst *CI, std::vector<Value *> &Args) {
|
||||
+ [=](CallInst *CI, std::vector<Value *> &Args) -> std::string {
|
||||
Info.PostProc(Args);
|
||||
// Order of args in OCL20:
|
||||
// object, 0-2 other args, 1-2 order, scope
|
||||
@@ -868,7 +921,22 @@ void OCLToSPIRVBase::transAtomicBuiltin(CallInst *CI,
|
||||
std::rotate(Args.begin() + 2, Args.begin() + OrderIdx,
|
||||
Args.end() - Offset);
|
||||
}
|
||||
- return getSPIRVFuncName(OCLSPIRVBuiltinMap::map(Info.UniqName));
|
||||
+
|
||||
+ llvm::Type* AtomicBuiltinsReturnType =
|
||||
+ CI->getCalledFunction()->getReturnType();
|
||||
+ auto IsFPType = [](llvm::Type *ReturnType) {
|
||||
+ return ReturnType->isHalfTy() || ReturnType->isFloatTy() ||
|
||||
+ ReturnType->isDoubleTy();
|
||||
+ };
|
||||
+ auto SPIRVFunctionName =
|
||||
+ getSPIRVFuncName(OCLSPIRVBuiltinMap::map(Info.UniqName));
|
||||
+ if (!IsFPType(AtomicBuiltinsReturnType))
|
||||
+ return SPIRVFunctionName;
|
||||
+ // Translate FP-typed atomic builtins.
|
||||
+ return llvm::StringSwitch<std::string>(SPIRVFunctionName)
|
||||
+ .Case("__spirv_AtomicIAdd", "__spirv_AtomicFAddEXT")
|
||||
+ .Case("__spirv_AtomicSMax", "__spirv_AtomicFMaxEXT")
|
||||
+ .Case("__spirv_AtomicSMin", "__spirv_AtomicFMinEXT");
|
||||
},
|
||||
&Attrs);
|
||||
}
|
||||
diff --git a/lib/SPIRV/OCLUtil.cpp b/lib/SPIRV/OCLUtil.cpp
|
||||
index 2de3f152..85155e39 100644
|
||||
--- a/lib/SPIRV/OCLUtil.cpp
|
||||
+++ b/lib/SPIRV/OCLUtil.cpp
|
||||
@@ -662,32 +662,6 @@ size_t getSPIRVAtomicBuiltinNumMemoryOrderArgs(Op OC) {
|
||||
return 1;
|
||||
}
|
||||
|
||||
-bool isComputeAtomicOCLBuiltin(StringRef DemangledName) {
|
||||
- if (!DemangledName.startswith(kOCLBuiltinName::AtomicPrefix) &&
|
||||
- !DemangledName.startswith(kOCLBuiltinName::AtomPrefix))
|
||||
- return false;
|
||||
-
|
||||
- return llvm::StringSwitch<bool>(DemangledName)
|
||||
- .EndsWith("add", true)
|
||||
- .EndsWith("sub", true)
|
||||
- .EndsWith("inc", true)
|
||||
- .EndsWith("dec", true)
|
||||
- .EndsWith("cmpxchg", true)
|
||||
- .EndsWith("min", true)
|
||||
- .EndsWith("max", true)
|
||||
- .EndsWith("and", true)
|
||||
- .EndsWith("or", true)
|
||||
- .EndsWith("xor", true)
|
||||
- .EndsWith("add_explicit", true)
|
||||
- .EndsWith("sub_explicit", true)
|
||||
- .EndsWith("or_explicit", true)
|
||||
- .EndsWith("xor_explicit", true)
|
||||
- .EndsWith("and_explicit", true)
|
||||
- .EndsWith("min_explicit", true)
|
||||
- .EndsWith("max_explicit", true)
|
||||
- .Default(false);
|
||||
-}
|
||||
-
|
||||
BarrierLiterals getBarrierLiterals(CallInst *CI) {
|
||||
auto N = CI->getNumArgOperands();
|
||||
assert(N == 1 || N == 2);
|
||||
diff --git a/lib/SPIRV/OCLUtil.h b/lib/SPIRV/OCLUtil.h
|
||||
index 4c05c672..c8577e9b 100644
|
||||
--- a/lib/SPIRV/OCLUtil.h
|
||||
+++ b/lib/SPIRV/OCLUtil.h
|
||||
@@ -394,10 +394,6 @@ size_t getAtomicBuiltinNumMemoryOrderArgs(StringRef Name);
|
||||
/// Get number of memory order arguments for spirv atomic builtin function.
|
||||
size_t getSPIRVAtomicBuiltinNumMemoryOrderArgs(Op OC);
|
||||
|
||||
-/// Return true for OpenCL builtins which do compute operations
|
||||
-/// (like add, sub, min, max, inc, dec, ...) atomically
|
||||
-bool isComputeAtomicOCLBuiltin(StringRef DemangledName);
|
||||
-
|
||||
/// Get OCL version from metadata opencl.ocl.version.
|
||||
/// \param AllowMulti Allows multiple operands if true.
|
||||
/// \return OCL version encoded as Major*10^5+Minor*10^3+Rev,
|
||||
diff --git a/test/negative/InvalidAtomicBuiltins.cl b/test/negative/InvalidAtomicBuiltins.cl
|
||||
index b8ec5b89..23dcc4e3 100644
|
||||
--- a/test/negative/InvalidAtomicBuiltins.cl
|
||||
+++ b/test/negative/InvalidAtomicBuiltins.cl
|
||||
@@ -1,7 +1,9 @@
|
||||
// Check that translator doesn't generate atomic instructions for atomic builtins
|
||||
// which are not defined in the spec.
|
||||
|
||||
-// RUN: %clang_cc1 -triple spir -O1 -cl-std=cl2.0 -fdeclare-opencl-builtins -finclude-default-header %s -emit-llvm-bc -o %t.bc
|
||||
+// To drop `fdeclare-opencl-builtins` option, since FP-typed atomic function
|
||||
+// TableGen definitions have not been introduced.
|
||||
+// RUN: %clang_cc1 -triple spir -O1 -cl-std=cl2.0 -finclude-default-header %s -emit-llvm-bc -o %t.bc
|
||||
// RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s
|
||||
// RUN: llvm-spirv %t.bc -o %t.spv
|
||||
// RUN: spirv-val %t.spv
|
||||
@@ -41,13 +43,9 @@ float __attribute__((overloadable)) atomic_fetch_xor(volatile generic atomic_flo
|
||||
double __attribute__((overloadable)) atomic_fetch_and(volatile generic atomic_double *object, double operand, memory_order order);
|
||||
double __attribute__((overloadable)) atomic_fetch_max(volatile generic atomic_double *object, double operand, memory_order order);
|
||||
double __attribute__((overloadable)) atomic_fetch_min(volatile generic atomic_double *object, double operand, memory_order order);
|
||||
-float __attribute__((overloadable)) atomic_fetch_add_explicit(volatile generic atomic_float *object, float operand, memory_order order);
|
||||
-float __attribute__((overloadable)) atomic_fetch_sub_explicit(volatile generic atomic_float *object, float operand, memory_order order);
|
||||
float __attribute__((overloadable)) atomic_fetch_or_explicit(volatile generic atomic_float *object, float operand, memory_order order);
|
||||
float __attribute__((overloadable)) atomic_fetch_xor_explicit(volatile generic atomic_float *object, float operand, memory_order order);
|
||||
double __attribute__((overloadable)) atomic_fetch_and_explicit(volatile generic atomic_double *object, double operand, memory_order order);
|
||||
-double __attribute__((overloadable)) atomic_fetch_max_explicit(volatile generic atomic_double *object, double operand, memory_order order);
|
||||
-double __attribute__((overloadable)) atomic_fetch_min_explicit(volatile generic atomic_double *object, double operand, memory_order order);
|
||||
|
||||
__kernel void test_atomic_fn(volatile __global float *p,
|
||||
volatile __global double *pp,
|
||||
@@ -86,11 +84,7 @@ __kernel void test_atomic_fn(volatile __global float *p,
|
||||
d = atomic_fetch_and(pp, val, order);
|
||||
d = atomic_fetch_min(pp, val, order);
|
||||
d = atomic_fetch_max(pp, val, order);
|
||||
- f = atomic_fetch_add_explicit(p, val, order);
|
||||
- f = atomic_fetch_sub_explicit(p, val, order);
|
||||
f = atomic_fetch_or_explicit(p, val, order);
|
||||
f = atomic_fetch_xor_explicit(p, val, order);
|
||||
d = atomic_fetch_and_explicit(pp, val, order);
|
||||
- d = atomic_fetch_min_explicit(pp, val, order);
|
||||
- d = atomic_fetch_max_explicit(pp, val, order);
|
||||
}
|
||||
diff --git a/test/transcoding/AtomicFAddEXTForOCL.ll b/test/transcoding/AtomicFAddEXTForOCL.ll
|
||||
new file mode 100644
|
||||
index 00000000..fb146fb9
|
||||
--- /dev/null
|
||||
+++ b/test/transcoding/AtomicFAddEXTForOCL.ll
|
||||
@@ -0,0 +1,64 @@
|
||||
+; RUN: llvm-as %s -o %t.bc
|
||||
+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_add -o %t.spv
|
||||
+; RUN: spirv-val %t.spv
|
||||
+; RUN: llvm-spirv -to-text %t.spv -o %t.spt
|
||||
+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
|
||||
+
|
||||
+; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc
|
||||
+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20
|
||||
+
|
||||
+; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc
|
||||
+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV
|
||||
+
|
||||
+target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
|
||||
+target triple = "spir-unknown-unknown"
|
||||
+
|
||||
+; CHECK-SPIRV: Capability AtomicFloat32AddEXT
|
||||
+; CHECK-SPIRV: Capability AtomicFloat64AddEXT
|
||||
+; CHECK-SPIRV: Extension "SPV_EXT_shader_atomic_float_add"
|
||||
+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32
|
||||
+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64
|
||||
+
|
||||
+
|
||||
+; Function Attrs: convergent norecurse nounwind
|
||||
+define dso_local spir_func void @test_atomic_float(float addrspace(1)* %a) local_unnamed_addr #0 {
|
||||
+entry:
|
||||
+ ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_32]]
|
||||
+ ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_add_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}})
|
||||
+ ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+fiif]]({{.*}})
|
||||
+ %call = tail call spir_func float @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2
|
||||
+ ret void
|
||||
+}
|
||||
+
|
||||
+; Function Attrs: convergent
|
||||
+declare spir_func float @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1
|
||||
+; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
|
||||
+
|
||||
+; Function Attrs: convergent norecurse nounwind
|
||||
+define dso_local spir_func void @test_atomic_double(double addrspace(1)* %a) local_unnamed_addr #0 {
|
||||
+entry:
|
||||
+ ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_64]]
|
||||
+ ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_add_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}})
|
||||
+ ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+diid]]({{.*}})
|
||||
+ %call = tail call spir_func double @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2
|
||||
+ ret void
|
||||
+}
|
||||
+; Function Attrs: convergent
|
||||
+declare spir_func double @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)*, double, i32) local_unnamed_addr #1
|
||||
+; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
|
||||
+
|
||||
+; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
|
||||
+; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
|
||||
+
|
||||
+attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
|
||||
+attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
|
||||
+attributes #2 = { convergent nounwind }
|
||||
+
|
||||
+!llvm.module.flags = !{!0}
|
||||
+!opencl.ocl.version = !{!1}
|
||||
+!opencl.spir.version = !{!1}
|
||||
+!llvm.ident = !{!2}
|
||||
+
|
||||
+!0 = !{i32 1, !"wchar_size", i32 4}
|
||||
+!1 = !{i32 2, i32 0}
|
||||
+!2 = !{!"clang version 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"}
|
||||
diff --git a/test/transcoding/AtomicFMaxEXTForOCL.ll b/test/transcoding/AtomicFMaxEXTForOCL.ll
|
||||
new file mode 100644
|
||||
index 00000000..1f2530d9
|
||||
--- /dev/null
|
||||
+++ b/test/transcoding/AtomicFMaxEXTForOCL.ll
|
||||
@@ -0,0 +1,64 @@
|
||||
+; RUN: llvm-as %s -o %t.bc
|
||||
+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_min_max -o %t.spv
|
||||
+; RUN: spirv-val %t.spv
|
||||
+; RUN: llvm-spirv -to-text %t.spv -o %t.spt
|
||||
+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
|
||||
+
|
||||
+; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc
|
||||
+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20
|
||||
+
|
||||
+; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc
|
||||
+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV
|
||||
+
|
||||
+target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
|
||||
+target triple = "spir-unknown-unknown"
|
||||
+
|
||||
+; CHECK-SPIRV: Capability AtomicFloat32MinMaxEXT
|
||||
+; CHECK-SPIRV: Capability AtomicFloat64MinMaxEXT
|
||||
+; CHECK-SPIRV: Extension "SPV_EXT_shader_atomic_float_min_max"
|
||||
+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32
|
||||
+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64
|
||||
+
|
||||
+; Function Attrs: convergent norecurse nounwind
|
||||
+define dso_local spir_func void @test_float(float addrspace(1)* %a) local_unnamed_addr #0 {
|
||||
+entry:
|
||||
+ ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_32]]
|
||||
+ ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_max_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}})
|
||||
+ ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+fiif]]({{.*}})
|
||||
+ %call = tail call spir_func float @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2
|
||||
+ ret void
|
||||
+}
|
||||
+
|
||||
+; Function Attrs: convergent
|
||||
+declare spir_func float @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1
|
||||
+; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
|
||||
+
|
||||
+; Function Attrs: convergent norecurse nounwind
|
||||
+define dso_local spir_func void @test_double(double addrspace(1)* %a) local_unnamed_addr #0 {
|
||||
+entry:
|
||||
+ ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_64]]
|
||||
+ ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_max_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}})
|
||||
+ ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+diid]]({{.*}})
|
||||
+ %call = tail call spir_func double @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2
|
||||
+ ret void
|
||||
+}
|
||||
+
|
||||
+; Function Attrs: convergent
|
||||
+declare spir_func double @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)*, double, i32) local_unnamed_addr #1
|
||||
+; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
|
||||
+
|
||||
+; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
|
||||
+; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
|
||||
+
|
||||
+attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
|
||||
+attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
|
||||
+attributes #2 = { convergent nounwind }
|
||||
+
|
||||
+!llvm.module.flags = !{!0}
|
||||
+!opencl.ocl.version = !{!1}
|
||||
+!opencl.spir.version = !{!1}
|
||||
+!llvm.ident = !{!2}
|
||||
+
|
||||
+!0 = !{i32 1, !"wchar_size", i32 4}
|
||||
+!1 = !{i32 2, i32 0}
|
||||
+!2 = !{!"clang version 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"}
|
||||
diff --git a/test/transcoding/AtomicFMinEXTForOCL.ll b/test/transcoding/AtomicFMinEXTForOCL.ll
|
||||
new file mode 100644
|
||||
index 00000000..6196b0f8
|
||||
--- /dev/null
|
||||
+++ b/test/transcoding/AtomicFMinEXTForOCL.ll
|
||||
@@ -0,0 +1,64 @@
|
||||
+; RUN: llvm-as %s -o %t.bc
|
||||
+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_min_max -o %t.spv
|
||||
+; RUN: spirv-val %t.spv
|
||||
+; RUN: llvm-spirv -to-text %t.spv -o %t.spt
|
||||
+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
|
||||
+
|
||||
+; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc
|
||||
+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20
|
||||
+
|
||||
+; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc
|
||||
+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV
|
||||
+
|
||||
+target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
|
||||
+target triple = "spir-unknown-unknown"
|
||||
+
|
||||
+; CHECK-SPIRV: Capability AtomicFloat32MinMaxEXT
|
||||
+; CHECK-SPIRV: Capability AtomicFloat64MinMaxEXT
|
||||
+; CHECK-SPIRV: Extension "SPV_EXT_shader_atomic_float_min_max"
|
||||
+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32
|
||||
+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64
|
||||
+
|
||||
+; Function Attrs: convergent norecurse nounwind
|
||||
+define dso_local spir_func void @test_float(float addrspace(1)* %a) local_unnamed_addr #0 {
|
||||
+entry:
|
||||
+ ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_32]]
|
||||
+ ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_min_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}})
|
||||
+ ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+fiif]]({{.*}})
|
||||
+ %call = tail call spir_func float @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2
|
||||
+ ret void
|
||||
+}
|
||||
+
|
||||
+; Function Attrs: convergent
|
||||
+declare spir_func float @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1
|
||||
+; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
|
||||
+
|
||||
+; Function Attrs: convergent norecurse nounwind
|
||||
+define dso_local spir_func void @test_double(double addrspace(1)* %a) local_unnamed_addr #0 {
|
||||
+entry:
|
||||
+ ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_64]]
|
||||
+ ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_min_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}})
|
||||
+ ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+diid]]({{.*}})
|
||||
+ %call = tail call spir_func double @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2
|
||||
+ ret void
|
||||
+}
|
||||
+
|
||||
+; Function Attrs: convergent
|
||||
+declare spir_func double @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)*, double, i32) local_unnamed_addr #1
|
||||
+; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
|
||||
+
|
||||
+; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
|
||||
+; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
|
||||
+
|
||||
+attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
|
||||
+attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
|
||||
+attributes #2 = { convergent nounwind }
|
||||
+
|
||||
+!llvm.module.flags = !{!0}
|
||||
+!opencl.ocl.version = !{!1}
|
||||
+!opencl.spir.version = !{!1}
|
||||
+!llvm.ident = !{!2}
|
||||
+
|
||||
+!0 = !{i32 1, !"wchar_size", i32 4}
|
||||
+!1 = !{i32 2, i32 0}
|
||||
+!2 = !{!"clang version 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"}
|
||||
--
|
||||
2.17.1
|
||||
|
||||
|
|
@ -1,35 +0,0 @@
|
|||
From ef27f1f99ad661c9604b7ff10efb1122466c508b Mon Sep 17 00:00:00 2001
|
||||
From: juanrod2 <>
|
||||
Date: Tue, 22 Dec 2020 08:33:08 +0800
|
||||
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]
|
||||
|
||||
Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com>
|
||||
|
||||
Cleaning a mutex inside ManagedStatic llvm class.
|
||||
---
|
||||
llvm/lib/Support/ManagedStatic.cpp | 6 +++++-
|
||||
1 file changed, 5 insertions(+), 1 deletion(-)
|
||||
|
||||
diff --git a/llvm/lib/Support/ManagedStatic.cpp b/llvm/lib/Support/ManagedStatic.cpp
|
||||
index 053493f72fb5..6571580ccecf 100644
|
||||
--- a/llvm/lib/Support/ManagedStatic.cpp
|
||||
+++ b/llvm/lib/Support/ManagedStatic.cpp
|
||||
@@ -76,8 +76,12 @@ void ManagedStaticBase::destroy() const {
|
||||
|
||||
/// llvm_shutdown - Deallocate and destroy all ManagedStatic variables.
|
||||
void llvm::llvm_shutdown() {
|
||||
- std::lock_guard<std::recursive_mutex> Lock(*getManagedStaticMutex());
|
||||
+ getManagedStaticMutex()->lock();
|
||||
|
||||
while (StaticList)
|
||||
StaticList->destroy();
|
||||
+
|
||||
+ getManagedStaticMutex()->unlock();
|
||||
+ delete ManagedStaticMutex;
|
||||
+ ManagedStaticMutex = nullptr;
|
||||
}
|
||||
--
|
||||
2.17.1
|
||||
|
||||
|
|
@ -1,49 +0,0 @@
|
|||
From a71ab6fb04b918b856f1dd802cfdb4a7ccd53290 Mon Sep 17 00:00:00 2001
|
||||
From: Feng Zou <feng.zou@intel.com>
|
||||
Date: Tue, 20 Oct 2020 11:29:04 +0800
|
||||
Subject: [PATCH 3/6] Remove repo name in LLVM IR
|
||||
|
||||
Upstream-Status: Backport [Taken from opencl-clang patches, https://github.com/intel/opencl-clang/blob/ocl-open-110/patches/llvm/0002-Remove-repo-name-in-LLVM-IR.patch]
|
||||
Signed-off-by: Feng Zou <feng.zou@intel.com>
|
||||
Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com>
|
||||
---
|
||||
llvm/cmake/modules/VersionFromVCS.cmake | 23 ++++++++++++-----------
|
||||
1 file changed, 12 insertions(+), 11 deletions(-)
|
||||
|
||||
diff --git a/llvm/cmake/modules/VersionFromVCS.cmake b/llvm/cmake/modules/VersionFromVCS.cmake
|
||||
index 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)
|
||||
else()
|
||||
set(remote "origin")
|
||||
endif()
|
||||
- execute_process(COMMAND ${GIT_EXECUTABLE} remote get-url ${remote}
|
||||
- WORKING_DIRECTORY ${path}
|
||||
- RESULT_VARIABLE git_result
|
||||
- OUTPUT_VARIABLE git_output
|
||||
- ERROR_QUIET)
|
||||
- if(git_result EQUAL 0)
|
||||
- string(STRIP "${git_output}" git_output)
|
||||
- set(${repository} ${git_output} PARENT_SCOPE)
|
||||
- else()
|
||||
- set(${repository} ${path} PARENT_SCOPE)
|
||||
- endif()
|
||||
+ # Do not show repo name in IR
|
||||
+ # execute_process(COMMAND ${GIT_EXECUTABLE} remote get-url ${remote}
|
||||
+ # WORKING_DIRECTORY ${path}
|
||||
+ # RESULT_VARIABLE git_result
|
||||
+ # OUTPUT_VARIABLE git_output
|
||||
+ # ERROR_QUIET)
|
||||
+ # if(git_result EQUAL 0)
|
||||
+ # string(STRIP "${git_output}" git_output)
|
||||
+ # set(${repository} ${git_output} PARENT_SCOPE)
|
||||
+ # else()
|
||||
+ # set(${repository} ${path} PARENT_SCOPE)
|
||||
+ # endif()
|
||||
endif()
|
||||
else()
|
||||
message(WARNING "Git not found. Version cannot be determined.")
|
||||
--
|
||||
2.17.1
|
||||
|
||||
|
|
@ -1,51 +0,0 @@
|
|||
From 546d9089fe5e21cccc671a0a89555cd4d5f8c817 Mon Sep 17 00:00:00 2001
|
||||
From: Naveen Saini <naveen.kumar.saini@intel.com>
|
||||
Date: Thu, 19 Aug 2021 15:52:24 +0800
|
||||
Subject: [PATCH 4/6] Remove __IMAGE_SUPPORT__ macro for SPIR since SPIR
|
||||
doesn't require image support
|
||||
|
||||
Upstream-Status: Backport [Taken from opencl-clang patches; https://github.com/intel/opencl-clang/blob/ocl-open-110/patches/clang/0002-Remove-__IMAGE_SUPPORT__-macro-for-SPIR.patch]
|
||||
|
||||
Signed-off-by: haonanya <haonan.yang@intel.com>
|
||||
Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com>
|
||||
---
|
||||
clang/lib/Frontend/InitPreprocessor.cpp | 3 ---
|
||||
clang/test/Preprocessor/predefined-macros.c | 2 --
|
||||
2 files changed, 5 deletions(-)
|
||||
|
||||
diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp
|
||||
index 5bb489c11909..cf3b48cb65d2 100644
|
||||
--- a/clang/lib/Frontend/InitPreprocessor.cpp
|
||||
+++ b/clang/lib/Frontend/InitPreprocessor.cpp
|
||||
@@ -1115,9 +1115,6 @@ static void InitializePredefinedMacros(const TargetInfo &TI,
|
||||
if (TI.getSupportedOpenCLOpts().isSupported(#Ext)) \
|
||||
Builder.defineMacro(#Ext);
|
||||
#include "clang/Basic/OpenCLExtensions.def"
|
||||
-
|
||||
- if (TI.getTriple().isSPIR())
|
||||
- Builder.defineMacro("__IMAGE_SUPPORT__");
|
||||
}
|
||||
|
||||
if (TI.hasInt128Type() && LangOpts.CPlusPlus && LangOpts.GNUMode) {
|
||||
diff --git a/clang/test/Preprocessor/predefined-macros.c b/clang/test/Preprocessor/predefined-macros.c
|
||||
index 6c80517ec4d4..b5e5d7e2d546 100644
|
||||
--- a/clang/test/Preprocessor/predefined-macros.c
|
||||
+++ b/clang/test/Preprocessor/predefined-macros.c
|
||||
@@ -186,14 +186,12 @@
|
||||
|
||||
// RUN: %clang_cc1 %s -E -dM -o - -x cl -triple spir-unknown-unknown \
|
||||
// RUN: | FileCheck -match-full-lines %s --check-prefix=CHECK-SPIR
|
||||
-// CHECK-SPIR-DAG: #define __IMAGE_SUPPORT__ 1
|
||||
// CHECK-SPIR-DAG: #define __SPIR__ 1
|
||||
// CHECK-SPIR-DAG: #define __SPIR32__ 1
|
||||
// CHECK-SPIR-NOT: #define __SPIR64__ 1
|
||||
|
||||
// RUN: %clang_cc1 %s -E -dM -o - -x cl -triple spir64-unknown-unknown \
|
||||
// RUN: | FileCheck -match-full-lines %s --check-prefix=CHECK-SPIR64
|
||||
-// CHECK-SPIR64-DAG: #define __IMAGE_SUPPORT__ 1
|
||||
// CHECK-SPIR64-DAG: #define __SPIR__ 1
|
||||
// CHECK-SPIR64-DAG: #define __SPIR64__ 1
|
||||
// CHECK-SPIR64-NOT: #define __SPIR32__ 1
|
||||
--
|
||||
2.17.1
|
||||
|
||||
|
|
@ -1,52 +0,0 @@
|
|||
From 747e48959e18ac8b586078a82472a0799d12925c Mon Sep 17 00:00:00 2001
|
||||
From: Raphael Isemann <teemperor@gmail.com>
|
||||
Date: Thu, 1 Apr 2021 18:41:44 +0200
|
||||
Subject: [PATCH 5/6] Avoid calling ParseCommandLineOptions in BackendUtil if
|
||||
possible
|
||||
|
||||
Calling `ParseCommandLineOptions` should only be called from `main` as the
|
||||
CommandLine setup code isn't thread-safe. As BackendUtil is part of the
|
||||
generic Clang FrontendAction logic, a process which has several threads executing
|
||||
Clang FrontendActions will randomly crash in the unsafe setup code.
|
||||
|
||||
This patch avoids calling the function unless either the debug-pass option or
|
||||
limit-float-precision option is set. Without these two options set the
|
||||
`ParseCommandLineOptions` call doesn't do anything beside parsing
|
||||
the command line `clang` which doesn't set any options.
|
||||
|
||||
See also D99652 where LLDB received a workaround for this crash.
|
||||
|
||||
Reviewed By: JDevlieghere
|
||||
|
||||
Differential Revision: https://reviews.llvm.org/D99740
|
||||
|
||||
Upstream-Status: Backport [Taken from opencl-clang patches; https://github.com/intel/opencl-clang/blob/ocl-open-110/patches/clang/0003-Avoid-calling-ParseCommandLineOptions-in-BackendUtil.patch]
|
||||
|
||||
Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com>
|
||||
---
|
||||
clang/lib/CodeGen/BackendUtil.cpp | 8 ++++++++
|
||||
1 file changed, 8 insertions(+)
|
||||
|
||||
diff --git a/clang/lib/CodeGen/BackendUtil.cpp b/clang/lib/CodeGen/BackendUtil.cpp
|
||||
index dce0940670a2..ab478090ed1c 100644
|
||||
--- a/clang/lib/CodeGen/BackendUtil.cpp
|
||||
+++ b/clang/lib/CodeGen/BackendUtil.cpp
|
||||
@@ -797,7 +797,15 @@ static void setCommandLineOpts(const CodeGenOptions &CodeGenOpts) {
|
||||
BackendArgs.push_back("-limit-float-precision");
|
||||
BackendArgs.push_back(CodeGenOpts.LimitFloatPrecision.c_str());
|
||||
}
|
||||
+ // Check for the default "clang" invocation that won't set any cl::opt values.
|
||||
+ // Skip trying to parse the command line invocation to avoid the issues
|
||||
+ // described below.
|
||||
+ if (BackendArgs.size() == 1)
|
||||
+ return;
|
||||
BackendArgs.push_back(nullptr);
|
||||
+ // FIXME: The command line parser below is not thread-safe and shares a global
|
||||
+ // state, so this call might crash or overwrite the options of another Clang
|
||||
+ // instance in the same process.
|
||||
llvm::cl::ParseCommandLineOptions(BackendArgs.size() - 1,
|
||||
BackendArgs.data());
|
||||
}
|
||||
--
|
||||
2.17.1
|
||||
|
||||
|
|
@ -1,353 +0,0 @@
|
|||
From a1b924d76cdacfa3f9dbb79a9e3edddcd75f61ca Mon Sep 17 00:00:00 2001
|
||||
From: Naveen Saini <naveen.kumar.saini@intel.com>
|
||||
Date: Thu, 19 Aug 2021 16:06:33 +0800
|
||||
Subject: [PATCH 6/6] [OpenCL] support cl_ext_float_atomics
|
||||
|
||||
Upstream-Status: Backport [Taken from opencl-clang patches; https://github.com/intel/opencl-clang/blob/ocl-open-110/patches/clang/0004-OpenCL-support-cl_ext_float_atomics.patch]
|
||||
|
||||
Signed-off-by: haonanya <haonan.yang@intel.com>
|
||||
Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com>
|
||||
---
|
||||
clang/lib/Headers/opencl-c-base.h | 25 ++++
|
||||
clang/lib/Headers/opencl-c.h | 195 ++++++++++++++++++++++++++
|
||||
clang/test/Headers/opencl-c-header.cl | 85 +++++++++++
|
||||
3 files changed, 305 insertions(+)
|
||||
|
||||
diff --git a/clang/lib/Headers/opencl-c-base.h b/clang/lib/Headers/opencl-c-base.h
|
||||
index afa900ab24d9..9a3ee8529acf 100644
|
||||
--- a/clang/lib/Headers/opencl-c-base.h
|
||||
+++ b/clang/lib/Headers/opencl-c-base.h
|
||||
@@ -62,6 +62,31 @@
|
||||
#endif
|
||||
#endif // defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ == CL_VERSION_2_0)
|
||||
|
||||
+#if (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200)
|
||||
+// For SPIR all extensions are supported.
|
||||
+#if defined(__SPIR__)
|
||||
+#define cl_ext_float_atomics
|
||||
+#ifdef cl_khr_fp16
|
||||
+#define __opencl_c_ext_fp16_global_atomic_load_store 1
|
||||
+#define __opencl_c_ext_fp16_local_atomic_load_store 1
|
||||
+#define __opencl_c_ext_fp16_global_atomic_add 1
|
||||
+#define __opencl_c_ext_fp16_local_atomic_add 1
|
||||
+#define __opencl_c_ext_fp16_global_atomic_min_max 1
|
||||
+#define __opencl_c_ext_fp16_local_atomic_min_max 1
|
||||
+#endif
|
||||
+#ifdef __opencl_c_fp64
|
||||
+#define __opencl_c_ext_fp64_global_atomic_add 1
|
||||
+#define __opencl_c_ext_fp64_local_atomic_add 1
|
||||
+#define __opencl_c_ext_fp64_global_atomic_min_max 1
|
||||
+#define __opencl_c_ext_fp64_local_atomic_min_max 1
|
||||
+#endif
|
||||
+#define __opencl_c_ext_fp32_global_atomic_add 1
|
||||
+#define __opencl_c_ext_fp32_local_atomic_add 1
|
||||
+#define __opencl_c_ext_fp32_global_atomic_min_max 1
|
||||
+#define __opencl_c_ext_fp32_local_atomic_min_max 1
|
||||
+#endif // defined(__SPIR__)
|
||||
+#endif // (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200)
|
||||
+
|
||||
// built-in scalar data types:
|
||||
|
||||
/**
|
||||
diff --git a/clang/lib/Headers/opencl-c.h b/clang/lib/Headers/opencl-c.h
|
||||
index 67d900eb1c3d..bda0f5c6df80 100644
|
||||
--- a/clang/lib/Headers/opencl-c.h
|
||||
+++ b/clang/lib/Headers/opencl-c.h
|
||||
@@ -14354,6 +14354,201 @@ intptr_t __ovld atomic_fetch_max_explicit(
|
||||
// defined(cl_khr_int64_extended_atomics)
|
||||
#endif // (__OPENCL_C_VERSION__ >= CL_VERSION_3_0)
|
||||
|
||||
+#if defined(cl_ext_float_atomics)
|
||||
+
|
||||
+#if defined(__opencl_c_ext_fp32_global_atomic_min_max)
|
||||
+float __ovld atomic_fetch_min(volatile __global atomic_float *object,
|
||||
+ float operand);
|
||||
+float __ovld atomic_fetch_max(volatile __global atomic_float *object,
|
||||
+ float operand);
|
||||
+float __ovld atomic_fetch_min_explicit(volatile __global atomic_float *object,
|
||||
+ float operand, memory_order order);
|
||||
+float __ovld atomic_fetch_max_explicit(volatile __global atomic_float *object,
|
||||
+ float operand, memory_order order);
|
||||
+float __ovld atomic_fetch_min_explicit(volatile __global atomic_float *object,
|
||||
+ float operand, memory_order order,
|
||||
+ memory_scope scope);
|
||||
+float __ovld atomic_fetch_max_explicit(volatile __global atomic_float *object,
|
||||
+ float operand, memory_order order,
|
||||
+ memory_scope scope);
|
||||
+#endif
|
||||
+#if defined(__opencl_c_ext_fp32_local_atomic_min_max)
|
||||
+float __ovld atomic_fetch_min(volatile __local atomic_float *object,
|
||||
+ float operand);
|
||||
+float __ovld atomic_fetch_max(volatile __local atomic_float *object,
|
||||
+ float operand);
|
||||
+float __ovld atomic_fetch_min_explicit(volatile __local atomic_float *object,
|
||||
+ float operand, memory_order order);
|
||||
+float __ovld atomic_fetch_max_explicit(volatile __local atomic_float *object,
|
||||
+ float operand, memory_order order);
|
||||
+float __ovld atomic_fetch_min_explicit(volatile __local atomic_float *object,
|
||||
+ float operand, memory_order order,
|
||||
+ memory_scope scope);
|
||||
+float __ovld atomic_fetch_max_explicit(volatile __local atomic_float *object,
|
||||
+ float operand, memory_order order,
|
||||
+ memory_scope scope);
|
||||
+#endif
|
||||
+#if defined(__opencl_c_ext_fp32_global_atomic_min_max) || \
|
||||
+ defined(__opencl_c_ext_fp32_local_atomic_min_max)
|
||||
+float __ovld atomic_fetch_min(volatile atomic_float *object, float operand);
|
||||
+float __ovld atomic_fetch_max(volatile atomic_float *object, float operand);
|
||||
+float __ovld atomic_fetch_min_explicit(volatile atomic_float *object,
|
||||
+ float operand, memory_order order);
|
||||
+float __ovld atomic_fetch_max_explicit(volatile atomic_float *object,
|
||||
+ float operand, memory_order order);
|
||||
+float __ovld atomic_fetch_min_explicit(volatile atomic_float *object,
|
||||
+ float operand, memory_order order,
|
||||
+ memory_scope scope);
|
||||
+float __ovld atomic_fetch_max_explicit(volatile atomic_float *object,
|
||||
+ float operand, memory_order order,
|
||||
+ memory_scope scope);
|
||||
+#endif
|
||||
+#if defined(__opencl_c_ext_fp64_global_atomic_min_max)
|
||||
+double __ovld atomic_fetch_min(volatile __global atomic_double *object,
|
||||
+ double operand);
|
||||
+double __ovld atomic_fetch_max(volatile __global atomic_double *object,
|
||||
+ double operand);
|
||||
+double __ovld atomic_fetch_min_explicit(volatile __global atomic_double *object,
|
||||
+ double operand, memory_order order);
|
||||
+double __ovld atomic_fetch_max_explicit(volatile __global atomic_double *object,
|
||||
+ double operand, memory_order order);
|
||||
+double __ovld atomic_fetch_min_explicit(volatile __global atomic_double *object,
|
||||
+ double operand, memory_order order,
|
||||
+ memory_scope scope);
|
||||
+double __ovld atomic_fetch_max_explicit(volatile __global atomic_double *object,
|
||||
+ double operand, memory_order order,
|
||||
+ memory_scope scope);
|
||||
+#endif
|
||||
+#if defined(__opencl_c_ext_fp64_local_atomic_min_max)
|
||||
+double __ovld atomic_fetch_min(volatile __local atomic_double *object,
|
||||
+ double operand);
|
||||
+double __ovld atomic_fetch_max(volatile __local atomic_double *object,
|
||||
+ double operand);
|
||||
+double __ovld atomic_fetch_min_explicit(volatile __local atomic_double *object,
|
||||
+ double operand, memory_order order);
|
||||
+double __ovld atomic_fetch_max_explicit(volatile __local atomic_double *object,
|
||||
+ double operand, memory_order order);
|
||||
+double __ovld atomic_fetch_min_explicit(volatile __local atomic_double *object,
|
||||
+ double operand, memory_order order,
|
||||
+ memory_scope scope);
|
||||
+double __ovld atomic_fetch_max_explicit(volatile __local atomic_double *object,
|
||||
+ double operand, memory_order order,
|
||||
+ memory_scope scope);
|
||||
+#endif
|
||||
+#if defined(__opencl_c_ext_fp64_global_atomic_min_max) || \
|
||||
+ defined(__opencl_c_ext_fp64_local_atomic_min_max)
|
||||
+double __ovld atomic_fetch_min(volatile atomic_double *object, double operand);
|
||||
+double __ovld atomic_fetch_max(volatile atomic_double *object, double operand);
|
||||
+double __ovld atomic_fetch_min_explicit(volatile atomic_double *object,
|
||||
+ double operand, memory_order order);
|
||||
+double __ovld atomic_fetch_max_explicit(volatile atomic_double *object,
|
||||
+ double operand, memory_order order);
|
||||
+double __ovld atomic_fetch_min_explicit(volatile atomic_double *object,
|
||||
+ double operand, memory_order order,
|
||||
+ memory_scope scope);
|
||||
+double __ovld atomic_fetch_max_explicit(volatile atomic_double *object,
|
||||
+ double operand, memory_order order,
|
||||
+ memory_scope scope);
|
||||
+#endif
|
||||
+
|
||||
+#if defined(__opencl_c_ext_fp32_global_atomic_add)
|
||||
+float __ovld atomic_fetch_add(volatile __global atomic_float *object,
|
||||
+ float operand);
|
||||
+float __ovld atomic_fetch_sub(volatile __global atomic_float *object,
|
||||
+ float operand);
|
||||
+float __ovld atomic_fetch_add_explicit(volatile __global atomic_float *object,
|
||||
+ float operand, memory_order order);
|
||||
+float __ovld atomic_fetch_sub_explicit(volatile __global atomic_float *object,
|
||||
+ float operand, memory_order order);
|
||||
+float __ovld atomic_fetch_add_explicit(volatile __global atomic_float *object,
|
||||
+ float operand, memory_order order,
|
||||
+ memory_scope scope);
|
||||
+float __ovld atomic_fetch_sub_explicit(volatile __global atomic_float *object,
|
||||
+ float operand, memory_order order,
|
||||
+ memory_scope scope);
|
||||
+#endif
|
||||
+#if defined(__opencl_c_ext_fp32_local_atomic_add)
|
||||
+float __ovld atomic_fetch_add(volatile __local atomic_float *object,
|
||||
+ float operand);
|
||||
+float __ovld atomic_fetch_sub(volatile __local atomic_float *object,
|
||||
+ float operand);
|
||||
+float __ovld atomic_fetch_add_explicit(volatile __local atomic_float *object,
|
||||
+ float operand, memory_order order);
|
||||
+float __ovld atomic_fetch_sub_explicit(volatile __local atomic_float *object,
|
||||
+ float operand, memory_order order);
|
||||
+float __ovld atomic_fetch_add_explicit(volatile __local atomic_float *object,
|
||||
+ float operand, memory_order order,
|
||||
+ memory_scope scope);
|
||||
+float __ovld atomic_fetch_sub_explicit(volatile __local atomic_float *object,
|
||||
+ float operand, memory_order order,
|
||||
+ memory_scope scope);
|
||||
+#endif
|
||||
+#if defined(__opencl_c_ext_fp32_global_atomic_add) || \
|
||||
+ defined(__opencl_c_ext_fp32_local_atomic_add)
|
||||
+float __ovld atomic_fetch_add(volatile atomic_float *object, float operand);
|
||||
+float __ovld atomic_fetch_sub(volatile atomic_float *object, float operand);
|
||||
+float __ovld atomic_fetch_add_explicit(volatile atomic_float *object,
|
||||
+ float operand, memory_order order);
|
||||
+float __ovld atomic_fetch_sub_explicit(volatile atomic_float *object,
|
||||
+ float operand, memory_order order);
|
||||
+float __ovld atomic_fetch_add_explicit(volatile atomic_float *object,
|
||||
+ float operand, memory_order order,
|
||||
+ memory_scope scope);
|
||||
+float __ovld atomic_fetch_sub_explicit(volatile atomic_float *object,
|
||||
+ float operand, memory_order order,
|
||||
+ memory_scope scope);
|
||||
+#endif
|
||||
+
|
||||
+#if defined(__opencl_c_ext_fp64_global_atomic_add)
|
||||
+double __ovld atomic_fetch_add(volatile __global atomic_double *object,
|
||||
+ double operand);
|
||||
+double __ovld atomic_fetch_sub(volatile __global atomic_double *object,
|
||||
+ double operand);
|
||||
+double __ovld atomic_fetch_add_explicit(volatile __global atomic_double *object,
|
||||
+ double operand, memory_order order);
|
||||
+double __ovld atomic_fetch_sub_explicit(volatile __global atomic_double *object,
|
||||
+ double operand, memory_order order);
|
||||
+double __ovld atomic_fetch_add_explicit(volatile __global atomic_double *object,
|
||||
+ double operand, memory_order order,
|
||||
+ memory_scope scope);
|
||||
+double __ovld atomic_fetch_sub_explicit(volatile __global atomic_double *object,
|
||||
+ double operand, memory_order order,
|
||||
+ memory_scope scope);
|
||||
+#endif
|
||||
+#if defined(__opencl_c_ext_fp64_local_atomic_add)
|
||||
+double __ovld atomic_fetch_add(volatile __local atomic_double *object,
|
||||
+ double operand);
|
||||
+double __ovld atomic_fetch_sub(volatile __local atomic_double *object,
|
||||
+ double operand);
|
||||
+double __ovld atomic_fetch_add_explicit(volatile __local atomic_double *object,
|
||||
+ double operand, memory_order order);
|
||||
+double __ovld atomic_fetch_sub_explicit(volatile __local atomic_double *object,
|
||||
+ double operand, memory_order order);
|
||||
+double __ovld atomic_fetch_add_explicit(volatile __local atomic_double *object,
|
||||
+ double operand, memory_order order,
|
||||
+ memory_scope scope);
|
||||
+double __ovld atomic_fetch_sub_explicit(volatile __local atomic_double *object,
|
||||
+ double operand, memory_order order,
|
||||
+ memory_scope scope);
|
||||
+#endif
|
||||
+#if defined(__opencl_c_ext_fp64_global_atomic_add) || \
|
||||
+ defined(__opencl_c_ext_fp64_local_atomic_add)
|
||||
+double __ovld atomic_fetch_add(volatile atomic_double *object, double operand);
|
||||
+double __ovld atomic_fetch_sub(volatile atomic_double *object, double operand);
|
||||
+double __ovld atomic_fetch_add_explicit(volatile atomic_double *object,
|
||||
+ double operand, memory_order order);
|
||||
+double __ovld atomic_fetch_sub_explicit(volatile atomic_double *object,
|
||||
+ double operand, memory_order order);
|
||||
+double __ovld atomic_fetch_add_explicit(volatile atomic_double *object,
|
||||
+ double operand, memory_order order,
|
||||
+ memory_scope scope);
|
||||
+double __ovld atomic_fetch_sub_explicit(volatile atomic_double *object,
|
||||
+ double operand, memory_order order,
|
||||
+ memory_scope scope);
|
||||
+#endif
|
||||
+
|
||||
+#endif // cl_ext_float_atomics
|
||||
+
|
||||
// atomic_store()
|
||||
|
||||
#if defined(__opencl_c_atomic_scope_device) && \
|
||||
diff --git a/clang/test/Headers/opencl-c-header.cl b/clang/test/Headers/opencl-c-header.cl
|
||||
index 2716076acdcf..6b3eca84e8b9 100644
|
||||
--- a/clang/test/Headers/opencl-c-header.cl
|
||||
+++ b/clang/test/Headers/opencl-c-header.cl
|
||||
@@ -98,3 +98,88 @@ global atomic_int z = ATOMIC_VAR_INIT(99);
|
||||
#pragma OPENCL EXTENSION cl_intel_planar_yuv : enable
|
||||
|
||||
// CHECK-MOD: Reading modules
|
||||
+
|
||||
+// For SPIR all extensions are supported.
|
||||
+#if defined(__SPIR__)
|
||||
+
|
||||
+#if (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200)
|
||||
+
|
||||
+#if __opencl_c_ext_fp16_global_atomic_load_store != 1
|
||||
+#error "Incorrectly defined __opencl_c_ext_fp16_global_atomic_load_store"
|
||||
+#endif
|
||||
+#if __opencl_c_ext_fp16_local_atomic_load_store != 1
|
||||
+#error "Incorrectly defined __opencl_c_ext_fp16_local_atomic_load_store"
|
||||
+#endif
|
||||
+#if __opencl_c_ext_fp16_global_atomic_add != 1
|
||||
+#error "Incorrectly defined __opencl_c_ext_fp16_global_atomic_add"
|
||||
+#endif
|
||||
+#if __opencl_c_ext_fp32_global_atomic_add != 1
|
||||
+#error "Incorrectly defined __opencl_c_ext_fp32_global_atomic_add"
|
||||
+#endif
|
||||
+#if __opencl_c_ext_fp16_local_atomic_add != 1
|
||||
+#error "Incorrectly defined __opencl_c_ext_fp16_local_atomic_add"
|
||||
+#endif
|
||||
+#if __opencl_c_ext_fp32_local_atomic_add != 1
|
||||
+#error "Incorrectly defined __opencl_c_ext_fp32_local_atomic_add"
|
||||
+#endif
|
||||
+#if __opencl_c_ext_fp16_global_atomic_min_max != 1
|
||||
+#error "Incorrectly defined __opencl_c_ext_fp16_global_atomic_min_max"
|
||||
+#endif
|
||||
+#if __opencl_c_ext_fp32_global_atomic_min_max != 1
|
||||
+#error "Incorrectly defined __opencl_c_ext_fp32_global_atomic_min_max"
|
||||
+#endif
|
||||
+#if __opencl_c_ext_fp16_local_atomic_min_max != 1
|
||||
+#error "Incorrectly defined __opencl_c_ext_fp16_local_atomic_min_max"
|
||||
+#endif
|
||||
+#if __opencl_c_ext_fp32_local_atomic_min_max != 1
|
||||
+#error "Incorrectly defined __opencl_c_ext_fp32_local_atomic_min_max"
|
||||
+#endif
|
||||
+
|
||||
+#else
|
||||
+#ifdef __opencl_c_ext_fp16_global_atomic_load_store
|
||||
+#error "Incorrectly __opencl_c_ext_fp16_global_atomic_load_store defined"
|
||||
+#endif
|
||||
+#ifdef __opencl_c_ext_fp16_local_atomic_load_store
|
||||
+#error "Incorrectly __opencl_c_ext_fp16_local_atomic_load_store defined"
|
||||
+#endif
|
||||
+#ifdef __opencl_c_ext_fp16_global_atomic_add
|
||||
+#error "Incorrectly __opencl_c_ext_fp16_global_atomic_add defined"
|
||||
+#endif
|
||||
+#ifdef __opencl_c_ext_fp32_global_atomic_add
|
||||
+#error "Incorrectly __opencl_c_ext_fp32_global_atomic_add defined"
|
||||
+#endif
|
||||
+#ifdef __opencl_c_ext_fp64_global_atomic_add
|
||||
+#error "Incorrectly __opencl_c_ext_fp64_global_atomic_add defined"
|
||||
+#endif
|
||||
+#ifdef __opencl_c_ext_fp16_local_atomic_add
|
||||
+#error "Incorrectly __opencl_c_ext_fp16_local_atomic_add defined"
|
||||
+#endif
|
||||
+#ifdef __opencl_c_ext_fp32_local_atomic_add
|
||||
+#error "Incorrectly __opencl_c_ext_fp32_local_atomic_add defined"
|
||||
+#endif
|
||||
+#ifdef __opencl_c_ext_fp64_local_atomic_add
|
||||
+#error "Incorrectly __opencl_c_ext_fp64_local_atomic_add defined"
|
||||
+#endif
|
||||
+#ifdef __opencl_c_ext_fp16_global_atomic_min_max
|
||||
+#error "Incorrectly __opencl_c_ext_fp16_global_atomic_min_max defined"
|
||||
+#endif
|
||||
+#ifdef __opencl_c_ext_fp32_global_atomic_min_max
|
||||
+#error "Incorrectly __opencl_c_ext_fp32_global_atomic_min_max defined"
|
||||
+#endif
|
||||
+#ifdef __opencl_c_ext_fp64_global_atomic_min_max
|
||||
+#error "Incorrectly __opencl_c_ext_fp64_global_atomic_min_max defined"
|
||||
+#endif
|
||||
+#ifdef __opencl_c_ext_fp16_local_atomic_min_max
|
||||
+#error "Incorrectly __opencl_c_ext_fp16_local_atomic_min_max defined"
|
||||
+#endif
|
||||
+#ifdef __opencl_c_ext_fp32_local_atomic_min_max
|
||||
+#error "Incorrectly __opencl_c_ext_fp32_local_atomic_min_max defined"
|
||||
+#endif
|
||||
+#ifdef __opencl_c_ext_fp64_local_atomic_min_max
|
||||
+#error "Incorrectly __opencl_c_ext_fp64_local_atomic_min_max defined"
|
||||
+#endif
|
||||
+
|
||||
+#endif //(defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200)
|
||||
+
|
||||
+#endif // defined(__SPIR__)
|
||||
+
|
||||
--
|
||||
2.17.1
|
||||
|
||||
|
|
@ -1,47 +0,0 @@
|
|||
From ef2b930a8e33078449737a93e7d522b2280ec58c Mon Sep 17 00:00:00 2001
|
||||
From: Naveen Saini <naveen.kumar.saini@intel.com>
|
||||
Date: Fri, 27 Aug 2021 11:39:16 +0800
|
||||
Subject: [PATCH 1/2] This patch is needed for ISPC for Gen only
|
||||
|
||||
Transformation of add to or is not safe for VC backend.
|
||||
|
||||
Upstream-Status: Backport [Taken from ispc,https://github.com/ispc/ispc/blob/v1.16.1/llvm_patches/11_0_11_1_disable-A-B-A-B-in-InstCombine.patch]
|
||||
|
||||
Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com>
|
||||
---
|
||||
.../lib/Transforms/InstCombine/InstCombineAddSub.cpp | 12 ++++++++----
|
||||
1 file changed, 8 insertions(+), 4 deletions(-)
|
||||
|
||||
diff --git a/llvm/lib/Transforms/InstCombine/InstCombineAddSub.cpp b/llvm/lib/Transforms/InstCombine/InstCombineAddSub.cpp
|
||||
index a7f5e0a7774d..bf02b0f70827 100644
|
||||
--- a/llvm/lib/Transforms/InstCombine/InstCombineAddSub.cpp
|
||||
+++ b/llvm/lib/Transforms/InstCombine/InstCombineAddSub.cpp
|
||||
@@ -15,6 +15,7 @@
|
||||
#include "llvm/ADT/APInt.h"
|
||||
#include "llvm/ADT/STLExtras.h"
|
||||
#include "llvm/ADT/SmallVector.h"
|
||||
+#include "llvm/ADT/Triple.h"
|
||||
#include "llvm/Analysis/InstructionSimplify.h"
|
||||
#include "llvm/Analysis/ValueTracking.h"
|
||||
#include "llvm/IR/Constant.h"
|
||||
@@ -1324,10 +1325,13 @@ Instruction *InstCombiner::visitAdd(BinaryOperator &I) {
|
||||
return BinaryOperator::CreateSRem(RHS, NewRHS);
|
||||
}
|
||||
}
|
||||
-
|
||||
- // A+B --> A|B iff A and B have no bits set in common.
|
||||
- if (haveNoCommonBitsSet(LHS, RHS, DL, &AC, &I, &DT))
|
||||
- return BinaryOperator::CreateOr(LHS, RHS);
|
||||
+
|
||||
+ // Disable this transformation for ISPC SPIR-V
|
||||
+ if (!Triple(I.getModule()->getTargetTriple()).isSPIR()) {
|
||||
+ // A+B --> A|B iff A and B have no bits set in common.
|
||||
+ if (haveNoCommonBitsSet(LHS, RHS, DL, &AC, &I, &DT))
|
||||
+ return BinaryOperator::CreateOr(LHS, RHS);
|
||||
+ }
|
||||
|
||||
// FIXME: We already did a check for ConstantInt RHS above this.
|
||||
// FIXME: Is this pattern covered by another fold? No regression tests fail on
|
||||
--
|
||||
2.17.1
|
||||
|
||||
|
|
@ -1,95 +0,0 @@
|
|||
From c20838176e8bea9e5a176c59c78bbce9051ec987 Mon Sep 17 00:00:00 2001
|
||||
From: Naveen Saini <naveen.kumar.saini@intel.com>
|
||||
Date: Fri, 27 Aug 2021 11:41:47 +0800
|
||||
Subject: [PATCH 2/2] [X86] When storing v1i1/v2i1/v4i1 to memory, make sure we
|
||||
store zeros in the rest of the byte
|
||||
|
||||
We can't store garbage in the unused bits. It possible that something like zextload from i1/i2/i4 is created to read the memory. Those zextloads would be legalized assuming the extra bits are 0.
|
||||
|
||||
I'm not sure that the code in lowerStore is executed for the v1i1/v2i1/v4i1 case. It looks like the DAG combine in combineStore may have converted them to v8i1 first. And I think we're missing some cases to avoid going to the stack in the first place. But I don't have time to investigate those things at the moment so I wanted to focus on the correctness issue.
|
||||
|
||||
Should fix PR48147.
|
||||
|
||||
Reviewed By: RKSimon
|
||||
|
||||
Differential Revision: https://reviews.llvm.org/D91294
|
||||
|
||||
Upstream-Status: Backport [https://github.com/llvm/llvm-project/commit/a4124e455e641db1e18d4221d2dacb31953fd13b]
|
||||
|
||||
Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com>
|
||||
---
|
||||
llvm/lib/Target/X86/X86ISelLowering.cpp | 19 ++++++++++++++-----
|
||||
llvm/lib/Target/X86/X86InstrAVX512.td | 3 ---
|
||||
2 files changed, 14 insertions(+), 8 deletions(-)
|
||||
|
||||
diff --git a/llvm/lib/Target/X86/X86ISelLowering.cpp b/llvm/lib/Target/X86/X86ISelLowering.cpp
|
||||
index 56690c3c555b..7e673a3163b7 100644
|
||||
--- a/llvm/lib/Target/X86/X86ISelLowering.cpp
|
||||
+++ b/llvm/lib/Target/X86/X86ISelLowering.cpp
|
||||
@@ -23549,17 +23549,22 @@ static SDValue LowerStore(SDValue Op, const X86Subtarget &Subtarget,
|
||||
// Without AVX512DQ, we need to use a scalar type for v2i1/v4i1/v8i1 stores.
|
||||
if (StoredVal.getValueType().isVector() &&
|
||||
StoredVal.getValueType().getVectorElementType() == MVT::i1) {
|
||||
- assert(StoredVal.getValueType().getVectorNumElements() <= 8 &&
|
||||
- "Unexpected VT");
|
||||
+ unsigned NumElts = StoredVal.getValueType().getVectorNumElements();
|
||||
+ assert(NumElts <= 8 && "Unexpected VT");
|
||||
assert(!St->isTruncatingStore() && "Expected non-truncating store");
|
||||
assert(Subtarget.hasAVX512() && !Subtarget.hasDQI() &&
|
||||
"Expected AVX512F without AVX512DQI");
|
||||
|
||||
+ // We must pad with zeros to ensure we store zeroes to any unused bits.
|
||||
StoredVal = DAG.getNode(ISD::INSERT_SUBVECTOR, dl, MVT::v16i1,
|
||||
DAG.getUNDEF(MVT::v16i1), StoredVal,
|
||||
DAG.getIntPtrConstant(0, dl));
|
||||
StoredVal = DAG.getBitcast(MVT::i16, StoredVal);
|
||||
StoredVal = DAG.getNode(ISD::TRUNCATE, dl, MVT::i8, StoredVal);
|
||||
+ // Make sure we store zeros in the extra bits.
|
||||
+ if (NumElts < 8)
|
||||
+ StoredVal = DAG.getZeroExtendInReg(StoredVal, dl,
|
||||
+ MVT::getIntegerVT(NumElts));
|
||||
|
||||
return DAG.getStore(St->getChain(), dl, StoredVal, St->getBasePtr(),
|
||||
St->getPointerInfo(), St->getOriginalAlign(),
|
||||
@@ -44133,17 +44138,21 @@ static SDValue combineStore(SDNode *N, SelectionDAG &DAG,
|
||||
if (VT == MVT::v1i1 && VT == StVT && Subtarget.hasAVX512() &&
|
||||
StoredVal.getOpcode() == ISD::SCALAR_TO_VECTOR &&
|
||||
StoredVal.getOperand(0).getValueType() == MVT::i8) {
|
||||
- return DAG.getStore(St->getChain(), dl, StoredVal.getOperand(0),
|
||||
+ SDValue Val = StoredVal.getOperand(0);
|
||||
+ // We must store zeros to the unused bits.
|
||||
+ Val = DAG.getZeroExtendInReg(Val, dl, MVT::i1);
|
||||
+ return DAG.getStore(St->getChain(), dl, Val,
|
||||
St->getBasePtr(), St->getPointerInfo(),
|
||||
St->getOriginalAlign(),
|
||||
St->getMemOperand()->getFlags());
|
||||
}
|
||||
|
||||
// Widen v2i1/v4i1 stores to v8i1.
|
||||
- if ((VT == MVT::v2i1 || VT == MVT::v4i1) && VT == StVT &&
|
||||
+ if ((VT == MVT::v1i1 || VT == MVT::v2i1 || VT == MVT::v4i1) && VT == StVT &&
|
||||
Subtarget.hasAVX512()) {
|
||||
unsigned NumConcats = 8 / VT.getVectorNumElements();
|
||||
- SmallVector<SDValue, 4> Ops(NumConcats, DAG.getUNDEF(VT));
|
||||
+ // We must store zeros to the unused bits.
|
||||
+ SmallVector<SDValue, 4> Ops(NumConcats, DAG.getConstant(0, dl, VT));
|
||||
Ops[0] = StoredVal;
|
||||
StoredVal = DAG.getNode(ISD::CONCAT_VECTORS, dl, MVT::v8i1, Ops);
|
||||
return DAG.getStore(St->getChain(), dl, StoredVal, St->getBasePtr(),
|
||||
diff --git a/llvm/lib/Target/X86/X86InstrAVX512.td b/llvm/lib/Target/X86/X86InstrAVX512.td
|
||||
index a3ad0b1c8dd6..aa1ccec02f2a 100644
|
||||
--- a/llvm/lib/Target/X86/X86InstrAVX512.td
|
||||
+++ b/llvm/lib/Target/X86/X86InstrAVX512.td
|
||||
@@ -2871,9 +2871,6 @@ def : Pat<(i64 (bitconvert (v64i1 VK64:$src))),
|
||||
|
||||
// Load/store kreg
|
||||
let Predicates = [HasDQI] in {
|
||||
- def : Pat<(store VK1:$src, addr:$dst),
|
||||
- (KMOVBmk addr:$dst, (COPY_TO_REGCLASS VK1:$src, VK8))>;
|
||||
-
|
||||
def : Pat<(v1i1 (load addr:$src)),
|
||||
(COPY_TO_REGCLASS (KMOVBkm addr:$src), VK1)>;
|
||||
def : Pat<(v2i1 (load addr:$src)),
|
||||
--
|
||||
2.17.1
|
||||
|
||||
|
|
@ -1,9 +1,6 @@
|
|||
FILESEXTRAPATHS:prepend:intel-x86-common := "${THISDIR}/files:"
|
||||
|
||||
SPIRV10_SRCREV = "fe4d6b767363a1995ccbfca27f79efb10dcfe110"
|
||||
SPIRV11_SRCREV = "ca3a50e6e3193e399d26446d4f74a90e2a531f3a"
|
||||
|
||||
SPIRV_SRCREV = "${@bb.utils.contains('LLVMVERSION', '10.0.1', '${SPIRV10_SRCREV}', '${SPIRV11_SRCREV}', d)}"
|
||||
SPIRV_SRCREV = "fe4d6b767363a1995ccbfca27f79efb10dcfe110"
|
||||
|
||||
SRC_URI_LLVM10_PATCHES = " \
|
||||
file://llvm10-0001-llvm-spirv-skip-building-tests.patch;patchdir=llvm/projects/llvm-spirv \
|
||||
|
|
@ -28,18 +25,6 @@ SRC_URI_LLVM10_PATCHES = " \
|
|||
file://llvm10-0015-ispc-10_0_vXi1calling_avx512_i8_i16.patch \
|
||||
"
|
||||
|
||||
SRC_URI_LLVM11_PATCHES = " \
|
||||
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 \
|
||||
file://llvm11-0007-ispc-11_0_11_1_disable-A-B-A-B-in-InstCombine.patch \
|
||||
file://llvm11-0008-ispc-11_0_11_1_packed_load_store_avx512.patch \
|
||||
"
|
||||
SRC_URI_LLVM12_PATCHES = " \
|
||||
file://llvm12-0001-Remove-__IMAGE_SUPPORT__-macro-for-SPIR-since-SPIR-d.patch \
|
||||
file://llvm12-0002-Avoid-calling-ParseCommandLineOptions-in-BackendUtil.patch \
|
||||
|
|
@ -51,12 +36,7 @@ SRC_URI_LLVM12_PATCHES = " \
|
|||
|
||||
SPIRV_LLVM10_SRC_URI = "git://github.com/KhronosGroup/SPIRV-LLVM-Translator.git;protocol=https;branch=llvm_release_100;destsuffix=git/llvm/projects/llvm-spirv;name=spirv"
|
||||
|
||||
SPIRV_LLVM11_SRC_URI = "git://github.com/KhronosGroup/SPIRV-LLVM-Translator.git;protocol=https;branch=llvm_release_110;destsuffix=git/llvm/projects/llvm-spirv;name=spirv"
|
||||
|
||||
|
||||
|
||||
SRC_URI:append:intel-x86-common = "${@bb.utils.contains('LLVMVERSION', '10.0.1', ' ${SPIRV_LLVM10_SRC_URI} ${SRC_URI_LLVM10_PATCHES} ', '', d)}"
|
||||
SRC_URI:append:intel-x86-common = "${@bb.utils.contains('LLVMVERSION', '11.1.0', ' ${SPIRV_LLVM11_SRC_URI} ${SRC_URI_LLVM11_PATCHES} ', '', d)}"
|
||||
SRC_URI:append:intel-x86-common = "${@bb.utils.contains('LLVMVERSION', '12.0.0', ' ${SRC_URI_LLVM12_PATCHES} ', '', d)}"
|
||||
|
||||
SRCREV_spirv = "${@bb.utils.contains_any('LLVMVERSION', [ '13.0.0', '12.0.0' ], '', '${SPIRV_SRCREV}', d)}"
|
||||
|
|
|
|||
|
|
@ -1,15 +0,0 @@
|
|||
require opencl-clang.inc
|
||||
|
||||
SRC_URI:append = " file://0001-don-t-redefine-LLVM_TABLEGEN_EXE.patch \
|
||||
"
|
||||
SRC_URI:append:class-native = " file://0002-make-sure-only-static-libraries-linked-for-native-bu.patch"
|
||||
|
||||
SRCREV = "c67648d41df00ea8ee9d701d17299b86f86f0321"
|
||||
|
||||
BRANCH = "ocl-open-110"
|
||||
|
||||
EXTRA_OECMAKE += "\
|
||||
-DLLVM_TABLEGEN_EXE=${STAGING_BINDIR_NATIVE}/llvm-tblgen \
|
||||
-DCMAKE_SKIP_RPATH=TRUE \
|
||||
-DPREFERRED_LLVM_VERSION="11.1.0" \
|
||||
"
|
||||
Loading…
Reference in New Issue
Block a user