Skip to content

Commit 61fe8e6

Browse files
steffenlarsensteffenlarsen
authored
[SYCL][E2E] Enable device_global tests (#10007)
This commit enables a selection of previously disabled paths for the SYCL E2E tests related to the device_global extension. Additionally, it splits each test into one case without the `device_image_scope` on its `device_global`s and one case with. --------- Signed-off-by: Larsen, Steffen <[email protected]>
1 parent b4d2b2c commit 61fe8e6

16 files changed

+271
-221
lines changed

‎sycl/test-e2e/DeviceGlobal/common.hpp

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,14 @@
1+
#pragma once
2+
3+
#include <sycl/sycl.hpp>
4+
5+
using namespace sycl;
6+
using namespace sycl::ext::oneapi::experimental;
7+
8+
// Property list that contains device_image_scope if USE_DEVICE_IMAGE_SCOPE is
9+
// defined.
10+
#ifdef USE_DEVICE_IMAGE_SCOPE
11+
using TestProperties = decltype(properties{device_image_scope});
12+
#else
13+
using TestProperties = decltype(properties{});
14+
#endif

‎sycl/test-e2e/DeviceGlobal/device_global_arrow.cpp

Lines changed: 10 additions & 64 deletions
Original file line numberDiff line numberDiff line change
@@ -1,70 +1,16 @@
1-
// TODO: device_global without the device_image_scope property is not currently
2-
// initialized on device. Enable the following test cases when it is
3-
// supported.
4-
// RUNx: %{build} -o %t.out
5-
// RUNx: %{run} %t.out
1+
// RUN: %{build} -o %t.out
2+
// RUN: %{run} %t.out
63
//
7-
// RUN: %{build} -fsycl-device-code-split=per_source -DUSE_DEVICE_IMAGE_SCOPE -o %t_dev_img_scope.out
8-
// RUN: %{run} %t_dev_img_scope.out
4+
// The HIP and OpenCL GPU backends do not currently support device_global
5+
// backend calls.
6+
// UNSUPPORTED: hip || (opencl && gpu)
97
//
10-
// CPU and accelerators are not currently guaranteed to support the required
11-
// extensions they are disabled until they are.
12-
// UNSUPPORTED: cpu, accelerator
8+
// Temporarily disabled for OpenCL CPU while we wait for CPU driver bump. Same
9+
// applies to the FPGA emulator.
10+
// UNSUPPORTED: opencl
1311
//
1412
// Tests operator-> on device_global.
15-
// NOTE: USE_DEVICE_IMAGE_SCOPE needs both kernels to be in the same image so
16-
// we set -fsycl-device-code-split=per_source.
1713

18-
#include <sycl/sycl.hpp>
14+
#include "device_global_arrow.hpp"
1915

20-
using namespace sycl;
21-
using namespace sycl::ext::oneapi::experimental;
22-
23-
struct StructWithMember {
24-
int x;
25-
int getX() { return x; }
26-
};
27-
28-
struct StructWithDeref {
29-
StructWithMember y[1];
30-
StructWithMember *operator->() { return y; }
31-
};
32-
33-
#ifdef USE_DEVICE_IMAGE_SCOPE
34-
device_global<StructWithMember *, decltype(properties{device_image_scope})>
35-
DeviceGlobalVar1;
36-
device_global<StructWithDeref, decltype(properties{device_image_scope})>
37-
DeviceGlobalVar2;
38-
#else
39-
device_global<StructWithMember *> DeviceGlobalVar1;
40-
device_global<StructWithDeref> DeviceGlobalVar2;
41-
#endif
42-
43-
int main() {
44-
queue Q;
45-
46-
StructWithMember *DGMem = malloc_device<StructWithMember>(1, Q);
47-
48-
Q.single_task([=]() {
49-
DeviceGlobalVar1 = DGMem;
50-
DeviceGlobalVar1->x = 1234;
51-
DeviceGlobalVar2->x = 4321;
52-
}).wait();
53-
54-
int Out[2] = {0, 0};
55-
{
56-
buffer<int, 1> OutBuf{Out, 2};
57-
Q.submit([&](handler &CGH) {
58-
auto OutAcc = OutBuf.get_access<access::mode::write>(CGH);
59-
CGH.single_task([=]() {
60-
OutAcc[0] = DeviceGlobalVar1->getX();
61-
OutAcc[1] = DeviceGlobalVar2->getX();
62-
});
63-
});
64-
}
65-
free(DGMem, Q);
66-
67-
assert(Out[0] == 1234 && "First value does not match.");
68-
assert(Out[1] == 4321 && "Second value does not match.");
69-
return 0;
70-
}
16+
int main() { return test(); }

‎sycl/test-e2e/DeviceGlobal/device_global_arrow.hpp

Lines changed: 45 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,45 @@
1+
#pragma once
2+
3+
#include "common.hpp"
4+
5+
struct StructWithMember {
6+
int x;
7+
int getX() { return x; }
8+
};
9+
10+
struct StructWithDeref {
11+
StructWithMember y[1];
12+
StructWithMember *operator->() { return y; }
13+
};
14+
15+
device_global<StructWithMember *, TestProperties> DeviceGlobalVar1;
16+
device_global<StructWithDeref, TestProperties> DeviceGlobalVar2;
17+
18+
int test() {
19+
queue Q;
20+
21+
StructWithMember *DGMem = malloc_device<StructWithMember>(1, Q);
22+
23+
Q.single_task([=]() {
24+
DeviceGlobalVar1 = DGMem;
25+
DeviceGlobalVar1->x = 1234;
26+
DeviceGlobalVar2->x = 4321;
27+
}).wait();
28+
29+
int Out[2] = {0, 0};
30+
{
31+
buffer<int, 1> OutBuf{Out, 2};
32+
Q.submit([&](handler &CGH) {
33+
auto OutAcc = OutBuf.get_access<access::mode::write>(CGH);
34+
CGH.single_task([=]() {
35+
OutAcc[0] = DeviceGlobalVar1->getX();
36+
OutAcc[1] = DeviceGlobalVar2->getX();
37+
});
38+
});
39+
}
40+
free(DGMem, Q);
41+
42+
assert(Out[0] == 1234 && "First value does not match.");
43+
assert(Out[1] == 4321 && "Second value does not match.");
44+
return 0;
45+
}

‎sycl/test-e2e/DeviceGlobal/device_global_arrow_dis.cpp

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,14 @@
1+
// RUN: %{build} -fsycl-device-code-split=per_source -DUSE_DEVICE_IMAGE_SCOPE -o %t.out
2+
// RUN: %{run} %t.out
3+
//
4+
// The HIP and OpenCL GPU backends do not currently support device_global
5+
// backend calls.
6+
// UNSUPPORTED: hip || (opencl && gpu)
7+
//
8+
// Tests operator-> on device_global with device_image_scope.
9+
// NOTE: USE_DEVICE_IMAGE_SCOPE needs both kernels to be in the same image so
10+
// we set -fsycl-device-code-split=per_source.
11+
12+
#include "device_global_arrow.hpp"
13+
14+
int main() { return test(); }

‎sycl/test-e2e/DeviceGlobal/device_global_device_only.cpp

Lines changed: 10 additions & 40 deletions
Original file line numberDiff line numberDiff line change
@@ -1,46 +1,16 @@
1-
// TODO: device_global without the device_image_scope property is not currently
2-
// initialized on device. Enable the following test cases when it is
3-
// supported.
4-
// RUNx: %{build} -o %t.out
5-
// RUNx: %{run} %t.out
1+
// RUN: %{build} -o %t.out
2+
// RUN: %{run} %t.out
63
//
7-
// RUN: %{build} -fsycl-device-code-split=per_source -DUSE_DEVICE_IMAGE_SCOPE -o %t_dev_img_scope.out
8-
// RUN: %{run} %t_dev_img_scope.out
4+
// The HIP and OpenCL GPU backends do not currently support device_global
5+
// backend calls.
6+
// UNSUPPORTED: hip || (opencl && gpu)
97
//
10-
// CPU and accelerators are not currently guaranteed to support the required
11-
// extensions they are disabled until they are.
12-
// UNSUPPORTED: cpu, accelerator
8+
// Temporarily disabled for OpenCL CPU while we wait for CPU driver bump. Same
9+
// applies to the FPGA emulator.
10+
// UNSUPPORTED: opencl
1311
//
1412
// Tests basic device_global access through device kernels.
15-
// NOTE: USE_DEVICE_IMAGE_SCOPE needs both kernels to be in the same image so
16-
// we set -fsycl-device-code-split=per_source.
1713

18-
#include <sycl/sycl.hpp>
14+
#include "device_global_device_only.hpp"
1915

20-
using namespace sycl;
21-
using namespace sycl::ext::oneapi::experimental;
22-
23-
#ifdef USE_DEVICE_IMAGE_SCOPE
24-
device_global<int[4], decltype(properties{device_image_scope})> DeviceGlobalVar;
25-
#else
26-
device_global<int[4]> DeviceGlobalVar;
27-
#endif
28-
29-
int main() {
30-
queue Q;
31-
32-
Q.single_task([=]() { DeviceGlobalVar.get()[0] = 42; });
33-
// Make sure that the write happens before subsequent read
34-
Q.wait();
35-
36-
int OutVal = 0;
37-
{
38-
buffer<int, 1> OutBuf(&OutVal, 1);
39-
Q.submit([&](handler &CGH) {
40-
auto OutAcc = OutBuf.get_access<access::mode::write>(CGH);
41-
CGH.single_task([=]() { OutAcc[0] = DeviceGlobalVar.get()[0]; });
42-
});
43-
}
44-
assert(OutVal == 42 && "Read value does not match.");
45-
return 0;
46-
}
16+
int main() { return test(); }

‎sycl/test-e2e/DeviceGlobal/device_global_device_only.hpp

Lines changed: 24 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,24 @@
1+
#pragma once
2+
3+
#include "common.hpp"
4+
5+
device_global<int[4], TestProperties> DeviceGlobalVar;
6+
7+
int test() {
8+
queue Q;
9+
10+
Q.single_task([=]() { DeviceGlobalVar.get()[0] = 42; });
11+
// Make sure that the write happens before subsequent read
12+
Q.wait();
13+
14+
int OutVal = 0;
15+
{
16+
buffer<int, 1> OutBuf(&OutVal, 1);
17+
Q.submit([&](handler &CGH) {
18+
auto OutAcc = OutBuf.get_access<access::mode::write>(CGH);
19+
CGH.single_task([=]() { OutAcc[0] = DeviceGlobalVar.get()[0]; });
20+
});
21+
}
22+
assert(OutVal == 42 && "Read value does not match.");
23+
return 0;
24+
}

‎sycl/test-e2e/DeviceGlobal/device_global_device_only_dis.cpp

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,15 @@
1+
// RUN: %{build} -fsycl-device-code-split=per_source -DUSE_DEVICE_IMAGE_SCOPE -o %t.out
2+
// RUN: %{run} %t.out
3+
//
4+
// The HIP and OpenCL GPU backends do not currently support device_global
5+
// backend calls.
6+
// UNSUPPORTED: hip || (opencl && gpu)
7+
//
8+
// Tests basic device_global with device_image_scope access through device
9+
// kernels.
10+
// NOTE: USE_DEVICE_IMAGE_SCOPE needs both kernels to be in the same image so
11+
// we set -fsycl-device-code-split=per_source.
12+
13+
#include "device_global_device_only.hpp"
14+
15+
int main() { return test(); }

‎sycl/test-e2e/DeviceGlobal/device_global_operator_passthrough.cpp

Lines changed: 10 additions & 44 deletions
Original file line numberDiff line numberDiff line change
@@ -1,50 +1,16 @@
1-
// TODO: device_global without the device_image_scope property is not currently
2-
// initialized on device. Enable the following test cases when it is
3-
// supported.
4-
// RUNx: %{build} -o %t.out
5-
// RUNx: %{run} %t.out
1+
// RUN: %{build} -o %t.out
2+
// RUN: %{run} %t.out
63
//
7-
// RUN: %{build} -fsycl-device-code-split=per_source -DUSE_DEVICE_IMAGE_SCOPE -o %t_dev_img_scope.out
8-
// RUN: %{run} %t_dev_img_scope.out
4+
// The HIP and OpenCL GPU backends do not currently support device_global
5+
// backend calls.
6+
// UNSUPPORTED: hip || (opencl && gpu)
97
//
10-
// CPU and accelerators are not currently guaranteed to support the required
11-
// extensions they are disabled until they are.
12-
// UNSUPPORTED: cpu, accelerator
8+
// Temporarily disabled for OpenCL CPU while we wait for CPU driver bump. Same
9+
// applies to the FPGA emulator.
10+
// UNSUPPORTED: opencl
1311
//
1412
// Tests the passthrough of operators on device_global.
15-
// NOTE: USE_DEVICE_IMAGE_SCOPE needs both kernels to be in the same image so
16-
// we set -fsycl-device-code-split=per_source.
1713

18-
#include <sycl/sycl.hpp>
14+
#include "device_global_operator_passthrough.hpp"
1915

20-
using namespace sycl;
21-
using namespace sycl::ext::oneapi::experimental;
22-
23-
#ifdef USE_DEVICE_IMAGE_SCOPE
24-
device_global<int, decltype(properties{device_image_scope})> DeviceGlobalVar;
25-
#else
26-
device_global<int> DeviceGlobalVar;
27-
#endif
28-
29-
int main() {
30-
queue Q;
31-
32-
Q.single_task([]() {
33-
DeviceGlobalVar = 2;
34-
DeviceGlobalVar += 3;
35-
DeviceGlobalVar = DeviceGlobalVar * DeviceGlobalVar;
36-
DeviceGlobalVar = DeviceGlobalVar - 3;
37-
DeviceGlobalVar = 25 - DeviceGlobalVar;
38-
}).wait();
39-
40-
int Out = 0;
41-
{
42-
buffer<int, 1> OutBuf{&Out, 1};
43-
Q.submit([&](handler &CGH) {
44-
auto OutAcc = OutBuf.get_access<access::mode::write>(CGH);
45-
CGH.single_task([=]() { OutAcc[0] = DeviceGlobalVar; });
46-
});
47-
}
48-
assert(Out == 3 && "Read value does not match.");
49-
return 0;
50-
}
16+
int main() { return test(); }

‎sycl/test-e2e/DeviceGlobal/device_global_operator_passthrough.hpp

Lines changed: 28 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,28 @@
1+
#pragma once
2+
3+
#include "common.hpp"
4+
5+
device_global<int, TestProperties> DeviceGlobalVar;
6+
7+
int test() {
8+
queue Q;
9+
10+
Q.single_task([]() {
11+
DeviceGlobalVar = 2;
12+
DeviceGlobalVar += 3;
13+
DeviceGlobalVar = DeviceGlobalVar * DeviceGlobalVar;
14+
DeviceGlobalVar = DeviceGlobalVar - 3;
15+
DeviceGlobalVar = 25 - DeviceGlobalVar;
16+
}).wait();
17+
18+
int Out = 0;
19+
{
20+
buffer<int, 1> OutBuf{&Out, 1};
21+
Q.submit([&](handler &CGH) {
22+
auto OutAcc = OutBuf.get_access<access::mode::write>(CGH);
23+
CGH.single_task([=]() { OutAcc[0] = DeviceGlobalVar; });
24+
});
25+
}
26+
assert(Out == 3 && "Read value does not match.");
27+
return 0;
28+
}

‎sycl/test-e2e/DeviceGlobal/device_global_operator_passthrough_dis.cpp

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,14 @@
1+
// RUN: %{build} -fsycl-device-code-split=per_source -DUSE_DEVICE_IMAGE_SCOPE -o %t.out
2+
// RUN: %{run} %t.out
3+
//
4+
// The HIP and OpenCL GPU backends do not currently support device_global
5+
// backend calls.
6+
// UNSUPPORTED: hip || (opencl && gpu)
7+
//
8+
// Tests the passthrough of operators on device_global with device_image_scope.
9+
// NOTE: USE_DEVICE_IMAGE_SCOPE needs both kernels to be in the same image so
10+
// we set -fsycl-device-code-split=per_source.
11+
12+
#include "device_global_operator_passthrough.hpp"
13+
14+
int main() { return test(); }

0 commit comments

Comments
 (0)