Skip to content

Commit c39327a

Browse files
author
Bjoern Knafla
authored
[SYCL][CUDA] Move interop tests (#1570)
Move LIT test code for buffers that uses OpenCL C online compilation to related interop test. Signed-off-by: Bjoern Knafla <[email protected]>
1 parent 0052d08 commit c39327a

File tree

2 files changed

+94
-82
lines changed

2 files changed

+94
-82
lines changed

‎sycl/test/basic_tests/buffer/buffer.cpp‎

Lines changed: 4 additions & 72 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,3 @@
1-
// REQUIRES: opencl
2-
31
// RUN: %clangxx %s -o %t1.out -lsycl -I %sycl_include
42
// RUN: env SYCL_DEVICE_TYPE=HOST %t1.out
53
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t2.out
@@ -8,9 +6,6 @@
86
// RUN: %GPU_RUN_PLACEHOLDER %t2.out
97
// RUN: %ACC_RUN_PLACEHOLDER %t2.out
108

11-
// TODO: Unexpected result and following assertion
12-
// XFAIL: cuda
13-
149
//==------------------- buffer.cpp - SYCL buffer basic test ----------------==//
1510
//
1611
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
@@ -397,8 +392,8 @@ int main() {
397392

398393
myQueue.submit([&](handler &cgh) {
399394
accessor<int, 2, access::mode::write, access::target::global_buffer,
400-
access::placeholder::false_t>
401-
B(Buffer, cgh, range<2>(20,20), id<2>(10,10));
395+
access::placeholder::false_t>
396+
B(Buffer, cgh, range<2>(20, 20), id<2>(10, 10));
402397
cgh.parallel_for<class bufferByRangeOffset>(
403398
range<2>{10, 5}, [=](id<2> index) { B[index] = 1; });
404399
});
@@ -494,9 +489,8 @@ int main() {
494489
myQueue.submit([&](handler &cgh) {
495490
auto B = b.get_access<access::mode::read_write>(cgh);
496491
cgh.parallel_for<class wb>(range<1>{10},
497-
[=](id<1> index) { B[index] = 0; });
492+
[=](id<1> index) { B[index] = 0; });
498493
});
499-
500494
}
501495
// Data is copied back because there is a user side ptr and write-back is
502496
// enabled
@@ -519,9 +513,8 @@ int main() {
519513
myQueue.submit([&](handler &cgh) {
520514
auto B = b.get_access<access::mode::read_write>(cgh);
521515
cgh.parallel_for<class notwb>(range<1>{10},
522-
[=](id<1> index) { B[index] = 0; });
516+
[=](id<1> index) { B[index] = 0; });
523517
});
524-
525518
}
526519
// Data is not copied back because write-back is canceled
527520
for (int i = 0; i < 10; i++)
@@ -551,33 +544,6 @@ int main() {
551544
assert(data1[i] == 0);
552545
}
553546

554-
{
555-
queue myQueue;
556-
if (!myQueue.is_host()) {
557-
std::vector<int> data1(10, -1);
558-
std::vector<int> data2(10, -2);
559-
{
560-
buffer<int, 1> a(data1.data(), range<1>(10));
561-
buffer<int, 1> b(data2);
562-
563-
program prog(myQueue.get_context());
564-
prog.build_with_source("kernel void override_source(global int* Acc) "
565-
"{Acc[get_global_id(0)] = 0; }\n");
566-
cl::sycl::kernel krn = prog.get_kernel("override_source");
567-
myQueue.submit([&](handler &cgh) {
568-
auto A = a.get_access<access::mode::read_write>(cgh);
569-
cgh.set_arg(0, A);
570-
auto B = b.get_access<access::mode::read_write>(cgh);
571-
cgh.parallel_for(cl::sycl::range<1>(10), krn);
572-
});
573-
} // Data is copied back
574-
for (int i = 0; i < 10; i++)
575-
assert(data2[i] == -2);
576-
for (int i = 0; i < 10; i++)
577-
assert(data1[i] == 0);
578-
}
579-
}
580-
581547
{
582548
std::vector<int> data1(10, -1);
583549
std::vector<int> data2(10, -2);
@@ -604,40 +570,6 @@ int main() {
604570
assert(data1[i] == 0);
605571
}
606572

607-
{
608-
queue myQueue;
609-
if (!myQueue.is_host()) {
610-
std::vector<int> data1(10, -1);
611-
std::vector<int> data2(10, -2);
612-
{
613-
buffer<int, 1> a(data1.data(), range<1>(10));
614-
buffer<int, 1> b(data2);
615-
accessor<int, 1, access::mode::read_write,
616-
access::target::global_buffer, access::placeholder::true_t>
617-
A(a);
618-
accessor<int, 1, access::mode::read_write,
619-
access::target::global_buffer, access::placeholder::true_t>
620-
B(b);
621-
622-
program prog(myQueue.get_context());
623-
prog.build_with_source("kernel void override_source_placeholder(global "
624-
"int* Acc) {Acc[get_global_id(0)] = 0; }\n");
625-
cl::sycl::kernel krn = prog.get_kernel("override_source_placeholder");
626-
627-
myQueue.submit([&](handler &cgh) {
628-
cgh.require(A);
629-
cgh.set_arg(0, A);
630-
cgh.require(B);
631-
cgh.parallel_for(cl::sycl::range<1>(10), krn);
632-
});
633-
} // Data is copied back
634-
for (int i = 0; i < 10; i++)
635-
assert(data2[i] == -2);
636-
for (int i = 0; i < 10; i++)
637-
assert(data1[i] == 0);
638-
}
639-
}
640-
641573
{
642574
int data[10];
643575
void *voidPtr = (void *)data;

‎sycl/test/basic_tests/buffer/buffer_interop.cpp‎

Lines changed: 90 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -5,7 +5,7 @@
55
// RUN: %GPU_RUN_PLACEHOLDER %t.out
66
// RUN: %ACC_RUN_PLACEHOLDER %t.out
77

8-
//==------------------- buffer.cpp - SYCL buffer basic test ----------------==//
8+
//==------------------- buffer_interop.cpp - SYCL buffer basic test ---------==//
99
//
1010
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
1111
// See https://llvm.org/LICENSE.txt for license information.
@@ -36,16 +36,16 @@ int main() {
3636
buffer<int, 1> Buffer{OpenCLBuffer, MyQueue.get_context()};
3737

3838
if (Buffer.get_range() != InteropRange) {
39-
assert(false);
40-
Failed = true;
39+
assert(false);
40+
Failed = true;
4141
}
4242
if (Buffer.get_size() != InteropSize) {
43-
assert(false);
44-
Failed = true;
43+
assert(false);
44+
Failed = true;
4545
}
4646
if (Buffer.get_count() != Size) {
47-
assert(false);
48-
Failed = true;
47+
assert(false);
48+
Failed = true;
4949
}
5050

5151
MyQueue.submit([&](handler &CGH) {
@@ -57,8 +57,7 @@ int main() {
5757
int Data[Size] = {10};
5858
std::vector<int> Result(Size, 0);
5959
{
60-
buffer<int, 1> BufferData{Data, range<1>{Size},
61-
{property::buffer::use_host_ptr()}};
60+
buffer<int, 1> BufferData{Data, range<1>{Size}, {property::buffer::use_host_ptr()}};
6261
BufferData.set_final_data(Result.begin());
6362
MyQueue.submit([&](handler &CGH) {
6463
auto Data = BufferData.get_access<access::mode::write>(CGH);
@@ -132,7 +131,7 @@ int main() {
132131
MyQueue.submit([&](handler &CGH) {
133132
auto B = Buffer.get_access<access::mode::write>(CGH);
134133
CGH.parallel_for<class HostAccess>(range<1>{Size},
135-
[=](id<1> Index) { B[Index] = 10; });
134+
[=](id<1> Index) { B[Index] = 10; });
136135
});
137136
auto Acc = Buffer.get_access<cl::sycl::access::mode::read>();
138137
for (size_t i = 0; i < Size; ++i) {
@@ -183,5 +182,86 @@ int main() {
183182
CHECK_OCL_CODE(clReleaseEvent(OpenCLEvent));
184183
}
185184

185+
{
186+
queue Queue;
187+
if (!Queue.is_host()) {
188+
std::vector<int> Data1(10, -1);
189+
std::vector<int> Data2(10, -2);
190+
{
191+
buffer<int, 1> BufferA(Data1.data(), range<1>(10));
192+
buffer<int, 1> BufferB(Data2);
193+
194+
program Program(Queue.get_context());
195+
Program.build_with_source("kernel void override_source(global int* Acc) "
196+
"{Acc[get_global_id(0)] = 0; }\n");
197+
cl::sycl::kernel Kernel = Program.get_kernel("override_source");
198+
Queue.submit([&](handler &CGH) {
199+
auto AccA = BufferA.get_access<access::mode::read_write>(CGH);
200+
CGH.set_arg(0, AccA);
201+
auto AccB = BufferB.get_access<access::mode::read_write>(CGH);
202+
CGH.parallel_for(cl::sycl::range<1>(10), Kernel);
203+
});
204+
} // Data is copied back
205+
for (int i = 0; i < 10; i++) {
206+
if (Data2[i] != -2) {
207+
std::cout << " Data2[" << i << "] is " << Data2[i] << " expected " << -2 << std::endl;
208+
assert(false);
209+
Failed = true;
210+
}
211+
}
212+
for (int i = 0; i < 10; i++) {
213+
if (Data1[i] != 0) {
214+
std::cout << " Data1[" << i << "] is " << Data1[i] << " expected " << 0 << std::endl;
215+
assert(false);
216+
Failed = true;
217+
}
218+
}
219+
}
220+
}
221+
222+
{
223+
queue Queue;
224+
if (!Queue.is_host()) {
225+
std::vector<int> Data1(10, -1);
226+
std::vector<int> Data2(10, -2);
227+
{
228+
buffer<int, 1> BufferA(Data1.data(), range<1>(10));
229+
buffer<int, 1> BufferB(Data2);
230+
accessor<int, 1, access::mode::read_write,
231+
access::target::global_buffer, access::placeholder::true_t>
232+
AccA(BufferA);
233+
accessor<int, 1, access::mode::read_write,
234+
access::target::global_buffer, access::placeholder::true_t>
235+
AccB(BufferB);
236+
237+
program Program(Queue.get_context());
238+
Program.build_with_source("kernel void override_source_placeholder(global "
239+
"int* Acc) {Acc[get_global_id(0)] = 0; }\n");
240+
cl::sycl::kernel Kernel = Program.get_kernel("override_source_placeholder");
241+
242+
Queue.submit([&](handler &CGH) {
243+
CGH.require(AccA);
244+
CGH.set_arg(0, AccA);
245+
CGH.require(AccB);
246+
CGH.parallel_for(cl::sycl::range<1>(10), Kernel);
247+
});
248+
} // Data is copied back
249+
for (int i = 0; i < 10; i++) {
250+
if (Data2[i] != -2) {
251+
std::cout << " Data2[" << i << "] is " << Data2[i] << " expected " << -2 << std::endl;
252+
assert(false);
253+
Failed = true;
254+
}
255+
}
256+
for (int i = 0; i < 10; i++) {
257+
if (Data1[i] != 0) {
258+
std::cout << " Data1[" << i << "] is " << Data1[i] << " expected " << 0 << std::endl;
259+
assert(false);
260+
Failed = true;
261+
}
262+
}
263+
}
264+
}
265+
186266
return Failed;
187267
}

0 commit comments

Comments
 (0)