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 43e5aa97d613a..217ba58f45dc2 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("DISABLE_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; @@ -3655,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/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 15df1b64e032d..3f10768b72b01 100644 --- a/clang/test/CodeGenSYCL/address-space-new.cpp +++ b/clang/test/CodeGenSYCL/address-space-new.cpp @@ -1,13 +1,24 @@ -// 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: 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-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-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; + + 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]; } 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 ); 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