Skip to content
This repository was archived by the owner on Mar 28, 2023. It is now read-only.

Commit 8300059

Browse files
romanovvladromanovvlad
authored
[SYCL] Add kernel_bundle_api.cpp test (#293)
1 parent 4ebf2e6 commit 8300059

File tree

1 file changed

+370
-0
lines changed

1 file changed

+370
-0
lines changed

‎SYCL/Basic/kernel_bundle/kernel_bundle_api.cpp

Lines changed: 370 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,370 @@
1+
// RUN: %clangxx -fsycl -fsycl-device-code-split=per_kernel -fsycl-targets=%sycl_triple %s -o %t.out
2+
// RUN: env SYCL_PI_TRACE=2 %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
3+
//
4+
// -fsycl-device-code-split is not supported for cuda
5+
// UNSUPPORTED: cuda
6+
7+
#include <CL/sycl.hpp>
8+
9+
#include <algorithm>
10+
#include <vector>
11+
12+
class Kernel1Name;
13+
class Kernel2Name;
14+
class Kernel3Name;
15+
16+
template <class TryBodyT>
17+
void checkException(TryBodyT TryBody, const std::string &ExpectedErrMsg) {
18+
bool ExceptionThrown = false;
19+
try {
20+
TryBody();
21+
} catch (std::exception &E) {
22+
std::cerr << "Caught: " << E.what() << std::endl;
23+
std::cerr << "Expect: " << ExpectedErrMsg << std::endl;
24+
const bool CorrectException =
25+
std::string(E.what()).find(ExpectedErrMsg) != std::string::npos;
26+
assert(CorrectException && "Test failed: caught exception is incorrect.");
27+
ExceptionThrown = true;
28+
}
29+
assert(ExceptionThrown && "Expected exception is not thrown");
30+
}
31+
32+
int main() {
33+
sycl::queue Q;
34+
sycl::queue Q2;
35+
36+
// No support for host device so far.
37+
if (Q.is_host() || Q2.is_host())
38+
return 0;
39+
40+
const sycl::context Ctx = Q.get_context();
41+
const sycl::device Dev = Q.get_device();
42+
43+
const sycl::context Ctx2 = Q2.get_context();
44+
const sycl::device Dev2 = Q2.get_device();
45+
46+
// The code is needed to just have device images in the executable
47+
if (0) {
48+
Q.submit([](sycl::handler &CGH) { CGH.single_task<Kernel1Name>([]() {}); });
49+
Q.submit([](sycl::handler &CGH) { CGH.single_task<Kernel2Name>([]() {}); });
50+
}
51+
52+
sycl::kernel_id Kernel1ID = sycl::get_kernel_id<Kernel1Name>();
53+
sycl::kernel_id Kernel2ID = sycl::get_kernel_id<Kernel2Name>();
54+
55+
{
56+
sycl::kernel_bundle KernelBundle1 =
57+
sycl::get_kernel_bundle<sycl::bundle_state::input>(Ctx, {Dev});
58+
59+
sycl::kernel_bundle KernelBundleCopy = KernelBundle1;
60+
assert(KernelBundleCopy == KernelBundle1);
61+
assert(!(KernelBundleCopy != KernelBundle1));
62+
assert(false == KernelBundle1.empty());
63+
assert(Ctx.get_platform().get_backend() == KernelBundle1.get_backend());
64+
assert(KernelBundle1.get_context() == Ctx);
65+
assert(KernelBundle1.get_devices() == (std::vector<sycl::device>){Dev});
66+
assert(KernelBundle1.has_kernel(Kernel1ID));
67+
assert(KernelBundle1.has_kernel(Kernel2ID));
68+
assert(KernelBundle1.has_kernel(Kernel1ID, Dev));
69+
assert(KernelBundle1.has_kernel(Kernel2ID, Dev));
70+
71+
assert(std::any_of(
72+
KernelBundle1.begin(), KernelBundle1.end(),
73+
[&Kernel1ID](
74+
const sycl::device_image<sycl::bundle_state::input> &DevImage) {
75+
return DevImage.has_kernel(Kernel1ID);
76+
}));
77+
78+
assert(std::any_of(
79+
KernelBundle1.begin(), KernelBundle1.end(),
80+
[&Kernel2ID](
81+
const sycl::device_image<sycl::bundle_state::input> &DevImage) {
82+
return DevImage.has_kernel(Kernel2ID);
83+
}));
84+
85+
assert(std::any_of(
86+
KernelBundle1.begin(), KernelBundle1.end(),
87+
[&Kernel1ID,
88+
&Dev](const sycl::device_image<sycl::bundle_state::input> &DevImage) {
89+
return DevImage.has_kernel(Kernel1ID, Dev);
90+
}));
91+
92+
assert(std::any_of(
93+
KernelBundle1.begin(), KernelBundle1.end(),
94+
[&Kernel2ID,
95+
&Dev](const sycl::device_image<sycl::bundle_state::input> &DevImage) {
96+
return DevImage.has_kernel(Kernel2ID, Dev);
97+
}));
98+
99+
assert(sycl::has_kernel_bundle<sycl::bundle_state::input>(Ctx, {Dev}));
100+
assert(sycl::has_kernel_bundle<sycl::bundle_state::input>(Ctx, {Dev},
101+
{Kernel2ID}));
102+
}
103+
104+
// The following check relies on "-fsycl-device-code-split=per_kernel" option,
105+
// so it is expected that each kernel in a separate device image.
106+
// Verify that get_kernel_bundle filters out device images based on vector
107+
// of kernel_id's and Selector.
108+
109+
{
110+
// Test get_kernel_bundle with filters, join and get_kernel_ids API.
111+
sycl::kernel_bundle KernelBundleInput1 =
112+
sycl::get_kernel_bundle<sycl::bundle_state::input>(Ctx, {Dev},
113+
{Kernel1ID});
114+
assert(KernelBundleInput1.has_kernel(Kernel1ID));
115+
assert(!KernelBundleInput1.has_kernel(Kernel2ID));
116+
117+
auto Selector =
118+
[&Kernel2ID](
119+
const sycl::device_image<sycl::bundle_state::input> &DevImage) {
120+
return DevImage.has_kernel(Kernel2ID);
121+
};
122+
123+
sycl::kernel_bundle KernelBundleInput2 =
124+
sycl::get_kernel_bundle<sycl::bundle_state::input>(Ctx, {Dev},
125+
Selector);
126+
assert(!KernelBundleInput2.has_kernel(Kernel1ID));
127+
assert(KernelBundleInput2.has_kernel(Kernel2ID));
128+
129+
sycl::kernel_bundle KernelBundleJoint =
130+
sycl::join(std::vector<sycl::kernel_bundle<sycl::bundle_state::input>>{
131+
KernelBundleInput1, KernelBundleInput2});
132+
133+
assert(KernelBundleJoint.has_kernel(Kernel1ID));
134+
assert(KernelBundleJoint.has_kernel(Kernel2ID));
135+
136+
std::vector<sycl::kernel_id> KernelIDs = KernelBundleJoint.get_kernel_ids();
137+
138+
assert(KernelIDs.size() == 2);
139+
}
140+
141+
{
142+
// Test compile, link, build
143+
sycl::kernel_bundle KernelBundleInput1 =
144+
sycl::get_kernel_bundle<sycl::bundle_state::input>(Ctx, {Dev},
145+
{Kernel1ID});
146+
147+
sycl::kernel_bundle KernelBundleInput2 =
148+
sycl::get_kernel_bundle<sycl::bundle_state::input>(Ctx, {Dev},
149+
{Kernel2ID});
150+
151+
sycl::kernel_bundle<sycl::bundle_state::object> KernelBundleObject1 =
152+
sycl::compile(KernelBundleInput1, KernelBundleInput1.get_devices());
153+
// CHECK:---> piProgramCreate
154+
// CHECK-NEXT: <unknown> : {{.*}}
155+
// CHECK-NEXT: <unknown> : {{.*}}
156+
// CHECK-NEXT: <unknown> : {{.*}}
157+
// CHECK-NEXT: <unknown> : {{.*}}
158+
// CHECK-NEXT: ) ---> pi_result : PI_SUCCESS
159+
// CHECK-NEXT: [out]<unknown> ** : {{.*}}[ [[PROGRAM_HANDLE1:[0-9a-fA-Fx]]]
160+
//
161+
// CHECK:---> piProgramCompile(
162+
// CHECK-Next: <unknown> : [[PROGRAM_HANDLE1]]
163+
164+
sycl::kernel_bundle<sycl::bundle_state::object> KernelBundleObject2 =
165+
sycl::compile(KernelBundleInput2, KernelBundleInput2.get_devices());
166+
// CHECK:---> piProgramCreate
167+
// CHECK-NEXT: <unknown> : {{.*}}
168+
// CHECK-NEXT: <unknown> : {{.*}}
169+
// CHECK-NEXT: <unknown> : {{.*}}
170+
// CHECK-NEXT: <unknown> : {{.*}}
171+
// CHECK-NEXT: ) ---> pi_result : PI_SUCCESS
172+
// CHECK-NEXT: [out]<unknown> ** : {{.*}}[ [[PROGRAM_HANDLE2:[0-9a-fA-Fx]]]
173+
//
174+
// CHECK:---> piProgramCompile(
175+
// CHECK-Next: <unknown> : [[PROGRAM_HANDLE2]]
176+
177+
sycl::kernel_bundle<sycl::bundle_state::executable> KernelBundleExecutable =
178+
sycl::link({KernelBundleObject1, KernelBundleObject2},
179+
KernelBundleObject1.get_devices());
180+
// CHECK:---> piProgramLink(
181+
// CHECK-NEXT: <unknown> : {{.*}}
182+
// CHECK-NEXT: <unknown> : {{.*}}
183+
// CHECK-NEXT: <unknown> : {{.*}}
184+
// CHECK-NEXT: <nullptr>
185+
// CHECK-NEXT: <unknown> : {{.*}}
186+
// CHECK-NEXT: <unknown> : {{.*}}
187+
// CHECK-NEXT: <nullptr>
188+
// CHECK-NEXT: <nullptr>
189+
// CHECK-NEXT: <unknown> : {{.*}}
190+
// CHECK-NEXT:---> pi_result : PI_SUCCESS
191+
// CHECK-NEXT: [out]<unknown> ** : {{.*}}
192+
// PI tracing doesn't allow checking for all input programs so far.
193+
194+
assert(KernelBundleExecutable.has_kernel(Kernel1ID));
195+
assert(KernelBundleExecutable.has_kernel(Kernel2ID));
196+
197+
sycl::kernel_bundle<sycl::bundle_state::executable>
198+
KernelBundleExecutable2 =
199+
sycl::build(KernelBundleInput1, KernelBundleInput1.get_devices());
200+
201+
// CHECK:---> piProgramCreate
202+
// CHECK-NEXT: <unknown> : {{.*}}
203+
// CHECK-NEXT: <unknown> : {{.*}}
204+
// CHECK-NEXT: <unknown> : {{.*}}
205+
// CHECK-NEXT: <unknown> : {{.*}}
206+
// CHECK-NEXT: ) ---> pi_result : PI_SUCCESS
207+
// CHECK-NEXT: [out]<unknown> ** : {{.*}}[ [[PROGRAM_HANDLE3:[0-9a-fA-Fx]]]
208+
//
209+
// CHECK:---> piProgramBuild(
210+
// CHECK-NEXT: <unknown> : [[PROGRAM_HANDLE3]]
211+
//
212+
// CHECK:---> piProgramRetain(
213+
// CHECK-NEXT: <unknown> : [[PROGRAM_HANDLE3]]
214+
// CHECK-NEXT:---> pi_result : PI_SUCCESS
215+
216+
// Version of link which finds intersection of associated devices between
217+
// input bundles
218+
sycl::kernel_bundle<sycl::bundle_state::executable>
219+
KernelBundleExecutable3 =
220+
sycl::link({KernelBundleObject1, KernelBundleObject2});
221+
}
222+
223+
{
224+
// Test handle::use_kernel_bundle APIs.
225+
sycl::kernel_id Kernel3ID = sycl::get_kernel_id<Kernel3Name>();
226+
227+
sycl::kernel_bundle KernelBundleExecutable =
228+
sycl::get_kernel_bundle<sycl::bundle_state::executable>(Ctx, {Dev},
229+
{Kernel3ID});
230+
// 3 SPIRV images - 3 calls to piextDeviceSelectBinary are expected
231+
// CHECK:---> piextDeviceSelectBinary
232+
// CHECK:---> piextDeviceSelectBinary
233+
// CHECK:---> piextDeviceSelectBinary
234+
// CHECK:---> piProgramCreate
235+
// CHECK-NEXT: <unknown> : {{.*}}
236+
// CHECK-NEXT: <unknown> : {{.*}}
237+
// CHECK-NEXT: <unknown> : {{.*}}
238+
// CHECK-NEXT: <unknown> : {{.*}}
239+
// CHECK-NEXT: ) ---> pi_result : PI_SUCCESS
240+
// CHECK-NEXT: [out]<unknown> ** : {{.*}}[ [[PROGRAM_HANDLE4:[0-9a-fA-Fx]]]
241+
//
242+
// CHECK:---> piProgramBuild(
243+
// CHECK-NEXT: <unknown> : [[PROGRAM_HANDLE4]]
244+
//
245+
// CHECK:---> piProgramRetain(
246+
// CHECK-NEXT: <unknown> : [[PROGRAM_HANDLE4]]
247+
// CHECK-NEXT:---> pi_result : PI_SUCCESS
248+
//
249+
// CHECK:---> piKernelCreate(
250+
// CHECK-NEXT: <unknown> : [[PROGRAM_HANDLE4]]
251+
// CHECK-NEXT:<const char *>: _ZTS11Kernel3Name
252+
// CHECK-NEXT: <unknown> : {{.*}}
253+
// CHECK-NEXT: ---> pi_result : PI_SUCCESS
254+
// CHECK-NEXT: [out]<unknown> ** : {{.*}}[ [[KERNEL_HANDLE:[0-9a-fA-Fx]]]
255+
//
256+
// CHECK:---> piKernelRetain(
257+
// CHECK-NEXT: <unknown> : [[KERNEL_HANDLE]]
258+
// CHECK-NEXT:---> pi_result : PI_SUCCESS
259+
//
260+
// CHECK:---> piEnqueueKernelLaunch(
261+
// CHECK-NEXT:<unknown> : {{.*}}
262+
// CHECK-NEXT:<unknown> : [[KERNEL_HANDLE]]
263+
//
264+
// CHECK:---> piKernelRelease(
265+
// CHECK-NEXT: <unknown> : [[KERNEL_HANDLE]]
266+
// CHECK-NEXT:---> pi_result : PI_SUCCESS
267+
268+
cl::sycl::buffer<int, 1> Buf(sycl::range<1>{1});
269+
270+
Q.submit([&](sycl::handler &CGH) {
271+
auto Acc = Buf.get_access<sycl::access::mode::write>(CGH);
272+
CGH.use_kernel_bundle(KernelBundleExecutable);
273+
CGH.single_task<Kernel3Name>([=]() { Acc[0] = 42; });
274+
});
275+
276+
{
277+
auto HostAcc = Buf.get_access<sycl::access::mode::write>();
278+
assert(HostAcc[0] == 42);
279+
}
280+
}
281+
282+
{
283+
// Error handling
284+
285+
std::cerr << "Empty list of devices for get_kernel_bundle" << std::endl;
286+
checkException(
287+
[&]() {
288+
sycl::get_kernel_bundle<sycl::bundle_state::input>(
289+
Ctx, std::vector<sycl::device>{});
290+
},
291+
"Not all devices are associated with the context or vector of devices "
292+
"is empty");
293+
294+
std::cerr << "Empty list of devices for compile" << std::endl;
295+
checkException(
296+
[&]() {
297+
sycl::kernel_bundle KernelBundleInput =
298+
sycl::get_kernel_bundle<sycl::bundle_state::input>(Ctx, {Dev});
299+
sycl::compile(KernelBundleInput, std::vector<sycl::device>{});
300+
},
301+
"Not all devices are in the set of associated "
302+
"devices for input bundle or vector of devices is empty");
303+
304+
std::cerr << "Mismatched contexts for link" << std::endl;
305+
checkException(
306+
[&]() {
307+
sycl::kernel_bundle KernelBundleObject1 =
308+
sycl::get_kernel_bundle<sycl::bundle_state::object>(Ctx, {Dev});
309+
310+
sycl::kernel_bundle KernelBundleObject2 =
311+
sycl::get_kernel_bundle<sycl::bundle_state::object>(Ctx2, {Dev2});
312+
313+
sycl::link({KernelBundleObject1, KernelBundleObject2}, {Dev2});
314+
},
315+
"Not all input bundles have the same associated context");
316+
317+
std::cerr << "Empty device list for link" << std::endl;
318+
checkException(
319+
[&]() {
320+
sycl::kernel_bundle KernelBundleObject1 =
321+
sycl::get_kernel_bundle<sycl::bundle_state::object>(Ctx, {Dev});
322+
323+
sycl::kernel_bundle KernelBundleObject2 =
324+
sycl::get_kernel_bundle<sycl::bundle_state::object>(Ctx, {Dev});
325+
326+
sycl::link({KernelBundleObject1, KernelBundleObject2},
327+
std::vector<sycl::device>{});
328+
},
329+
"Not all devices are in the set of associated devices for input "
330+
"bundles or vector of devices is empty");
331+
332+
std::cerr << "Mismatched contexts for join" << std::endl;
333+
checkException(
334+
[&]() {
335+
sycl::kernel_bundle KernelBundleObject1 =
336+
sycl::get_kernel_bundle<sycl::bundle_state::object>(Ctx);
337+
338+
sycl::kernel_bundle KernelBundleObject2 =
339+
sycl::get_kernel_bundle<sycl::bundle_state::object>(Ctx2);
340+
341+
sycl::join(
342+
std::vector<sycl::kernel_bundle<sycl::bundle_state::object>>{
343+
KernelBundleObject1, KernelBundleObject2});
344+
},
345+
"Not all input bundles have the same associated context");
346+
347+
std::cerr << "Not found kernel" << std::endl;
348+
checkException(
349+
[&]() {
350+
sycl::kernel_id Kernel3ID = sycl::get_kernel_id<Kernel3Name>();
351+
sycl::kernel_bundle KernelBundleExecutable =
352+
sycl::get_kernel_bundle<sycl::bundle_state::executable>(
353+
Ctx, {Dev}, {Kernel3ID});
354+
355+
KernelBundleExecutable.get_kernel(Kernel1ID);
356+
},
357+
"The kernel bundle does not contain the kernel identified by kernelId");
358+
359+
std::cerr << "Empty devices for has_kernel_bundle" << std::endl;
360+
checkException(
361+
[&]() {
362+
sycl::has_kernel_bundle<sycl::bundle_state::input>(
363+
Ctx, std::vector<sycl::device>{});
364+
},
365+
"Not all devices are associated with the context or vector of devices "
366+
"is empty");
367+
}
368+
369+
return 0;
370+
}

0 commit comments

Comments
 (0)