Skip to content

Commit 10e58e4

Browse files
author
iclsrc
committed
Merge from 'sycl' to 'sycl-web' (#4)
2 parents 29838d4 + bb68eef commit 10e58e4

File tree

10 files changed

+373
-159
lines changed

10 files changed

+373
-159
lines changed

‎clang/lib/CodeGen/CodeGenModule.cpp

+4-1
Original file line numberDiff line numberDiff line change
@@ -4034,7 +4034,10 @@ LangAS CodeGenModule::getStringLiteralAddressSpace() const {
40344034
// const char *getLiteral() n{
40354035
// return "AB";
40364036
// }
4037-
return LangAS::opencl_private;
4037+
// Use global address space to avoid illegal casts from constant to generic.
4038+
// Private address space is not used here because in SPIR-V global values
4039+
// cannot have private address space.
4040+
return LangAS::opencl_global;
40384041
if (auto AS = getTarget().getConstantAddressSpace())
40394042
return AS.getValue();
40404043
return LangAS::Default;

‎clang/test/CodeGenSYCL/address-space-new.cpp

+5-5
Original file line numberDiff line numberDiff line change
@@ -29,7 +29,7 @@ void test() {
2929
(void)bars;
3030
// 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
3131

32-
// CHECK: @[[STR:[.a-zA-Z0-9_]+]] = private unnamed_addr constant [14 x i8] c"Hello, world!\00", align 1
32+
// CHECK: @[[STR:[.a-zA-Z0-9_]+]] = private unnamed_addr addrspace(1) constant [14 x i8] c"Hello, world!\00", align 1
3333

3434
// CHECK: %i.ascast = addrspacecast i32* %i to i32 addrspace(4)*
3535
// CHECK: %[[ARR:[a-zA-Z0-9]+]] = alloca [42 x i32]
@@ -69,7 +69,7 @@ void test() {
6969
// CHECK: %cmp{{[0-9]+}} = icmp ult i32 addrspace(4)* %[[VALAPTR]], %[[ADDPTRCAST]]
7070

7171
const char *str = "Hello, world!";
72-
// 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
72+
// 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
7373

7474
i = str[0];
7575

@@ -85,11 +85,11 @@ void test() {
8585
// CHECK: [[CONDFALSE]]:
8686

8787
// CHECK: [[CONDEND]]:
88-
// 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]] ]
88+
// 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]] ]
8989

9090
const char *select_null = i > 2 ? "Yet another Hello world" : nullptr;
9191
(void)select_null;
92-
// 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
92+
// 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
9393

9494
const char *select_str_trivial1 = true ? str : "Another hello world!";
9595
(void)select_str_trivial1;
@@ -98,7 +98,7 @@ void test() {
9898

9999
const char *select_str_trivial2 = false ? str : "Another hello world!";
100100
(void)select_str_trivial2;
101-
// 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)** %{{.*}}
101+
// 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)** %{{.*}}
102102
//
103103
//
104104
Y yy;

‎clang/test/CodeGenSYCL/address-space-of-returns.cpp

+1-1
Original file line numberDiff line numberDiff line change
@@ -7,7 +7,7 @@ struct A {
77
const char *ret_char() {
88
return "N";
99
}
10-
// CHECK: ret i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([2 x i8], [2 x i8]* @.str, i64 0, i64 0) to i8 addrspace(4)*)
10+
// 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)*)
1111

1212
const char *ret_arr() {
1313
const static char Arr[36] = "Carrots, cabbage, radish, potatoes!";

‎clang/test/CodeGenSYCL/static-var-address-space.cpp

+29
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,29 @@
1+
// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
2+
#include "Inputs/sycl.hpp"
3+
struct C {
4+
static int c;
5+
};
6+
7+
template <typename T>
8+
struct D {
9+
static T d;
10+
};
11+
12+
template <typename T>
13+
void test() {
14+
// CHECK: @_ZZ4testIiEvvE1a = linkonce_odr addrspace(1) constant i32 0, comdat, align 4
15+
static const int a = 0;
16+
// CHECK: @_ZZ4testIiEvvE1b = linkonce_odr addrspace(1) constant i32 0, comdat, align 4
17+
static const T b = T(0);
18+
// CHECK: @_ZN1C1cE = external addrspace(1) global i32, align 4
19+
C::c = 10;
20+
const C struct_c;
21+
// CHECK: @_ZN1DIiE1dE = external addrspace(1) global i32, align 4
22+
D<int>::d = 11;
23+
const D<int> struct_d;
24+
}
25+
26+
int main() {
27+
cl::sycl::kernel_single_task<class fake_kernel>([]() { test<int>(); });
28+
return 0;
29+
}

‎clang/test/CodeGenSYCL/unique-stable-name.cpp

+19-19
Original file line numberDiff line numberDiff line change
@@ -1,12 +1,12 @@
11
// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -fsycl -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
2-
// CHECK: @[[INT:[^\w]+]] = private unnamed_addr constant [[INT_SIZE:\[[0-9]+ x i8\]]] c"_ZTSi\00"
3-
// CHECK: @[[LAMBDA_X:[^\w]+]] = private unnamed_addr constant [[LAMBDA_X_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE42_5clEvEUlvE46_16\00"
4-
// CHECK: @[[MACRO_X:[^\w]+]] = private unnamed_addr constant [[MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE42_5clEvEUlvE52_7m28_18\00"
5-
// CHECK: @[[MACRO_Y:[^\w]+]] = private unnamed_addr constant [[MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE42_5clEvEUlvE52_7m28_41\00"
6-
// CHECK: @[[MACRO_MACRO_X:[^\w]+]] = private unnamed_addr constant [[MACRO_MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE42_5clEvEUlvE55_7m28_18m33_4\00"
7-
// CHECK: @[[MACRO_MACRO_Y:[^\w]+]] = private unnamed_addr constant [[MACRO_MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE42_5clEvEUlvE55_7m28_41m33_4\00"
8-
// CHECK: @[[LAMBDA_IN_DEP_INT:[^\w]+]] = private unnamed_addr constant [[DEP_INT_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ28lambda_in_dependent_functionIiEvvEUlvE23_12\00",
9-
// 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",
2+
// CHECK: @[[INT:[^\w]+]] = private unnamed_addr addrspace(1) constant [[INT_SIZE:\[[0-9]+ x i8\]]] c"_ZTSi\00"
3+
// CHECK: @[[LAMBDA_X:[^\w]+]] = private unnamed_addr addrspace(1) constant [[LAMBDA_X_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE42_5clEvEUlvE46_16\00"
4+
// CHECK: @[[MACRO_X:[^\w]+]] = private unnamed_addr addrspace(1) constant [[MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE42_5clEvEUlvE52_7m28_18\00"
5+
// CHECK: @[[MACRO_Y:[^\w]+]] = private unnamed_addr addrspace(1) constant [[MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE42_5clEvEUlvE52_7m28_41\00"
6+
// 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"
7+
// CHECK: @[[MACRO_MACRO_Y:[^\w]+]] = private unnamed_addr addrspace(1) constant [[MACRO_MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE42_5clEvEUlvE55_7m28_41m33_4\00"
8+
// 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",
9+
// 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",
1010

1111
extern "C" void printf(const char *) {}
1212

@@ -41,36 +41,36 @@ int main() {
4141
kernel_single_task<class kernel>(
4242
[]() {
4343
printf(__builtin_unique_stable_name(int));
44-
// CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]]* @[[INT]]
44+
// CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]] addrspace(1)* @[[INT]]
4545

4646
auto x = [](){};
4747
printf(__builtin_unique_stable_name(x));
4848
printf(__builtin_unique_stable_name(decltype(x)));
49-
// CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[LAMBDA_X_SIZE]], [[LAMBDA_X_SIZE]]* @[[LAMBDA_X]]
50-
// CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[LAMBDA_X_SIZE]], [[LAMBDA_X_SIZE]]* @[[LAMBDA_X]]
49+
// 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]]
50+
// 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]]
5151

5252
DEF_IN_MACRO();
53-
// CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]]* @[[MACRO_X]]
54-
// CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]]* @[[MACRO_Y]]
53+
// CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]] addrspace(1)* @[[MACRO_X]]
54+
// CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]] addrspace(1)* @[[MACRO_Y]]
5555
MACRO_CALLS_MACRO();
56-
// CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]]* @[[MACRO_MACRO_X]]
57-
// CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]]* @[[MACRO_MACRO_Y]]
56+
// 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]]
57+
// 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]]
5858

5959
template_param<int>();
6060
// CHECK: define linkonce_odr spir_func void @_Z14template_paramIiEvv
61-
// CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]]* @[[INT]]
61+
// CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]] addrspace(1)* @[[INT]]
6262

6363
template_param<decltype(x)>();
6464
// CHECK: define internal spir_func void @"_Z14template_paramIZZ4mainENK3
65-
// CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[LAMBDA_X_SIZE]], [[LAMBDA_X_SIZE]]* @[[LAMBDA_X]]
65+
// 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]]
6666

6767
lambda_in_dependent_function<int>();
6868
// CHECK: define linkonce_odr spir_func void @_Z28lambda_in_dependent_functionIiEvv
69-
// CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[DEP_INT_SIZE]], [[DEP_INT_SIZE]]* @[[LAMBDA_IN_DEP_INT]]
69+
// 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]]
7070

7171
lambda_in_dependent_function<decltype(x)>();
7272
// CHECK: define internal spir_func void @"_Z28lambda_in_dependent_functionIZZ4mainENK3$_0clEvEUlvE_Evv
73-
// CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[DEP_LAMBDA_SIZE]], [[DEP_LAMBDA_SIZE]]* @[[LAMBDA_IN_DEP_X]]
73+
// 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]]
7474

7575
});
7676
}

‎sycl/include/CL/sycl/INTEL/esimd/detail/esimd_util.hpp

+2-2
Original file line numberDiff line numberDiff line change
@@ -77,7 +77,7 @@ constexpr unsigned int ElemsPerAddrDecoding(unsigned int ElemsPerAddrEncoded) {
7777
return (1 << ElemsPerAddrEncoded);
7878
}
7979

80-
namespace details {
80+
namespace detail {
8181

8282
/// type traits
8383
template <typename T> struct is_esimd_vector {
@@ -236,7 +236,7 @@ template <> struct word_type<int> { using type = short; };
236236
template <> struct word_type<uchar> { using type = ushort; };
237237
template <> struct word_type<uint> { using type = ushort; };
238238

239-
} // namespace details
239+
} // namespace detail
240240
} // namespace gpu
241241
} // namespace INTEL
242242
} // namespace sycl

0 commit comments

Comments
 (0)