Skip to content

Commit ff52284

Browse files
authored
[SYCL][CUDA] Test case covers unnamed lambda and hierarchical kernels (intel#260)
Signed-off-by: JackAKirk <jack.kirk@codeplay.com>
1 parent bbd0476 commit ff52284

File tree

1 file changed

+36
-4
lines changed

1 file changed

+36
-4
lines changed

SYCL/InorderQueue/in_order_kernels.cpp

Lines changed: 36 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,13 +1,11 @@
1-
// UNSUPPORTED: cuda
2-
// CUDA does not support unnamed lambdas.
31
//
42
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsycl-unnamed-lambda %s -o %t.out
53
// RUN: %HOST_RUN_PLACEHOLDER %t.out
64
// RUN: %ACC_RUN_PLACEHOLDER %t.out
75
// RUN: %CPU_RUN_PLACEHOLDER %t.out
86
// RUN: %GPU_RUN_PLACEHOLDER %t.out
97

10-
//==------ oq_kernels.cpp - SYCL ordered queue kernel shortcut test --------==//
8+
// SYCL ordered queue kernel shortcut test
119
//
1210
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
1311
// See https://llvm.org/LICENSE.txt for license information.
@@ -42,28 +40,62 @@ int main() {
4240
A[i]++;
4341
});
4442

43+
q.single_task([=]() {
44+
for (int i = 0; i < N; i++) {
45+
A[i]++;
46+
}
47+
});
48+
4549
q.single_task<class Bar>([=]() {
4650
for (int i = 0; i < N; i++) {
4751
A[i]++;
4852
}
4953
});
5054

5155
id<1> offset(0);
56+
q.parallel_for(range<1>{N}, offset, [=](id<1> ID) {
57+
auto i = ID[0];
58+
A[i]++;
59+
});
60+
5261
q.parallel_for<class Baz>(range<1>{N}, offset, [=](id<1> ID) {
5362
auto i = ID[0];
5463
A[i]++;
5564
});
5665

5766
nd_range<1> NDR(range<1>{N}, range<1>{2});
67+
q.parallel_for(NDR, [=](nd_item<1> Item) {
68+
auto i = Item.get_global_id(0);
69+
A[i]++;
70+
});
71+
5872
q.parallel_for<class NDFoo>(NDR, [=](nd_item<1> Item) {
5973
auto i = Item.get_global_id(0);
6074
A[i]++;
6175
});
6276

77+
q.submit([&](handler &cgh) {
78+
cgh.parallel_for_work_group<class WkGrp>(
79+
range<1>{N / 2}, range<1>{2}, [=](group<1> myGroup) {
80+
auto j = myGroup.get_id(0);
81+
myGroup.parallel_for_work_item(
82+
[&](h_item<1> it) { A[(j * 2) + it.get_local_id(0)]++; });
83+
});
84+
});
85+
86+
q.submit([&](handler &cgh) {
87+
cgh.parallel_for_work_group(
88+
range<1>{N / 2}, range<1>{2}, [=](group<1> myGroup) {
89+
auto j = myGroup.get_id(0);
90+
myGroup.parallel_for_work_item(
91+
[&](h_item<1> it) { A[(j * 2) + it.get_local_id(0)]++; });
92+
});
93+
});
94+
6395
q.wait();
6496

6597
for (int i = 0; i < N; i++) {
66-
if (A[i] != 6)
98+
if (A[i] != 11)
6799
return 1;
68100
}
69101
}

0 commit comments

Comments
 (0)