Skip to content

Commit a5ad7a1

Browse files
garimagubader
garimagu
authored and
bader
committed
[SYCL] Add Clang support for OpenCL image classes in SYCL mode. (#270)
Add SYCL internal classes utilizing OpenCL classes to generate LLVM bitcode compatible with OpenCL toolchain. Signed-off-by: Garima Gupta <garima.gupta@intel.com>
1 parent 85fa8a6 commit a5ad7a1

File tree

8 files changed

+479
-6
lines changed

8 files changed

+479
-6
lines changed

‎clang/lib/Sema/Sema.cpp

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -267,6 +267,15 @@ void Sema::Initialize() {
267267
if (getLangOpts().SYCLIsDevice) {
268268
addImplicitTypedef("__ocl_event_t", Context.OCLEventTy);
269269
addImplicitTypedef("__ocl_sampler_t", Context.OCLSamplerTy);
270+
#ifdef SEMA_STRINGIZE
271+
#error "Undefine SEMA_STRINGIZE macro."
272+
#endif
273+
#define SEMA_STRINGIZE(s) #s
274+
#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
275+
addImplicitTypedef(SEMA_STRINGIZE(__ocl_##ImgType##_##Suffix##_t), \
276+
Context.SingletonId);
277+
#include "clang/Basic/OpenCLImageTypes.def"
278+
#undef SEMA_STRINGIZE
270279
}
271280

272281
// Initialize predefined OpenCL types and supported extensions and (optional)

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

Lines changed: 145 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -135,11 +135,86 @@ class accessor {
135135
template <typename... T>
136136
void use(T... args) const {}
137137
_ImplT<dimensions> impl;
138+
138139
private:
139140
void __init(__global dataT *Ptr, range<dimensions> AccessRange,
140141
range<dimensions> MemRange, id<dimensions> Offset) {}
141142
};
142143

144+
template <int dimensions, access::mode accessmode, access::target accesstarget>
145+
struct opencl_image_type;
146+
147+
#define IMAGETY_DEFINE(dim, accessmode, amsuffix, Target, ifarray_) \
148+
template <> \
149+
struct opencl_image_type<dim, access::mode::accessmode, \
150+
access::target::Target> { \
151+
using type = __ocl_image##dim##d_##ifarray_##amsuffix##_t; \
152+
};
153+
154+
#define IMAGETY_READ_3_DIM_IMAGE \
155+
IMAGETY_DEFINE(1, read, ro, image, ) \
156+
IMAGETY_DEFINE(2, read, ro, image, ) \
157+
IMAGETY_DEFINE(3, read, ro, image, )
158+
159+
#define IMAGETY_WRITE_3_DIM_IMAGE \
160+
IMAGETY_DEFINE(1, write, wo, image, ) \
161+
IMAGETY_DEFINE(2, write, wo, image, ) \
162+
IMAGETY_DEFINE(3, write, wo, image, )
163+
164+
#define IMAGETY_READ_2_DIM_IARRAY \
165+
IMAGETY_DEFINE(1, read, ro, image_array, array_) \
166+
IMAGETY_DEFINE(2, read, ro, image_array, array_)
167+
168+
#define IMAGETY_WRITE_2_DIM_IARRAY \
169+
IMAGETY_DEFINE(1, write, wo, image_array, array_) \
170+
IMAGETY_DEFINE(2, write, wo, image_array, array_)
171+
172+
IMAGETY_READ_3_DIM_IMAGE
173+
IMAGETY_WRITE_3_DIM_IMAGE
174+
175+
IMAGETY_READ_2_DIM_IARRAY
176+
IMAGETY_WRITE_2_DIM_IARRAY
177+
178+
template <int dim, access::mode accessmode, access::target accesstarget>
179+
struct _ImageImplT {
180+
#ifdef __SYCL_DEVICE_ONLY__
181+
typename opencl_image_type<dim, accessmode, accesstarget>::type MImageObj;
182+
#else
183+
range<dim> AccessRange;
184+
range<dim> MemRange;
185+
id<dim> Offset;
186+
#endif
187+
};
188+
189+
template <typename dataT, int dimensions, access::mode accessmode>
190+
class accessor<dataT, dimensions, accessmode, access::target::image, access::placeholder::false_t> {
191+
public:
192+
void use(void) const {}
193+
template <typename... T>
194+
void use(T... args) {}
195+
template <typename... T>
196+
void use(T... args) const {}
197+
_ImageImplT<dimensions, accessmode, access::target::image> impl;
198+
#ifdef __SYCL_DEVICE_ONLY__
199+
void __init(typename opencl_image_type<dimensions, accessmode, access::target::image>::type ImageObj) { impl.MImageObj = ImageObj; }
200+
#endif
201+
};
202+
203+
template <typename dataT, int dimensions, access::mode accessmode>
204+
class accessor<dataT, dimensions, accessmode, access::target::host_image, access::placeholder::false_t> {
205+
public:
206+
void use(void) const {}
207+
template <typename... T>
208+
void use(T... args) {}
209+
template <typename... T>
210+
void use(T... args) const {}
211+
_ImageImplT<dimensions, accessmode, access::target::host_image> impl;
212+
};
213+
214+
// TODO: Add support for image_array accessor.
215+
// template <typename dataT, int dimensions, access::mode accessmode>
216+
//class accessor<dataT, dimensions, accessmode, access::target::image_array, access::placeholder::false_t>
217+
143218
class kernel {};
144219
class context {};
145220
class device {};
@@ -241,13 +316,81 @@ class buffer {
241316
accessor<T, dimensions, mode, access::target::host_buffer,
242317
access::placeholder::false_t>
243318
get_access() {
244-
accessor<T, dimensions, mode, access::target::host_buffer,
245-
access::placeholder::false_t>{};
319+
return accessor<T, dimensions, mode, access::target::host_buffer,
320+
access::placeholder::false_t>{};
246321
}
247322

248323
template <typename Destination>
249324
void set_final_data(Destination finalData = nullptr) {}
250325
};
251326

327+
enum class image_channel_order : unsigned int {
328+
a,
329+
r,
330+
rx,
331+
rg,
332+
rgx,
333+
ra,
334+
rgb,
335+
rgbx,
336+
rgba,
337+
argb,
338+
bgra,
339+
intensity,
340+
luminance,
341+
abgr
342+
};
343+
344+
enum class image_channel_type : unsigned int {
345+
snorm_int8,
346+
snorm_int16,
347+
unorm_int8,
348+
unorm_int16,
349+
unorm_short_565,
350+
unorm_short_555,
351+
unorm_int_101010,
352+
signed_int8,
353+
signed_int16,
354+
signed_int32,
355+
unsigned_int8,
356+
unsigned_int16,
357+
unsigned_int32,
358+
fp16,
359+
fp32
360+
};
361+
362+
template <int dimensions = 1, typename AllocatorT = int>
363+
class image {
364+
public:
365+
image(image_channel_order Order, image_channel_type Type,
366+
const range<dimensions> &Range, const property_list &PropList = {}) {}
367+
368+
/* -- common interface members -- */
369+
370+
image(const image &rhs) = default;
371+
372+
image(image &&rhs) = default;
373+
374+
image &operator=(const image &rhs) = default;
375+
376+
image &operator=(image &&rhs) = default;
377+
378+
~image() = default;
379+
380+
template <typename dataT, access::mode accessmode>
381+
accessor<dataT, dimensions, accessmode,
382+
access::target::image, access::placeholder::false_t>
383+
get_access(handler &commandGroupHandler) {
384+
return accessor<dataT, dimensions, accessmode, access::target::image, access::placeholder::false_t>{};
385+
}
386+
387+
template <typename dataT, access::mode accessmode>
388+
accessor<dataT, dimensions, accessmode,
389+
access::target::host_image, access::placeholder::false_t>
390+
get_access() {
391+
return accessor<dataT, dimensions, accessmode, access::target::host_image, access::placeholder::false_t>{};
392+
}
393+
};
394+
252395
} // namespace sycl
253396
} // namespace cl

‎clang/test/CodeGenSYCL/image_accessor.cpp

Lines changed: 111 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,111 @@
1+
// 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 %t.ll
2+
// RUN: FileCheck < %t.ll --enable-var-scope %s --check-prefix=CHECK-1DRO
3+
// RUN: FileCheck < %t.ll --enable-var-scope %s --check-prefix=CHECK-2DRO
4+
// RUN: FileCheck < %t.ll --enable-var-scope %s --check-prefix=CHECK-3DRO
5+
// RUN: FileCheck < %t.ll --enable-var-scope %s --check-prefix=CHECK-1DWO
6+
// RUN: FileCheck < %t.ll --enable-var-scope %s --check-prefix=CHECK-2DWO
7+
// RUN: FileCheck < %t.ll --enable-var-scope %s --check-prefix=CHECK-3DWO
8+
//
9+
// CHECK-1DRO: %opencl.image1d_ro_t = type opaque
10+
// CHECK-1DRO: define spir_kernel void @{{.*}}(%opencl.image1d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
11+
// CHECK-1DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image1d_ro_t addrspace(1)* %{{[0-9]+}})
12+
//
13+
// CHECK-2DRO: %opencl.image2d_ro_t = type opaque
14+
// CHECK-2DRO: define spir_kernel void @{{.*}}(%opencl.image2d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
15+
// CHECK-2DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image2d_ro_t addrspace(1)* %{{[0-9]+}})
16+
//
17+
// CHECK-3DRO: %opencl.image3d_ro_t = type opaque
18+
// CHECK-3DRO: define spir_kernel void @{{.*}}(%opencl.image3d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
19+
// CHECK-3DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image3d_ro_t addrspace(1)* %{{[0-9]+}})
20+
//
21+
// CHECK-1DWO: %opencl.image1d_wo_t = type opaque
22+
// CHECK-1DWO: define spir_kernel void @{{.*}}(%opencl.image1d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
23+
// CHECK-1DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image1d_wo_t addrspace(1)* %{{[0-9]+}})
24+
//
25+
// CHECK-2DWO: %opencl.image2d_wo_t = type opaque
26+
// CHECK-2DWO: define spir_kernel void @{{.*}}(%opencl.image2d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
27+
// CHECK-2DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image2d_wo_t addrspace(1)* %{{[0-9]+}})
28+
//
29+
// CHECK-3DWO: %opencl.image3d_wo_t = type opaque
30+
// CHECK-3DWO: define spir_kernel void @{{.*}}(%opencl.image3d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
31+
// CHECK-3DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image3d_wo_t addrspace(1)* %{{[0-9]+}})
32+
//
33+
// TODO: Add tests for the image_array opencl datatype support.
34+
#include "sycl.hpp"
35+
36+
int main() {
37+
38+
{
39+
cl::sycl::image<1> MyImage1d(cl::sycl::image_channel_order::rgbx, cl::sycl::image_channel_type::unorm_short_565, cl::sycl::range<1>(3));
40+
cl::sycl::queue Q;
41+
Q.submit([&](cl::sycl::handler &cgh) {
42+
auto Acc = MyImage1d.get_access<int, cl::sycl::access::mode::read>(cgh);
43+
44+
cgh.single_task<class image_accessor1dro>([=]() {
45+
Acc.use();
46+
});
47+
});
48+
}
49+
50+
{
51+
cl::sycl::image<2> MyImage2d(cl::sycl::image_channel_order::rgbx, cl::sycl::image_channel_type::unorm_short_565, cl::sycl::range<2>(3, 2));
52+
cl::sycl::queue Q;
53+
Q.submit([&](cl::sycl::handler &cgh) {
54+
auto Acc = MyImage2d.get_access<int, cl::sycl::access::mode::read>(cgh);
55+
56+
cgh.single_task<class image_accessor2dro>([=]() {
57+
Acc.use();
58+
});
59+
});
60+
}
61+
62+
{
63+
cl::sycl::image<3> MyImage3d(cl::sycl::image_channel_order::rgbx, cl::sycl::image_channel_type::unorm_short_565, cl::sycl::range<3>(3, 2, 4));
64+
cl::sycl::queue Q;
65+
Q.submit([&](cl::sycl::handler &cgh) {
66+
auto Acc = MyImage3d.get_access<int, cl::sycl::access::mode::read>(cgh);
67+
68+
cgh.single_task<class image_accessor3dro>([=]() {
69+
Acc.use();
70+
});
71+
});
72+
}
73+
74+
{
75+
cl::sycl::image<1> MyImage1d(cl::sycl::image_channel_order::rgbx, cl::sycl::image_channel_type::unorm_short_565, cl::sycl::range<1>(3));
76+
cl::sycl::queue Q;
77+
Q.submit([&](cl::sycl::handler &cgh) {
78+
auto Acc = MyImage1d.get_access<int, cl::sycl::access::mode::write>(cgh);
79+
80+
cgh.single_task<class image_accessor1dwo>([=]() {
81+
Acc.use();
82+
});
83+
});
84+
}
85+
86+
{
87+
cl::sycl::image<2> MyImage2d(cl::sycl::image_channel_order::rgbx, cl::sycl::image_channel_type::unorm_short_565, cl::sycl::range<2>(3, 2));
88+
cl::sycl::queue Q;
89+
Q.submit([&](cl::sycl::handler &cgh) {
90+
auto Acc = MyImage2d.get_access<int, cl::sycl::access::mode::write>(cgh);
91+
92+
cgh.single_task<class image_accessor2dwo>([=]() {
93+
Acc.use();
94+
});
95+
});
96+
}
97+
98+
{
99+
cl::sycl::image<3> MyImage3d(cl::sycl::image_channel_order::rgbx, cl::sycl::image_channel_type::unorm_short_565, cl::sycl::range<3>(3, 2, 4));
100+
cl::sycl::queue Q;
101+
Q.submit([&](cl::sycl::handler &cgh) {
102+
auto Acc = MyImage3d.get_access<int, cl::sycl::access::mode::write>(cgh);
103+
104+
cgh.single_task<class image_accessor3dwo>([=]() {
105+
Acc.use();
106+
});
107+
});
108+
}
109+
110+
return 0;
111+
}

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

Lines changed: 63 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -51,9 +51,9 @@ struct id {
5151

5252
template <int dim>
5353
struct _ImplT {
54-
range<dim> AccessRange;
55-
range<dim> MemRange;
56-
id<dim> Offset;
54+
range<dim> AccessRange;
55+
range<dim> MemRange;
56+
id<dim> Offset;
5757
};
5858

5959
template <typename dataT, access::target accessTarget>
@@ -81,7 +81,7 @@ class accessor {
8181

8282
public:
8383
void use(void) const {}
84-
void use(void*) const {}
84+
void use(void *) const {}
8585
_ImplT<dimensions> impl;
8686

8787
private:
@@ -90,6 +90,65 @@ class accessor {
9090
range<dimensions> MemRange, id<dimensions> Offset) {}
9191
};
9292

93+
template <int dimensions, access::mode accessmode, access::target accesstarget>
94+
struct opencl_image_type;
95+
96+
#define IMAGETY_DEFINE(dim, accessmode, amsuffix, Target, ifarray_) \
97+
template <> \
98+
struct opencl_image_type<dim, access::mode::accessmode, \
99+
access::target::Target> { \
100+
using type = __ocl_image##dim##d_##ifarray_##amsuffix##_t; \
101+
};
102+
103+
#define IMAGETY_READ_3_DIM_IMAGE \
104+
IMAGETY_DEFINE(1, read, ro, image, ) \
105+
IMAGETY_DEFINE(2, read, ro, image, ) \
106+
IMAGETY_DEFINE(3, read, ro, image, )
107+
108+
#define IMAGETY_WRITE_3_DIM_IMAGE \
109+
IMAGETY_DEFINE(1, write, wo, image, ) \
110+
IMAGETY_DEFINE(2, write, wo, image, ) \
111+
IMAGETY_DEFINE(3, write, wo, image, )
112+
113+
#define IMAGETY_READ_2_DIM_IARRAY \
114+
IMAGETY_DEFINE(1, read, ro, image_array, array_) \
115+
IMAGETY_DEFINE(2, read, ro, image_array, array_)
116+
117+
#define IMAGETY_WRITE_2_DIM_IARRAY \
118+
IMAGETY_DEFINE(1, write, wo, image_array, array_) \
119+
IMAGETY_DEFINE(2, write, wo, image_array, array_)
120+
121+
IMAGETY_READ_3_DIM_IMAGE
122+
IMAGETY_WRITE_3_DIM_IMAGE
123+
124+
IMAGETY_READ_2_DIM_IARRAY
125+
IMAGETY_WRITE_2_DIM_IARRAY
126+
127+
template <int dim, access::mode accessmode, access::target accesstarget>
128+
struct _ImageImplT {
129+
#ifdef __SYCL_DEVICE_ONLY__
130+
typename opencl_image_type<dim, accessmode, accesstarget>::type MImageObj;
131+
#else
132+
range<dim> AccessRange;
133+
range<dim> MemRange;
134+
id<dim> Offset;
135+
#endif
136+
};
137+
138+
template <typename dataT, int dimensions, access::mode accessmode>
139+
class accessor<dataT, dimensions, accessmode, access::target::image, access::placeholder::false_t> {
140+
public:
141+
void use(void) const {}
142+
template <typename... T>
143+
void use(T... args) {}
144+
template <typename... T>
145+
void use(T... args) const {}
146+
_ImageImplT<dimensions, accessmode, access::target::image> impl;
147+
#ifdef __SYCL_DEVICE_ONLY__
148+
void __init(typename opencl_image_type<dimensions, accessmode, access::target::image>::type ImageObj) { impl.MImageObj = ImageObj; }
149+
#endif
150+
};
151+
93152
struct sampler_impl {
94153
#ifdef __SYCL_DEVICE_ONLY__
95154
__ocl_sampler_t m_Sampler;

0 commit comments

Comments
 (0)