Skip to content

Commit 4609257

Browse files
vladimirlazvladimirlaz
committed
[SYCL] Initialize the offset of the accessor class
Signed-off-by: Vladimir Lazarev <[email protected]>
1 parent e2b3b46 commit 4609257

File tree

9 files changed

+144
-129
lines changed

9 files changed

+144
-129
lines changed

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 97 additions & 112 deletions
Large diffs are not rendered by default.

‎clang/test/CodeGenSYCL/Inputs/CL/sycl.hpp

Lines changed: 4 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -123,6 +123,7 @@ class property_list {
123123

124124
template <int dim>
125125
struct id {
126+
template<typename ...T> id(T...args) {} // fake constructor
126127
};
127128

128129
template <int dim>
@@ -137,6 +138,7 @@ struct nd_range {
137138
template <int dim>
138139
struct _ImplT {
139140
range<dim> Range;
141+
id<dim> Offset;
140142
};
141143

142144
template <typename dataT, int dimensions, access::mode accessmode,
@@ -146,9 +148,8 @@ template <typename dataT, int dimensions, access::mode accessmode,
146148

147149
public:
148150

149-
void __set_pointer(__global dataT *Ptr) { }
150-
void __set_range(range<dimensions> Range) {
151-
__impl.Range = Range;
151+
void __init(__global dataT *Ptr, range<dimensions> Range,
152+
id<dimensions> Offset) {
152153
}
153154
void use(void) const {}
154155
template <typename ...T> void use(T...args) { }

‎clang/test/CodeGenSYCL/Inputs/sycl.hpp

Lines changed: 4 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -123,6 +123,7 @@ class property_list {
123123

124124
template <int dim>
125125
struct id {
126+
template<typename ...T> id(T...args) {} // fake constructor
126127
};
127128

128129
template <int dim>
@@ -137,6 +138,7 @@ struct nd_range {
137138
template <int dim>
138139
struct _ImplT {
139140
range<dim> Range;
141+
id<dim> Offset;
140142
};
141143

142144
template <typename dataT, int dimensions, access::mode accessmode,
@@ -146,9 +148,8 @@ template <typename dataT, int dimensions, access::mode accessmode,
146148

147149
public:
148150

149-
void __set_pointer(__global dataT *Ptr) { }
150-
void __set_range(range<dimensions> Range) {
151-
__impl.Range = Range;
151+
void __init(__global dataT *Ptr, range<dimensions> Range,
152+
id<dimensions> Offset) {
152153
}
153154
void use(void) const {}
154155
template <typename ...T> void use(T...args) { }

‎clang/test/CodeGenSYCL/integration_header.cpp

Lines changed: 11 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -24,18 +24,22 @@
2424
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
2525
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 2014, 4 },
2626
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 4 },
27-
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 2016, 5 },
2827
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 5 },
28+
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 2016, 6 },
29+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 6 },
30+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 7 },
2931
// CHECK-EMPTY:
3032
// CHECK-NEXT: //--- ::second_namespace::second_kernel<char>
3133
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
3234
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 2016, 4 },
3335
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 4 },
36+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 5 },
3437
// CHECK-EMPTY:
3538
// CHECK-NEXT: //--- ::third_kernel<1, int, ::point<X> >
3639
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
3740
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 2016, 4 },
3841
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 4 },
42+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 5 },
3943
// CHECK-EMPTY:
4044
// CHECK-NEXT: };
4145
//
@@ -81,8 +85,12 @@ enum class address_space : int {
8185
struct range {
8286
};
8387

88+
struct id {
89+
};
90+
8491
struct _ImplT {
8592
range Range;
93+
id Offset;
8694
};
8795

8896
template <typename dataT, int dimensions, access::mode accessmode,
@@ -92,6 +100,8 @@ class accessor {
92100

93101
public:
94102
void use(void) const {}
103+
void __init(dataT *Ptr, range Range, id Offset) {
104+
}
95105

96106
_ImplT __impl; // compiler looks for this field
97107

‎clang/test/CodeGenSYCL/kernel-with-id.cpp

Lines changed: 8 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -38,9 +38,14 @@ template <int dim>
3838
struct range {
3939
};
4040

41+
template <int dim>
42+
struct id {
43+
};
44+
4145
template <int dim>
4246
struct _ImplT {
4347
range<dim> Range;
48+
id<dim> Offset;
4449
};
4550

4651
template <typename dataT, int dimensions, access::mode accessmode,
@@ -50,9 +55,8 @@ class accessor {
5055

5156
public:
5257

53-
void __set_pointer(__global dataT *Ptr) { }
54-
void __set_range(range<dimensions> Range) {
55-
__impl.Range = Range;
58+
void __init(__global dataT *Ptr, range<dimensions> Range,
59+
id<dimensions> Offset) {
5660
}
5761
void use(void) const {}
5862
void use(void*) const {}
@@ -69,8 +73,7 @@ __attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
6973

7074
int main() {
7175
cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write> accessorA;
72-
// CHECK: call spir_func void @{{.*}}__set_pointer{{.*}}(%"class.cl::sycl::accessor"* %{{.*}}, i32 addrspace(1)* %{{.*}})
73-
// CHECK: call spir_func void @{{.*}}__set_range{{.*}}(%"class.cl::sycl::accessor"* %{{.*}}, %"struct.cl::sycl::range"* byval align 1 %{{.*}})
76+
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.cl::sycl::accessor"* %{{.*}}, i32 addrspace(1)* %{{.*}}, %"struct.cl::sycl::range"* byval align 1 %{{.*}}, %"struct.cl::sycl::id"* byval align 1 %{{.*}})
7477
kernel<class kernel_function>(
7578
[=]() {
7679
accessorA.use();

‎clang/test/CodeGenSYCL/struct_kernel_param.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -5,6 +5,7 @@
55
// CHECK-NEXT: //--- MyKernel
66
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 2014, 0 },
77
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 8, 16 },
8+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 8, 24 },
89
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 24, 32 },
910
// CHECK-EMPTY:
1011
// CHECK-NEXT:};

‎clang/test/SemaSYCL/Inputs/sycl.hpp

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3,6 +3,10 @@
33

44
// Shared code for SYCL tests
55

6+
#ifndef __SYCL_DEVICE_ONLY__
7+
#define __global
8+
#endif
9+
610
namespace cl {
711
namespace sycl {
812
namespace access {
@@ -41,9 +45,14 @@ template <int dim>
4145
struct range {
4246
};
4347

48+
template <int dim>
49+
struct id {
50+
};
51+
4452
template <int dim>
4553
struct _ImplT {
4654
range<dim> Range;
55+
id<dim> Offset;
4756
};
4857

4958
template <typename dataT, int dimensions, access::mode accessmode,
@@ -54,6 +63,11 @@ class accessor {
5463
public:
5564
void use(void) const {}
5665
void use(void*) const {}
66+
void __init(__global dataT *Ptr, range<dimensions> Range,
67+
id<dimensions> Offset) {
68+
}
69+
70+
5771
_ImplT<dimensions> __impl;
5872
};
5973

‎clang/test/SemaSYCL/built-in-type-kernel-arg.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -21,5 +21,5 @@ int main() {
2121
});
2222
return 0;
2323
}
24-
// CHECK: kernel_function 'void (__global int *__global, range<1>)
25-
// CHECK: kernel_local_acc 'void (__local int *__local, range<1>)
24+
// CHECK: kernel_function 'void (__global int *, range<1>, id<1>)
25+
// CHECK: kernel_local_acc 'void (__local int *, range<1>, id<1>)

‎clang/test/SemaSYCL/fake-accessors.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -51,6 +51,6 @@ int main() {
5151
});
5252
return 0;
5353
}
54-
// CHECK: fake_accessors 'void (__global int *__global, range<1>, foo::cl::sycl::accessor, accessor)
55-
// CHECK: accessor_typedef 'void (__global int *__global, range<1>, foo::cl::sycl::accessor, accessor)
56-
// CHECK: accessor_alias 'void (__global int *__global, range<1>, foo::cl::sycl::accessor, accessor)
54+
// CHECK: fake_accessors 'void (__global int *, range<1>, id<1>, foo::cl::sycl::accessor, accessor)
55+
// CHECK: accessor_typedef 'void (__global int *, range<1>, id<1>, foo::cl::sycl::accessor, accessor)
56+
// CHECK: accessor_alias 'void (__global int *, range<1>, id<1>, foo::cl::sycl::accessor, accessor)

0 commit comments

Comments
 (0)