Skip to content

Commit 580ab4d

Browse files
Seanst98SYCL Unbound Team
Seanst98
and
SYCL Unbound Team
authored
[SYCL][Bindless][4/4] Add tests for experimental implementation of SYCL bindless images extension (#10500)
# Experimental Implementation of SYCL Bindless Images Extension This commit stands as the fourth, and final, commit of four to make code review easier, covering the additional tests for bindless images to the e2e test suite. Implementing [revision 4 of the bindless images extension proposal](#9842). This will not compile or run until [PR3](#10454) has been merged. However, it can be reviewed simultaneously with PR3. ## Overview 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. ## Previous PRs * [1/4] [libclc](#9808) * [2/4] [PI/UR](#10112) * [3/4] [SYCL API](#10454) * [4/4] Tests <--- This one ## Authors 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 Co-authored-by: SYCL Unbound Team <sycl.unbound@codeplay.com>
1 parent 531aabf commit 580ab4d

27 files changed

+5427
-0
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: 204 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,204 @@
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 <iostream>
8+
#include <sycl/sycl.hpp>
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 imgMem(desc, dev, ctxt);
44+
45+
// Extension: query for bindless image support -- device aspects
46+
bool bindlessSupport = dev.has(sycl::aspect::ext_oneapi_bindless_images);
47+
bool bindlessSharedUsmSupport =
48+
dev.has(sycl::aspect::ext_oneapi_bindless_images_shared_usm);
49+
bool usm1dSupport =
50+
dev.has(sycl::aspect::ext_oneapi_bindless_images_1d_usm);
51+
bool usm2dSupport =
52+
dev.has(sycl::aspect::ext_oneapi_bindless_images_2d_usm);
53+
54+
#ifdef VERBOSE_PRINT
55+
std::cout << "bindless_images_support: " << bindlessSupport
56+
<< "\nbindless_images_shared_usm_support: "
57+
<< bindlessSharedUsmSupport
58+
<< "\nbindless_images_1d_usm_support: " 1dS
59+
<< "\nbindless_images_2d_usm_support: " << S << "\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 pitchAlign = dev.get_info<
68+
sycl::ext::oneapi::experimental::info::device::image_row_pitch_align>();
69+
auto maxPitch = dev.get_info<sycl::ext::oneapi::experimental::info::device::
70+
max_image_linear_row_pitch>();
71+
auto maxWidth = dev.get_info<sycl::ext::oneapi::experimental::info::device::
72+
max_image_linear_width>();
73+
auto maxheight = dev.get_info<sycl::ext::oneapi::experimental::info::
74+
device::max_image_linear_height>();
75+
76+
#ifdef VERBOSE_PRINT
77+
std::cout << "image_row_pitch_align: " << pitchAlign
78+
<< "\nmax_image_linear_row_pitch: " << maxPitch
79+
<< "\nmax_image_linear_width: " << maxWidth
80+
<< "\nmax_image_linear_height: " << maxheight << "\n";
81+
#endif
82+
83+
// Extension: query for bindless image mipmaps support -- aspects & info
84+
bool mipmapSupport = dev.has(sycl::aspect::ext_oneapi_mipmap);
85+
bool mipmapAnisotropySupport =
86+
dev.has(sycl::aspect::ext_oneapi_mipmap_anisotropy);
87+
float mipmapMaxAnisotropy = dev.get_info<
88+
sycl::ext::oneapi::experimental::info::device::mipmap_max_anisotropy>();
89+
bool mipmapLevelReferenceSupport =
90+
dev.has(sycl::aspect::ext_oneapi_mipmap_level_reference);
91+
92+
#ifdef VERBOSE_PRINT
93+
std::cout << "mipmapSupport: " << mipmapSupport
94+
<< "\nmipmapAnisotropySupport: " << mipmapAnisotropySupport
95+
<< "\nmipmapMaxAnisotropy: " << mipmapMaxAnisotropy
96+
<< "\nmipmapLevelReferenceSupport: "
97+
<< mipmapLevelReferenceSupport << "\n";
98+
#endif
99+
100+
// Extension: query for bindless image interop support -- device aspects
101+
bool interopMemoryImportSupport =
102+
dev.has(sycl::aspect::ext_oneapi_interop_memory_import);
103+
bool interopMemoryExportSupport =
104+
dev.has(sycl::aspect::ext_oneapi_interop_memory_export);
105+
bool interopSemaphoreImportSupport =
106+
dev.has(sycl::aspect::ext_oneapi_interop_semaphore_import);
107+
bool interopSemaphoreExportSupport =
108+
dev.has(sycl::aspect::ext_oneapi_interop_semaphore_export);
109+
110+
#ifdef VERBOSE_PRINT
111+
std::cout << "interopMemoryImportSupport: " << interopMemoryImportSupport
112+
<< "\ninteropMemoryExportSupport: " << interopMemoryExportSupport
113+
<< "\ninteropSemaphoreImportSupport: "
114+
<< interopSemaphoreImportSupport
115+
<< "\ninteropSemaphoreExportSupport: "
116+
<< interopSemaphoreExportSupport << "\n";
117+
#endif
118+
119+
auto rangeMem = imgMem.get_range();
120+
auto range = sycl::ext::oneapi::experimental::get_image_range(
121+
imgMem.get_handle(), dev, ctxt);
122+
if (rangeMem != range) {
123+
printString("handle and mem object disagree on image dimensions!\n");
124+
validated = false;
125+
}
126+
if (range[0] == width) {
127+
printString("width is correct!\n");
128+
} else {
129+
printString("width is NOT correct!\n");
130+
validated = false;
131+
}
132+
if (range[1] == height) {
133+
printString("height is correct!\n");
134+
} else {
135+
printString("height is NOT correct!\n");
136+
validated = false;
137+
}
138+
if (range[2] == depth) {
139+
printString("depth is correct!\n");
140+
} else {
141+
printString("depth is NOT correct!\n");
142+
validated = false;
143+
}
144+
145+
auto type = imgMem.get_type();
146+
if (type == sycl::ext::oneapi::experimental::image_type::standard) {
147+
printString("image type is correct!\n");
148+
} else {
149+
printString("image type is NOT correct!\n");
150+
validated = false;
151+
}
152+
153+
auto ctypeMem = imgMem.get_channel_type();
154+
auto ctype = sycl::ext::oneapi::experimental::get_image_channel_type(
155+
imgMem.get_handle(), dev, ctxt);
156+
if (ctypeMem != ctype) {
157+
printString("handle and mem object disagree on image channel type!\n");
158+
validated = false;
159+
}
160+
if (ctype == sycl::image_channel_type::signed_int32) {
161+
printString("channel type is correct!\n");
162+
} else {
163+
printString("channel type is NOT correct!\n");
164+
validated = false;
165+
}
166+
167+
auto corder = imgMem.get_channel_order();
168+
if (corder == sycl::image_channel_order::r) {
169+
printString("channel order is correct!\n");
170+
} else {
171+
printString("channel order is NOT correct!\n");
172+
validated = false;
173+
}
174+
175+
auto numchannelsMem = imgMem.get_num_channels();
176+
auto numchannels = sycl::ext::oneapi::experimental::get_image_num_channels(
177+
imgMem.get_handle(), dev, ctxt);
178+
if (numchannelsMem != numchannels) {
179+
printString("handle and mem object disagree on number of channels!\n");
180+
validated = false;
181+
}
182+
if (numchannels == 1) {
183+
printString("num channels is correct!\n");
184+
} else {
185+
printString("num channels is NOT correct!\n");
186+
validated = false;
187+
}
188+
189+
} catch (sycl::exception e) {
190+
std::cerr << "SYCL exception caught! : " << e.what() << "\n";
191+
return 1;
192+
} catch (...) {
193+
std::cerr << "Unknown exception caught!\n";
194+
return 2;
195+
}
196+
197+
if (validated) {
198+
std::cout << "Test Passed!\n";
199+
return 0;
200+
}
201+
202+
std::cout << "Test Failed!" << std::endl;
203+
return 3;
204+
}

‎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 <iostream>
8+
#include <sycl/sycl.hpp>
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 numLevels = 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, numLevels);
50+
51+
// Extension: allocate mipmap memory on device
52+
sycl::ext::oneapi::experimental::image_mem mipMem(desc, dev, ctxt);
53+
54+
// Extension: retrieve level 0
55+
sycl::ext::oneapi::experimental::image_mem_handle imgMem1 =
56+
mipMem.get_mip_level_mem_handle(0);
57+
58+
// Extension: copy over data to device at level 0
59+
q.ext_oneapi_copy(dataIn1.data(), imgMem1, desc);
60+
61+
// Extension: copy data to device at level 1
62+
q.ext_oneapi_copy(dataIn2.data(), mipMem.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)numLevels, 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(mipMem, 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(mipMem.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+
return 1;
112+
} catch (...) {
113+
std::cerr << "Unknown exception caught!\n";
114+
return 2;
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 3;
142+
}

0 commit comments

Comments
 (0)