Skip to content

Commit 159e9f6

Browse files
SYCL Unbound TeamSeanst98
SYCL Unbound Team
authored and
Seanst98
committed
[SYCL][Bindless][4/4] Add experimental implementation of SYCL bindless images extension
This commit stands as the fourth, and final, commit of four to make code review easier, mostly covering the changes made to the e2e tests with the additional tests for bindless images. The bindless images extension provides a new interface for allocating, creating, and accessing images in SYCL. Image memory allocation is seperated from image handle creation, and image handles can be passed to kernels without requesting access through accessors. This approach provides much more flexibility to the user, as well as enabling programs to implement features that were impossible to implement using standard SYCL images, such as a texture atlas. In addition to providing a new interface for images, this extension also provides initial experimental support for importing external memory into SYCL. Co-authored-by: Isaac Ault <isaac.ault@codeplay.com> Co-authored-by: Hugh Bird <hugh.bird@codeplay.com> Co-authored-by: Duncan Brawley <duncan.brawley@codeplay.com> Co-authored-by: Przemek Malon <przemek.malon@codeplay.com> Co-authored-by: Chedy Najjar <chedy.najjar@codeplay.com> Co-authored-by: Sean Stirling <sean.stirling@codeplay.com> Co-authored-by: Peter Zuzek <peter@codeplay.com> Implement revision 4 of the bindless images extension proposal: #9842
1 parent 99e8dbe commit 159e9f6

27 files changed

+5464
-4
lines changed

‎sycl/test-e2e/CMakeLists.txt‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -24,6 +24,8 @@ endif() # Standalone.
2424
find_package(Threads REQUIRED)
2525
set(SYCL_THREADS_LIB ${CMAKE_THREAD_LIBS_INIT})
2626

27+
find_package(Vulkan)
28+
2729
if(NOT LLVM_LIT)
2830
find_program(LLVM_LIT
2931
NAMES llvm-lit lit.py lit

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

Lines changed: 206 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,206 @@
1+
// REQUIRES: linux
2+
// REQUIRES: cuda
3+
4+
// RUN: %clangxx -fsycl -fsycl-targets=%{sycl_triple} %s -o %t.out
5+
// RUN: %t.out
6+
7+
#include <CL/sycl.hpp>
8+
#include <iostream>
9+
10+
// Uncomment to print additional test information
11+
// #define VERBOSE_PRINT
12+
13+
void printString(std::string name) {
14+
#ifdef VERBOSE_PRINT
15+
std::cout << name;
16+
#endif
17+
}
18+
19+
int main() {
20+
21+
sycl::device dev;
22+
sycl::queue q(dev);
23+
auto ctxt = q.get_context();
24+
25+
size_t height = 13;
26+
size_t width = 7;
27+
size_t depth = 11;
28+
29+
bool validated = true;
30+
31+
try {
32+
// Submit dummy kernel to let the runtime decide the backend (CUDA)
33+
// Without this, the default Level Zero backend is active
34+
q.submit([&](sycl::handler &cgh) { cgh.single_task([]() {}); });
35+
36+
// Extension: image descriptor - can use the same for both images
37+
sycl::ext::oneapi::experimental::image_descriptor desc(
38+
{width, height, depth}, sycl::image_channel_order::r,
39+
sycl::image_channel_type::signed_int32);
40+
41+
// Extension: returns the device pointer to the allocated memory
42+
// Input images memory
43+
sycl::ext::oneapi::experimental::image_mem img_mem_0(desc, dev, ctxt);
44+
45+
// Extension: query for bindless image support -- device aspects
46+
bool bindless_support = dev.has(sycl::aspect::ext_oneapi_bindless_images);
47+
bool bindless_shared_usm_support =
48+
dev.has(sycl::aspect::ext_oneapi_bindless_images_shared_usm);
49+
bool usm_1d_support =
50+
dev.has(sycl::aspect::ext_oneapi_bindless_images_1d_usm);
51+
bool usm_2d_support =
52+
dev.has(sycl::aspect::ext_oneapi_bindless_images_2d_usm);
53+
54+
#ifdef VERBOSE_PRINT
55+
std::cout << "bindless_images_support: " << bindless_support
56+
<< "\nbindless_images_shared_usm_support: "
57+
<< bindless_shared_usm_support
58+
<< "\nbindless_images_1d_usm_support: " << usm_1d_support
59+
<< "\nbindless_images_2d_usm_support: " << usm_2d_support << "\n";
60+
#endif
61+
62+
// Extension: get pitch alignment information from device -- device info
63+
// Make sure our pitch alignment queries work properly
64+
// These can be different depending on the device so we cannot test that the
65+
// values are correct
66+
// But we should at least see that the query itself works
67+
auto pitch_align = dev.get_info<
68+
sycl::ext::oneapi::experimental::info::device::image_pitch_align>();
69+
auto max_pitch = dev.get_info<sycl::ext::oneapi::experimental::info::
70+
device::max_image_linear_pitch>();
71+
auto max_width = dev.get_info<sycl::ext::oneapi::experimental::info::
72+
device::max_image_linear_width>();
73+
auto max_height = dev.get_info<sycl::ext::oneapi::experimental::info::
74+
device::max_image_linear_height>();
75+
76+
#ifdef VERBOSE_PRINT
77+
std::cout << "image_pitch_align: " << pitch_align
78+
<< "\nmax_image_linear_pitch: " << max_pitch
79+
<< "\nmax_image_linear_width: " << max_width
80+
<< "\nmax_image_linear_height: " << max_height << "\n";
81+
#endif
82+
83+
// Extension: query for bindless image mipmaps support -- aspects & info
84+
bool mipmap_support = dev.has(sycl::aspect::ext_oneapi_mipmap);
85+
bool mipmap_anisotropy_support =
86+
dev.has(sycl::aspect::ext_oneapi_mipmap_anisotropy);
87+
float mipmap_max_anisotropy = dev.get_info<
88+
sycl::ext::oneapi::experimental::info::device::mipmap_max_anisotropy>();
89+
bool mipmap_level_reference_support =
90+
dev.has(sycl::aspect::ext_oneapi_mipmap_level_reference);
91+
92+
#ifdef VERBOSE_PRINT
93+
std::cout << "mipmap_support: " << mipmap_support
94+
<< "\nmipmap_anisotropy_support: " << mipmap_anisotropy_support
95+
<< "\nmipmap_max_anisotropy: " << mipmap_max_anisotropy
96+
<< "\nmipmap_level_reference_support: "
97+
<< mipmap_level_reference_support << "\n";
98+
#endif
99+
100+
// Extension: query for bindless image interop support -- device aspects
101+
bool interop_memory_import_support =
102+
dev.has(sycl::aspect::ext_oneapi_interop_memory_import);
103+
bool interop_memory_export_support =
104+
dev.has(sycl::aspect::ext_oneapi_interop_memory_export);
105+
bool interop_semaphore_import_support =
106+
dev.has(sycl::aspect::ext_oneapi_interop_semaphore_import);
107+
bool interop_semaphore_export_support =
108+
dev.has(sycl::aspect::ext_oneapi_interop_semaphore_export);
109+
110+
#ifdef VERBOSE_PRINT
111+
std::cout << "interop_memory_import_support: "
112+
<< interop_memory_import_support
113+
<< "\ninterop_memory_export_support: "
114+
<< interop_memory_export_support
115+
<< "\ninterop_semaphore_import_support: "
116+
<< interop_semaphore_import_support
117+
<< "\ninterop_semaphore_export_support: "
118+
<< interop_semaphore_export_support << "\n";
119+
#endif
120+
121+
auto rangeMem = img_mem_0.get_range();
122+
auto range = sycl::ext::oneapi::experimental::get_image_range(
123+
img_mem_0.get_handle(), dev, ctxt);
124+
if (rangeMem != range) {
125+
printString("handle and mem object disagree on image dimensions!\n");
126+
validated = false;
127+
}
128+
if (range[0] == width) {
129+
printString("width is correct!\n");
130+
} else {
131+
printString("width is NOT correct!\n");
132+
validated = false;
133+
}
134+
if (range[1] == height) {
135+
printString("height is correct!\n");
136+
} else {
137+
printString("height is NOT correct!\n");
138+
validated = false;
139+
}
140+
if (range[2] == depth) {
141+
printString("depth is correct!\n");
142+
} else {
143+
printString("depth is NOT correct!\n");
144+
validated = false;
145+
}
146+
147+
auto type = img_mem_0.get_type();
148+
if (type == sycl::ext::oneapi::experimental::image_type::standard) {
149+
printString("image type is correct!\n");
150+
} else {
151+
printString("image type is NOT correct!\n");
152+
validated = false;
153+
}
154+
155+
auto ctypeMem = img_mem_0.get_channel_type();
156+
auto ctype = sycl::ext::oneapi::experimental::get_image_channel_type(
157+
img_mem_0.get_handle(), dev, ctxt);
158+
if (ctypeMem != ctype) {
159+
printString("handle and mem object disagree on image channel type!\n");
160+
validated = false;
161+
}
162+
if (ctype == sycl::image_channel_type::signed_int32) {
163+
printString("channel type is correct!\n");
164+
} else {
165+
printString("channel type is NOT correct!\n");
166+
validated = false;
167+
}
168+
169+
auto corder = img_mem_0.get_channel_order();
170+
if (corder == sycl::image_channel_order::r) {
171+
printString("channel order is correct!\n");
172+
} else {
173+
printString("channel order is NOT correct!\n");
174+
validated = false;
175+
}
176+
177+
auto numchannelsMem = img_mem_0.get_num_channels();
178+
auto numchannels = sycl::ext::oneapi::experimental::get_image_num_channels(
179+
img_mem_0.get_handle(), dev, ctxt);
180+
if (numchannelsMem != numchannels) {
181+
printString("handle and mem object disagree on number of channels!\n");
182+
validated = false;
183+
}
184+
if (numchannels == 1) {
185+
printString("num channels is correct!\n");
186+
} else {
187+
printString("num channels is NOT correct!\n");
188+
validated = false;
189+
}
190+
191+
} catch (sycl::exception e) {
192+
std::cerr << "SYCL exception caught! : " << e.what() << "\n";
193+
exit(-1);
194+
} catch (...) {
195+
std::cerr << "Unknown exception caught!\n";
196+
exit(-1);
197+
}
198+
199+
if (validated) {
200+
std::cout << "Test Passed!\n";
201+
return 0;
202+
}
203+
204+
std::cout << "Test Failed!\n";
205+
return 1;
206+
}

‎sycl/test-e2e/bindless_images/mipmap/mipmap_read_1D.cpp‎

Lines changed: 142 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,142 @@
1+
// REQUIRES: linux
2+
// REQUIRES: cuda
3+
4+
// RUN: %clangxx -fsycl -fsycl-targets=%{sycl_triple} %s -o %t.out
5+
// RUN: %t.out
6+
7+
#include <CL/sycl.hpp>
8+
#include <iostream>
9+
10+
// Uncomment to print additional test information
11+
// #define VERBOSE_PRINT
12+
13+
class image_addition;
14+
15+
int main() {
16+
17+
sycl::device dev;
18+
sycl::queue q(dev);
19+
auto ctxt = q.get_context();
20+
21+
// declare image data
22+
constexpr size_t N = 16;
23+
std::vector<float> out(N);
24+
std::vector<float> expected(N);
25+
std::vector<sycl::float4> dataIn1(N);
26+
std::vector<sycl::float4> dataIn2(N / 2);
27+
std::vector<sycl::float4> copyOut(N / 2);
28+
int j = 0;
29+
for (int i = 0; i < N; i++) {
30+
expected[i] = i + (j + 10);
31+
if (i % 2)
32+
j++;
33+
dataIn1[i] = sycl::float4(i, i, i, i);
34+
if (i < (N / 2)) {
35+
dataIn2[i] = sycl::float4(i + 10, i + 10, i + 10, i + 10);
36+
copyOut[i] = sycl::float4{0, 0, 0, 0};
37+
}
38+
}
39+
40+
try {
41+
42+
size_t width = N;
43+
unsigned int num_levels = 2;
44+
45+
// Extension: image descriptor -- number of levels
46+
sycl::ext::oneapi::experimental::image_descriptor desc(
47+
{width}, sycl::image_channel_order::rgba,
48+
sycl::image_channel_type::fp32,
49+
sycl::ext::oneapi::experimental::image_type::mipmap, num_levels);
50+
51+
// Extension: allocate mipmap memory on device
52+
sycl::ext::oneapi::experimental::image_mem mip_mem(desc, dev, ctxt);
53+
54+
// Extension: retrieve level 0
55+
sycl::ext::oneapi::experimental::image_mem_handle img_mem1 =
56+
mip_mem.get_mip_level_mem_handle(0);
57+
58+
// Extension: copy over data to device at level 0
59+
q.ext_oneapi_copy(dataIn1.data(), img_mem1, desc);
60+
61+
// Extension: copy data to device at level 1
62+
q.ext_oneapi_copy(dataIn2.data(), mip_mem.get_mip_level_mem_handle(1),
63+
desc.get_mip_level_desc(1));
64+
q.wait_and_throw();
65+
66+
// Extension: define a sampler object -- extended mipmap attributes
67+
sycl::ext::oneapi::experimental::bindless_image_sampler samp(
68+
sycl::addressing_mode::mirrored_repeat,
69+
sycl::coordinate_normalization_mode::normalized,
70+
sycl::filtering_mode::nearest, sycl::filtering_mode::nearest, 0.0f,
71+
(float)num_levels, 8.0f);
72+
73+
// Extension: create a sampled image handle to represent the mipmap
74+
sycl::ext::oneapi::experimental::sampled_image_handle mipHandle =
75+
sycl::ext::oneapi::experimental::create_image(mip_mem, samp, desc, dev,
76+
ctxt);
77+
78+
sycl::buffer<float, 1> buf((float *)out.data(), N);
79+
q.submit([&](sycl::handler &cgh) {
80+
auto outAcc = buf.get_access<sycl::access_mode::write>(cgh, N);
81+
82+
cgh.parallel_for<image_addition>(N, [=](sycl::id<1> id) {
83+
float sum = 0;
84+
float x = float(id[0] + 0.5) / (float)N;
85+
// Extension: read mipmap level 0 with anisotropic filtering and level 1
86+
// with LOD
87+
sycl::float4 px1 =
88+
sycl::ext::oneapi::experimental::read_image<sycl::float4>(
89+
mipHandle, x, 0.0f, 0.0f);
90+
sycl::float4 px2 =
91+
sycl::ext::oneapi::experimental::read_image<sycl::float4>(mipHandle,
92+
x, 1.0f);
93+
94+
sum = px1[0] + px2[0];
95+
outAcc[id] = sum;
96+
});
97+
});
98+
99+
q.wait_and_throw();
100+
101+
// Extension: copy data from device
102+
q.ext_oneapi_copy(mip_mem.get_mip_level_mem_handle(1), copyOut.data(),
103+
desc.get_mip_level_desc(1));
104+
q.wait_and_throw();
105+
106+
// Extension: cleanup
107+
sycl::ext::oneapi::experimental::destroy_image_handle(mipHandle, dev, ctxt);
108+
109+
} catch (sycl::exception e) {
110+
std::cerr << "SYCL exception caught! : " << e.what() << "\n";
111+
exit(-1);
112+
} catch (...) {
113+
std::cerr << "Unknown exception caught!\n";
114+
exit(-1);
115+
}
116+
117+
// collect and validate output
118+
bool validated = true;
119+
for (int i = 0; i < N; i++) {
120+
bool mismatch = false;
121+
if (out[i] != expected[i]) {
122+
mismatch = true;
123+
validated = false;
124+
}
125+
126+
if (mismatch) {
127+
#ifdef VERBOSE_PRINT
128+
std::cout << "Result mismatch! Expected: " << expected[i]
129+
<< ", Actual: " << out[i] << std::endl;
130+
#else
131+
break;
132+
#endif
133+
}
134+
}
135+
if (validated) {
136+
std::cout << "Test passed!" << std::endl;
137+
return 0;
138+
}
139+
140+
std::cout << "Test failed!" << std::endl;
141+
return 1;
142+
}

0 commit comments

Comments
 (0)