Skip to content

Commit 44bc0b4

Browse files
JackAKirkJackAKirk
authored
[Bindless] Add O0 bindless tests for fetch/sample_image that fail on L0 && windows (#18485)
These tests fail with ``` SYCL exception caught! : The program was built for 1 devices Build program log for 'Intel(R) Arc(TM) A770 Graphics': error: Inconsistent use of image! error: backend compiler failed build. ``` on L0 backend, only on windows. This error goes away if `fetch_image` (using either sampled or unsampled image)/ `sample_image` calls are commented out. This only occurs on L0 backend, pointing to issues with spirv implementations for this backend. See #18919 --------- Signed-off-by: JackAKirk <jack.kirk@codeplay.com>
1 parent 0c4e3ac commit 44bc0b4

File tree

9 files changed

+403
-354
lines changed

9 files changed

+403
-354
lines changed

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

Lines changed: 2 additions & 140 deletions
Original file line numberDiff line numberDiff line change
@@ -3,147 +3,9 @@
33
// RUN: %{build} -o %t.out
44
// RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out
55

6-
#include <iostream>
7-
#include <sycl/detail/core.hpp>
8-
9-
#include <sycl/ext/oneapi/bindless_images.hpp>
10-
116
// Uncomment to print additional test information
127
// #define VERBOSE_PRINT
138

14-
class image_addition;
15-
16-
int main() {
17-
18-
sycl::device dev;
19-
sycl::queue q(dev);
20-
auto ctxt = q.get_context();
21-
22-
// declare image data
23-
constexpr size_t width = 512;
24-
std::vector<float> out(width);
25-
std::vector<float> expected(width);
26-
std::vector<sycl::float4> dataIn1(width);
27-
std::vector<sycl::float4> dataIn2(width);
28-
float exp = 512;
29-
for (int i = 0; i < width; i++) {
30-
expected[i] = exp;
31-
dataIn1[i] = sycl::float4(i, i, i, i);
32-
dataIn2[i] = sycl::float4(width - i, width - i, width - i, width - i);
33-
}
34-
35-
try {
36-
// Extension: image descriptor - can use the same for both images
37-
sycl::ext::oneapi::experimental::image_descriptor desc(
38-
{width}, 4, sycl::image_channel_type::fp32);
39-
40-
// Extension: allocate memory on device and create the handle
41-
sycl::ext::oneapi::experimental::image_mem imgMem0(desc, dev, ctxt);
42-
sycl::ext::oneapi::experimental::image_mem imgMem1(desc, dev, ctxt);
43-
44-
// std::hash specialization to ensure `image_mem` follows common reference
45-
// semantics
46-
assert(std::hash<sycl::ext::oneapi::experimental::image_mem>{}(imgMem0) !=
47-
std::hash<sycl::ext::oneapi::experimental::image_mem>{}(imgMem1));
48-
49-
// We're able to use move semantics
50-
// Move construct
51-
sycl::ext::oneapi::experimental::image_mem imgMem0MoveConstruct(
52-
std::move(imgMem0));
53-
// Move assign
54-
sycl::ext::oneapi::experimental::image_mem imgMem0MoveAssign;
55-
imgMem0MoveAssign = std::move(imgMem0MoveConstruct);
56-
57-
// We're able to use copy semantics
58-
// Copy construct
59-
sycl::ext::oneapi::experimental::image_mem imgMem1CopyConstruct(imgMem1);
60-
// Copy assign
61-
sycl::ext::oneapi::experimental::image_mem imgMem1CopyAssign;
62-
imgMem1CopyAssign = imgMem1CopyConstruct;
63-
64-
// Equality operators to ensure `image_mem` follows common reference
65-
// semantics
66-
assert(imgMem0MoveAssign != imgMem1CopyAssign);
67-
assert(imgMem1 == imgMem1CopyAssign);
68-
69-
// We can default construct image handles
70-
sycl::ext::oneapi::experimental::unsampled_image_handle imgHandle1;
71-
72-
// Extension: create the image and return the handle
73-
sycl::ext::oneapi::experimental::unsampled_image_handle tmpHandle =
74-
sycl::ext::oneapi::experimental::create_image(imgMem0MoveAssign, desc,
75-
dev, ctxt);
76-
sycl::ext::oneapi::experimental::unsampled_image_handle imgHandle2 =
77-
sycl::ext::oneapi::experimental::create_image(imgMem1CopyAssign, desc,
78-
dev, ctxt);
79-
80-
// Default constructed image handles are not valid until we assign a valid
81-
// raw handle to the struct
82-
imgHandle1.raw_handle = tmpHandle.raw_handle;
83-
84-
// Extension: copy over data to device
85-
q.ext_oneapi_copy(dataIn1.data(), imgMem0MoveAssign.get_handle(), desc);
86-
q.ext_oneapi_copy(dataIn2.data(), imgMem1CopyAssign.get_handle(), desc);
87-
88-
q.wait_and_throw();
89-
90-
sycl::buffer<float, 1> buf((float *)out.data(), width);
91-
q.submit([&](sycl::handler &cgh) {
92-
auto outAcc = buf.get_access<sycl::access_mode::write>(cgh, width);
93-
94-
cgh.parallel_for<image_addition>(width, [=](sycl::id<1> id) {
95-
float sum = 0;
96-
// Extension: fetch image data from handle
97-
sycl::float4 px1 =
98-
sycl::ext::oneapi::experimental::fetch_image<sycl::float4>(
99-
imgHandle1, int(id[0]));
100-
sycl::float4 px2 =
101-
sycl::ext::oneapi::experimental::fetch_image<sycl::float4>(
102-
imgHandle2, int(id[0]));
103-
104-
sum = px1[0] + px2[0];
105-
outAcc[id] = sum;
106-
});
107-
});
108-
109-
q.wait_and_throw();
110-
111-
// Extension: cleanup
112-
sycl::ext::oneapi::experimental::destroy_image_handle(imgHandle1, dev,
113-
ctxt);
114-
sycl::ext::oneapi::experimental::destroy_image_handle(imgHandle2, dev,
115-
ctxt);
116-
} catch (sycl::exception e) {
117-
std::cerr << "SYCL exception caught! : " << e.what() << "\n";
118-
return 1;
119-
} catch (...) {
120-
std::cerr << "Unknown exception caught!\n";
121-
return 2;
122-
}
123-
124-
// collect and validate output
125-
bool validated = true;
126-
for (int i = 0; i < width; i++) {
127-
bool mismatch = false;
128-
if (out[i] != expected[i]) {
129-
mismatch = true;
130-
validated = false;
131-
}
132-
133-
if (mismatch) {
134-
#ifdef VERBOSE_PRINT
135-
std::cout << "Result mismatch! Expected: " << expected[i]
136-
<< ", Actual: " << out[i] << std::endl;
137-
#else
138-
break;
139-
#endif
140-
}
141-
}
142-
if (validated) {
143-
std::cout << "Test passed!" << std::endl;
144-
return 0;
145-
}
9+
#include "read_1D.hpp"
14610

147-
std::cout << "Test failed!" << std::endl;
148-
return 3;
149-
}
11+
int main() { return test(); }

‎sycl/test-e2e/bindless_images/read_1D.hpp

Lines changed: 141 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,141 @@
1+
#include <iostream>
2+
#include <sycl/detail/core.hpp>
3+
4+
#include <sycl/ext/oneapi/bindless_images.hpp>
5+
6+
class image_addition;
7+
8+
int test() {
9+
10+
sycl::device dev;
11+
sycl::queue q(dev);
12+
auto ctxt = q.get_context();
13+
14+
// declare image data
15+
constexpr size_t width = 512;
16+
std::vector<float> out(width);
17+
std::vector<float> expected(width);
18+
std::vector<sycl::float4> dataIn1(width);
19+
std::vector<sycl::float4> dataIn2(width);
20+
float exp = 512;
21+
for (int i = 0; i < width; i++) {
22+
expected[i] = exp;
23+
dataIn1[i] = sycl::float4(i, i, i, i);
24+
dataIn2[i] = sycl::float4(width - i, width - i, width - i, width - i);
25+
}
26+
27+
try {
28+
// Extension: image descriptor - can use the same for both images
29+
sycl::ext::oneapi::experimental::image_descriptor desc(
30+
{width}, 4, sycl::image_channel_type::fp32);
31+
32+
// Extension: allocate memory on device and create the handle
33+
sycl::ext::oneapi::experimental::image_mem imgMem0(desc, dev, ctxt);
34+
sycl::ext::oneapi::experimental::image_mem imgMem1(desc, dev, ctxt);
35+
36+
// std::hash specialization to ensure `image_mem` follows common reference
37+
// semantics
38+
assert(std::hash<sycl::ext::oneapi::experimental::image_mem>{}(imgMem0) !=
39+
std::hash<sycl::ext::oneapi::experimental::image_mem>{}(imgMem1));
40+
41+
// We're able to use move semantics
42+
// Move construct
43+
sycl::ext::oneapi::experimental::image_mem imgMem0MoveConstruct(
44+
std::move(imgMem0));
45+
// Move assign
46+
sycl::ext::oneapi::experimental::image_mem imgMem0MoveAssign;
47+
imgMem0MoveAssign = std::move(imgMem0MoveConstruct);
48+
49+
// We're able to use copy semantics
50+
// Copy construct
51+
sycl::ext::oneapi::experimental::image_mem imgMem1CopyConstruct(imgMem1);
52+
// Copy assign
53+
sycl::ext::oneapi::experimental::image_mem imgMem1CopyAssign;
54+
imgMem1CopyAssign = imgMem1CopyConstruct;
55+
56+
// Equality operators to ensure `image_mem` follows common reference
57+
// semantics
58+
assert(imgMem0MoveAssign != imgMem1CopyAssign);
59+
assert(imgMem1 == imgMem1CopyAssign);
60+
61+
// We can default construct image handles
62+
sycl::ext::oneapi::experimental::unsampled_image_handle imgHandle1;
63+
64+
// Extension: create the image and return the handle
65+
sycl::ext::oneapi::experimental::unsampled_image_handle tmpHandle =
66+
sycl::ext::oneapi::experimental::create_image(imgMem0MoveAssign, desc,
67+
dev, ctxt);
68+
sycl::ext::oneapi::experimental::unsampled_image_handle imgHandle2 =
69+
sycl::ext::oneapi::experimental::create_image(imgMem1CopyAssign, desc,
70+
dev, ctxt);
71+
72+
// Default constructed image handles are not valid until we assign a valid
73+
// raw handle to the struct
74+
imgHandle1.raw_handle = tmpHandle.raw_handle;
75+
76+
// Extension: copy over data to device
77+
q.ext_oneapi_copy(dataIn1.data(), imgMem0MoveAssign.get_handle(), desc);
78+
q.ext_oneapi_copy(dataIn2.data(), imgMem1CopyAssign.get_handle(), desc);
79+
80+
q.wait_and_throw();
81+
82+
sycl::buffer<float, 1> buf((float *)out.data(), width);
83+
q.submit([&](sycl::handler &cgh) {
84+
auto outAcc = buf.get_access<sycl::access_mode::write>(cgh, width);
85+
86+
cgh.parallel_for<image_addition>(width, [=](sycl::id<1> id) {
87+
float sum = 0;
88+
// Extension: fetch image data from handle
89+
sycl::float4 px1 =
90+
sycl::ext::oneapi::experimental::fetch_image<sycl::float4>(
91+
imgHandle1, int(id[0]));
92+
sycl::float4 px2 =
93+
sycl::ext::oneapi::experimental::fetch_image<sycl::float4>(
94+
imgHandle2, int(id[0]));
95+
96+
sum = px1[0] + px2[0];
97+
outAcc[id] = sum;
98+
});
99+
});
100+
101+
q.wait_and_throw();
102+
103+
// Extension: cleanup
104+
sycl::ext::oneapi::experimental::destroy_image_handle(imgHandle1, dev,
105+
ctxt);
106+
sycl::ext::oneapi::experimental::destroy_image_handle(imgHandle2, dev,
107+
ctxt);
108+
} catch (sycl::exception e) {
109+
std::cerr << "SYCL exception caught! : " << e.what() << "\n";
110+
return 1;
111+
} catch (...) {
112+
std::cerr << "Unknown exception caught!\n";
113+
return 2;
114+
}
115+
116+
// collect and validate output
117+
bool validated = true;
118+
for (int i = 0; i < width; i++) {
119+
bool mismatch = false;
120+
if (out[i] != expected[i]) {
121+
mismatch = true;
122+
validated = false;
123+
}
124+
125+
if (mismatch) {
126+
#ifdef VERBOSE_PRINT
127+
std::cout << "Result mismatch! Expected: " << expected[i]
128+
<< ", Actual: " << out[i] << std::endl;
129+
#else
130+
break;
131+
#endif
132+
}
133+
}
134+
if (validated) {
135+
std::cout << "Test passed!" << std::endl;
136+
return 0;
137+
}
138+
139+
std::cout << "Test failed!" << std::endl;
140+
return 3;
141+
}

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

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,12 @@
1+
// REQUIRES: aspect-ext_oneapi_bindless_images
2+
// XFAIL: level_zero && windows
3+
// XFAIL-TRACKER: https://github.com/intel/llvm/issues/18919
4+
// RUN: %{build} %O0 -o %t.out
5+
// RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out
6+
7+
// Uncomment to print additional test information
8+
// #define VERBOSE_PRINT
9+
10+
#include "read_1D.hpp"
11+
12+
int main() { return test(); }

0 commit comments

Comments
 (0)