Skip to content

Commit b62febc

Browse files
przemektmalonprzemektmalon
authored
[SYCL][Bindless][E2E] Add 1D host USM tests (#17374)
Adds 1D Host USM backed image tests. This patch also fixes device aspect queries for 1D & 2D USM backed image sampling support.
1 parent e1f723e commit b62febc

File tree

5 files changed

+230
-4
lines changed

5 files changed

+230
-4
lines changed

‎sycl/source/detail/device_impl.cpp

+2-4
Original file line numberDiff line numberDiff line change
@@ -671,17 +671,15 @@ bool device_impl::has(aspect Aspect) const {
671671
ur_bool_t support = false;
672672
bool call_successful =
673673
getAdapter()->call_nocheck<UrApiKind::urDeviceGetInfo>(
674-
MDevice,
675-
UR_DEVICE_INFO_BINDLESS_SAMPLED_IMAGE_FETCH_1D_USM_SUPPORT_EXP,
674+
MDevice, UR_DEVICE_INFO_BINDLESS_SAMPLE_1D_USM_SUPPORT_EXP,
676675
sizeof(ur_bool_t), &support, nullptr) == UR_RESULT_SUCCESS;
677676
return call_successful && support;
678677
}
679678
case aspect::ext_oneapi_bindless_images_sample_2d_usm: {
680679
ur_bool_t support = false;
681680
bool call_successful =
682681
getAdapter()->call_nocheck<UrApiKind::urDeviceGetInfo>(
683-
MDevice,
684-
UR_DEVICE_INFO_BINDLESS_SAMPLED_IMAGE_FETCH_2D_USM_SUPPORT_EXP,
682+
MDevice, UR_DEVICE_INFO_BINDLESS_SAMPLE_2D_USM_SUPPORT_EXP,
685683
sizeof(ur_bool_t), &support, nullptr) == UR_RESULT_SUCCESS;
686684
return call_successful && support;
687685
}

‎sycl/test-e2e/bindless_images/sampled_fetch/fetch_1D_USM_host.cpp

+102
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,102 @@
1+
// REQUIRES: aspect-ext_oneapi_bindless_images
2+
// REQUIRES: aspect-ext_oneapi_bindless_sampled_image_fetch_1d_usm
3+
// UNSUPPORTED: target-amd
4+
// UNSUPPORTED-INTENDED: Sampled fetch not currently supported on AMD
5+
6+
// RUN: %{build} -o %t.out
7+
// RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out
8+
9+
#include <iostream>
10+
#include <sycl/detail/core.hpp>
11+
#include <sycl/ext/oneapi/bindless_images.hpp>
12+
#include <sycl/usm.hpp>
13+
14+
class kernel_sampled_fetch;
15+
16+
// Uncomment to print additional test information
17+
// #define VERBOSE_PRINT
18+
19+
int main() {
20+
sycl::device dev;
21+
sycl::queue q(dev);
22+
auto ctxt = q.get_context();
23+
24+
// Declare image size, and expected output and actual output vectors
25+
constexpr size_t width = 32;
26+
constexpr size_t widthInBytes = width * sizeof(float);
27+
std::vector<float> out(width);
28+
std::vector<float> expected(width);
29+
for (int i = 0; i < width; ++i) {
30+
expected[i] = static_cast<float>(i);
31+
}
32+
33+
namespace syclexp = sycl::ext::oneapi::experimental;
34+
35+
try {
36+
// Extension: image descriptor
37+
syclexp::image_descriptor desc({width}, 1, sycl::image_channel_type::fp32);
38+
39+
// Extension: Image creation requires a sampler, but it will have no effect
40+
// on the result, as we will use `fetch_image` in the kernel.
41+
syclexp::bindless_image_sampler samp(
42+
sycl::addressing_mode::repeat,
43+
sycl::coordinate_normalization_mode::normalized,
44+
sycl::filtering_mode::linear);
45+
46+
// Allocate Host USM and initialize with expected data
47+
float *imgMem = sycl::malloc_host<float>(width, q);
48+
memcpy(imgMem, expected.data(), widthInBytes);
49+
50+
// Extension: create the image backed by Host USM and return the handle
51+
auto imgHandle = syclexp::create_image(imgMem, 0, samp, desc, q);
52+
53+
// Create a buffer to output the result from `fetch_image`
54+
sycl::buffer outBuf(out.data(), sycl::range{width});
55+
q.submit([&](sycl::handler &cgh) {
56+
sycl::accessor outAcc{outBuf, cgh, sycl::write_only};
57+
58+
cgh.parallel_for<kernel_sampled_fetch>(width, [=](sycl::id<1> id) {
59+
// Extension: fetch data from sampled image handle
60+
outAcc[id] = syclexp::fetch_image<float>(imgHandle, int(id[0]));
61+
});
62+
});
63+
64+
q.wait_and_throw();
65+
66+
// Extension: cleanup
67+
syclexp::destroy_image_handle(imgHandle, dev, ctxt);
68+
sycl::free(imgMem, ctxt);
69+
} catch (sycl::exception e) {
70+
std::cerr << "SYCL exception caught! : " << e.what() << "\n";
71+
return 1;
72+
} catch (...) {
73+
std::cerr << "Unknown exception caught!\n";
74+
return 2;
75+
}
76+
77+
// collect and validate output
78+
bool validated = true;
79+
for (int i = 0; i < width; i++) {
80+
bool mismatch = false;
81+
if (out[i] != expected[i]) {
82+
mismatch = true;
83+
validated = false;
84+
}
85+
86+
if (mismatch) {
87+
#ifdef VERBOSE_PRINT
88+
std::cout << "Result mismatch! Expected: " << expected[i]
89+
<< ", Actual: " << out[i] << std::endl;
90+
#else
91+
break;
92+
#endif
93+
}
94+
}
95+
if (validated) {
96+
std::cout << "Test passed!" << std::endl;
97+
return 0;
98+
}
99+
100+
std::cout << "Test failed!" << std::endl;
101+
return 3;
102+
}

‎sycl/test-e2e/bindless_images/sampled_fetch/fetch_1D_USM.cpp renamed to ‎sycl/test-e2e/bindless_images/sampled_fetch/fetch_1D_USM_shared.cpp

File renamed without changes.

‎sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D_USM.cpp renamed to ‎sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D_USM_device.cpp

File renamed without changes.

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

+126
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,126 @@
1+
// REQUIRES: aspect-ext_oneapi_bindless_images
2+
// REQUIRES: aspect-ext_oneapi_bindless_images_sample_1d_usm
3+
4+
// UNSUPPORTED: hip
5+
// UNSUPPORTED-INTENDED: Host USM backed image support is not yet enabled in UR
6+
// adapter. Also, when provionally enabled, the test crashes upon image
7+
// creation, whereas Device USM backed images do not crash. This issue is
8+
// undetermined.
9+
10+
// RUN: %{build} -o %t.out
11+
// RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out
12+
13+
#include <cmath>
14+
#include <iostream>
15+
#include <sycl/detail/core.hpp>
16+
17+
#include <sycl/ext/oneapi/bindless_images.hpp>
18+
#include <sycl/usm.hpp>
19+
20+
// Uncomment to print additional test information
21+
// #define VERBOSE_PRINT
22+
23+
class sample_host_usm_image_kernel;
24+
25+
int main() {
26+
27+
sycl::device dev;
28+
sycl::queue q(dev);
29+
auto ctxt = q.get_context();
30+
31+
// declare image data
32+
size_t width = 32;
33+
size_t widthInBytes = width * sizeof(float);
34+
std::vector<float> out(width);
35+
std::vector<float> expected(width);
36+
for (int i = 0; i < width; ++i) {
37+
expected[i] = static_cast<float>(i);
38+
}
39+
40+
try {
41+
sycl::ext::oneapi::experimental::bindless_image_sampler samp(
42+
sycl::addressing_mode::clamp,
43+
sycl::coordinate_normalization_mode::normalized,
44+
sycl::filtering_mode::linear);
45+
46+
// Extension: image descriptor
47+
sycl::ext::oneapi::experimental::image_descriptor desc(
48+
{width}, 1, sycl::image_channel_type::fp32);
49+
50+
// Host USM allocation
51+
float *imgMem = sycl::malloc_host<float>(width, ctxt);
52+
53+
if (imgMem == nullptr) {
54+
std::cerr << "Error allocating host USM!" << std::endl;
55+
return 1;
56+
}
57+
58+
// Initialize input data
59+
for (int i = 0; i < width; ++i) {
60+
imgMem[i] = static_cast<float>(i);
61+
}
62+
63+
// Extension: create the image and return the handle
64+
sycl::ext::oneapi::experimental::sampled_image_handle imgHandle =
65+
sycl::ext::oneapi::experimental::create_image(imgMem, 0 /* pitch */,
66+
samp, desc, dev, ctxt);
67+
68+
sycl::buffer<float, 1> buf((float *)out.data(), sycl::range<1>{width});
69+
q.submit([&](sycl::handler &cgh) {
70+
auto outAcc =
71+
buf.get_access<sycl::access_mode::write>(cgh, sycl::range<1>{width});
72+
73+
cgh.parallel_for<sample_host_usm_image_kernel>(
74+
sycl::nd_range<1>{{width}, {width}}, [=](sycl::nd_item<1> it) {
75+
size_t dim0 = it.get_local_id(0);
76+
77+
// Normalize coordinates -- +0.5 to look towards centre of pixel
78+
float fdim0 = float(dim0 + 0.5f) / (float)width;
79+
80+
// Extension: sample image data from handle
81+
float px = sycl::ext::oneapi::experimental::sample_image<float>(
82+
imgHandle, (float)fdim0);
83+
84+
outAcc[sycl::id<1>{dim0}] = px;
85+
});
86+
});
87+
88+
q.wait_and_throw();
89+
90+
// Extension: cleanup
91+
sycl::ext::oneapi::experimental::destroy_image_handle(imgHandle, dev, ctxt);
92+
sycl::free(imgMem, ctxt);
93+
} catch (sycl::exception e) {
94+
std::cerr << "SYCL exception caught! : " << e.what() << "\n";
95+
return 1;
96+
} catch (...) {
97+
std::cerr << "Unknown exception caught!\n";
98+
return 2;
99+
}
100+
101+
// collect and validate output
102+
bool validated = true;
103+
for (int i = 0; i < width; i++) {
104+
bool mismatch = false;
105+
if (out[i] != expected[i]) {
106+
mismatch = true;
107+
validated = false;
108+
}
109+
110+
if (mismatch) {
111+
#ifdef VERBOSE_PRINT
112+
std::cout << "Result mismatch! Expected: " << expected[i]
113+
<< ", Actual: " << out[i] << std::endl;
114+
#else
115+
break;
116+
#endif
117+
}
118+
}
119+
if (validated) {
120+
std::cout << "Test passed!" << std::endl;
121+
return 0;
122+
}
123+
124+
std::cout << "Test failed!" << std::endl;
125+
return 3;
126+
}

0 commit comments

Comments
 (0)