Skip to content

Commit 3eef33b

Browse files
cppchedycppchedy
authored
[SYCL][E2E] increase test coverage for subregions copies (#14046)
* Updated 2D subregion copy test so that it follows the other subregion tests. * Updated the 3D and 2D subregion tests to cover test cases for copying with non-zero offsets from device to host.
1 parent fe2b47f commit 3eef33b

File tree

3 files changed

+58
-27
lines changed

3 files changed

+58
-27
lines changed

‎sycl/plugins/unified_runtime/CMakeLists.txt

Lines changed: 7 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -126,7 +126,13 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT)
126126

127127
fetch_adapter_source(cuda
128128
${UNIFIED_RUNTIME_REPO}
129-
${UNIFIED_RUNTIME_TAG}
129+
# commit 08acafb6b36c9ff12753922365df150f677655e7
130+
# Merge: 1e9b1b49 8e5a33b6
131+
# Author: Callum Fare <[email protected]>
132+
# Date: Mon Jun 24 10:56:12 2024 +0100
133+
# Merge pull request #1714 from cppchedy/chedy/fix-copy-in-bytes-bug
134+
# [CUDA][Bindless][Exp] Fix subregion copies
135+
08acafb6b36c9ff12753922365df150f677655e7
130136
)
131137

132138
fetch_adapter_source(hip

‎sycl/test-e2e/bindless_images/read_write_2D_subregion.cpp

Lines changed: 26 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -24,14 +24,14 @@ int main() {
2424
size_t height = 32;
2525
size_t N = width * height;
2626
std::vector<float> out(N);
27-
std::vector<float> expected(N / 4);
28-
std::vector<float> dataIn1(N / 4);
29-
std::vector<float> dataIn2(N / 4);
30-
for (int i = 0; i < width / 2; i++) {
31-
for (int j = 0; j < height / 2; j++) {
32-
expected[j + ((height / 2) * i)] = j * 3;
33-
dataIn1[j + ((height / 2) * i)] = j;
34-
dataIn2[j + ((height / 2) * i)] = j * 2;
27+
std::vector<float> expected(N);
28+
std::vector<float> dataIn1(N);
29+
std::vector<float> dataIn2(N);
30+
for (int i = 0; i < width; i++) {
31+
for (int j = 0; j < height; j++) {
32+
expected[j + ((height)*i)] = j * 3;
33+
dataIn1[j + ((height)*i)] = j;
34+
dataIn2[j + ((height)*i)] = j * 2;
3535
}
3636
}
3737

@@ -55,29 +55,29 @@ int main() {
5555

5656
// Extension: copy over data to device (four subregions/quadrants)
5757
sycl::range copyExtent1 = {width / 2, height / 2, 1};
58-
sycl::range srcExtent = {width / 2, height / 2, 0};
58+
sycl::range srcExtent = {width, height, 0};
5959

6060
q.ext_oneapi_copy(dataIn1.data(), {0, 0, 0}, srcExtent,
6161
imgMem0.get_handle(), {0, 0, 0}, desc, copyExtent1);
62-
q.ext_oneapi_copy(dataIn1.data(), {0, 0, 0}, srcExtent,
62+
q.ext_oneapi_copy(dataIn1.data(), {width / 2, 0, 0}, srcExtent,
6363
imgMem0.get_handle(), {width / 2, 0, 0}, desc,
6464
copyExtent1);
65-
q.ext_oneapi_copy(dataIn1.data(), {0, 0, 0}, srcExtent,
65+
q.ext_oneapi_copy(dataIn1.data(), {0, height / 2, 0}, srcExtent,
6666
imgMem0.get_handle(), {0, height / 2, 0}, desc,
6767
copyExtent1);
68-
q.ext_oneapi_copy(dataIn1.data(), {0, 0, 0}, srcExtent,
68+
q.ext_oneapi_copy(dataIn1.data(), {width / 2, height / 2, 0}, srcExtent,
6969
imgMem0.get_handle(), {width / 2, height / 2, 0}, desc,
7070
copyExtent1);
7171

7272
q.ext_oneapi_copy(dataIn2.data(), {0, 0, 0}, srcExtent,
7373
imgMem1.get_handle(), {0, 0, 0}, desc, copyExtent1);
74-
q.ext_oneapi_copy(dataIn2.data(), {0, 0, 0}, srcExtent,
74+
q.ext_oneapi_copy(dataIn2.data(), {width / 2, 0, 0}, srcExtent,
7575
imgMem1.get_handle(), {width / 2, 0, 0}, desc,
7676
copyExtent1);
77-
q.ext_oneapi_copy(dataIn2.data(), {0, 0, 0}, srcExtent,
77+
q.ext_oneapi_copy(dataIn2.data(), {0, height / 2, 0}, srcExtent,
7878
imgMem1.get_handle(), {0, height / 2, 0}, desc,
7979
copyExtent1);
80-
q.ext_oneapi_copy(dataIn2.data(), {0, 0, 0}, srcExtent,
80+
q.ext_oneapi_copy(dataIn2.data(), {width / 2, height / 2, 0}, srcExtent,
8181
imgMem1.get_handle(), {width / 2, height / 2, 0}, desc,
8282
copyExtent1);
8383

@@ -104,13 +104,18 @@ int main() {
104104
});
105105
q.wait_and_throw();
106106

107-
// Extension: copy data from device to host (two sub-regions)
108-
sycl::range copyExtent2 = {width, height / 2, 1};
109-
sycl::range destExtent = {width, height, 0};
107+
// Extension: copy data from device to host (four subregions/quadrants)
108+
auto destExtent = srcExtent;
110109
q.ext_oneapi_copy(imgMem2.get_handle(), {0, 0, 0}, desc, out.data(),
111-
{0, 0, 0}, destExtent, copyExtent2);
110+
{0, 0, 0}, destExtent, copyExtent1);
111+
q.ext_oneapi_copy(imgMem2.get_handle(), {width / 2, 0, 0}, desc, out.data(),
112+
{width / 2, 0, 0}, destExtent, copyExtent1);
112113
q.ext_oneapi_copy(imgMem2.get_handle(), {0, height / 2, 0}, desc,
113-
out.data(), {0, height / 2, 0}, destExtent, copyExtent2);
114+
out.data(), {0, height / 2, 0}, destExtent, copyExtent1);
115+
q.ext_oneapi_copy(imgMem2.get_handle(), {width / 2, height / 2, 0}, desc,
116+
out.data(), {width / 2, height / 2, 0}, destExtent,
117+
copyExtent1);
118+
114119
q.wait_and_throw();
115120

116121
// Extension: cleanup
@@ -128,7 +133,7 @@ int main() {
128133
bool validated = true;
129134
for (int i = 0; i < N; i++) {
130135
bool mismatch = false;
131-
if (out[i] != expected[i % (N / 4)]) {
136+
if (out[i] != expected[i]) {
132137
mismatch = true;
133138
validated = false;
134139
}

‎sycl/test-e2e/bindless_images/read_write_3D_subregion.cpp

Lines changed: 25 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -124,13 +124,32 @@ int main() {
124124

125125
q.wait_and_throw();
126126

127-
// Extension: copy data from device to host (two sub-regions)
128-
sycl::range copyExtent3 = {width, height, depth / 2};
129-
sycl::range destExtent = {width, height, depth};
127+
// Extension: copy data from device to host (8 sub-regions)
128+
129+
sycl::range<3> destExtent2 = srcExtent1;
130+
130131
q.ext_oneapi_copy(imgMem2.get_handle(), {0, 0, 0}, desc, out.data(),
131-
{0, 0, 0}, destExtent, copyExtent3);
132+
{0, 0, 0}, destExtent2, copyExtent1);
133+
134+
q.ext_oneapi_copy(imgMem2.get_handle(), {width / 2, 0, 0}, desc, out.data(),
135+
{width / 2, 0, 0}, destExtent2, copyExtent1);
136+
q.ext_oneapi_copy(imgMem2.get_handle(), {0, height / 2, 0}, desc,
137+
out.data(), {0, height / 2, 0}, destExtent2, copyExtent1);
132138
q.ext_oneapi_copy(imgMem2.get_handle(), {0, 0, depth / 2}, desc, out.data(),
133-
{0, 0, depth / 2}, destExtent, copyExtent3);
139+
{0, 0, depth / 2}, destExtent2, copyExtent1);
140+
q.ext_oneapi_copy(imgMem2.get_handle(), {width / 2, height / 2, 0}, desc,
141+
out.data(), {width / 2, height / 2, 0}, destExtent2,
142+
copyExtent1);
143+
q.ext_oneapi_copy(imgMem2.get_handle(), {0, height / 2, depth / 2}, desc,
144+
out.data(), {0, height / 2, depth / 2}, destExtent2,
145+
copyExtent1);
146+
q.ext_oneapi_copy(imgMem2.get_handle(), {width / 2, 0, depth / 2}, desc,
147+
out.data(), {width / 2, 0, depth / 2}, destExtent2,
148+
copyExtent1);
149+
q.ext_oneapi_copy(imgMem2.get_handle(), {width / 2, height / 2, depth / 2},
150+
desc, out.data(), {width / 2, height / 2, depth / 2},
151+
destExtent2, copyExtent1);
152+
134153
q.wait_and_throw();
135154

136155
// Extension: cleanup
@@ -163,6 +182,7 @@ int main() {
163182
#endif
164183
}
165184
}
185+
166186
if (validated) {
167187
std::cout << "Test passed!" << std::endl;
168188
return 0;

0 commit comments

Comments
 (0)