From 9be97686f0f56f6e6bbe1da7e4f25ffa84fd77f0 Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Tue, 18 Jun 2019 13:46:19 +0300 Subject: [PATCH 1/4] [SYCL] Move static variables to global address space by default Static variables without address space now reside in global address space, unless they have an explicit address space qualifier in source code. Signed-off-by: Andrew Savonichev --- clang/lib/CodeGen/CodeGenModule.cpp | 8 ++++++++ clang/test/CodeGenSYCL/address-space-new.cpp | 10 ++++++++-- 2 files changed, 16 insertions(+), 2 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 43e5aa97d613a..3bc612fdae88b 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -3630,6 +3630,14 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) { return AddrSpace; } + if (LangOpts.SYCLIsDevice) { + if (getenv("ENABLE_INFER_AS")) { + if (!D || D->getType().getAddressSpace() == LangAS::Default) { + return LangAS::opencl_global; + } + } + } + if (LangOpts.CUDA && LangOpts.CUDAIsDevice) { if (D && D->hasAttr()) return LangAS::cuda_constant; diff --git a/clang/test/CodeGenSYCL/address-space-new.cpp b/clang/test/CodeGenSYCL/address-space-new.cpp index 15df1b64e032d..385a82037cbec 100644 --- a/clang/test/CodeGenSYCL/address-space-new.cpp +++ b/clang/test/CodeGenSYCL/address-space-new.cpp @@ -1,13 +1,19 @@ -// 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-prefix=CHECK-DEFAULT +// 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-prefix=CHECK-LEGACY // RUN: ENABLE_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-prefix=CHECK-NEW 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 + int i = 0; int *pptr = &i; - // CHECK-DEFAULT: store i32* %i, i32** %pptr + // 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 + + *pptr = foo; } From 0dc80bdf4323201bd7efbe0a2657a091c53608e6 Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Tue, 18 Jun 2019 13:43:46 +0300 Subject: [PATCH 2/4] [SYCL] Return __constant pointers from 'constant' multi_ptr We cannot return a generic pointer like in 'global', 'local' and 'private' cases, because 'constant' and 'generic' address spaces do not overlap. This patch returns __constant pointer as is, and assumes that a user have to deal with __constant pointer limitations (ie. such pointers cannot be casted to plain (default) pointers). Signed-off-by: Andrew Savonichev --- sycl/include/CL/sycl/multi_ptr.hpp | 40 ++++++++++++++++++++++-------- 1 file changed, 29 insertions(+), 11 deletions(-) diff --git a/sycl/include/CL/sycl/multi_ptr.hpp b/sycl/include/CL/sycl/multi_ptr.hpp index fbc845d710c78..62c4a32f92709 100644 --- a/sycl/include/CL/sycl/multi_ptr.hpp +++ b/sycl/include/CL/sycl/multi_ptr.hpp @@ -69,17 +69,37 @@ template class multi_ptr { m_Pointer = nullptr; return *this; } - ElementType &operator*() const { - return *(reinterpret_cast(m_Pointer)); + +#ifdef __SYCL_ENABLE_INFER_AS__ + using ReturnPtr = + typename std::conditional::type; + using ReturnRef = + typename std::conditional::type; + using ReturnConstRef = + typename std::conditional::type; +#else + using ReturnPtr = ElementType *; + using ReturnRef = ElementType &; + using ReturnConstRef = const ElementType &; +#endif + + ReturnRef operator*() const { + return *reinterpret_cast(m_Pointer); } - ElementType *operator->() const { - return reinterpret_cast(m_Pointer); + + ReturnPtr operator->() const { + return reinterpret_cast(m_Pointer); } - ElementType &operator[](difference_type index) { - return *(reinterpret_cast(m_Pointer + index)); + + ReturnRef operator[](difference_type index) { + return reinterpret_cast(m_Pointer)[index]; } - ElementType operator[](difference_type index) const { - return *(reinterpret_cast(m_Pointer + index)); + + ReturnConstRef operator[](difference_type index) const { + return reinterpret_cast(m_Pointer)[index]; } // Only if Space == global_space @@ -181,9 +201,7 @@ template class multi_ptr { pointer_t get() const { return m_Pointer; } // Implicit conversion to the underlying pointer type - operator ElementType *() const { - return reinterpret_cast(m_Pointer); - } + operator ReturnPtr() const { return reinterpret_cast(m_Pointer); } // Implicit conversion to a multi_ptr // Only available when ElementType is not const-qualified From 6ea08ec33de43847efa35fdb479de1933a6b8049 Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Tue, 18 Jun 2019 13:51:29 +0300 Subject: [PATCH 3/4] [SYCL] Turn new address space rules by default This patch enables new rules for address space handling in SYCL by default. Previous rules for address space handling can still be enabled by DISABLE_INFER_AS=1 environment variable. This approach is known as "generic pointers by default", and it essentially makes all pointers without an explicit address space qualifier to be pointers in generic address space. For example: void foo() { int* ip = foo(); } void* bar(int* ip); template void swap(T&, T&) {} All pointers (and references) in the example above are CodeGen'ed as pointers in addrspace(4) in LLVM IR (where '4' is a number for generic address space). Concrete address space (global, local and private) can originate from: 1) an address-of operation applied to an automatic variable: int i = 42; (void) &i // <--- this is a private pointer 2) Static variable 3) SYCL global, local, or constant buffers In all three cases, once this pointer is used, it is immediately addrspacecast'ed to generic, because a user does not (and should not) specify address space for pointers in source code. Signed-off-by: Andrew Savonichev --- clang/lib/Basic/Targets/SPIR.h | 2 +- clang/lib/CodeGen/BackendUtil.cpp | 4 +- clang/lib/CodeGen/CGCall.cpp | 2 +- clang/lib/CodeGen/CGExpr.cpp | 2 +- clang/lib/CodeGen/CodeGenModule.cpp | 2 +- clang/lib/Frontend/InitPreprocessor.cpp | 2 + clang/test/CodeGenSYCL/address-space-new.cpp | 4 +- .../address-space-parameter-conversions.cpp | 117 ++++++++++++------ clang/test/CodeGenSYCL/address-space-swap.cpp | 36 ++++++ .../test/CodeGenSYCL/basic-kernel-wrapper.cpp | 11 +- .../CodeGenSYCL/debug-info-srcpos-kernel.cpp | 6 +- clang/test/CodeGenSYCL/device-functions.cpp | 6 +- clang/test/CodeGenSYCL/sampler.cpp | 7 +- clang/test/CodeGenSYCL/spir-calling-conv.cpp | 9 +- clang/test/SemaSYCL/spir-enum.cpp | 6 +- 15 files changed, 155 insertions(+), 61 deletions(-) create mode 100644 clang/test/CodeGenSYCL/address-space-swap.cpp diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h index 9a2895e24d000..9468320fa6594 100644 --- a/clang/lib/Basic/Targets/SPIR.h +++ b/clang/lib/Basic/Targets/SPIR.h @@ -63,7 +63,7 @@ class LLVM_LIBRARY_VISIBILITY SPIRTargetInfo : public TargetInfo { VLASupported = false; LongWidth = LongAlign = 64; if (Triple.getEnvironment() == llvm::Triple::SYCLDevice && - getenv("ENABLE_INFER_AS")) { + !getenv("DISABLE_INFER_AS")) { AddrSpaceMap = &SYCLAddrSpaceMap; } else { AddrSpaceMap = &SPIRAddrSpaceMap; diff --git a/clang/lib/CodeGen/BackendUtil.cpp b/clang/lib/CodeGen/BackendUtil.cpp index ce3be99226466..7d9a54df4f34c 100644 --- a/clang/lib/CodeGen/BackendUtil.cpp +++ b/clang/lib/CodeGen/BackendUtil.cpp @@ -830,7 +830,7 @@ void EmitAssemblyHelper::EmitAssembly(BackendAction Action, case Backend_EmitBC: if (LangOpts.SYCLIsDevice) { - if (!getenv("ENABLE_INFER_AS")) + if (getenv("DISABLE_INFER_AS")) PerModulePasses.add(createASFixerPass()); PerModulePasses.add(createDeadCodeEliminationPass()); } @@ -1230,7 +1230,7 @@ void EmitAssemblyHelper::EmitAssemblyWithNewPassManager( case Backend_EmitBC: if (LangOpts.SYCLIsDevice) { - if (!getenv("ENABLE_INFER_AS")) + if (getenv("DISABLE_INFER_AS")) CodeGenPasses.add(createASFixerPass()); CodeGenPasses.add(createDeadCodeEliminationPass()); } diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp index 59b33b516ea48..7359f39be6164 100644 --- a/clang/lib/CodeGen/CGCall.cpp +++ b/clang/lib/CodeGen/CGCall.cpp @@ -4256,7 +4256,7 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo, deactivateArgCleanupsBeforeCall(*this, CallArgs); // Addrspace cast to generic if necessary - if (getenv("ENABLE_INFER_AS")) { + if (!getenv("DISABLE_INFER_AS")) { for (unsigned i = 0; i < IRFuncTy->getNumParams(); ++i) { if (auto *PtrTy = dyn_cast(IRCallArgs[i]->getType())) { auto *ExpectedPtrType = diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp index 61760966c80a2..9d2a88ecc008f 100644 --- a/clang/lib/CodeGen/CGExpr.cpp +++ b/clang/lib/CodeGen/CGExpr.cpp @@ -1745,7 +1745,7 @@ void CodeGenFunction::EmitStoreOfScalar(llvm::Value *Value, Address Addr, return; } - if (getenv("ENABLE_INFER_AS")) { + if (!getenv("DISABLE_INFER_AS")) { if (auto *PtrTy = dyn_cast(Value->getType())) { auto *ExpectedPtrType = cast(Addr.getType()->getElementType()); diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 3bc612fdae88b..e7447b67d0805 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -3631,7 +3631,7 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) { } if (LangOpts.SYCLIsDevice) { - if (getenv("ENABLE_INFER_AS")) { + if (!getenv("DISABLE_INFER_AS")) { if (!D || D->getType().getAddressSpace() == LangAS::Default) { return LangAS::opencl_global; } diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp index 5aff3cae1b6a6..893914be14493 100644 --- a/clang/lib/Frontend/InitPreprocessor.cpp +++ b/clang/lib/Frontend/InitPreprocessor.cpp @@ -1063,6 +1063,8 @@ static void InitializePredefinedMacros(const TargetInfo &TI, // SYCL device compiler which doesn't produce host binary. if (LangOpts.SYCLIsDevice) { Builder.defineMacro("__SYCL_DEVICE_ONLY__", "1"); + if (!getenv("DISABLE_INFER_AS")) + Builder.defineMacro("__SYCL_ENABLE_INFER_AS__", "1"); } // OpenCL definitions. diff --git a/clang/test/CodeGenSYCL/address-space-new.cpp b/clang/test/CodeGenSYCL/address-space-new.cpp index 385a82037cbec..9d6f833772093 100644 --- a/clang/test/CodeGenSYCL/address-space-new.cpp +++ b/clang/test/CodeGenSYCL/address-space-new.cpp @@ -1,5 +1,5 @@ -// 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-prefix=CHECK-LEGACY -// RUN: ENABLE_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-prefix=CHECK-NEW +// 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-prefix=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-prefix=CHECK-NEW void test() { diff --git a/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp b/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp index 671f781fc5242..6210544c41b78 100644 --- a/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp +++ b/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp @@ -1,14 +1,19 @@ -// 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: 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-DAG: define spir_func void @[[RAW_REF:[a-zA-Z0-9_]+]](i32* dereferenceable(4) % +// 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-DAG: define spir_func void @[[RAW_REF2:[a-zA-Z0-9_]+]](i32* dereferenceable(4) % +// 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(__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* % +// 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-DAG: define spir_func void @[[RAW_PTR2:[a-zA-Z0-9_]+]](i32* % +// 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)* % @@ -21,17 +26,23 @@ void usages() { __attribute__((address_space(1))) int *GLOB; // CHECK-DAG: [[LOC:%[a-zA-Z0-9]+]] = alloca i32 addrspace(3)* __local int *LOC; - // CHECK-DAG: [[NoAS:%[a-zA-Z0-9]+]] = alloca i32* + // 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* + __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 @new.[[RAW_REF]](i32 addrspace(4)* [[GLOB_CAST]]) + // 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-DAG: call spir_func void @new.[[RAW_REF2]](i32 addrspace(4)* [[GLOB_CAST2]]) + // 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]] @@ -39,36 +50,48 @@ void usages() { 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 @new.[[RAW_REF2]](i32 addrspace(4)* [[LOC_CAST2]]) + // 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-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load i32*, i32** [[NoAS]] - // CHECK-DAG: call spir_func void @[[RAW_REF]](i32* dereferenceable(4) [[NoAS_LOAD]]) + // 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-DAG: [[NoAS_LOAD2:%[a-zA-Z0-9]+]] = load i32*, i32** [[NoAS]] - // CHECK-DAG: call spir_func void @[[RAW_REF2]](i32* dereferenceable(4) [[NoAS_LOAD2]]) + // 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-DAG: call spir_func void @new.[[RAW_PTR]](i32 addrspace(4)* [[GLOB_CAST3]]) + // 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-DAG: call spir_func void @new.[[RAW_PTR2]](i32 addrspace(4)* [[GLOB_CAST4]]) + // 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-DAG: call spir_func void @new.[[RAW_PTR2]](i32 addrspace(4)* [[LOC_CAST4]]) + // 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-DAG: [[NoAS_LOAD3:%[a-zA-Z0-9]+]] = load i32*, i32** [[NoAS]] - // CHECK-DAG: call spir_func void @[[RAW_PTR]](i32* [[NoAS_LOAD3]]) + // 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-DAG: [[NoAS_LOAD4:%[a-zA-Z0-9]+]] = load i32*, i32** [[NoAS]] - // CHECK-DAG: call spir_func void @[[RAW_PTR2]](i32* [[NoAS_LOAD4]]) + // 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); @@ -77,14 +100,21 @@ void usages() { 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*, i32** [[NoAS]] - // CHECK-DAG: call spir_func void [[AS0_TMPL:@[a-zA-Z0-9_]+]](i32* [[NoAS_LOAD5]]) + // 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-DAG: define linkonce_odr spir_func void [[AS0_TMPL]](i32* % +// 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; @@ -108,46 +138,55 @@ void usages2() { bar(*PRIV_NUM); // CHECK-DAG: [[PRIV_NUM_LOAD:%[a-zA-Z0-9]+]] = load i32*, i32** [[PRIV_NUM]] - // CHECK-DAG: call spir_func void @[[RAW_REF]](i32* dereferenceable(4) [[PRIV_NUM_LOAD]]) + // 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-DAG: call spir_func void @[[RAW_REF]](i32* dereferenceable(4) [[PRIV_NUM2_LOAD]]) + // 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-DAG: call spir_func void @[[RAW_REF]](i32* dereferenceable(4) [[PRIV_LOAD]]) + // 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-DAG: call spir_func void @new.[[RAW_REF]](i32 addrspace(4)* [[GLOB_NUM_CAST]]) + // 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-DAG: call spir_func void @new.[[RAW_REF]](i32 addrspace(4)* [[GLOB_CAST]]) + // 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-DAG: call spir_func void @new.[[RAW_REF]](i32 addrspace(4)* [[CONST_NUM_CAST]]) + // 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-DAG: call spir_func void @new.[[RAW_REF]](i32 addrspace(4)* [[CONST_CAST]]) + // 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-DAG: call spir_func void @new.[[RAW_REF2]](i32 addrspace(4)* [[LOCAL_NUM_CAST]]) + // 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-DAG: call spir_func void @new.[[RAW_REF2]](i32 addrspace(4)* [[LOCAL_CAST]]) + // 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-DAG: define spir_func void @new.[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) - -// CHECK-DAG: define spir_func void @new.[[RAW_REF2]](i32 addrspace(4)* dereferenceable(4) - -// CHECK-DAG: define spir_func void @new.[[RAW_PTR]](i32 addrspace(4)* - -// CHECK-DAG: define spir_func void @new.[[RAW_PTR2]](i32 addrspace(4)* +// 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) { diff --git a/clang/test/CodeGenSYCL/address-space-swap.cpp b/clang/test/CodeGenSYCL/address-space-swap.cpp new file mode 100644 index 0000000000000..422c2891249a9 --- /dev/null +++ b/clang/test/CodeGenSYCL/address-space-swap.cpp @@ -0,0 +1,36 @@ +// RUN: %clang --sycl -S -emit-llvm -x c++ %s -o - | FileCheck %s +#include + + +void test() { + static int foo = 0x42; +// CHECK: @[[FOO:[a-zA-Z0-9_]+]] = internal addrspace(1) global i32 66, align 4 + int i = 43; +// CHECK: %[[I:[a-zA-Z0-9_]+]] = alloca i32, align 4 + + int* p1 = &foo; + int* p2 = &i; +// CHECK: %[[P1:[a-zA-Z0-9_]+]] = alloca i32 addrspace(4)*, align 8 +// CHECK: %[[P2:[a-zA-Z0-9_]+]] = alloca i32 addrspace(4)*, align 8 +// CHECK: %[[P1GEN:[a-zA-Z0-9_]+]] = addrspacecast i32 addrspace(4)** %[[P1]] to i32 addrspace(4)* addrspace(4)* +// CHECK: %[[P2GEN:[a-zA-Z0-9_]+]] = addrspacecast i32 addrspace(4)** %[[P2]] to i32 addrspace(4)* addrspace(4)* + + std::swap(p1, p2); +// CHECK: call spir_func void @_ZSt4swap{{.*}}(i32 addrspace(4)* addrspace(4)* dereferenceable(8) %[[P1GEN]], i32 addrspace(4)* addrspace(4)* dereferenceable(8) %[[P2GEN]]) + + std::swap(foo, i); +// CHECK: %[[ICAST:[a-zA-Z0-9_]+]] = addrspacecast i32* %[[I]] to i32 addrspace(4)* +// CHECK: call spir_func void @_ZSt4swap{{.*}}(i32 addrspace(4)* dereferenceable(4) addrspacecast (i32 addrspace(1)* @[[FOO]] to i32 addrspace(4)*), i32 addrspace(4)* dereferenceable(4) %[[ICAST]]) +} + + +template +__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { + kernelFunc(); +} + + +int main() { + kernel_single_task([]() { test(); }); + return 0; +} diff --git a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp index e5d58b91f0c1a..5ba24f924057b 100644 --- a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp +++ b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp @@ -1,4 +1,5 @@ -// 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 +// 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 // This test checks that compiler generates correct kernel wrapper for basic // case. @@ -37,7 +38,11 @@ int main() { // CHECK: [[MEM_LOAD:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG]].addr // Check accessor __init method call -// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor"* [[ACCESSOR]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval align 4 [[ACC_RANGE]], %"struct.{{.*}}.cl::sycl::range"* byval align 4 [[MEM_RANGE]], %"struct.{{.*}}.cl::sycl::id"* byval align 4 [[OFFSET]]) +// CHECK-OLD: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor"* [[ACCESSOR]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval align 4 [[ACC_RANGE]], %"struct.{{.*}}.cl::sycl::range"* byval align 4 [[MEM_RANGE]], %"struct.{{.*}}.cl::sycl::id"* byval align 4 [[OFFSET]]) +// 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 [[ACC_RANGE]], %"struct.{{.*}}.cl::sycl::range"* byval align 4 [[MEM_RANGE]], %"struct.{{.*}}.cl::sycl::id"* byval align 4 [[OFFSET]]) // Check lambda "()" operator call -// CHECK: call spir_func void @{{.*}}(%"class.{{.*}}.anon"* [[ANON]]) +// 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]]) diff --git a/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp b/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp index 2d2d8cc7dc4b9..51390aa8c7a15 100644 --- a/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp +++ b/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp @@ -1,4 +1,5 @@ -// RUN: %clang --sycl %s -S -emit-llvm -g -o - | FileCheck %s +// RUN: DISABLE_INFER_AS=1 %clang --sycl %s -S -emit-llvm -g -o - | FileCheck %s --check-prefixes CHECK,CHECK-OLD +// RUN: %clang --sycl %s -S -emit-llvm -g -o - | FileCheck %s --check-prefixes CHECK,CHECK-NEW // // Verify the SYCL kernel routine is marked artificial. // @@ -21,7 +22,8 @@ int main() { return 0; } -// CHECK: define{{.*}} spir_kernel {{.*}}void @_ZTSZ4mainE15kernel_function(i32*{{.*}}){{.*}} !dbg [[KERNEL:![0-9]+]] {{.*}}{ +// CHECK-OLD: define{{.*}} spir_kernel {{.*}}void @_ZTSZ4mainE15kernel_function(i32*{{.*}}){{.*}} !dbg [[KERNEL:![0-9]+]] {{.*}}{ +// CHECK-NEW: define{{.*}} spir_kernel {{.*}}void @_ZTSZ4mainE15kernel_function(i32 addrspace(4)*{{.*}}){{.*}} !dbg [[KERNEL:![0-9]+]] {{.*}}{ // CHECK: [[FILE:![0-9]+]] = !DIFile(filename: "{{.*}}debug-info-srcpos-kernel.cpp"{{.*}}) // CHECK: [[KERNEL]] = {{.*}}!DISubprogram(name: "_ZTSZ4mainE15kernel_function" // CHECK-SAME: scope: [[FILE]] diff --git a/clang/test/CodeGenSYCL/device-functions.cpp b/clang/test/CodeGenSYCL/device-functions.cpp index 114e3430fbc05..bb43e379bbde8 100644 --- a/clang/test/CodeGenSYCL/device-functions.cpp +++ b/clang/test/CodeGenSYCL/device-functions.cpp @@ -1,4 +1,5 @@ -// 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 +// 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 template T bar(T arg); @@ -22,6 +23,7 @@ int main() { return 0; } // CHECK: define spir_kernel void @_ZTSZ4mainE11fake_kernel() -// CHECK: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}.anon"* %this) +// 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 spir_func void @_Z3foov() // CHECK: define linkonce_odr spir_func i32 @_Z3barIiET_S0_(i32 %arg) diff --git a/clang/test/CodeGenSYCL/sampler.cpp b/clang/test/CodeGenSYCL/sampler.cpp index 8d75cc4508c7f..b7c4dd475e2f3 100644 --- a/clang/test/CodeGenSYCL/sampler.cpp +++ b/clang/test/CodeGenSYCL/sampler.cpp @@ -1,4 +1,5 @@ -// 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 +// 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 // 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 @@ -8,7 +9,9 @@ // 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-NEXT: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.{{.*}}.cl::sycl::sampler"* [[GEP]], %opencl.sampler_t addrspace(2)* [[LOAD_SAMPLER_ARG]]) +// 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]]) // #include "sycl.hpp" diff --git a/clang/test/CodeGenSYCL/spir-calling-conv.cpp b/clang/test/CodeGenSYCL/spir-calling-conv.cpp index daace615a635c..114b81a12d6e3 100644 --- a/clang/test/CodeGenSYCL/spir-calling-conv.cpp +++ b/clang/test/CodeGenSYCL/spir-calling-conv.cpp @@ -1,4 +1,5 @@ -// 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 +// 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 template __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { @@ -9,9 +10,11 @@ int main() { // CHECK: define spir_kernel void @_ZTSZ4mainE15kernel_function() - // CHECK: call spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}.anon"* %0) + // 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: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}anon"* %this) + // 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) kernel_single_task([]() {}); return 0; diff --git a/clang/test/SemaSYCL/spir-enum.cpp b/clang/test/SemaSYCL/spir-enum.cpp index 89f785afdd38b..e9d943d7d7271 100644 --- a/clang/test/SemaSYCL/spir-enum.cpp +++ b/clang/test/SemaSYCL/spir-enum.cpp @@ -1,4 +1,5 @@ -// 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 +// 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 template __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { @@ -23,7 +24,8 @@ int main() { // CHECK: define spir_kernel void @_ZTSZ4test9enum_typeE15kernel_function(i32 %_arg_) // CHECK: getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* - // CHECK: call spir_func void @"_ZZ4test9enum_typeENK3$_0clEv"(%"class.{{.*}}.anon"* %0) + // 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) test( enum_type::B ); From 02774e12f97530c83cf56e9f915d6c27fe24c312 Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Fri, 21 Jun 2019 19:53:24 +0300 Subject: [PATCH 4/4] [SYCL] Put string literals to private address space Although OpenCL specification explicitly states that a string literal should reside in constant address space, it does not work for SYCL with "generic by default" address space rules. For example: const char *getLiteral() { return "A"; } void func(bool AorB) { char B[] = {'B', '\0'}; const char* C = AorB ? A : B; } If `A' reside in constant address space, it cannot be returned from a function `getLiteral', because it returns a generic const char*. Signed-off-by: Andrew Savonichev --- clang/lib/CodeGen/CodeGenModule.cpp | 8 ++++++++ clang/test/CodeGenSYCL/address-space-new.cpp | 13 +++++++++---- 2 files changed, 17 insertions(+), 4 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index e7447b67d0805..217ba58f45dc2 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -3663,6 +3663,14 @@ 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 we keep a literal string in constant address space, the following code + // becomes illegal: + // + // const char *getLiteral() n{ + // return "AB"; + // } + return LangAS::opencl_private; if (auto AS = getTarget().getConstantAddressSpace()) return AS.getValue(); return LangAS::Default; diff --git a/clang/test/CodeGenSYCL/address-space-new.cpp b/clang/test/CodeGenSYCL/address-space-new.cpp index 9d6f833772093..3f10768b72b01 100644 --- a/clang/test/CodeGenSYCL/address-space-new.cpp +++ b/clang/test/CodeGenSYCL/address-space-new.cpp @@ -1,19 +1,24 @@ -// 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-prefix=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-prefix=CHECK-NEW - +// 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 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: @[[STR:[.a-zA-Z0-9_]+]] = private unnamed_addr constant [14 x i8] c"Hello, world!\00", align 1 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 - *pptr = foo; + + const char *str = "Hello, world!"; + // CHECK-LEGACY: store i8* getelementptr inbounds ([14 x i8], [14 x i8]* @[[STR]], i64 0, i64 0), i8** %{{.*}}, 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)** %{{.*}}, align 8 + + i = str[0]; }