From 3f8bc66c390c63092cfa090fad23eae396e5a4fa Mon Sep 17 00:00:00 2001 From: Vlad Romanov Date: Fri, 29 Nov 2019 22:06:14 +0300 Subject: [PATCH 1/3] [SYCL] Remove ASFixer and DISABLE_INFER_AS support Signed-off-by: Vlad Romanov --- clang/lib/Basic/Targets/SPIR.h | 3 +- clang/lib/CodeGen/BackendUtil.cpp | 5 - clang/lib/CodeGen/CGCall.cpp | 20 +- clang/lib/CodeGen/CGExpr.cpp | 18 +- clang/lib/CodeGen/CGStmt.cpp | 35 +- clang/lib/CodeGen/CMakeLists.txt | 1 - clang/lib/CodeGen/CodeGenModule.cpp | 8 +- clang/lib/Frontend/InitPreprocessor.cpp | 2 - clang/test/CodeGenSYCL/address-space-new.cpp | 89 +-- .../address-space-parameter-conversions.cpp | 201 ------ .../test/CodeGenSYCL/basic-kernel-wrapper.cpp | 14 +- clang/test/CodeGenSYCL/device-functions.cpp | 9 +- clang/test/CodeGenSYCL/sampler.cpp | 11 +- clang/test/CodeGenSYCL/spir-calling-conv.cpp | 12 +- clang/test/SemaSYCL/spir-enum.cpp | 9 +- llvm/include/llvm/InitializePasses.h | 1 - llvm/include/llvm/SYCL/ASFixer.h | 25 - llvm/lib/CMakeLists.txt | 1 - llvm/lib/LLVMBuild.txt | 1 - llvm/lib/SYCL/ASFixer.cpp | 602 ------------------ llvm/lib/SYCL/CMakeLists.txt | 9 - llvm/lib/SYCL/LLVMBuild.txt | 21 - .../address-space-parameter-conversions.ll | 67 -- llvm/test/SYCL/ASFixer/binary-operators.ll | 120 ---- .../ASFixer/numbers-and-local-variables.ll | 93 --- llvm/tools/opt/CMakeLists.txt | 1 - llvm/tools/opt/opt.cpp | 1 - 27 files changed, 81 insertions(+), 1298 deletions(-) delete mode 100644 clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp delete mode 100644 llvm/include/llvm/SYCL/ASFixer.h delete mode 100644 llvm/lib/SYCL/ASFixer.cpp delete mode 100644 llvm/lib/SYCL/CMakeLists.txt delete mode 100644 llvm/lib/SYCL/LLVMBuild.txt delete mode 100644 llvm/test/SYCL/ASFixer/address-space-parameter-conversions.ll delete mode 100644 llvm/test/SYCL/ASFixer/binary-operators.ll delete mode 100644 llvm/test/SYCL/ASFixer/numbers-and-local-variables.ll diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h index dfc37bb5ad68f..475c750f3cf86 100644 --- a/clang/lib/Basic/Targets/SPIR.h +++ b/clang/lib/Basic/Targets/SPIR.h @@ -64,8 +64,7 @@ class LLVM_LIBRARY_VISIBILITY SPIRTargetInfo : public TargetInfo { TLSSupported = false; VLASupported = false; LongWidth = LongAlign = 64; - if (Triple.getEnvironment() == llvm::Triple::SYCLDevice && - !getenv("DISABLE_INFER_AS")) { + if (Triple.getEnvironment() == llvm::Triple::SYCLDevice) { AddrSpaceMap = &SYCLAddrSpaceMap; } else { AddrSpaceMap = &SPIRAddrSpaceMap; diff --git a/clang/lib/CodeGen/BackendUtil.cpp b/clang/lib/CodeGen/BackendUtil.cpp index 47a0209e67395..d6a2c36169862 100644 --- a/clang/lib/CodeGen/BackendUtil.cpp +++ b/clang/lib/CodeGen/BackendUtil.cpp @@ -39,7 +39,6 @@ #include "llvm/Passes/PassBuilder.h" #include "llvm/Passes/PassPlugin.h" #include "llvm/Passes/StandardInstrumentations.h" -#include "llvm/SYCL/ASFixer.h" #include "llvm/Support/BuryPointer.h" #include "llvm/Support/CommandLine.h" #include "llvm/Support/Debug.h" @@ -854,8 +853,6 @@ void EmitAssemblyHelper::EmitAssembly(BackendAction Action, case Backend_EmitBC: if (LangOpts.SYCLIsDevice) { - if (getenv("DISABLE_INFER_AS")) - PerModulePasses.add(createASFixerPass()); PerModulePasses.add(createDeadCodeEliminationPass()); } if (CodeGenOpts.PrepareForThinLTO && !CodeGenOpts.DisableLLVMPasses) { @@ -1282,8 +1279,6 @@ void EmitAssemblyHelper::EmitAssemblyWithNewPassManager( case Backend_EmitBC: if (LangOpts.SYCLIsDevice) { - if (getenv("DISABLE_INFER_AS")) - CodeGenPasses.add(createASFixerPass()); CodeGenPasses.add(createDeadCodeEliminationPass()); } if (CodeGenOpts.PrepareForThinLTO && !CodeGenOpts.DisableLLVMPasses) { diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp index 266700b5a8f45..3d292f84c79b4 100644 --- a/clang/lib/CodeGen/CGCall.cpp +++ b/clang/lib/CodeGen/CGCall.cpp @@ -4316,17 +4316,15 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo, deactivateArgCleanupsBeforeCall(*this, CallArgs); // Addrspace cast to generic if necessary - if (!getenv("DISABLE_INFER_AS")) { - for (unsigned i = 0; i < IRFuncTy->getNumParams(); ++i) { - if (auto *PtrTy = dyn_cast(IRCallArgs[i]->getType())) { - auto *ExpectedPtrType = - cast(IRFuncTy->getParamType(i)); - unsigned ValueAS = PtrTy->getAddressSpace(); - unsigned ExpectedAS = ExpectedPtrType->getAddressSpace(); - if (ValueAS != ExpectedAS) { - IRCallArgs[i] = Builder.CreatePointerBitCastOrAddrSpaceCast( - IRCallArgs[i], ExpectedPtrType); - } + for (unsigned i = 0; i < IRFuncTy->getNumParams(); ++i) { + if (auto *PtrTy = dyn_cast(IRCallArgs[i]->getType())) { + auto *ExpectedPtrType = + cast(IRFuncTy->getParamType(i)); + unsigned ValueAS = PtrTy->getAddressSpace(); + unsigned ExpectedAS = ExpectedPtrType->getAddressSpace(); + if (ValueAS != ExpectedAS) { + IRCallArgs[i] = Builder.CreatePointerBitCastOrAddrSpaceCast( + IRCallArgs[i], ExpectedPtrType); } } } diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp index ae653705084e7..6becd35976d85 100644 --- a/clang/lib/CodeGen/CGExpr.cpp +++ b/clang/lib/CodeGen/CGExpr.cpp @@ -1748,16 +1748,14 @@ void CodeGenFunction::EmitStoreOfScalar(llvm::Value *Value, Address Addr, return; } - if (!getenv("DISABLE_INFER_AS")) { - if (auto *PtrTy = dyn_cast(Value->getType())) { - auto *ExpectedPtrType = - cast(Addr.getType()->getElementType()); - unsigned ValueAS = PtrTy->getAddressSpace(); - unsigned ExpectedAS = ExpectedPtrType->getAddressSpace(); - if (ValueAS != ExpectedAS) { - Value = - Builder.CreatePointerBitCastOrAddrSpaceCast(Value, ExpectedPtrType); - } + if (auto *PtrTy = dyn_cast(Value->getType())) { + auto *ExpectedPtrType = + cast(Addr.getType()->getElementType()); + unsigned ValueAS = PtrTy->getAddressSpace(); + unsigned ExpectedAS = ExpectedPtrType->getAddressSpace(); + if (ValueAS != ExpectedAS) { + Value = + Builder.CreatePointerBitCastOrAddrSpaceCast(Value, ExpectedPtrType); } } llvm::StoreInst *Store = Builder.CreateStore(Value, Addr, Volatile); diff --git a/clang/lib/CodeGen/CGStmt.cpp b/clang/lib/CodeGen/CGStmt.cpp index ffe0320242f73..f3a4e98edc3a5 100644 --- a/clang/lib/CodeGen/CGStmt.cpp +++ b/clang/lib/CodeGen/CGStmt.cpp @@ -1110,16 +1110,13 @@ void CodeGenFunction::EmitReturnStmt(const ReturnStmt &S) { // rather than the value. RValue Result = EmitReferenceBindingToExpr(RV); llvm::Value *Val = Result.getScalarVal(); - if (!getenv("DISABLE_INFER_AS")) { - if (auto *PtrTy = dyn_cast(Val->getType())) { - auto *ExpectedPtrType = - cast(ReturnValue.getType()->getElementType()); - unsigned ValueAS = PtrTy->getAddressSpace(); - unsigned ExpectedAS = ExpectedPtrType->getAddressSpace(); - if (ValueAS != ExpectedAS) { - Val = - Builder.CreatePointerBitCastOrAddrSpaceCast(Val, ExpectedPtrType); - } + if (auto *PtrTy = dyn_cast(Val->getType())) { + auto *ExpectedPtrType = + cast(ReturnValue.getType()->getElementType()); + unsigned ValueAS = PtrTy->getAddressSpace(); + unsigned ExpectedAS = ExpectedPtrType->getAddressSpace(); + if (ValueAS != ExpectedAS) { + Val = Builder.CreatePointerBitCastOrAddrSpaceCast(Val, ExpectedPtrType); } } Builder.CreateStore(Val, ReturnValue); @@ -1128,16 +1125,14 @@ void CodeGenFunction::EmitReturnStmt(const ReturnStmt &S) { case TEK_Scalar: { llvm::Value *Val = EmitScalarExpr(RV); - if (!getenv("DISABLE_INFER_AS")) { - if (auto *PtrTy = dyn_cast(Val->getType())) { - auto *ExpectedPtrType = - cast(ReturnValue.getType()->getElementType()); - unsigned ValueAS = PtrTy->getAddressSpace(); - unsigned ExpectedAS = ExpectedPtrType->getAddressSpace(); - if (ValueAS != ExpectedAS) { - Val = Builder.CreatePointerBitCastOrAddrSpaceCast( - Val, ExpectedPtrType); - } + if (auto *PtrTy = dyn_cast(Val->getType())) { + auto *ExpectedPtrType = + cast(ReturnValue.getType()->getElementType()); + unsigned ValueAS = PtrTy->getAddressSpace(); + unsigned ExpectedAS = ExpectedPtrType->getAddressSpace(); + if (ValueAS != ExpectedAS) { + Val = + Builder.CreatePointerBitCastOrAddrSpaceCast(Val, ExpectedPtrType); } } Builder.CreateStore(Val, ReturnValue); diff --git a/clang/lib/CodeGen/CMakeLists.txt b/clang/lib/CodeGen/CMakeLists.txt index 3a8f1b0beb405..22f7ef97d9349 100644 --- a/clang/lib/CodeGen/CMakeLists.txt +++ b/clang/lib/CodeGen/CMakeLists.txt @@ -24,7 +24,6 @@ set(LLVM_LINK_COMPONENTS Support Target TransformUtils - ASFixer ) # In a standard Clang+LLVM build, we need to generate intrinsics before diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 003992abb9765..343f95b942814 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -3775,10 +3775,8 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) { if (Scope && Scope->isWorkGroup()) return LangAS::opencl_local; - if (!getenv("DISABLE_INFER_AS")) { - if (!D || D->getType().getAddressSpace() == LangAS::Default) { - return LangAS::opencl_global; - } + if (!D || D->getType().getAddressSpace() == LangAS::Default) { + return LangAS::opencl_global; } } @@ -3807,7 +3805,7 @@ LangAS CodeGenModule::getStringLiteralAddressSpace() const { // OpenCL v1.2 s6.5.3: a string literal is in the constant address space. if (LangOpts.OpenCL) return LangAS::opencl_constant; - if (LangOpts.SYCLIsDevice && !getenv("DISABLE_INFER_AS")) + if (LangOpts.SYCLIsDevice) // If we keep a literal string in constant address space, the following code // becomes illegal: // diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp index 9792a8c27e600..17b8640890db8 100644 --- a/clang/lib/Frontend/InitPreprocessor.cpp +++ b/clang/lib/Frontend/InitPreprocessor.cpp @@ -1089,8 +1089,6 @@ static void InitializePredefinedMacros(const TargetInfo &TI, if (LangOpts.SYCLIsDevice) { Builder.defineMacro("__SYCL_DEVICE_ONLY__", "1"); Builder.defineMacro("SYCL_EXTERNAL", "__attribute__((sycl_device))"); - if (!getenv("DISABLE_INFER_AS")) - Builder.defineMacro("__SYCL_ENABLE_INFER_AS__", "1"); } if (LangOpts.SYCLUnnamedLambda) Builder.defineMacro("__SYCL_UNNAMED_LAMBDA__", "1"); diff --git a/clang/test/CodeGenSYCL/address-space-new.cpp b/clang/test/CodeGenSYCL/address-space-new.cpp index 64e3601b48e55..018aac0d52e7b 100644 --- a/clang/test/CodeGenSYCL/address-space-new.cpp +++ b/clang/test/CodeGenSYCL/address-space-new.cpp @@ -1,5 +1,4 @@ -// RUN: DISABLE_INFER_AS=1 %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -emit-llvm -x c++ %s -o - | FileCheck %s --check-prefixes=CHECK,CHECK-LEGACY -// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -emit-llvm -x c++ %s -o - | FileCheck %s --check-prefixes=CHECK,CHECK-NEW +// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -emit-llvm -x c++ %s -o - | FileCheck %s struct SpaceWaster { int i, j; @@ -19,8 +18,7 @@ void baz(Y &y) { void test() { static const int foo = 0x42; - // CHECK-LEGACY: @_ZZ4testvE3foo = internal constant i32 66, align 4 - // CHECK-NEW: @_ZZ4testvE3foo = internal addrspace(1) constant i32 66, align 4 + // CHECK: @_ZZ4testvE3foo = internal addrspace(1) constant i32 66, align 4 // Intentionally leave a part of an array uninitialized. This triggers a // different code path contrary to a fully initialized array. @@ -29,60 +27,49 @@ void test() { 11, 12, 13, 14, 15, 16, 17, 18, 19, 20 }; (void)bars; - // CHECK-LEGACY: @_ZZ4testvE4bars = internal constant <{ [21 x i32], [235 x i32] }> <{ [21 x i32] [i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15, i32 16, i32 17, i32 18, i32 19, i32 20], [235 x i32] zeroinitializer }>, align 4 - // CHECK-NEW: @_ZZ4testvE4bars = internal addrspace(1) constant <{ [21 x i32], [235 x i32] }> <{ [21 x i32] [i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15, i32 16, i32 17, i32 18, i32 19, i32 20], [235 x i32] zeroinitializer }>, align 4 + // CHECK: @_ZZ4testvE4bars = internal addrspace(1) constant <{ [21 x i32], [235 x i32] }> <{ [21 x i32] [i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15, i32 16, i32 17, i32 18, i32 19, i32 20], [235 x i32] zeroinitializer }>, align 4 // CHECK: @[[STR:[.a-zA-Z0-9_]+]] = private unnamed_addr constant [14 x i8] c"Hello, world!\00", align 1 - // CHECK-NEW: %i.ascast = addrspacecast i32* %i to i32 addrspace(4)* + // CHECK: %i.ascast = addrspacecast i32* %i to i32 addrspace(4)* // CHECK: %[[ARR:[a-zA-Z0-9]+]] = alloca [42 x i32] int i = 0; int *pptr = &i; - // CHECK-LEGACY: store i32* %i, i32** %pptr - // CHECK-NEW: %[[GEN:[0-9]+]] = addrspacecast i32* %i to i32 addrspace(4)* - // CHECK-NEW: store i32 addrspace(4)* %[[GEN]], i32 addrspace(4)** %pptr + // CHECK: %[[GEN:[0-9]+]] = addrspacecast i32* %i to i32 addrspace(4)* + // CHECK: store i32 addrspace(4)* %[[GEN]], i32 addrspace(4)** %pptr bool is_i_ptr = (pptr == &i); - // CHECK-LEGACY: %[[VALPPTR:[0-9]+]] = load i32*, i32** %pptr - // CHECK-LEGACY: %cmp{{[0-9]*}} = icmp eq i32* %[[VALPPTR]], %i - // CHECK-NEW: %[[VALPPTR:[0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)** %pptr - // CHECK-NEW: %cmp{{[0-9]*}} = icmp eq i32 addrspace(4)* %[[VALPPTR]], %i.ascast + // CHECK: %[[VALPPTR:[0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)** %pptr + // CHECK: %cmp{{[0-9]*}} = icmp eq i32 addrspace(4)* %[[VALPPTR]], %i.ascast *pptr = foo; int var23 = 23; char *cp = (char *)&var23; *cp = 41; // CHECK: store i32 23, i32* %[[VAR:[a-zA-Z0-9]+]] - // CHECK-LEGACY: [[VARCAST:[a-zA-Z0-9]+]] = bitcast i32* %[[VAR]] to i8* - // CHECK-LEGACY: store i8* %[[VARCAST]], i8** %{{.*}} - // CHECK-NEW: [[VARAS:[a-zA-Z0-9]+]] = addrspacecast i32* %[[VAR]] to i32 addrspace(4)* - // CHECK-NEW: [[VARCAST:[a-zA-Z0-9]+]] = bitcast i32 addrspace(4)* %[[VARAS]] to i8 addrspace(4)* - // CHECK-NEW: store i8 addrspace(4)* %[[VARCAST]], i8 addrspace(4)** %{{.*}} + // CHECK: [[VARAS:[a-zA-Z0-9]+]] = addrspacecast i32* %[[VAR]] to i32 addrspace(4)* + // CHECK: [[VARCAST:[a-zA-Z0-9]+]] = bitcast i32 addrspace(4)* %[[VARAS]] to i8 addrspace(4)* + // CHECK: store i8 addrspace(4)* %[[VARCAST]], i8 addrspace(4)** %{{.*}} int arr[42]; char *cpp = (char *)arr; *cpp = 43; // CHECK: %[[ARRDECAY:[a-zA-Z0-9]+]] = getelementptr inbounds [42 x i32], [42 x i32]* %[[ARR]], i64 0, i64 0 - // CHECK-LEGACY: [[ARRCAST:[a-zA-Z0-9]+]] = bitcast i32* %[[ARRDECAY]] to i8* - // CHECK-LEGACY: store i8* %[[ARRCAST]], i8** %{{.*}} - // CHECK-NEW: %[[ARRAS:[a-zA-Z0-9]+]] = addrspacecast i32* %[[ARRDECAY]] to i32 addrspace(4)* - // CHECK-NEW: %[[ARRCAST:[a-zA-Z0-9]+]] = bitcast i32 addrspace(4)* %[[ARRAS]] to i8 addrspace(4)* - // CHECK-NEW: store i8 addrspace(4)* %[[ARRCAST]], i8 addrspace(4)** %{{.*}} + // CHECK: %[[ARRAS:[a-zA-Z0-9]+]] = addrspacecast i32* %[[ARRDECAY]] to i32 addrspace(4)* + // CHECK: %[[ARRCAST:[a-zA-Z0-9]+]] = bitcast i32 addrspace(4)* %[[ARRAS]] to i8 addrspace(4)* + // CHECK: store i8 addrspace(4)* %[[ARRCAST]], i8 addrspace(4)** %{{.*}} int *aptr = arr + 10; if (aptr < arr + sizeof(arr)) *aptr = 44; - // CHECK-LEGACY: %[[VALAPTR:[0-9]+]] = load i32*, i32** %aptr - // CHECK-NEW: %[[VALAPTR:[0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)** %aptr + // CHECK: %[[VALAPTR:[0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)** %aptr // CHECK: %[[ARRDCY2:[a-zA-Z0-9]+]] = getelementptr inbounds [42 x i32], [42 x i32]* %[[ARR]], i64 0, i64 0 // CHECK: %[[ADDPTR:[a-zA-Z0-9.]+]] = getelementptr inbounds i32, i32* %[[ARRDCY2]], i64 168 - // CHECK-LEGACY: %cmp{{[0-9]+}} = icmp ult i32* %[[VALAPTR]], %[[ADDPTR]] - // CHECK-NEW: %[[ADDPTRCAST:[a-zA-Z0-9.]+]] = addrspacecast i32* %[[ADDPTR]] to i32 addrspace(4)* - // CHECK-NEW: %cmp{{[0-9]+}} = icmp ult i32 addrspace(4)* %[[VALAPTR]], %[[ADDPTRCAST]] + // CHECK: %[[ADDPTRCAST:[a-zA-Z0-9.]+]] = addrspacecast i32* %[[ADDPTR]] to i32 addrspace(4)* + // CHECK: %cmp{{[0-9]+}} = icmp ult i32 addrspace(4)* %[[VALAPTR]], %[[ADDPTRCAST]] const char *str = "Hello, world!"; - // CHECK-LEGACY: store i8* getelementptr inbounds ([14 x i8], [14 x i8]* @[[STR]], i64 0, i64 0), i8** %[[STRVAL:[a-zA-Z0-9]+]], align 8 - // CHECK-NEW: store i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([14 x i8], [14 x i8]* @[[STR]], i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)** %[[STRVAL:[a-zA-Z0-9]+]], align 8 + // CHECK: store i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([14 x i8], [14 x i8]* @[[STR]], i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)** %[[STRVAL:[a-zA-Z0-9]+]], align 8 i = str[0]; @@ -92,48 +79,35 @@ void test() { // CHECK: br i1 %[[COND]], label %[[CONDTRUE:[.a-zA-Z0-9]+]], label %[[CONDFALSE:[.a-zA-Z0-9]+]] // CHECK: [[CONDTRUE]]: - // CHECK-LEGACY-NEXT: %[[VALTRUE:[a-zA-Z0-9]+]] = load i8*, i8** %[[STRVAL]] - // CHECK-LEGACY-NEXT: br label %[[CONDEND:[.a-zA-Z0-9]+]] - // CHECK-NEW-NEXT: %[[VALTRUE:[a-zA-Z0-9]+]] = load i8 addrspace(4)*, i8 addrspace(4)** %[[STRVAL]] - // CHECK-NEW-NEXT: br label %[[CONDEND:[.a-zA-Z0-9]+]] + // CHECK-NEXT: %[[VALTRUE:[a-zA-Z0-9]+]] = load i8 addrspace(4)*, i8 addrspace(4)** %[[STRVAL]] + // CHECK-NEXT: br label %[[CONDEND:[.a-zA-Z0-9]+]] // CHECK: [[CONDFALSE]]: - // CHECK-LEGACY-NEXT: br label %[[CONDEND]] - // CHECK-LEGACY: [[CONDEND]]: - // CHECK-NEW: [[CONDEND]]: - // CHECK-NEW-NEXT: phi i8 addrspace(4)* [ %[[VALTRUE]], %[[CONDTRUE]] ], [ addrspacecast (i8* getelementptr inbounds ([21 x i8], [21 x i8]* @{{.*}}, i64 0, i64 0) to i8 addrspace(4)*), %[[CONDFALSE]] ] - // CHECK-LEGACY-NEXT: phi i8* [ %[[VALTRUE]], %[[CONDTRUE]] ], [ getelementptr inbounds ([21 x i8], [21 x i8]* @{{.*}}, i64 0, i64 0), %[[CONDFALSE]] ] + // CHECK: [[CONDEND]]: + // CHECK-NEXT: phi i8 addrspace(4)* [ %[[VALTRUE]], %[[CONDTRUE]] ], [ addrspacecast (i8* getelementptr inbounds ([21 x i8], [21 x i8]* @{{.*}}, i64 0, i64 0) to i8 addrspace(4)*), %[[CONDFALSE]] ] const char *select_null = i > 2 ? "Yet another Hello world" : nullptr; (void)select_null; - // CHECK-LEGACY: select i1 %{{.*}}, i8* getelementptr inbounds ([24 x i8], [24 x i8]* @{{.*}}, i64 0, i64 0), i8* null - // CHECK-NEW: select i1 %{{.*}}, i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([24 x i8], [24 x i8]* @{{.*}}, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* null + // CHECK: select i1 %{{.*}}, i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([24 x i8], [24 x i8]* @{{.*}}, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* null const char *select_str_trivial1 = true ? str : "Another hello world!"; (void)select_str_trivial1; - // CHECK-LEGACY: %[[TRIVIALTRUE:[a-zA-Z0-9]+]] = load i8*, i8** %[[STRVAL]] - // CHECK-LEGACY: store i8* %[[TRIVIALTRUE]], i8** %{{.*}}, align 8 - // CHECK-NEW: %[[TRIVIALTRUE:[a-zA-Z0-9]+]] = load i8 addrspace(4)*, i8 addrspace(4)** %[[STRVAL]] - // CHECK-NEW: store i8 addrspace(4)* %[[TRIVIALTRUE]], i8 addrspace(4)** %{{.*}}, align 8 + // CHECK: %[[TRIVIALTRUE:[a-zA-Z0-9]+]] = load i8 addrspace(4)*, i8 addrspace(4)** %[[STRVAL]] + // CHECK: store i8 addrspace(4)* %[[TRIVIALTRUE]], i8 addrspace(4)** %{{.*}}, align 8 const char *select_str_trivial2 = false ? str : "Another hello world!"; (void)select_str_trivial2; - // CHECK-LEGACY: store i8* getelementptr inbounds ([21 x i8], [21 x i8]* @{{.*}}, i64 0, i64 0), i8** %{{.*}} - // CHECK-NEW: store i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([21 x i8], [21 x i8]* @{{.*}}, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)** %{{.*}} + // CHECK: store i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([21 x i8], [21 x i8]* @{{.*}}, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)** %{{.*}} // // Y yy; baz(yy); // CHECK: define spir_func void @{{.*}}baz{{.*}} - // CHECK-LEGACY: %[[FIRST:[a-zA-Z0-9]+]] = bitcast %struct.{{.*}}.Y* %{{.*}} to i8* - // CHECK-LEGACY: %[[OFFSET:[a-zA-Z0-9]+]].ptr = getelementptr inbounds i8, i8* %[[FIRST]], i64 8 - // CHECK-LEGACY: %[[SECOND:[a-zA-Z0-9]+]] = bitcast i8* %[[OFFSET]].ptr to %struct.{{.*}}.HasX* - // CHECK-LEGACY: call spir_func void @{{.*}}bar{{.*}}(%struct.{{.*}}.HasX* dereferenceable(4) %[[SECOND]]) - // CHECK-NEW: %[[FIRST:[a-zA-Z0-9]+]] = bitcast %struct.{{.*}}.Y addrspace(4)* %{{.*}} to i8 addrspace(4)* - // CHECK-NEW: %[[OFFSET:[a-zA-Z0-9]+]].ptr = getelementptr inbounds i8, i8 addrspace(4)* %[[FIRST]], i64 8 - // CHECK-NEW: %[[SECOND:[a-zA-Z0-9]+]] = bitcast i8 addrspace(4)* %[[OFFSET]].ptr to %struct.{{.*}}.HasX addrspace(4)* - // CHECK-NEW: call spir_func void @{{.*}}bar{{.*}}(%struct.{{.*}}.HasX addrspace(4)* dereferenceable(4) %[[SECOND]]) + // CHECK: %[[FIRST:[a-zA-Z0-9]+]] = bitcast %struct.{{.*}}.Y addrspace(4)* %{{.*}} to i8 addrspace(4)* + // CHECK: %[[OFFSET:[a-zA-Z0-9]+]].ptr = getelementptr inbounds i8, i8 addrspace(4)* %[[FIRST]], i64 8 + // CHECK: %[[SECOND:[a-zA-Z0-9]+]] = bitcast i8 addrspace(4)* %[[OFFSET]].ptr to %struct.{{.*}}.HasX addrspace(4)* + // CHECK: call spir_func void @{{.*}}bar{{.*}}(%struct.{{.*}}.HasX addrspace(4)* dereferenceable(4) %[[SECOND]]) } @@ -147,6 +121,3 @@ int main() { kernel_single_task([]() { test(); }); return 0; } - -// TODO: SYCL specific fail - analyze and enable -// XFAIL: windows-msvc diff --git a/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp b/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp deleted file mode 100644 index 682e9455bd466..0000000000000 --- a/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp +++ /dev/null @@ -1,201 +0,0 @@ -// RUN: DISABLE_INFER_AS=1 %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -emit-llvm -x c++ %s -o - | opt -asfix -S -o - | FileCheck %s --check-prefixes CHECK,CHECK-OLD -// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -emit-llvm -x c++ %s -o - | opt -asfix -S -o - | FileCheck %s --check-prefixes CHECK,CHECK-NEW -void bar(int & Data) {} -// CHECK-OLD-DAG: define spir_func void @[[RAW_REF:[a-zA-Z0-9_]+]](i32* dereferenceable(4) % -// CHECK-NEW-DAG: define spir_func void @[[RAW_REF:[a-zA-Z0-9_]+]](i32 addrspace(4)* dereferenceable(4) % -void bar2(int & Data) {} -// CHECK-OLD-DAG: define spir_func void @[[RAW_REF2:[a-zA-Z0-9_]+]](i32* dereferenceable(4) % -// CHECK-NEW-DAG: define spir_func void @[[RAW_REF2:[a-zA-Z0-9_]+]](i32 addrspace(4)* dereferenceable(4) % -void bar(__attribute__((ocl_local)) int &Data) {} -// CHECK-DAG: define spir_func void [[LOC_REF:@[a-zA-Z0-9_]+]](i32 addrspace(3)* dereferenceable(4) % -void foo(int * Data) {} -// CHECK-OLD-DAG: define spir_func void @[[RAW_PTR:[a-zA-Z0-9_]+]](i32* % -// CHECK-NEW-DAG: define spir_func void @[[RAW_PTR:[a-zA-Z0-9_]+]](i32 addrspace(4)* % -void foo2(int * Data) {} -// CHECK-OLD-DAG: define spir_func void @[[RAW_PTR2:[a-zA-Z0-9_]+]](i32* % -// CHECK-NEW-DAG: define spir_func void @[[RAW_PTR2:[a-zA-Z0-9_]+]](i32 addrspace(4)* % -void foo(__attribute__((address_space(3))) int * Data) {} -// CHECK-DAG: define spir_func void [[LOC_PTR:@[a-zA-Z0-9_]+]](i32 addrspace(3)* % - -template -void tmpl(T t){} -// See Check Lines below. - -void usages() { - // CHECK-DAG: [[GLOB:%[a-zA-Z0-9]+]] = alloca i32 addrspace(1)* - __attribute__((address_space(1))) int *GLOB; - // CHECK-DAG: [[LOC:%[a-zA-Z0-9]+]] = alloca i32 addrspace(3)* - __attribute__((ocl_local)) int *LOC; - // CHECK-OLD-DAG: [[NoAS:%[a-zA-Z0-9]+]] = alloca i32* - // CHECK-NEW-DAG: [[NoAS:%[a-zA-Z0-9]+]] = alloca i32 addrspace(4)* - int *NoAS; - - // CHECK-DAG: [[PRIV:%[a-zA-Z0-9]+]] = alloca i32* - __attribute__((ocl_private)) int *PRIV; - - bar(*GLOB); - // CHECK-DAG: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB]] - // CHECK-DAG: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(1)* [[GLOB_LOAD]] to i32 addrspace(4)* - // CHECK-OLD-DAG: call spir_func void @new.[[RAW_REF]](i32 addrspace(4)* [[GLOB_CAST]]) - // CHECK-NEW-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[GLOB_CAST]]) - bar2(*GLOB); - // CHECK-DAG: [[GLOB_LOAD2:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB]] - // CHECK-DAG: [[GLOB_CAST2:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(1)* [[GLOB_LOAD2]] to i32 addrspace(4)* - // CHECK-OLD-DAG: call spir_func void @new.[[RAW_REF2]](i32 addrspace(4)* [[GLOB_CAST2]]) - // CHECK-NEW-DAG: call spir_func void @[[RAW_REF2]](i32 addrspace(4)* dereferenceable(4) [[GLOB_CAST2]]) - - bar(*LOC); - // CHECK-DAG: [[LOC_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOC]] - // CHECK-DAG: call spir_func void [[LOC_REF]](i32 addrspace(3)* dereferenceable(4) [[LOC_LOAD]]) - bar2(*LOC); - // CHECK-DAG: [[LOC_LOAD2:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOC]] - // CHECK-DAG: [[LOC_CAST2:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(3)* [[LOC_LOAD2]] to i32 addrspace(4)* - // CHECK-OLD-DAG: call spir_func void @new.[[RAW_REF2]](i32 addrspace(4)* [[LOC_CAST2]]) - // CHECK-NEW-DAG: call spir_func void @[[RAW_REF2]](i32 addrspace(4)* dereferenceable(4) [[LOC_CAST2]]) - - bar(*NoAS); - // CHECK-OLD-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load i32*, i32** [[NoAS]] - // CHECK-OLD-DAG: call spir_func void @[[RAW_REF]](i32* dereferenceable(4) [[NoAS_LOAD]]) - // CHECK-NEW-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)** [[NoAS]] - // CHECK-NEW-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[NoAS_LOAD]]) - bar2(*NoAS); - // CHECK-OLD-DAG: [[NoAS_LOAD2:%[a-zA-Z0-9]+]] = load i32*, i32** [[NoAS]] - // CHECK-OLD-DAG: call spir_func void @[[RAW_REF2]](i32* dereferenceable(4) [[NoAS_LOAD2]]) - // CHECK-NEW-DAG: [[NoAS_LOAD2:%[a-zA-Z0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)** [[NoAS]] - // CHECK-NEW-DAG: call spir_func void @[[RAW_REF2]](i32 addrspace(4)* dereferenceable(4) [[NoAS_LOAD2]]) - - foo(GLOB); - // CHECK-DAG: [[GLOB_LOAD3:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB]] - // CHECK-DAG: [[GLOB_CAST3:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(1)* [[GLOB_LOAD3]] to i32 addrspace(4)* - // CHECK-OLD-DAG: call spir_func void @new.[[RAW_PTR]](i32 addrspace(4)* [[GLOB_CAST3]]) - // CHECK-NEW-DAG: call spir_func void @[[RAW_PTR]](i32 addrspace(4)* [[GLOB_CAST3]]) - foo2(GLOB); - // CHECK-DAG: [[GLOB_LOAD4:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB]] - // CHECK-DAG: [[GLOB_CAST4:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(1)* [[GLOB_LOAD4]] to i32 addrspace(4)* - // CHECK-OLD-DAG: call spir_func void @new.[[RAW_PTR2]](i32 addrspace(4)* [[GLOB_CAST4]]) - // CHECK-NEW-DAG: call spir_func void @[[RAW_PTR2]](i32 addrspace(4)* [[GLOB_CAST4]]) - foo(LOC); - // CHECK-DAG: [[LOC_LOAD3:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOC]] - // CHECK-DAG: call spir_func void [[LOC_PTR]](i32 addrspace(3)* [[LOC_LOAD3]]) - foo2(LOC); - // CHECK-DAG: [[LOC_LOAD4:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOC]] - // CHECK-DAG: [[LOC_CAST4:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(3)* [[LOC_LOAD4]] to i32 addrspace(4)* - // CHECK-OLD-DAG: call spir_func void @new.[[RAW_PTR2]](i32 addrspace(4)* [[LOC_CAST4]]) - // CHECK-NEW-DAG: call spir_func void @[[RAW_PTR2]](i32 addrspace(4)* [[LOC_CAST4]]) - foo(NoAS); - // CHECK-OLD-DAG: [[NoAS_LOAD3:%[a-zA-Z0-9]+]] = load i32*, i32** [[NoAS]] - // CHECK-OLD-DAG: call spir_func void @[[RAW_PTR]](i32* [[NoAS_LOAD3]]) - // CHECK-NEW-DAG: [[NoAS_LOAD3:%[a-zA-Z0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)** [[NoAS]] - // CHECK-NEW-DAG: call spir_func void @[[RAW_PTR]](i32 addrspace(4)* [[NoAS_LOAD3]]) - foo2(NoAS); - // CHECK-OLD-DAG: [[NoAS_LOAD4:%[a-zA-Z0-9]+]] = load i32*, i32** [[NoAS]] - // CHECK-OLD-DAG: call spir_func void @[[RAW_PTR2]](i32* [[NoAS_LOAD4]]) - // CHECK-NEW-DAG: [[NoAS_LOAD4:%[a-zA-Z0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)** [[NoAS]] - // CHECK-NEW-DAG: call spir_func void @[[RAW_PTR2]](i32 addrspace(4)* [[NoAS_LOAD4]]) - - // Ensure that we still get 3 different template instantiations. - tmpl(GLOB); - // CHECK-DAG: [[GLOB_LOAD4:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB]] - // CHECK-DAG: call spir_func void [[GLOB_TMPL:@[a-zA-Z0-9_]+]](i32 addrspace(1)* [[GLOB_LOAD4]]) - tmpl(LOC); - // CHECK-DAG: [[LOC_LOAD5:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOC]] - // CHECK-DAG: call spir_func void [[LOC_TMPL:@[a-zA-Z0-9_]+]](i32 addrspace(3)* [[LOC_LOAD5]]) - tmpl(PRIV); - // CHECK-DAG: [[PRIV_LOAD5:%[a-zA-Z0-9]+]] = load i32*, i32** [[PRIV]] - // CHECK-DAG: call spir_func void [[PRIV_TMPL:@[a-zA-Z0-9_]+]](i32* [[PRIV_LOAD5]]) - tmpl(NoAS); - // CHECK-OLD-DAG: [[NoAS_LOAD5:%[a-zA-Z0-9]+]] = load i32*, i32** [[NoAS]] - // CHECK-OLD-DAG: call spir_func void [[AS0_TMPL:@[a-zA-Z0-9_]+]](i32* [[NoAS_LOAD5]]) - // CHECK-NEW-DAG: [[NoAS_LOAD5:%[a-zA-Z0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)** [[NoAS]] - // CHECK-NEW-DAG: call spir_func void [[GEN_TMPL:@[a-zA-Z0-9_]+]](i32 addrspace(4)* [[NoAS_LOAD5]]) -} - -// CHECK-DAG: define linkonce_odr spir_func void [[GLOB_TMPL]](i32 addrspace(1)* % -// CHECK-DAG: define linkonce_odr spir_func void [[LOC_TMPL]](i32 addrspace(3)* % -// CHECK-OLD-DAG: define linkonce_odr spir_func void [[AS0_TMPL]](i32* % -// CHECK-NEW-DAG: define linkonce_odr spir_func void [[PRIV_TMPL]](i32* % -// CHECK-NEW-DAG: define linkonce_odr spir_func void [[GEN_TMPL]](i32 addrspace(4)* % - -void usages2() { - __attribute__((address_space(0))) int *PRIV_NUM; - // CHECK-DAG: [[PRIV_NUM:%[a-zA-Z0-9_]+]] = alloca i32* - __attribute__((address_space(0))) int *PRIV_NUM2; - // CHECK-DAG: [[PRIV_NUM2:%[a-zA-Z0-9_]+]] = alloca i32* - __attribute__((ocl_private)) int *PRIV; - // CHECK-DAG: [[PRIV:%[a-zA-Z0-9_]+]] = alloca i32* - __attribute__((address_space(1))) int *GLOB_NUM; - // CHECK-DAG: [[GLOB_NUM:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(1)* - __attribute__((ocl_global)) int *GLOB; - // CHECK-DAG: [[GLOB:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(1)* - __attribute__((address_space(2))) int *CONST_NUM; - // CHECK-DAG: [[CONST_NUM:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(2)* - __attribute__((ocl_constant)) int *CONST; - // CHECK-DAG: [[CONST:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(2)* - __attribute__((address_space(3))) int *LOCAL_NUM; - // CHECK-DAG: [[LOCAL_NUM:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(3)* - __attribute__((ocl_local)) int *LOCAL; - // CHECK-DAG: [[LOCAL:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(3)* - - bar(*PRIV_NUM); - // CHECK-DAG: [[PRIV_NUM_LOAD:%[a-zA-Z0-9]+]] = load i32*, i32** [[PRIV_NUM]] - // CHECK-OLD-DAG: call spir_func void @[[RAW_REF]](i32* dereferenceable(4) [[PRIV_NUM_LOAD]]) - // CHECK-NEW-DAG: [[PRIV_NUM_ASCAST:%[a-zA-Z0-9]+]] = addrspacecast i32* [[PRIV_NUM_LOAD]] to i32 addrspace(4)* - // CHECK-NEW-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[PRIV_NUM_ASCAST]]) - bar(*PRIV_NUM2); - // CHECK-DAG: [[PRIV_NUM2_LOAD:%[a-zA-Z0-9]+]] = load i32*, i32** [[PRIV_NUM2]] - // CHECK-OLD-DAG: call spir_func void @[[RAW_REF]](i32* dereferenceable(4) [[PRIV_NUM2_LOAD]]) - // CHECK-NEW-DAG: [[PRIV_NUM2_ASCAST:%[a-zA-Z0-9]+]] = addrspacecast i32* [[PRIV_NUM2_LOAD]] to i32 addrspace(4)* - // CHECK-NEW-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[PRIV_NUM2_ASCAST]]) - bar(*PRIV); - // CHECK-DAG: [[PRIV_LOAD:%[a-zA-Z0-9]+]] = load i32*, i32** [[PRIV]] - // CHECK-OLD-DAG: call spir_func void @[[RAW_REF]](i32* dereferenceable(4) [[PRIV_LOAD]]) - // CHECK-NEW-DAG: [[PRIV_ASCAST:%[a-zA-Z0-9]+]] = addrspacecast i32* [[PRIV_LOAD]] to i32 addrspace(4)* - // CHECK-NEW-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[PRIV_ASCAST]]) - bar(*GLOB_NUM); - // CHECK-DAG: [[GLOB_NUM_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB_NUM]] - // CHECK-DAG: [[GLOB_NUM_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(1)* [[GLOB_NUM_LOAD]] to i32 addrspace(4)* - // CHECK-OLD-DAG: call spir_func void @new.[[RAW_REF]](i32 addrspace(4)* [[GLOB_NUM_CAST]]) - // CHECK-NEW-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[GLOB_NUM_CAST]]) - bar(*GLOB); - // CHECK-DAG: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB]] - // CHECK-DAG: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(1)* [[GLOB_LOAD]] to i32 addrspace(4)* - // CHECK-OLD-DAG: call spir_func void @new.[[RAW_REF]](i32 addrspace(4)* [[GLOB_CAST]]) - // CHECK-NEW-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[GLOB_CAST]]) - bar(*CONST_NUM); - // CHECK-DAG: [[CONST_NUM_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(2)*, i32 addrspace(2)** [[CONST_NUM]] - // CHECK-DAG: [[CONST_NUM_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(2)* [[CONST_NUM_LOAD]] to i32 addrspace(4)* - // CHECK-OLD-DAG: call spir_func void @new.[[RAW_REF]](i32 addrspace(4)* [[CONST_NUM_CAST]]) - // CHECK-NEW-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[CONST_NUM_CAST]]) - bar(*CONST); - // CHECK-DAG: [[CONST_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(2)*, i32 addrspace(2)** [[CONST]] - // CHECK-DAG: [[CONST_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(2)* [[CONST_LOAD]] to i32 addrspace(4)* - // CHECK-OLD-DAG: call spir_func void @new.[[RAW_REF]](i32 addrspace(4)* [[CONST_CAST]]) - // CHECK-NEW-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[CONST_CAST]]) - bar2(*LOCAL_NUM); - // CHECK-DAG: [[LOCAL_NUM_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOCAL_NUM]] - // CHECK-DAG: [[LOCAL_NUM_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(3)* [[LOCAL_NUM_LOAD]] to i32 addrspace(4)* - // CHECK-OLD-DAG: call spir_func void @new.[[RAW_REF2]](i32 addrspace(4)* [[LOCAL_NUM_CAST]]) - // CHECK-NEW-DAG: call spir_func void @[[RAW_REF2]](i32 addrspace(4)* dereferenceable(4) [[LOCAL_NUM_CAST]]) - bar2(*LOCAL); - // CHECK-DAG: [[LOCAL_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOCAL]] - // CHECK-DAG: [[LOCAL_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(3)* [[LOCAL_LOAD]] to i32 addrspace(4)* - // CHECK-OLD-DAG: call spir_func void @new.[[RAW_REF2]](i32 addrspace(4)* [[LOCAL_CAST]]) - // CHECK-NEW-DAG: call spir_func void @[[RAW_REF2]](i32 addrspace(4)* dereferenceable(4) [[LOCAL_CAST]]) -} - -// CHECK-OLD-DAG: define spir_func void @new.[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) -// CHECK-OLD-DAG: define spir_func void @new.[[RAW_REF2]](i32 addrspace(4)* dereferenceable(4) -// CHECK-OLD-DAG: define spir_func void @new.[[RAW_PTR]](i32 addrspace(4)* -// CHECK-OLD-DAG: define spir_func void @new.[[RAW_PTR2]](i32 addrspace(4)* - -template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { - kernelFunc(); -} -int main() { - kernel_single_task([]() { usages();usages2(); }); - return 0; -} - -// TODO: SYCL specific fail - analyze and enable -// XFAIL: windows-msvc diff --git a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp index b7da2c8255848..cfeeec7d692c5 100644 --- a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp +++ b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp @@ -1,5 +1,4 @@ -// RUN: DISABLE_INFER_AS=1 %clang_cc1 -I %S/Inputs -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -S -emit-llvm %s -o - | FileCheck %s --check-prefixes CHECK,CHECK-OLD -// RUN: %clang_cc1 -I %S/Inputs -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -S -emit-llvm %s -o - | FileCheck %s --check-prefixes CHECK,CHECK-NEW +// RUN: %clang_cc1 -I %S/Inputs -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -S -emit-llvm %s -o - | FileCheck %s // This test checks that compiler generates correct kernel wrapper for basic // case. @@ -48,13 +47,10 @@ int main() { // Check accessor __init method call // CHECK-OLD: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor"* [[ACCESSOR]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]]) -// CHECK-NEW: [[ACCESSORCAST:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[ACCESSOR]] to %"class{{.*}}accessor" addrspace(4)* -// CHECK-NEW: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACCESSORCAST]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]]) +// CHECK: [[ACCESSORCAST:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[ACCESSOR]] to %"class{{.*}}accessor" addrspace(4)* +// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACCESSORCAST]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]]) // Check lambda "()" operator call // CHECK-OLD: call spir_func void @{{.*}}(%"class.{{.*}}.anon"* [[ANON]]) -// CHECK-NEW: [[ANONCAST:%[0-9]+]] = addrspacecast %"class{{.*}}anon"* {{.*}} to %"class{{.*}}anon" addrspace(4)* -// CHECK-NEW: call spir_func void @{{.*}}(%"class.{{.*}}.anon" addrspace(4)* [[ANONCAST]]) - -// TODO: SYCL specific fail - analyze and enable -// XFAIL: windows-msvc +// CHECK: [[ANONCAST:%[0-9]+]] = addrspacecast %"class{{.*}}anon"* {{.*}} to %"class{{.*}}anon" addrspace(4)* +// CHECK: call spir_func void @{{.*}}(%"class.{{.*}}.anon" addrspace(4)* [[ANONCAST]]) diff --git a/clang/test/CodeGenSYCL/device-functions.cpp b/clang/test/CodeGenSYCL/device-functions.cpp index 474aa53268c3a..90bb7c2565253 100644 --- a/clang/test/CodeGenSYCL/device-functions.cpp +++ b/clang/test/CodeGenSYCL/device-functions.cpp @@ -1,5 +1,4 @@ -// RUN: DISABLE_INFER_AS=1 %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -S -emit-llvm -x c++ %s -o - | FileCheck %s --check-prefixes CHECK,CHECK-OLD -// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -S -emit-llvm -x c++ %s -o - | FileCheck %s --check-prefixes CHECK,CHECK-NEW +// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -S -emit-llvm -x c++ %s -o - | FileCheck %s template T bar(T arg); @@ -23,10 +22,6 @@ int main() { return 0; } // CHECK: define spir_kernel void @_ZTSZ4mainE11fake_kernel() -// CHECK-OLD: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}.anon"* %this) -// CHECK-NEW: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}.anon" addrspace(4)* %this) +// CHECK: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}.anon" addrspace(4)* %this) // CHECK: define spir_func void @_Z3foov() // CHECK: define linkonce_odr spir_func i32 @_Z3barIiET_S0_(i32 %arg) - -// TODO: SYCL specific fail - analyze and enable -// XFAIL: windows-msvc diff --git a/clang/test/CodeGenSYCL/sampler.cpp b/clang/test/CodeGenSYCL/sampler.cpp index 5e7ba43494ad2..9ceef6f27795b 100644 --- a/clang/test/CodeGenSYCL/sampler.cpp +++ b/clang/test/CodeGenSYCL/sampler.cpp @@ -1,5 +1,4 @@ -// RUN: DISABLE_INFER_AS=1 %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -I %S/Inputs -fsycl-is-device -disable-llvm-passes -emit-llvm -x c++ %s -o - | FileCheck --enable-var-scope %s --check-prefixes CHECK,CHECK-OLD -// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -I %S/Inputs -fsycl-is-device -disable-llvm-passes -emit-llvm -x c++ %s -o - | FileCheck --enable-var-scope %s --check-prefixes CHECK,CHECK-NEW +// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -I %S/Inputs -fsycl-is-device -disable-llvm-passes -emit-llvm -x c++ %s -o - | FileCheck --enable-var-scope %s // CHECK: define spir_kernel void @{{[a-zA-Z0-9_]+}}(%opencl.sampler_t addrspace(2)* [[SAMPLER_ARG:%[a-zA-Z0-9_]+]]) // CHECK-NEXT: entry: // CHECK-NEXT: [[SAMPLER_ARG]].addr = alloca %opencl.sampler_t addrspace(2)*, align 8 @@ -9,9 +8,8 @@ // CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* [[BITCAST]]) #4 // CHECK-NEXT: [[GEP:%[0-9]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[ANON]], i32 0, i32 0 // CHECK-NEXT: [[LOAD_SAMPLER_ARG:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG]].addr, align 8 -// CHECK-OLD-NEXT: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.{{.*}}.cl::sycl::sampler"* [[GEP]], %opencl.sampler_t addrspace(2)* [[LOAD_SAMPLER_ARG]]) -// CHECK-NEW-NEXT: [[GEPCAST:%[0-9]+]] = addrspacecast %"class{{.*}}.cl::sycl::sampler"* [[GEP]] to %"class{{.*}}.cl::sycl::sampler" addrspace(4)* -// CHECK-NEW-NEXT: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.{{.*}}.cl::sycl::sampler" addrspace(4)* [[GEPCAST]], %opencl.sampler_t addrspace(2)* [[LOAD_SAMPLER_ARG]]) +// CHECK-NEXT: [[GEPCAST:%[0-9]+]] = addrspacecast %"class{{.*}}.cl::sycl::sampler"* [[GEP]] to %"class{{.*}}.cl::sycl::sampler" addrspace(4)* +// CHECK-NEXT: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.{{.*}}.cl::sycl::sampler" addrspace(4)* [[GEPCAST]], %opencl.sampler_t addrspace(2)* [[LOAD_SAMPLER_ARG]]) // #include "sycl.hpp" @@ -28,6 +26,3 @@ int main() { return 0; } - -// TODO: SYCL specific fail - analyze and enable -// XFAIL: windows-msvc diff --git a/clang/test/CodeGenSYCL/spir-calling-conv.cpp b/clang/test/CodeGenSYCL/spir-calling-conv.cpp index 46decdc1a4258..0a37ad888bae1 100644 --- a/clang/test/CodeGenSYCL/spir-calling-conv.cpp +++ b/clang/test/CodeGenSYCL/spir-calling-conv.cpp @@ -1,5 +1,4 @@ -// RUN: DISABLE_INFER_AS=1 %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -S -emit-llvm -x c++ %s -o - | FileCheck %s --check-prefixes CHECK,CHECK-OLD -// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -S -emit-llvm -x c++ %s -o - | FileCheck %s --check-prefixes CHECK,CHECK-NEW +// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -S -emit-llvm -x c++ %s -o - | FileCheck %s template __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { @@ -10,15 +9,10 @@ int main() { // CHECK: define spir_kernel void @_ZTSZ4mainE15kernel_function() - // CHECK-OLD: call spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}.anon"* %0) - // CHECK-NEW: call spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}.anon" addrspace(4)* %2) + // CHECK: call spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}.anon" addrspace(4)* %2) - // CHECK-OLD: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}anon"* %this) - // CHECK-NEW: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}anon" addrspace(4)* %this) + // CHECK: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}anon" addrspace(4)* %this) kernel_single_task([]() {}); return 0; } - -// TODO: SYCL specific fail - analyze and enable -// XFAIL: windows-msvc diff --git a/clang/test/SemaSYCL/spir-enum.cpp b/clang/test/SemaSYCL/spir-enum.cpp index a8fc23f2d2f31..d3eb0c8823b80 100644 --- a/clang/test/SemaSYCL/spir-enum.cpp +++ b/clang/test/SemaSYCL/spir-enum.cpp @@ -1,5 +1,4 @@ -// RUN: DISABLE_INFER_AS=1 %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-optzns -disable-llvm-passes -S -emit-llvm -x c++ %s -o - | FileCheck %s --check-prefixes CHECK,CHECK-OLD -// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-optzns -disable-llvm-passes -S -emit-llvm -x c++ %s -o - | FileCheck %s --check-prefixes CHECK,CHECK-NEW +// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-optzns -disable-llvm-passes -S -emit-llvm -x c++ %s -o - | FileCheck %s template __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { @@ -24,13 +23,9 @@ int main() { // CHECK: define spir_kernel void @_ZTSZ4test9enum_typeE15kernel_function(i32 %_arg_) // CHECK: getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* - // CHECK-OLD: call spir_func void @"_ZZ4test9enum_typeENK3$_0clEv"(%"class.{{.*}}.anon"* %0) - // CHECK-NEW: call spir_func void @"_ZZ4test9enum_typeENK3$_0clEv"(%"class.{{.*}}.anon" addrspace(4)* %4) + // CHECK: call spir_func void @"_ZZ4test9enum_typeENK3$_0clEv"(%"class.{{.*}}.anon" addrspace(4)* %4) test( enum_type::B ); return 0; } - -// TODO: SYCL specific fail - analyze and enable -// XFAIL: windows-msvc diff --git a/llvm/include/llvm/InitializePasses.h b/llvm/include/llvm/InitializePasses.h index ecca66b9af243..dbae32e843936 100644 --- a/llvm/include/llvm/InitializePasses.h +++ b/llvm/include/llvm/InitializePasses.h @@ -422,7 +422,6 @@ void initializeWinEHPreparePass(PassRegistry&); void initializeWriteBitcodePassPass(PassRegistry&); void initializeWriteThinLTOBitcodePass(PassRegistry&); void initializeXRayInstrumentationPass(PassRegistry&); -void initializeASFixerPass(PassRegistry&); } // end namespace llvm diff --git a/llvm/include/llvm/SYCL/ASFixer.h b/llvm/include/llvm/SYCL/ASFixer.h deleted file mode 100644 index 9dcf322df581a..0000000000000 --- a/llvm/include/llvm/SYCL/ASFixer.h +++ /dev/null @@ -1,25 +0,0 @@ -//===- ASFixer.h - SYCL address spaces fixer pass -------------------------===// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// -// -// SYCL address spaces fixer pass -// -//===----------------------------------------------------------------------===// - -#ifndef LLVM_SYCL_ASFIXER_H -#define LLVM_SYCL_ASFIXER_H - -#include "llvm/IR/Module.h" -#include "llvm/Pass.h" - -namespace llvm { - -ModulePass *createASFixerPass(); - -} - -#endif diff --git a/llvm/lib/CMakeLists.txt b/llvm/lib/CMakeLists.txt index 30beb2012774c..5a41d65b0cf37 100644 --- a/llvm/lib/CMakeLists.txt +++ b/llvm/lib/CMakeLists.txt @@ -32,4 +32,3 @@ if (LLVM_INCLUDE_TESTS) add_subdirectory(Testing) endif() add_subdirectory(WindowsManifest) -add_subdirectory(SYCL) diff --git a/llvm/lib/LLVMBuild.txt b/llvm/lib/LLVMBuild.txt index 4ebb1033acde1..4c039176267c3 100644 --- a/llvm/lib/LLVMBuild.txt +++ b/llvm/lib/LLVMBuild.txt @@ -39,7 +39,6 @@ subdirectories = Remarks Passes ProfileData - SYCL Support TableGen TextAPI diff --git a/llvm/lib/SYCL/ASFixer.cpp b/llvm/lib/SYCL/ASFixer.cpp deleted file mode 100644 index 8f5f36470bf7a..0000000000000 --- a/llvm/lib/SYCL/ASFixer.cpp +++ /dev/null @@ -1,602 +0,0 @@ -//===- ASFixer.cpp - Address spaces fixer pass implementation -------------===// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception// -//===----------------------------------------------------------------------===// -// -// This implements the address spaces fixer pass. -// -//===----------------------------------------------------------------------===// - -#include "llvm/ADT/SetVector.h" -#include "llvm/ADT/SmallSet.h" -#include "llvm/ADT/StringRef.h" -#include "llvm/IR/ConstantFolder.h" -#include "llvm/IR/Function.h" -#include "llvm/IR/Instructions.h" -#include "llvm/IR/Module.h" -#include "llvm/IR/Operator.h" -#include "llvm/Pass.h" -#include "llvm/Support/raw_ostream.h" -#include "llvm/SYCL/ASFixer.h" -#include "llvm/Transforms/Utils/Cloning.h" - -#include -#include -#include - -using namespace llvm; - -namespace { - -typedef std::pair ValueValuePair; -typedef std::stack> WorkListType; -typedef DenseMap ValueToValueMap; -typedef DenseMap TypeToTypeMap; - -typedef SmallVector, 32> OperandsVector; - -// Contains Instruction stub and it's future operands -typedef DenseMap UserToOperandsMap; - -typedef DenseMap FunctionToNewTypeMap; - -enum SPIRAddressSpace { - SPIRAS_Private, - SPIRAS_Global, - SPIRAS_Constant, - SPIRAS_Local, - SPIRAS_Generic, -}; - -static CallInst *createCallInstStub(FunctionType *FTy, Function *Func) { - SmallVector Args; - for (auto &Arg : FTy->params()) { - Args.push_back(UndefValue::get(Arg)); - } - CallInst *Result = CallInst::Create(Func, Args, ""); - Result->setCallingConv(Func->getCallingConv()); - return Result; -} - -static Function *createNewFunction(Function *ExistingF, FunctionType *NewFTy, - ValueToValueMap &VMap, - UserToOperandsMap &UOpMap) { - - Function *Replacement = - Function::Create(NewFTy, ExistingF->getLinkage(), - "new." + ExistingF->getName(), ExistingF->getParent()); - Replacement->setCallingConv(ExistingF->getCallingConv()); - ValueToValueMapTy CloneMap; - auto ExistFArgIt = ExistingF->arg_begin(); - auto RArgIt = Replacement->arg_begin(); - auto RArgItE = Replacement->arg_end(); - for (; RArgIt != RArgItE; ++RArgIt, ++ExistFArgIt) { - CloneMap[&*ExistFArgIt] = &*RArgIt; - for (auto &ArgUse : ExistFArgIt->uses()) { - Value *Usr = ArgUse.getUser(); - Value *NewVal = VMap[Usr]; - if (NewVal) { - UOpMap[cast(NewVal)].push_back( - std::make_pair(ArgUse.getOperandNo(), RArgIt)); - } - } - } - - SmallVector Returns; - CloneFunctionInto(Replacement, ExistingF, CloneMap, true, Returns, ""); - assert(Replacement && "CloneFunctionInto failed"); - - // Remap new instructions to clones. - SmallVector NewInsts; - for (auto It : CloneMap) { - if (isa(It.first)) { - Instruction *OldInstClone = cast(It.second); - Value *NewVal = VMap[const_cast(It.first)]; - if (NewVal) { - VMap[OldInstClone] = NewVal; - auto NewInst = cast(NewVal); - if (!NewInst->getParent()) { - NewInst->insertBefore(OldInstClone); - } - NewInsts.push_back(NewInst); - } - VMap.erase(const_cast(It.first)); - } - } - // If new instruction uses some old instructions - - // it should use old instructions clones - for (auto NI : NewInsts) { - for (auto &OpIt : NI->operands()) { - Value *OpVal = OpIt; - if (CloneMap[OpVal]) - OpIt.set(CloneMap[OpVal]); - } - } - return Replacement; -} - -static Value *createTypeUserStub(Value *OldValUser, Value *NewVal) { - Type *NewTy = NewVal->getType(); - auto *UndefType = UndefValue::get(NewTy); - - auto *OldInst = cast(OldValUser); - - if (auto *Alloca = dyn_cast(OldInst)) { - return new AllocaInst(NewTy, Alloca->getType()->getAddressSpace(), - Alloca->getArraySize(), "new." + Alloca->getName()); - } - - if (auto *Store = dyn_cast(OldInst)) { - Value *DestPtr = UndefValue::get(PointerType::getUnqual(NewTy)); - auto SI = new StoreInst(UndefType, DestPtr, Store->isVolatile(), - MaybeAlign(Store->getAlignment()), - Store->getOrdering(), Store->getSyncScopeID()); - return SI; - } - - if (auto *GEP = dyn_cast(OldInst)) { - SmallVector IdxList(GEP->idx_begin(), GEP->idx_end()); - - return GetElementPtrInst::Create(NewTy->getPointerElementType(), UndefType, - IdxList, OldInst->getName()); - } - - if (auto *Load = dyn_cast(OldInst)) { - return new LoadInst(NewTy->getPointerElementType(), UndefType, - Load->getName(), Load->isVolatile()); - } - if (auto *Select = dyn_cast(OldInst)) { - return SelectInst::Create(Select->getCondition(), UndefType, UndefType, - OldInst->getName()); - } - - if (auto *Phi = dyn_cast(OldInst)) { - auto *NewPhi = - PHINode::Create(NewTy, Phi->getNumIncomingValues(), OldInst->getName()); - for (auto *BB : Phi->blocks()) { - NewPhi->addIncoming(UndefType, BB); - } - return NewPhi; - } - - if (auto *Ret = dyn_cast(OldInst)) { - return ReturnInst::Create(Ret->getContext(), UndefType); - } - - if (auto *BCast = dyn_cast(OldInst)) { - unsigned AS = NewTy->getPointerAddressSpace(); - return new BitCastInst( - UndefType, - BCast->getDestTy()->getPointerElementType()->getPointerTo(AS), - OldInst->getName()); - } - - if (auto *Shuffle = dyn_cast(OldInst)) - return Shuffle->clone(); - - if (auto *BinOp = dyn_cast(OldInst)) - return BinOp->clone(); - - llvm_unreachable("Unsupported instruction."); -} - -static Value *getAllocaOrArgValue(Function *F, const unsigned ArgNo) { - assert(ArgNo < F->arg_size() && "Invalid ArgNo"); - auto Arg = F->arg_begin() + ArgNo; - for (auto *ArgUser : Arg->users()) { - if (auto *Store = dyn_cast(ArgUser)) { - Value *POperand = Store->getPointerOperand(); - assert(isa(POperand) && "Expected alloca for argument"); - return POperand; - } - } - return Arg; -} - -static bool checkFunctionArgument(FunctionType *FTy, Type *Ty, - const unsigned ArgNo) { - assert(ArgNo < FTy->getNumParams() && "Invalid ArgNo"); - return (FTy->getParamType(ArgNo) == Ty); -} - -static bool checkFunctionRetType(FunctionType *FTy, Type *Ty) { - return (FTy->getReturnType() == Ty); -} - -static FunctionType *createNewFuncReplacementType(FunctionType *FTy, Type *Ty, - const unsigned ArgNo) { - SmallVector Args; - for (auto P : FTy->params()) { - Args.push_back(P); - } - Args[ArgNo] = Ty; - return FunctionType::get(FTy->getReturnType(), Args, FTy->isVarArg()); -} - -static FunctionType *createNewFuncReplacementType(FunctionType *FTy, Type *Ty) { - return FunctionType::get(Ty, FTy->params(), FTy->isVarArg()); -} - -static bool valueIsReplacement(ValueToValueMap &VMap, Value *V) { - for (auto It : VMap) { - if (It.second == V) - return true; - } - return false; -} - -static void collectTypeReplacementData(Type *OldTy, Type *NewTy, - ValueToValueMap &VMap, - WorkListType &WorkList, - UserToOperandsMap &UOpMap, - FunctionToNewTypeMap &FTyMap) { - - while (!WorkList.empty()) { - Value *OldVal = WorkList.top().first; - Value *NewVal = WorkList.top().second; - WorkList.pop(); - - for (const auto &U : OldVal->uses()) { - auto OpNo = U.getOperandNo(); - Value *OldValUser = U.getUser(); - - if (CallInst *Call = dyn_cast(OldValUser)) { - auto F = Call->getCalledFunction(); - assert(F && "Indirect function call?"); - FunctionType *&FuncReplacementType = FTyMap[Call->getCalledFunction()]; - if (!FuncReplacementType) - FuncReplacementType = F->getFunctionType(); - if (!checkFunctionArgument(FuncReplacementType, NewVal->getType(), - OpNo)) { - FuncReplacementType = createNewFuncReplacementType( - FuncReplacementType, NewVal->getType(), OpNo); - } - - Function *FuncStub = Function::Create( - FuncReplacementType, F->getLinkage(), "new." + F->getName()); - Value *&CallInstStub = VMap[Call]; - if (!CallInstStub) { - CallInstStub = createCallInstStub(FuncReplacementType, FuncStub); - } else { - cast(CallInstStub) - ->setCalledFunction(FuncReplacementType, FuncStub); - } - auto Arg = getAllocaOrArgValue(F, OpNo); - if (isa(Arg)) { - auto *&NewValUser = VMap[Arg]; - if (!NewValUser) { - NewValUser = createTypeUserStub(Arg, NewVal); - WorkList.push(std::make_pair(Arg, NewValUser)); - } - } else { - WorkList.push( - std::make_pair(Arg, getAllocaOrArgValue(FuncStub, OpNo))); - } - UOpMap[cast(CallInstStub)].push_back( - std::make_pair(OpNo, NewVal)); - // TODO: Case when call used next - continue; - } - - if (auto *Ret = dyn_cast(OldValUser)) { - auto F = Ret->getFunction(); - FunctionType *&FuncReplacementType = FTyMap[F]; - if (!FuncReplacementType) - FuncReplacementType = F->getFunctionType(); - if (!checkFunctionRetType(FuncReplacementType, NewVal->getType())) { - FuncReplacementType = createNewFuncReplacementType( - FuncReplacementType, NewVal->getType()); - } - } - - // TODO: Don't handle this case separately - if (auto *Store = dyn_cast(OldValUser)) { - if (OpNo == 1) { - - auto *UndefType = - UndefValue::get(NewVal->getType()->getPointerElementType()); - auto *&NewValUser = VMap[OldValUser]; - if (!NewValUser) { - NewValUser = new StoreInst( - UndefType, UndefValue::get(NewVal->getType()), - Store->isVolatile(), MaybeAlign(Store->getAlignment()), - Store->getOrdering(), Store->getSyncScopeID()); - WorkList.push(std::make_pair(OldValUser, NewValUser)); - } - UOpMap[cast(NewValUser)].push_back( - std::make_pair(OpNo, NewVal)); - continue; - } - } - - // We are cloning shuffle vectors and binary operators, - // so these new instructions contains references to - // old instructions and we can see new instructions as old instructions - // here, so we should skip these instructions - // to don't create clone for clone. - if (valueIsReplacement(VMap, OldValUser)) - continue; - - auto *&NewValUser = VMap[OldValUser]; - if (!NewValUser) { - NewValUser = createTypeUserStub(OldValUser, NewVal); - WorkList.push(std::make_pair(OldValUser, NewValUser)); - } - UOpMap[cast(NewValUser)].push_back(std::make_pair(OpNo, NewVal)); - } - } -} - -static void traceAddressSpace(AddrSpaceCastInst *AS, - ValueToValueMap &VMap, UserToOperandsMap &UOpMap, - FunctionToNewTypeMap &FTyMap) { - - Type *OldTy = AS->getType(); - Type *NewTy = - AS->getSrcTy()->getPointerElementType()->getPointerTo(SPIRAS_Generic); - AddrSpaceCastInst *NewAS = - new AddrSpaceCastInst(AS->getPointerOperand(), NewTy, "", AS); - WorkListType WorkList; - for (const auto &U : AS->uses()) { - Value *User = U.getUser(); - if (auto Call = dyn_cast(User)) { - WorkList.push(std::make_pair(AS, NewAS)); - auto F = Call->getCalledFunction(); - assert(F && "No function info."); - for (auto &Arg : F->args()) { - if (Arg.getType() == OldTy) { - auto ActArg = Call->getArgOperand(Arg.getArgNo()); - assert(ActArg && "No argument info."); - if (!isa(ActArg)) { - AddrSpaceCastInst *AddAS = - new AddrSpaceCastInst(ActArg, NewTy, "", Call); - Call->setArgOperand(Arg.getArgNo(), AddAS); - WorkList.push(std::make_pair(AddAS, AddAS)); - } - } - } - } - } - collectTypeReplacementData(OldTy, NewTy, VMap, WorkList, UOpMap, FTyMap); -} - -static void doReplace(ValueToValueMap &VMap, UserToOperandsMap &UOpMap, - FunctionToNewTypeMap &FTyMap) { - - SmallVector, 32> Calls; - for (auto It : VMap) { - if (It.second) { - if (auto Call = dyn_cast(It.first)) { - auto NewCall = cast(It.second); - Calls.push_back(std::make_pair(Call, NewCall)); - } - } - } - - // Create all functions - ValueToValueMap FMap; - for (auto It : Calls) { - auto Call = It.first; - auto NewCall = It.second; - Function *ExistingF = Call->getCalledFunction(); - FunctionType *NewFTy = FTyMap[ExistingF]; - Value *&NewFunc = FMap[ExistingF]; - if (NewFTy && NewFTy != ExistingF->getFunctionType()) { - if (!NewFunc) { - assert(NewFTy && "Forgot function?"); - NewFunc = createNewFunction(ExistingF, NewFTy, VMap, UOpMap); - } - } else { - NewFunc = ExistingF; - } - NewCall->setCalledFunction(cast(NewFunc)->getFunctionType(), - NewFunc); - NewCall->setCallingConv(cast(NewFunc)->getCallingConv()); - NewCall->setDebugLoc(Call->getDebugLoc()); - } - - // Set all operands - for (auto it : UOpMap) { - auto Operands = it.second; - User *Usr = it.first; - for (auto OpIt : Operands) { - Usr->setOperand(OpIt.first, OpIt.second); - } - } - - // New instructions can use some old instructions, so - // we need to set correspondig operands - for (auto It : VMap) { - if (It.second) { - if (auto Usr = dyn_cast(It.first)) { - auto NewUsr = cast(It.second); - for (auto &Op : NewUsr->operands()) { - if (isa(Op.get())) { - NewUsr->setOperand(Op.getOperandNo(), - Usr->getOperand(Op.getOperandNo())); - } - } - } - } - } - - // Replace old instructions - for (auto it : VMap) { - if (it.second) { - if (auto NewInst = dyn_cast(it.second)) { - auto OldInst = cast(it.first); - if (!NewInst->getParent()) { - NewInst->insertBefore(OldInst); - } - OldInst->mutateType(NewInst->getType()); - OldInst->replaceAllUsesWith(NewInst); - if (OldInst->use_empty()) { - OldInst->eraseFromParent(); - } - } - } - } -} - -static Type *createNewStructType(Type *NewElTy, StructType *OldTy) { - SmallVector Elements; - for (auto T : OldTy->elements()) { - if (T->isPointerTy() && - (T->getPointerElementType() == NewElTy->getPointerElementType()) && - (T->getPointerAddressSpace() != NewElTy->getPointerAddressSpace())) { - Elements.push_back(NewElTy); - } else { - Elements.push_back(T); - } - } - auto NewStruct = - StructType::create(OldTy->getContext(), Elements, - "new." + std::string(OldTy->getStructName())); - return NewStruct; -} - -static AllocaInst *createAllocaReplacement(AllocaInst *OldAlloca, Type *NewTy) { - auto NewAlloca = new AllocaInst( - NewTy, OldAlloca->getType()->getAddressSpace(), OldAlloca->getArraySize(), - "new." + OldAlloca->getName(), OldAlloca); - NewAlloca->setAlignment(MaybeAlign(OldAlloca->getAlignment())); - return NewAlloca; -} - -static bool needToReplaceAlloca(AllocaInst *Alloca, - ValueToValueMap &VMap, - UserToOperandsMap &UOpMap, - TypeToTypeMap &TMap) { - SmallSet Seen; - std::stack WorkList; - Type *AllocType = Alloca->getAllocatedType(); - WorkList.push(Alloca); - while (!WorkList.empty()) { - Value *Val = WorkList.top(); - WorkList.pop(); - for (const auto &U : Val->uses()) { - Value *Usr = U.getUser(); - Value *NextUsr = nullptr; - if (CallInst *Call = dyn_cast(Usr)) { - if (VMap[Call]) { - auto F = Call->getCalledFunction(); - assert(F && "No function info."); - NextUsr = - getAllocaOrArgValue(F, U.getOperandNo()); - } - } else { - NextUsr = Usr; - } - if (NextUsr && !Seen.count(NextUsr)) { - WorkList.push(NextUsr); - Seen.insert(NextUsr); - } - // TODO: Try only with stores, maybe add more cases later - auto *&NewVal = VMap[Usr]; - auto Store = dyn_cast(Usr); - if (Store && NewVal) { - auto NewStore = cast(NewVal); - auto StoreValType = NewStore->getValueOperand()->getType(); - auto Operands = UOpMap[NewStore]; - if (Operands.size() == 1 && Operands[0].first == 0) { - if (StoreValType->isPointerTy()) { - auto *&NewStructTy = TMap[AllocType]; - if (!NewStructTy) { - NewStructTy = createNewStructType(StoreValType, - cast(AllocType)); - } - return true; - } - } - } - } - } - return false; -} - -static bool structContainsPointers(StructType *Struct) { - // TODO: more general case, for example when struct - // contains struct which contains pointers. - for (auto *T : Struct->elements()) { - if (T->isPointerTy()) - return true; - } - return false; -} - -struct ASFixer : public ModulePass { - static char ID; - ASFixer() : ModulePass(ID) {} - bool runOnModule(Module &M) override { - bool Changed = false; - ValueToValueMap VMap; - UserToOperandsMap UOpMap; - FunctionToNewTypeMap FTyMap; - // We find and replace all address space casts to private - // address space - for (auto &F : M.functions()) { - for (auto &BB : F) { - for (auto &I : BB) { - auto AS = dyn_cast(&I); - if (AS && AS->getDestAddressSpace() == SPIRAS_Private && - AS->getSrcAddressSpace() != SPIRAS_Generic) { - traceAddressSpace(AS, VMap, UOpMap, FTyMap); - Changed = true; - } - } - } - } - // Pointer with changed address space can be stored - // into structure so we need to check structures with - // pointers and replace it if needed. - // As described in SYCL spec structures with pointers - // can't be passed as kernel argument so we check - // allocas of structures with pointers. - std::vector> BadAllocas; - TypeToTypeMap TMap; - if (Changed) { - for (auto &F : M.functions()) { - for (auto &BB : F) { - for (auto &I : BB) { - if (auto Alloca = dyn_cast(&I)) { - Type *AllocType = Alloca->getAllocatedType(); - if (auto StructTy = dyn_cast(AllocType)) { - if (structContainsPointers(StructTy)) { - if (needToReplaceAlloca(Alloca, VMap, UOpMap, TMap)) { - auto NewStructTy = TMap[StructTy]; - AllocaInst *AllocaReplacement = - createAllocaReplacement(Alloca, NewStructTy); - BadAllocas.push_back( - std::make_pair(Alloca, AllocaReplacement)); - } - } - } - } - } - } - } - } - for (auto It : BadAllocas) { - WorkListType W; - W.push(It); - collectTypeReplacementData(It.first->getType(), It.second->getType(), - VMap, W, UOpMap, FTyMap); - } - doReplace(VMap, UOpMap, FTyMap); - return Changed; - } - - virtual llvm::StringRef getPassName() const { return "ASFixer"; } -}; -} // namespace - -namespace llvm { -void initializeASFixerPass(PassRegistry &Registry); -} - -INITIALIZE_PASS(ASFixer, "asfix", "Fix SYCL address spaces", false, false) -ModulePass *llvm::createASFixerPass() { return new ASFixer(); } - -char ASFixer::ID = 0; diff --git a/llvm/lib/SYCL/CMakeLists.txt b/llvm/lib/SYCL/CMakeLists.txt deleted file mode 100644 index 6ecc4ba5c32d2..0000000000000 --- a/llvm/lib/SYCL/CMakeLists.txt +++ /dev/null @@ -1,9 +0,0 @@ -add_llvm_library(LLVMASFixer - ASFixer.cpp - - ADDITIONAL_HEADER_DIRS - ${LLVM_MAIN_INCLUDE_DIR}/llvm/SYCL - - DEPENDS - intrinsics_gen -) diff --git a/llvm/lib/SYCL/LLVMBuild.txt b/llvm/lib/SYCL/LLVMBuild.txt deleted file mode 100644 index 28dfad9a30cd2..0000000000000 --- a/llvm/lib/SYCL/LLVMBuild.txt +++ /dev/null @@ -1,21 +0,0 @@ -;===- ./lib/SYCL/LLVMBuild.txt ------------------------------*- Conf -*-----===; -; -; Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -; See https://llvm.org/LICENSE.txt for license information. -; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -; -;===------------------------------------------------------------------------===; -; -; This is an LLVMBuild description file for the components in this subdirectory. -; -; For more information on the LLVMBuild system, please see: -; -; http://llvm.org/docs/LLVMBuild.html -; -;===------------------------------------------------------------------------===; - -[component_0] -type = Library -name = ASFixer -parent = Libraries -required_libraries = Passes Core Support TransformUtils diff --git a/llvm/test/SYCL/ASFixer/address-space-parameter-conversions.ll b/llvm/test/SYCL/ASFixer/address-space-parameter-conversions.ll deleted file mode 100644 index f3d8032b04522..0000000000000 --- a/llvm/test/SYCL/ASFixer/address-space-parameter-conversions.ll +++ /dev/null @@ -1,67 +0,0 @@ -; RUN: opt -asfix %s -S -o - | FileCheck %s -; ModuleID = 'address-space-parameter-conversions.cpp' -source_filename = "address-space-parameter-conversions.cpp" -target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" -target triple = "spir64-unknown-linux-sycldevice" - -; Function Attrs: noinline nounwind optnone -define dso_local spir_func void @_Z3fooPi(i32* %Data) #1 { -entry: - %Data.addr = alloca i32*, align 8 - store i32* %Data, i32** %Data.addr, align 8 - ret void -} - -; Function Attrs: noinline nounwind optnone -define dso_local spir_func void @_Z3fooPU3AS3i(i32 addrspace(3)* %Data) #1 { -entry: - %Data.addr = alloca i32 addrspace(3)*, align 8 - store i32 addrspace(3)* %Data, i32 addrspace(3)** %Data.addr, align 8 - ret void -} - -; Function Attrs: noinline nounwind optnone -define dso_local spir_func void @_Z6usagesv() #1 { -entry: - %GLOB = alloca i32 addrspace(1)*, align 8 - %LOC = alloca i32 addrspace(3)*, align 8 - %NoAS = alloca i32*, align 8 -; CHECK: %[[GLOB:.*]] = alloca i32 addrspace(1)*, align 8 -; CHECK: %[[LOC:.*]] = alloca i32 addrspace(3)*, align 8 -; CHECK: %[[NoAS:.*]] = alloca i32*, align 8 - %0 = load i32 addrspace(1)*, i32 addrspace(1)** %GLOB, align 8 - %1 = addrspacecast i32 addrspace(1)* %0 to i32* -; CHECK: %[[GLOB_LOAD:.*]] = load i32 addrspace(1)*, i32 addrspace(1)** %[[GLOB]], align 8 -; CHECK: %[[NEW_CAST:.*]] = addrspacecast i32 addrspace(1)* %[[GLOB_LOAD]] to i32 addrspace(4)* -; CHECK: %[[OLD_CAST:.*]] = addrspacecast i32 addrspace(1)* %[[GLOB_LOAD]] to i32* -; CHECK: call spir_func void @new.[[FOO:.*]](i32 addrspace(4)* %[[NEW_CAST]]) -; CHECK-NOT: call spir_func void @[[FOO]](i32* %[[OLD_CAST]]) - call spir_func void @_Z3fooPi(i32* %1) - %2 = load i32 addrspace(3)*, i32 addrspace(3)** %LOC, align 8 -; CHECK: %[[LOC_LOAD:.*]] = load i32 addrspace(3)*, i32 addrspace(3)** %[[LOC]], align 8 -; CHECK: call spir_func void @[[BAR:.*]](i32 addrspace(3)* %[[LOC_LOAD]]) - call spir_func void @_Z3fooPU3AS3i(i32 addrspace(3)* %2) - %3 = load i32*, i32** %NoAS, align 8 -; CHECK: %[[NoAS_LOAD:.*]] = load i32*, i32** %[[NoAS]], align 8 -; CHECK: call spir_func void @[[FOO]](i32* %[[NoAS_LOAD]]) - call spir_func void @_Z3fooPi(i32* %3) - ret void -} - -; CHECK: define dso_local spir_func void @new.[[FOO]](i32 addrspace(4)* -; CHECK: %[[PAR_ALLOC:.*]] = alloca i32 addrspace(4)* -; CHECK: store i32 addrspace(4)* %{{.*}}, i32 addrspace(4)** %[[PAR_ALLOC]], align 8 - -attributes #0 = { noinline optnone "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="true" "no-frame-pointer-elim-non-leaf" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" } -attributes #1 = { noinline nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="true" "no-frame-pointer-elim-non-leaf" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } - -!llvm.module.flags = !{!0} -!opencl.spir.version = !{!1} -!spirv.Source = !{!2} -!llvm.ident = !{!3} - -!0 = !{i32 1, !"wchar_size", i32 4} -!1 = !{i32 1, i32 2} -!2 = !{i32 4, i32 100000} -!3 = !{!"clang version 8.0.0"} -!4 = !{} diff --git a/llvm/test/SYCL/ASFixer/binary-operators.ll b/llvm/test/SYCL/ASFixer/binary-operators.ll deleted file mode 100644 index 0598f9b809570..0000000000000 --- a/llvm/test/SYCL/ASFixer/binary-operators.ll +++ /dev/null @@ -1,120 +0,0 @@ -; RUN: opt -asfix %s -S -o %t.out -; FileCheck %s --input-file %t.out -; FileCheck %s --input-file %t.out -check-prefix=CHECK-MUL -; FileCheck %s --input-file %t.out -check-prefix=CHECK-ADD -; ModuleID = 'bugpoint-reduced-simplified.ll' -source_filename = "scan.cpp" -target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" -target triple = "spir64-unknown-linux-sycldevice" - -%"struct.std::plus" = type { i8 } -%"struct.std::multiplies" = type { i8 } - -$_ZZZ8par_scanIiSt4plusIiEEvRN2cl4sycl6bufferIT_Li1ESaIS5_EEERNS3_5queueEENKUlRNS3_7handlerEE_clESC_ENKUlNS3_7nd_itemILi1EEEE_clESF_ = comdat any - -$_ZZZ8par_scanIlSt10multipliesIlEEvRN2cl4sycl6bufferIT_Li1ESaIS5_EEERNS3_5queueEENKUlRNS3_7handlerEE_clESC_ENKUlNS3_7nd_itemILi1EEEE_clESF_ = comdat any - -$_ZNKSt4plusIiEclERKiS2_ = comdat any - -$_ZNKSt10multipliesIlEclERKlS2_ = comdat any - -; Function Attrs: noinline -define linkonce_odr dso_local spir_func void @_ZZZ8par_scanIiSt4plusIiEEvRN2cl4sycl6bufferIT_Li1ESaIS5_EEERNS3_5queueEENKUlRNS3_7handlerEE_clESC_ENKUlNS3_7nd_itemILi1EEEE_clESF_(i32 addrspace(3)* %par1, i32 addrspace(3)* %par2) #0 comdat align 2 { - %1 = alloca i32 addrspace(3)*, align 8 - %2 = alloca i32 addrspace(3)*, align 8 - %3 = alloca %"struct.std::plus", align 1 - store i32 addrspace(3)* %par1, i32 addrspace(3)** %1, align 8 - store i32 addrspace(3)* %par2, i32 addrspace(3)** %2, align 8 - %4 = load i32 addrspace(3)*, i32 addrspace(3)** %1, align 8 - %5 = load i32 addrspace(3)*, i32 addrspace(3)** %2, align 8 -; CHECK: %[[CAST1:.*]] = addrspacecast i32 addrspace(3)* %{{.*}} to i32 addrspace(4)* - %6 = addrspacecast i32 addrspace(3)* %4 to i32* -; CHECK: %[[CAST2:.*]] = addrspacecast i32 addrspace(3)* %{{.*}} to i32 addrspace(4)* - %7 = addrspacecast i32 addrspace(3)* %5 to i32* -; CHECK: %{{.*}} = call spir_func i32 @new.[[PLUS:.*]](%"struct.std::plus"* %{{.*}}, i32 addrspace(4)* %[[CAST1]], i32 addrspace(4)* %[[CAST2]]) - %8 = call spir_func i32 @_ZNKSt4plusIiEclERKiS2_(%"struct.std::plus"* %3, i32* dereferenceable(4) %6, i32* dereferenceable(4) %7) - ret void -} - -; Function Attrs: noinline -define linkonce_odr dso_local spir_func void @_ZZZ8par_scanIlSt10multipliesIlEEvRN2cl4sycl6bufferIT_Li1ESaIS5_EEERNS3_5queueEENKUlRNS3_7handlerEE_clESC_ENKUlNS3_7nd_itemILi1EEEE_clESF_(i64 addrspace(3)* %par1, i64 addrspace(3)* %par2) #0 comdat align 2 { - %1 = alloca i64 addrspace(3)*, align 8 - %2 = alloca i64 addrspace(3)*, align 8 - %3 = alloca %"struct.std::multiplies", align 1 - store i64 addrspace(3)* %par1, i64 addrspace(3)** %1, align 8 - store i64 addrspace(3)* %par2, i64 addrspace(3)** %2, align 8 - %4 = load i64 addrspace(3)*, i64 addrspace(3)** %1, align 8 - %5 = load i64 addrspace(3)*, i64 addrspace(3)** %2, align 8 -; CHECK: %[[CAST1:.*]] = addrspacecast i64 addrspace(3)* %{{.*}} to i64 addrspace(4)* - %6 = addrspacecast i64 addrspace(3)* %4 to i64* -; CHECK: %[[CAST2:.*]] = addrspacecast i64 addrspace(3)* %{{.*}} to i64 addrspace(4)* - %7 = addrspacecast i64 addrspace(3)* %5 to i64* -; CHECK: %{{.*}} = call spir_func i64 @new.[[MUL:.*]](%"struct.std::multiplies"* %{{.*}}, i64 addrspace(4)* %[[CAST1]], i64 addrspace(4)* %[[CAST2]]) - %8 = call spir_func i64 @_ZNKSt10multipliesIlEclERKlS2_(%"struct.std::multiplies"* %3, i64* dereferenceable(8) %6, i64* dereferenceable(8) %7) - ret void -} - -; CHECK-ADD: define linkonce_odr dso_local spir_func i32 @new.[[PLUS]](%"struct.std::plus"*, i32 addrspace(4)* dereferenceable(4), i32 addrspace(4)* dereferenceable(4)) #1 align 2 { -; Function Attrs: noinline nounwind optnone -define linkonce_odr dso_local spir_func i32 @_ZNKSt4plusIiEclERKiS2_(%"struct.std::plus"*, i32* dereferenceable(4), i32* dereferenceable(4)) #1 comdat align 2 { - %4 = alloca %"struct.std::plus"*, align 8 -; CHECK-ADD: %[[ALLOC1:.*]] = alloca i32 addrspace(4)* - %5 = alloca i32*, align 8 -; CHECK-ADD: %[[ALLOC2:.*]] = alloca i32 addrspace(4)* - %6 = alloca i32*, align 8 - store %"struct.std::plus"* %0, %"struct.std::plus"** %4, align 8 -; CHECK-ADD: store i32 addrspace(4)* %{{.*}}, i32 addrspace(4)** %[[ALLOC1]], align 8 - store i32* %1, i32** %5, align 8 -; CHECK-ADD: store i32 addrspace(4)* %{{.*}}, i32 addrspace(4)** %[[ALLOC2]], align 8 - store i32* %2, i32** %6, align 8 - %7 = load %"struct.std::plus"*, %"struct.std::plus"** %4, align 8 -; CHECK-ADD: %[[LOAD1:.*]] = load i32 addrspace(4)*, i32 addrspace(4)** %[[ALLOC1]] - %8 = load i32*, i32** %5, align 8 -; CHECK-ADD: %[[NEXT_LOAD1:.*]] = load i32, i32 addrspace(4)* %[[LOAD1]] - %9 = load i32, i32* %8, align 4 -; CHECK-ADD: %[[LOAD2:.*]] = load i32 addrspace(4)*, i32 addrspace(4)** %[[ALLOC2]] - %10 = load i32*, i32** %6, align 8 -; CHECK-ADD: %[[NEXT_LOAD2:.*]] = load i32, i32 addrspace(4)* %[[LOAD2]] - %11 = load i32, i32* %10, align 4 -; CHECK-ADD: %[[ADD:.*]] = add nsw i32 %[[NEXT_LOAD1]], %[[NEXT_LOAD2]] - %12 = add nsw i32 %9, %11 -; CHECK-ADD: %[[ADD2:.*]] = add nsw i32 %[[ADD]], 1 - %13 = add nsw i32 %12, 1 -; CHECK-ADD: ret i32 %[[ADD2]] - ret i32 %13 -} - -; CHECK-MUL: define linkonce_odr dso_local spir_func i64 @new.[[MUL]](%"struct.std::multiplies"*, i64 addrspace(4)* dereferenceable(8), i64 addrspace(4)* dereferenceable(8)) #1 align 2 { -; Function Attrs: noinline nounwind optnone -define linkonce_odr dso_local spir_func i64 @_ZNKSt10multipliesIlEclERKlS2_(%"struct.std::multiplies"*, i64* dereferenceable(8), i64* dereferenceable(8)) #1 comdat align 2 { - %4 = alloca %"struct.std::multiplies"*, align 8 -; CHECK-MUL: %[[ALLOC1:.*]] = alloca i64 addrspace(4)* - %5 = alloca i64*, align 8 -; CHECK-MUL: %[[ALLOC2:.*]] = alloca i64 addrspace(4)* - %6 = alloca i64*, align 8 - store %"struct.std::multiplies"* %0, %"struct.std::multiplies"** %4, align 8 -; CHECK-MUL: store i64 addrspace(4)* %{{.*}}, i64 addrspace(4)** %[[ALLOC1]], align 8 - store i64* %1, i64** %5, align 8 -; CHECK-MUL: store i64 addrspace(4)* %{{.*}}, i64 addrspace(4)** %[[ALLOC2]], align 8 - store i64* %2, i64** %6, align 8 - %7 = load %"struct.std::multiplies"*, %"struct.std::multiplies"** %4, align 8 -; CHECK-MUL: %[[LOAD1:.*]] = load i64 addrspace(4)*, i64 addrspace(4)** %[[ALLOC1]] - %8 = load i64*, i64** %5, align 8 -; CHECK-MUL: %[[NEXT_LOAD1:.*]] = load i64, i64 addrspace(4)* %[[LOAD1]] - %9 = load i64, i64* %8, align 8 -; CHECK-MUL: %[[LOAD2:.*]] = load i64 addrspace(4)*, i64 addrspace(4)** %[[ALLOC2]] - %10 = load i64*, i64** %6, align 8 -; CHECK-MUL: %[[NEXT_LOAD2:.*]] = load i64, i64 addrspace(4)* %[[LOAD2]] - %11 = load i64, i64* %10, align 8 -; CHECK-MUL: %[[MUL:.*]] = mul nsw i64 %[[NEXT_LOAD1]], %[[NEXT_LOAD2]] - %12 = mul nsw i64 %9, %11 -; CHECK-MUL: ret i64 %[[MUL]] - ret i64 %12 -} - -attributes #0 = { noinline "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="true" "no-frame-pointer-elim-non-leaf" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } -attributes #1 = { noinline nounwind "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="true" "no-frame-pointer-elim-non-leaf" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } - -!llvm.ident = !{!0} - -!0 = !{!"clang version 8.0.0"} diff --git a/llvm/test/SYCL/ASFixer/numbers-and-local-variables.ll b/llvm/test/SYCL/ASFixer/numbers-and-local-variables.ll deleted file mode 100644 index dd09550b4c8be..0000000000000 --- a/llvm/test/SYCL/ASFixer/numbers-and-local-variables.ll +++ /dev/null @@ -1,93 +0,0 @@ -; RUN: opt -asfix %s -S -o - | FileCheck %s -; Compiled from: -; -; void foo(int * Data) { -; int a = 10; -; *Data = 1 + a; -; *Data = 10; -; } -; -; void usages() { -; __attribute__((address_space(1))) int *GLOB; -; foo(GLOB); -; } -; ; ModuleID = 'new_test.cpp' -source_filename = "new_test.cpp" -target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" -target triple = "spir64-unknown-linux-sycldevice" - -%class.anon = type { i8 } - -; Function Attrs: argmemonly nounwind -declare void @llvm.lifetime.start.p0i8(i64, i8* nocapture) #1 - -; Function Attrs: argmemonly nounwind -declare void @llvm.lifetime.end.p0i8(i64, i8* nocapture) #1 - -; Function Attrs: nounwind -define spir_func void @_Z6usagesv() #3 { -entry: - %GLOB = alloca i32 addrspace(1)*, align 8 - %0 = bitcast i32 addrspace(1)** %GLOB to i8* - call void @llvm.lifetime.start.p0i8(i64 8, i8* %0) #4 - %1 = load i32 addrspace(1)*, i32 addrspace(1)** %GLOB, align 8, !tbaa !5 -; CHECK: %[[CAST:.*]] = addrspacecast i32 addrspace(1)* %{{.*}} to i32 addrspace(4)* - %2 = addrspacecast i32 addrspace(1)* %1 to i32* -; CHECK: call spir_func void @new.[[FOO:.*]](i32 addrspace(4)* %[[CAST]]) - call spir_func void @_Z3fooPi(i32* %2) - %3 = bitcast i32 addrspace(1)** %GLOB to i8* - call void @llvm.lifetime.end.p0i8(i64 8, i8* %3) #4 - ret void -} - -;CHECK: define spir_func void @new.[[FOO]](i32 addrspace(4)* %{{.*}}) -; Function Attrs: nounwind -define spir_func void @_Z3fooPi(i32* %Data) #3 { -entry: -; CHECK: %[[DATA_ADDR:.*]] = alloca i32 addrspace(4)* - %Data.addr = alloca i32*, align 8 -; CHECK: %[[A:.*]] = alloca i32, align 4 - %a = alloca i32, align 4 -; CHECK: store i32 addrspace(4)* %{{.*}}, i32 addrspace(4)** %[[DATA_ADDR]], align 8 - store i32* %Data, i32** %Data.addr, align 8, !tbaa !5 - %0 = bitcast i32* %a to i8* - call void @llvm.lifetime.start.p0i8(i64 4, i8* %0) #4 - store i32 10, i32* %a, align 4, !tbaa !9 - %1 = load i32, i32* %a, align 4, !tbaa !9 -; CHECK: %[[ADD:.*]] = add nsw i32 1, %{{.*}} - %add = add nsw i32 1, %1 -; CHECK: %[[DATA_LOAD:.*]] = load i32 addrspace(4)*, i32 addrspace(4)** %[[DATA_ADDR]] - %2 = load i32*, i32** %Data.addr, align 8, !tbaa !5 -; CHECK: store i32 %[[ADD]], i32 addrspace(4)* %[[DATA_LOAD]], align 4 - store i32 %add, i32* %2, align 4, !tbaa !9 -; CHECK: %[[DATA_LOAD:.*]] = load i32 addrspace(4)*, i32 addrspace(4)** %[[DATA_ADDR]] - %3 = load i32*, i32** %Data.addr, align 8, !tbaa !5 -; CHECK: store i32 10, i32 addrspace(4)* %[[DATA_LOAD]], align 4 - store i32 10, i32* %3, align 4, !tbaa !9 - %4 = bitcast i32* %a to i8* - call void @llvm.lifetime.end.p0i8(i64 4, i8* %4) #4 - ret void -} - -attributes #0 = { nounwind "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" } -attributes #1 = { argmemonly nounwind } -attributes #2 = { inlinehint nounwind "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } -attributes #3 = { nounwind "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } -attributes #4 = { nounwind } - -!llvm.module.flags = !{!0} -!opencl.spir.version = !{!1} -!spirv.Source = !{!2} -!llvm.ident = !{!3} - -!0 = !{i32 1, !"wchar_size", i32 4} -!1 = !{i32 1, i32 2} -!2 = !{i32 4, i32 100000} -!3 = !{!"clang version 8.0.0"} -!4 = !{} -!5 = !{!6, !6, i64 0} -!6 = !{!"any pointer", !7, i64 0} -!7 = !{!"omnipotent char", !8, i64 0} -!8 = !{!"Simple C++ TBAA"} -!9 = !{!10, !10, i64 0} -!10 = !{!"int", !7, i64 0} diff --git a/llvm/tools/opt/CMakeLists.txt b/llvm/tools/opt/CMakeLists.txt index caca16c183237..90730e324c7d4 100644 --- a/llvm/tools/opt/CMakeLists.txt +++ b/llvm/tools/opt/CMakeLists.txt @@ -22,7 +22,6 @@ set(LLVM_LINK_COMPONENTS TransformUtils Vectorize Passes - ASFixer ) add_llvm_tool(opt diff --git a/llvm/tools/opt/opt.cpp b/llvm/tools/opt/opt.cpp index 737fb67259b00..092932237fd67 100644 --- a/llvm/tools/opt/opt.cpp +++ b/llvm/tools/opt/opt.cpp @@ -565,7 +565,6 @@ int main(int argc, char **argv) { initializeExpandReductionsPass(Registry); initializeWasmEHPreparePass(Registry); initializeWriteBitcodePassPass(Registry); - initializeASFixerPass(Registry); initializeHardwareLoopsPass(Registry); #ifdef BUILD_EXAMPLES From d5a14302490e72694eb391e67ac6e5e951006b88 Mon Sep 17 00:00:00 2001 From: Vlad Romanov Date: Mon, 2 Dec 2019 17:35:14 +0300 Subject: [PATCH 2/3] remove XFAIL from anyremove tests on asfixer Signed-off-by: Vlad Romanov --- .../address-space-parameter-conversions.cpp | 164 ++++++++++++++++++ clang/test/CodeGenSYCL/bool-vectors.cpp | 2 +- 2 files changed, 165 insertions(+), 1 deletion(-) create mode 100644 clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp diff --git a/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp b/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp new file mode 100644 index 0000000000000..55ea1a917ad2c --- /dev/null +++ b/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp @@ -0,0 +1,164 @@ +// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -emit-llvm -x c++ %s -o - | FileCheck %s +void bar(int & Data) {} +// CHECK-DAG: define spir_func void @[[RAW_REF:[a-zA-Z0-9_]+]](i32 addrspace(4)* dereferenceable(4) % +void bar2(int & Data) {} +// CHECK-DAG: define spir_func void @[[RAW_REF2:[a-zA-Z0-9_]+]](i32 addrspace(4)* dereferenceable(4) % +void bar(__attribute__((ocl_local)) int &Data) {} +// CHECK-DAG: define spir_func void [[LOC_REF:@[a-zA-Z0-9_]+]](i32 addrspace(3)* dereferenceable(4) % +void foo(int * Data) {} +// CHECK-DAG: define spir_func void @[[RAW_PTR:[a-zA-Z0-9_]+]](i32 addrspace(4)* % +void foo2(int * Data) {} +// CHECK-DAG: define spir_func void @[[RAW_PTR2:[a-zA-Z0-9_]+]](i32 addrspace(4)* % +void foo(__attribute__((address_space(3))) int * Data) {} +// CHECK-DAG: define spir_func void [[LOC_PTR:@[a-zA-Z0-9_]+]](i32 addrspace(3)* % + +template +void tmpl(T t){} +// See Check Lines below. + +void usages() { + // CHECK-DAG: [[GLOB:%[a-zA-Z0-9]+]] = alloca i32 addrspace(1)* + __attribute__((address_space(1))) int *GLOB; + // CHECK-DAG: [[LOC:%[a-zA-Z0-9]+]] = alloca i32 addrspace(3)* + __attribute__((ocl_local)) int *LOC; + // CHECK-DAG: [[NoAS:%[a-zA-Z0-9]+]] = alloca i32 addrspace(4)* + int *NoAS; + + // CHECK-DAG: [[PRIV:%[a-zA-Z0-9]+]] = alloca i32* + __attribute__((ocl_private)) int *PRIV; + + bar(*GLOB); + // CHECK-DAG: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB]] + // CHECK-DAG: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(1)* [[GLOB_LOAD]] to i32 addrspace(4)* + // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[GLOB_CAST]]) + bar2(*GLOB); + // CHECK-DAG: [[GLOB_LOAD2:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB]] + // CHECK-DAG: [[GLOB_CAST2:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(1)* [[GLOB_LOAD2]] to i32 addrspace(4)* + // CHECK-DAG: call spir_func void @[[RAW_REF2]](i32 addrspace(4)* dereferenceable(4) [[GLOB_CAST2]]) + + bar(*LOC); + // CHECK-DAG: [[LOC_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOC]] + // CHECK-DAG: call spir_func void [[LOC_REF]](i32 addrspace(3)* dereferenceable(4) [[LOC_LOAD]]) + bar2(*LOC); + // CHECK-DAG: [[LOC_LOAD2:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOC]] + // CHECK-DAG: [[LOC_CAST2:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(3)* [[LOC_LOAD2]] to i32 addrspace(4)* + // CHECK-DAG: call spir_func void @[[RAW_REF2]](i32 addrspace(4)* dereferenceable(4) [[LOC_CAST2]]) + + bar(*NoAS); + // CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)** [[NoAS]] + // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[NoAS_LOAD]]) + bar2(*NoAS); + // CHECK-DAG: [[NoAS_LOAD2:%[a-zA-Z0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)** [[NoAS]] + // CHECK-DAG: call spir_func void @[[RAW_REF2]](i32 addrspace(4)* dereferenceable(4) [[NoAS_LOAD2]]) + + foo(GLOB); + // CHECK-DAG: [[GLOB_LOAD3:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB]] + // CHECK-DAG: [[GLOB_CAST3:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(1)* [[GLOB_LOAD3]] to i32 addrspace(4)* + // CHECK-DAG: call spir_func void @[[RAW_PTR]](i32 addrspace(4)* [[GLOB_CAST3]]) + foo2(GLOB); + // CHECK-DAG: [[GLOB_LOAD4:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB]] + // CHECK-DAG: [[GLOB_CAST4:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(1)* [[GLOB_LOAD4]] to i32 addrspace(4)* + // CHECK-DAG: call spir_func void @[[RAW_PTR2]](i32 addrspace(4)* [[GLOB_CAST4]]) + foo(LOC); + // CHECK-DAG: [[LOC_LOAD3:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOC]] + // CHECK-DAG: call spir_func void [[LOC_PTR]](i32 addrspace(3)* [[LOC_LOAD3]]) + foo2(LOC); + // CHECK-DAG: [[LOC_LOAD4:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOC]] + // CHECK-DAG: [[LOC_CAST4:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(3)* [[LOC_LOAD4]] to i32 addrspace(4)* + // CHECK-DAG: call spir_func void @[[RAW_PTR2]](i32 addrspace(4)* [[LOC_CAST4]]) + foo(NoAS); + // CHECK-DAG: [[NoAS_LOAD3:%[a-zA-Z0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)** [[NoAS]] + // CHECK-DAG: call spir_func void @[[RAW_PTR]](i32 addrspace(4)* [[NoAS_LOAD3]]) + foo2(NoAS); + // CHECK-DAG: [[NoAS_LOAD4:%[a-zA-Z0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)** [[NoAS]] + // CHECK-DAG: call spir_func void @[[RAW_PTR2]](i32 addrspace(4)* [[NoAS_LOAD4]]) + + // Ensure that we still get 3 different template instantiations. + tmpl(GLOB); + // CHECK-DAG: [[GLOB_LOAD4:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB]] + // CHECK-DAG: call spir_func void [[GLOB_TMPL:@[a-zA-Z0-9_]+]](i32 addrspace(1)* [[GLOB_LOAD4]]) + tmpl(LOC); + // CHECK-DAG: [[LOC_LOAD5:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOC]] + // CHECK-DAG: call spir_func void [[LOC_TMPL:@[a-zA-Z0-9_]+]](i32 addrspace(3)* [[LOC_LOAD5]]) + tmpl(PRIV); + // CHECK-DAG: [[PRIV_LOAD5:%[a-zA-Z0-9]+]] = load i32*, i32** [[PRIV]] + // CHECK-DAG: call spir_func void [[PRIV_TMPL:@[a-zA-Z0-9_]+]](i32* [[PRIV_LOAD5]]) + tmpl(NoAS); + // CHECK-DAG: [[NoAS_LOAD5:%[a-zA-Z0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)** [[NoAS]] + // CHECK-DAG: call spir_func void [[GEN_TMPL:@[a-zA-Z0-9_]+]](i32 addrspace(4)* [[NoAS_LOAD5]]) +} + +// CHECK-DAG: define linkonce_odr spir_func void [[GLOB_TMPL]](i32 addrspace(1)* % +// CHECK-DAG: define linkonce_odr spir_func void [[LOC_TMPL]](i32 addrspace(3)* % +// CHECK-DAG: define linkonce_odr spir_func void [[PRIV_TMPL]](i32* % +// CHECK-DAG: define linkonce_odr spir_func void [[GEN_TMPL]](i32 addrspace(4)* % + +void usages2() { + __attribute__((address_space(0))) int *PRIV_NUM; + // CHECK-DAG: [[PRIV_NUM:%[a-zA-Z0-9_]+]] = alloca i32* + __attribute__((address_space(0))) int *PRIV_NUM2; + // CHECK-DAG: [[PRIV_NUM2:%[a-zA-Z0-9_]+]] = alloca i32* + __attribute__((ocl_private)) int *PRIV; + // CHECK-DAG: [[PRIV:%[a-zA-Z0-9_]+]] = alloca i32* + __attribute__((address_space(1))) int *GLOB_NUM; + // CHECK-DAG: [[GLOB_NUM:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(1)* + __attribute__((ocl_global)) int *GLOB; + // CHECK-DAG: [[GLOB:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(1)* + __attribute__((address_space(2))) int *CONST_NUM; + // CHECK-DAG: [[CONST_NUM:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(2)* + __attribute__((ocl_constant)) int *CONST; + // CHECK-DAG: [[CONST:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(2)* + __attribute__((address_space(3))) int *LOCAL_NUM; + // CHECK-DAG: [[LOCAL_NUM:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(3)* + __attribute__((ocl_local)) int *LOCAL; + // CHECK-DAG: [[LOCAL:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(3)* + + bar(*PRIV_NUM); + // CHECK-DAG: [[PRIV_NUM_LOAD:%[a-zA-Z0-9]+]] = load i32*, i32** [[PRIV_NUM]] + // CHECK-DAG: [[PRIV_NUM_ASCAST:%[a-zA-Z0-9]+]] = addrspacecast i32* [[PRIV_NUM_LOAD]] to i32 addrspace(4)* + // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[PRIV_NUM_ASCAST]]) + bar(*PRIV_NUM2); + // CHECK-DAG: [[PRIV_NUM2_LOAD:%[a-zA-Z0-9]+]] = load i32*, i32** [[PRIV_NUM2]] + // CHECK-DAG: [[PRIV_NUM2_ASCAST:%[a-zA-Z0-9]+]] = addrspacecast i32* [[PRIV_NUM2_LOAD]] to i32 addrspace(4)* + // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[PRIV_NUM2_ASCAST]]) + bar(*PRIV); + // CHECK-DAG: [[PRIV_LOAD:%[a-zA-Z0-9]+]] = load i32*, i32** [[PRIV]] + // CHECK-DAG: [[PRIV_ASCAST:%[a-zA-Z0-9]+]] = addrspacecast i32* [[PRIV_LOAD]] to i32 addrspace(4)* + // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[PRIV_ASCAST]]) + bar(*GLOB_NUM); + // CHECK-DAG: [[GLOB_NUM_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB_NUM]] + // CHECK-DAG: [[GLOB_NUM_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(1)* [[GLOB_NUM_LOAD]] to i32 addrspace(4)* + // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[GLOB_NUM_CAST]]) + bar(*GLOB); + // CHECK-DAG: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB]] + // CHECK-DAG: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(1)* [[GLOB_LOAD]] to i32 addrspace(4)* + // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[GLOB_CAST]]) + bar(*CONST_NUM); + // CHECK-DAG: [[CONST_NUM_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(2)*, i32 addrspace(2)** [[CONST_NUM]] + // CHECK-DAG: [[CONST_NUM_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(2)* [[CONST_NUM_LOAD]] to i32 addrspace(4)* + // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[CONST_NUM_CAST]]) + bar(*CONST); + // CHECK-DAG: [[CONST_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(2)*, i32 addrspace(2)** [[CONST]] + // CHECK-DAG: [[CONST_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(2)* [[CONST_LOAD]] to i32 addrspace(4)* + // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[CONST_CAST]]) + bar2(*LOCAL_NUM); + // CHECK-DAG: [[LOCAL_NUM_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOCAL_NUM]] + // CHECK-DAG: [[LOCAL_NUM_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(3)* [[LOCAL_NUM_LOAD]] to i32 addrspace(4)* + // CHECK-DAG: call spir_func void @[[RAW_REF2]](i32 addrspace(4)* dereferenceable(4) [[LOCAL_NUM_CAST]]) + bar2(*LOCAL); + // CHECK-DAG: [[LOCAL_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOCAL]] + // CHECK-DAG: [[LOCAL_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(3)* [[LOCAL_LOAD]] to i32 addrspace(4)* + // CHECK-DAG: call spir_func void @[[RAW_REF2]](i32 addrspace(4)* dereferenceable(4) [[LOCAL_CAST]]) +} + +template +__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { + kernelFunc(); +} +int main() { + kernel_single_task([]() { usages();usages2(); }); + return 0; +} + +// TODO: SYCL specific fail - analyze and enable +// XFAIL: windows-msvc diff --git a/clang/test/CodeGenSYCL/bool-vectors.cpp b/clang/test/CodeGenSYCL/bool-vectors.cpp index 8fb2313e46c6e..b4f64de7bcba1 100644 --- a/clang/test/CodeGenSYCL/bool-vectors.cpp +++ b/clang/test/CodeGenSYCL/bool-vectors.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -emit-llvm -x c++ %s -o - | opt -asfix -S -o - | FileCheck %s +// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -emit-llvm -x c++ %s -o - | FileCheck %s template __attribute__((sycl_kernel)) void kernel(Func kernelFunc) { From dea4b0e9b32bcf1f0d0dc9b2d136cb7f9525e8c1 Mon Sep 17 00:00:00 2001 From: Vlad Romanov Date: Mon, 2 Dec 2019 22:50:25 +0300 Subject: [PATCH 3/3] remove XFAIL from one more test Signed-off-by: Vlad Romanov --- clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp | 3 --- 1 file changed, 3 deletions(-) diff --git a/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp b/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp index 55ea1a917ad2c..1d5beced187bd 100644 --- a/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp +++ b/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp @@ -159,6 +159,3 @@ int main() { kernel_single_task([]() { usages();usages2(); }); return 0; } - -// TODO: SYCL specific fail - analyze and enable -// XFAIL: windows-msvc