From 5c5d0ae0115e4b62d8f4cbeb146b05c4ff0814c7 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Thu, 24 Sep 2020 13:12:29 +0300 Subject: [PATCH 1/6] Change adress space for global variables Signed-off-by: Aleksander Fadeev --- clang/lib/CodeGen/CGDecl.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/clang/lib/CodeGen/CGDecl.cpp b/clang/lib/CodeGen/CGDecl.cpp index 87070a3504a5..102a5a84647a 100644 --- a/clang/lib/CodeGen/CGDecl.cpp +++ b/clang/lib/CodeGen/CGDecl.cpp @@ -1144,6 +1144,8 @@ Address CodeGenModule::createUnnamedGlobalFrom(const VarDecl &D, llvm::GlobalVariable *InsertBefore = nullptr; unsigned AS = getContext().getTargetAddressSpace(getStringLiteralAddressSpace()); + if (AS == 0) + AS = 1; std::string Name; if (D.hasGlobalStorage()) Name = getMangledName(&D).str() + ".const"; From 8b7064561913914b6a696d7d758e82c5bcebd94d Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Thu, 24 Sep 2020 15:39:40 +0300 Subject: [PATCH 2/6] Add better solution Signed-off-by: Aleksander Fadeev --- clang/lib/CodeGen/CGDecl.cpp | 2 -- clang/lib/CodeGen/CodeGenModule.cpp | 2 +- 2 files changed, 1 insertion(+), 3 deletions(-) diff --git a/clang/lib/CodeGen/CGDecl.cpp b/clang/lib/CodeGen/CGDecl.cpp index 102a5a84647a..87070a3504a5 100644 --- a/clang/lib/CodeGen/CGDecl.cpp +++ b/clang/lib/CodeGen/CGDecl.cpp @@ -1144,8 +1144,6 @@ Address CodeGenModule::createUnnamedGlobalFrom(const VarDecl &D, llvm::GlobalVariable *InsertBefore = nullptr; unsigned AS = getContext().getTargetAddressSpace(getStringLiteralAddressSpace()); - if (AS == 0) - AS = 1; std::string Name; if (D.hasGlobalStorage()) Name = getMangledName(&D).str() + ".const"; diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index a8da0b910092..ec076a224ac5 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -4012,7 +4012,7 @@ LangAS CodeGenModule::getStringLiteralAddressSpace() const { // const char *getLiteral() n{ // return "AB"; // } - return LangAS::opencl_private; + return LangAS::opencl_global; if (auto AS = getTarget().getConstantAddressSpace()) return AS.getValue(); return LangAS::Default; From 878f996ce394e0c98462f52733c83ce9600e0638 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Thu, 24 Sep 2020 17:07:51 +0300 Subject: [PATCH 3/6] Change tests Signed-off-by: Aleksander Fadeev --- clang/test/CodeGenSYCL/address-space-new.cpp | 10 ++--- .../CodeGenSYCL/address-space-of-returns.cpp | 2 +- clang/test/CodeGenSYCL/unique-stable-name.cpp | 38 +++++++++---------- 3 files changed, 25 insertions(+), 25 deletions(-) diff --git a/clang/test/CodeGenSYCL/address-space-new.cpp b/clang/test/CodeGenSYCL/address-space-new.cpp index 1caf5d49dd20..788d7e0e2f13 100644 --- a/clang/test/CodeGenSYCL/address-space-new.cpp +++ b/clang/test/CodeGenSYCL/address-space-new.cpp @@ -29,7 +29,7 @@ void test() { (void)bars; // 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: @[[STR:[.a-zA-Z0-9_]+]] = private unnamed_addr addrspace(1) constant [14 x i8] c"Hello, world!\00", align 1 // CHECK: %i.ascast = addrspacecast i32* %i to i32 addrspace(4)* // CHECK: %[[ARR:[a-zA-Z0-9]+]] = alloca [42 x i32] @@ -69,7 +69,7 @@ void test() { // CHECK: %cmp{{[0-9]+}} = icmp ult i32 addrspace(4)* %[[VALAPTR]], %[[ADDPTRCAST]] const char *str = "Hello, world!"; - // 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 + // CHECK: store i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([14 x i8], [14 x i8] addrspace(1)* @[[STR]], i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)** %[[STRVAL:[a-zA-Z0-9]+]], align 8 i = str[0]; @@ -85,11 +85,11 @@ void test() { // CHECK: [[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]] ] + // CHECK-NEXT: phi i8 addrspace(4)* [ %[[VALTRUE]], %[[CONDTRUE]] ], [ addrspacecast (i8 addrspace(1)* getelementptr inbounds ([21 x i8], [21 x i8] addrspace(1)* @{{.*}}, i64 0, i64 0) to i8 addrspace(4)*), %[[CONDFALSE]] ] const char *select_null = i > 2 ? "Yet another Hello world" : nullptr; (void)select_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 + // CHECK: select i1 %{{.*}}, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([24 x i8], [24 x i8] addrspace(1)* @{{.*}}, 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; @@ -98,7 +98,7 @@ void test() { const char *select_str_trivial2 = false ? str : "Another hello world!"; (void)select_str_trivial2; - // 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)** %{{.*}} + // CHECK: store i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([21 x i8], [21 x i8] addrspace(1)* @{{.*}}, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)** %{{.*}} // // Y yy; diff --git a/clang/test/CodeGenSYCL/address-space-of-returns.cpp b/clang/test/CodeGenSYCL/address-space-of-returns.cpp index 24bd762bb28d..3b56e34bd5be 100644 --- a/clang/test/CodeGenSYCL/address-space-of-returns.cpp +++ b/clang/test/CodeGenSYCL/address-space-of-returns.cpp @@ -7,7 +7,7 @@ struct A { const char *ret_char() { return "N"; } -// CHECK: ret i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([2 x i8], [2 x i8]* @.str, i64 0, i64 0) to i8 addrspace(4)*) +// CHECK: ret i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([2 x i8], [2 x i8] addrspace(1)* @.str, i64 0, i64 0) to i8 addrspace(4)*) const char *ret_arr() { const static char Arr[36] = "Carrots, cabbage, radish, potatoes!"; diff --git a/clang/test/CodeGenSYCL/unique-stable-name.cpp b/clang/test/CodeGenSYCL/unique-stable-name.cpp index 66ca499e6cda..64b06f0a2fd7 100644 --- a/clang/test/CodeGenSYCL/unique-stable-name.cpp +++ b/clang/test/CodeGenSYCL/unique-stable-name.cpp @@ -1,12 +1,12 @@ // RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -fsycl -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s -// CHECK: @[[INT:[^\w]+]] = private unnamed_addr constant [[INT_SIZE:\[[0-9]+ x i8\]]] c"_ZTSi\00" -// CHECK: @[[LAMBDA_X:[^\w]+]] = private unnamed_addr constant [[LAMBDA_X_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE42_5clEvEUlvE46_16\00" -// CHECK: @[[MACRO_X:[^\w]+]] = private unnamed_addr constant [[MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE42_5clEvEUlvE52_7m28_18\00" -// CHECK: @[[MACRO_Y:[^\w]+]] = private unnamed_addr constant [[MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE42_5clEvEUlvE52_7m28_41\00" -// CHECK: @[[MACRO_MACRO_X:[^\w]+]] = private unnamed_addr constant [[MACRO_MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE42_5clEvEUlvE55_7m28_18m33_4\00" -// CHECK: @[[MACRO_MACRO_Y:[^\w]+]] = private unnamed_addr constant [[MACRO_MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE42_5clEvEUlvE55_7m28_41m33_4\00" -// CHECK: @[[LAMBDA_IN_DEP_INT:[^\w]+]] = private unnamed_addr constant [[DEP_INT_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ28lambda_in_dependent_functionIiEvvEUlvE23_12\00", -// CHECK: @[[LAMBDA_IN_DEP_X:[^\w]+]] = private unnamed_addr constant [[DEP_LAMBDA_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ28lambda_in_dependent_functionIZZ4mainENKUlvE42_5clEvEUlvE46_16EvvEUlvE23_12\00", +// CHECK: @[[INT:[^\w]+]] = private unnamed_addr addrspace(1) constant [[INT_SIZE:\[[0-9]+ x i8\]]] c"_ZTSi\00" +// CHECK: @[[LAMBDA_X:[^\w]+]] = private unnamed_addr addrspace(1) constant [[LAMBDA_X_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE42_5clEvEUlvE46_16\00" +// CHECK: @[[MACRO_X:[^\w]+]] = private unnamed_addr addrspace(1) constant [[MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE42_5clEvEUlvE52_7m28_18\00" +// CHECK: @[[MACRO_Y:[^\w]+]] = private unnamed_addr addrspace(1) constant [[MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE42_5clEvEUlvE52_7m28_41\00" +// CHECK: @[[MACRO_MACRO_X:[^\w]+]] = private unnamed_addr addrspace(1) constant [[MACRO_MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE42_5clEvEUlvE55_7m28_18m33_4\00" +// CHECK: @[[MACRO_MACRO_Y:[^\w]+]] = private unnamed_addr addrspace(1) constant [[MACRO_MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE42_5clEvEUlvE55_7m28_41m33_4\00" +// CHECK: @[[LAMBDA_IN_DEP_INT:[^\w]+]] = private unnamed_addr addrspace(1) constant [[DEP_INT_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ28lambda_in_dependent_functionIiEvvEUlvE23_12\00", +// CHECK: @[[LAMBDA_IN_DEP_X:[^\w]+]] = private unnamed_addr addrspace(1) constant [[DEP_LAMBDA_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ28lambda_in_dependent_functionIZZ4mainENKUlvE42_5clEvEUlvE46_16EvvEUlvE23_12\00", extern "C" void printf(const char *) {} @@ -41,36 +41,36 @@ int main() { kernel_single_task( []() { printf(__builtin_unique_stable_name(int)); - // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]]* @[[INT]] + // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]] addrspace(1)* @[[INT]] auto x = [](){}; printf(__builtin_unique_stable_name(x)); printf(__builtin_unique_stable_name(decltype(x))); - // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[LAMBDA_X_SIZE]], [[LAMBDA_X_SIZE]]* @[[LAMBDA_X]] - // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[LAMBDA_X_SIZE]], [[LAMBDA_X_SIZE]]* @[[LAMBDA_X]] + // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[LAMBDA_X_SIZE]], [[LAMBDA_X_SIZE]] addrspace(1)* @[[LAMBDA_X]] + // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[LAMBDA_X_SIZE]], [[LAMBDA_X_SIZE]] addrspace(1)* @[[LAMBDA_X]] DEF_IN_MACRO(); - // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]]* @[[MACRO_X]] - // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]]* @[[MACRO_Y]] + // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]] addrspace(1)* @[[MACRO_X]] + // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]] addrspace(1)* @[[MACRO_Y]] MACRO_CALLS_MACRO(); - // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]]* @[[MACRO_MACRO_X]] - // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]]* @[[MACRO_MACRO_Y]] + // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]] addrspace(1)* @[[MACRO_MACRO_X]] + // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]] addrspace(1)* @[[MACRO_MACRO_Y]] template_param(); // CHECK: define linkonce_odr spir_func void @_Z14template_paramIiEvv - // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]]* @[[INT]] + // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]] addrspace(1)* @[[INT]] template_param(); // CHECK: define internal spir_func void @"_Z14template_paramIZZ4mainENK3 - // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[LAMBDA_X_SIZE]], [[LAMBDA_X_SIZE]]* @[[LAMBDA_X]] + // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[LAMBDA_X_SIZE]], [[LAMBDA_X_SIZE]] addrspace(1)* @[[LAMBDA_X]] lambda_in_dependent_function(); // CHECK: define linkonce_odr spir_func void @_Z28lambda_in_dependent_functionIiEvv - // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[DEP_INT_SIZE]], [[DEP_INT_SIZE]]* @[[LAMBDA_IN_DEP_INT]] + // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[DEP_INT_SIZE]], [[DEP_INT_SIZE]] addrspace(1)* @[[LAMBDA_IN_DEP_INT]] lambda_in_dependent_function(); // CHECK: define internal spir_func void @"_Z28lambda_in_dependent_functionIZZ4mainENK3$_0clEvEUlvE_Evv - // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[DEP_LAMBDA_SIZE]], [[DEP_LAMBDA_SIZE]]* @[[LAMBDA_IN_DEP_X]] + // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[DEP_LAMBDA_SIZE]], [[DEP_LAMBDA_SIZE]] addrspace(1)* @[[LAMBDA_IN_DEP_X]] }); } From d5d72900fe71cefde519b2e5df96a656bcafc260 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Tue, 29 Sep 2020 11:13:32 +0300 Subject: [PATCH 4/6] Add test and comments Signed-off-by: Aleksander Fadeev --- clang/lib/CodeGen/CodeGenModule.cpp | 6 +++ .../CodeGenSYCL/static-var-address-space.cpp | 37 +++++++++++++++++++ 2 files changed, 43 insertions(+) create mode 100644 clang/test/CodeGenSYCL/static-var-address-space.cpp diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index ec076a224ac5..3d48f90390a0 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -4012,6 +4012,12 @@ LangAS CodeGenModule::getStringLiteralAddressSpace() const { // const char *getLiteral() n{ // return "AB"; // } + // because there is a addressspacecast to generic address space in IR, + // but adressspacecast from constant to generic forbitten because of + // constant address space is not part of generic address space. + // The private adress space doesn't suit here because a IR is translated + // in SPIRV in SYCLIsDevice mode, and all global Value shouldn't + // be private in IR for rigth translation to SPIRV. return LangAS::opencl_global; if (auto AS = getTarget().getConstantAddressSpace()) return AS.getValue(); diff --git a/clang/test/CodeGenSYCL/static-var-address-space.cpp b/clang/test/CodeGenSYCL/static-var-address-space.cpp new file mode 100644 index 000000000000..7c8ff3df1ec2 --- /dev/null +++ b/clang/test/CodeGenSYCL/static-var-address-space.cpp @@ -0,0 +1,37 @@ + +// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s + +struct C { + static int c; +}; + +template +struct D { + static T d; +}; + +template +void test() { + // CHECK: @_ZZ4testIiEvvE1a = linkonce_odr addrspace(1) constant i32 0, comdat, align 4 + static const int a = 0; + // CHECK: @_ZZ4testIiEvvE1b = linkonce_odr addrspace(1) constant i32 0, comdat, align 4 + static const T b = T(0); + // CHECK: @_ZN1C1cE = external addrspace(1) global i32, align 4 + C::c = 10; + const C struct_c; + // CHECK: @_ZN1DIiE1dE = external addrspace(1) global i32, align 4 + D::d = 11; + const D struct_d; + +} + + +template +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { + kernelFunc(); +} + +int main() { + kernel_single_task([]() { test(); }); + return 0; +} From 9b50de712bf99b33075ea6ea19bdab127db7c02b Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Tue, 29 Sep 2020 11:18:11 +0300 Subject: [PATCH 5/6] Add formatting Signed-off-by: Aleksander Fadeev --- clang/lib/CodeGen/CodeGenModule.cpp | 4 ++-- clang/test/CodeGenSYCL/static-var-address-space.cpp | 6 ++---- 2 files changed, 4 insertions(+), 6 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 3d48f90390a0..0256223a407c 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -4013,9 +4013,9 @@ LangAS CodeGenModule::getStringLiteralAddressSpace() const { // return "AB"; // } // because there is a addressspacecast to generic address space in IR, - // but adressspacecast from constant to generic forbitten because of + // but adressspacecast from constant to generic forbitten because of // constant address space is not part of generic address space. - // The private adress space doesn't suit here because a IR is translated + // The private adress space doesn't suit here because a IR is translated // in SPIRV in SYCLIsDevice mode, and all global Value shouldn't // be private in IR for rigth translation to SPIRV. return LangAS::opencl_global; diff --git a/clang/test/CodeGenSYCL/static-var-address-space.cpp b/clang/test/CodeGenSYCL/static-var-address-space.cpp index 7c8ff3df1ec2..d9c664d50875 100644 --- a/clang/test/CodeGenSYCL/static-var-address-space.cpp +++ b/clang/test/CodeGenSYCL/static-var-address-space.cpp @@ -5,12 +5,12 @@ struct C { static int c; }; -template +template struct D { static T d; }; -template +template void test() { // CHECK: @_ZZ4testIiEvvE1a = linkonce_odr addrspace(1) constant i32 0, comdat, align 4 static const int a = 0; @@ -22,10 +22,8 @@ void test() { // CHECK: @_ZN1DIiE1dE = external addrspace(1) global i32, align 4 D::d = 11; const D struct_d; - } - template __attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); From 52eeb920ec753f99048b77e6119a678b0973d34d Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Wed, 30 Sep 2020 13:38:55 +0300 Subject: [PATCH 6/6] Change comment and take single_task from sycl.hpp Signed-off-by: Aleksander Fadeev --- clang/lib/CodeGen/CodeGenModule.cpp | 9 +++------ clang/test/CodeGenSYCL/static-var-address-space.cpp | 10 ++-------- 2 files changed, 5 insertions(+), 14 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 0256223a407c..3df6120840cc 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -4012,12 +4012,9 @@ LangAS CodeGenModule::getStringLiteralAddressSpace() const { // const char *getLiteral() n{ // return "AB"; // } - // because there is a addressspacecast to generic address space in IR, - // but adressspacecast from constant to generic forbitten because of - // constant address space is not part of generic address space. - // The private adress space doesn't suit here because a IR is translated - // in SPIRV in SYCLIsDevice mode, and all global Value shouldn't - // be private in IR for rigth translation to SPIRV. + // Use global address space to avoid illegal casts from constant to generic. + // Private address space is not used here because in SPIR-V global values + // cannot have private address space. return LangAS::opencl_global; if (auto AS = getTarget().getConstantAddressSpace()) return AS.getValue(); diff --git a/clang/test/CodeGenSYCL/static-var-address-space.cpp b/clang/test/CodeGenSYCL/static-var-address-space.cpp index d9c664d50875..f5d6b7b27041 100644 --- a/clang/test/CodeGenSYCL/static-var-address-space.cpp +++ b/clang/test/CodeGenSYCL/static-var-address-space.cpp @@ -1,6 +1,5 @@ - // RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s - +#include "Inputs/sycl.hpp" struct C { static int c; }; @@ -24,12 +23,7 @@ void test() { const D struct_d; } -template -__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { - kernelFunc(); -} - int main() { - kernel_single_task([]() { test(); }); + cl::sycl::kernel_single_task([]() { test(); }); return 0; }