Skip to content

Commit 7b0d944

Browse files
cperkinsintelcperkinsintel
committed
making same range odor changes to DX12 bindless image tests
1 parent 2a290e0 commit 7b0d944

File tree

5 files changed

+56
-23
lines changed

5 files changed

+56
-23
lines changed

‎sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_2D_arithmetic.cpp‎

Lines changed: 16 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -461,7 +461,10 @@ int runTest(
461461
sycl::image_channel_type syclType = syclOverride.has_value()
462462
? syclOverride.value()
463463
: getSyclChannelType<T>();
464-
syclexp::image_descriptor imgDesc(sycl::range<2>(width, height), channels,
464+
// Bindless images use X,Y,Z order where X is dimension 0, Y is dimension 1.
465+
// This differs from SYCL 2020 "fastest-incrementing" convention.
466+
// See sycl_ext_oneapi_bindless_images.asciidoc for details.
467+
syclexp::image_descriptor imgDesc(sycl::range<2>(height, width), channels,
465468
syclType);
466469

467470
auto imgMemA = syclexp::map_external_image_memory(
@@ -498,9 +501,12 @@ int runTest(
498501

499502
kernelEvent = q.submit([&](sycl::handler &h) {
500503
h.depends_on(waitEvents);
501-
h.parallel_for(sycl::range<2>(width, height), [=](sycl::item<2> item) {
502-
int x = item.get_id(0);
503-
int y = item.get_id(1);
504+
// Range order matches image descriptor: (height, width)
505+
h.parallel_for(sycl::range<2>(height, width), [=](sycl::item<2> item) {
506+
// Extract x from dimension 1, y from dimension 0 to match bindless
507+
// coordinate order
508+
int x = item.get_id(1);
509+
int y = item.get_id(0);
504510

505511
bool isUnorm = (syclType == sycl::image_channel_type::unorm_int8);
506512
using Vec4 = sycl::vec<float, 4>;
@@ -561,9 +567,12 @@ int runTest(
561567

562568
kernelEvent = q.submit([&](sycl::handler &h) {
563569
h.depends_on(waitEvents);
564-
h.parallel_for(sycl::range<2>(width, height), [=](sycl::item<2> item) {
565-
int x = item.get_id(0);
566-
int y = item.get_id(1);
570+
// Range order matches image descriptor: (height, width)
571+
h.parallel_for(sycl::range<2>(height, width), [=](sycl::item<2> item) {
572+
// Extract x from dimension 1, y from dimension 0 to match bindless
573+
// coordinate order
574+
int x = item.get_id(1);
575+
int y = item.get_id(0);
567576

568577
bool isUnorm = (syclType == sycl::image_channel_type::unorm_int8);
569578
using Vec4 = sycl::vec<float, 4>;

‎sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_2D_read.cpp‎

Lines changed: 10 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -220,7 +220,10 @@ int runTest(
220220
sycl::image_channel_type syclType = syclOverride.has_value()
221221
? syclOverride.value()
222222
: getSyclChannelType<T>();
223-
syclexp::image_descriptor imgDesc(sycl::range<2>(width, height), channels,
223+
// Bindless images use X,Y,Z order where X is dimension 0, Y is dimension 1.
224+
// This differs from SYCL 2020 "fastest-incrementing" convention.
225+
// See sycl_ext_oneapi_bindless_images.asciidoc for details.
226+
syclexp::image_descriptor imgDesc(sycl::range<2>(height, width), channels,
224227
syclType);
225228

226229
syclexp::image_mem_handle devHandle = syclexp::map_external_image_memory(
@@ -275,9 +278,12 @@ int runTest(
275278
h.depends_on(dependencyEvent);
276279

277280
sycl::accessor outAcc(checkBuf, h, sycl::write_only);
278-
h.parallel_for(sycl::range<2>(width, height), [=](sycl::item<2> item) {
279-
int x = item.get_id(0);
280-
int y = item.get_id(1);
281+
// Range order matches image descriptor: (height, width)
282+
h.parallel_for(sycl::range<2>(height, width), [=](sycl::item<2> item) {
283+
// Extract x from dimension 1, y from dimension 0 to match bindless
284+
// coordinate order
285+
int x = item.get_id(1);
286+
int y = item.get_id(0);
281287
size_t base = (y * width + x) * channels;
282288

283289
if (useSampled) {

‎sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_2D_write_unsampled.cpp‎

Lines changed: 10 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -201,7 +201,10 @@ int runTest(
201201
sycl::image_channel_type syclType = syclOverride.has_value()
202202
? syclOverride.value()
203203
: getSyclChannelType<T>();
204-
syclexp::image_descriptor imgDesc(sycl::range<2>(width, height), channels,
204+
// Bindless images use X,Y,Z order where X is dimension 0, Y is dimension 1.
205+
// This differs from SYCL 2020 "fastest-incrementing" convention.
206+
// See sycl_ext_oneapi_bindless_images.asciidoc for details.
207+
syclexp::image_descriptor imgDesc(sycl::range<2>(height, width), channels,
205208
syclType);
206209

207210
syclexp::image_mem_handle devHandle = syclexp::map_external_image_memory(
@@ -232,9 +235,12 @@ int runTest(
232235
if (useSemaphores)
233236
h.depends_on(dependencyEvent);
234237

235-
h.parallel_for(sycl::range<2>(width, height), [=](sycl::item<2> item) {
236-
int x = item.get_id(0);
237-
int y = item.get_id(1);
238+
// Range order matches image descriptor: (height, width)
239+
h.parallel_for(sycl::range<2>(height, width), [=](sycl::item<2> item) {
240+
// Extract x from dimension 1, y from dimension 0 to match bindless
241+
// coordinate order
242+
int x = item.get_id(1);
243+
int y = item.get_id(0);
238244

239245
size_t index = y * width + x;
240246
size_t totalPixels = width * height;

‎sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_3D_read.cpp‎

Lines changed: 10 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -213,7 +213,10 @@ int runTest(
213213
sycl::image_channel_type syclType = syclOverride.has_value()
214214
? syclOverride.value()
215215
: getSyclChannelType<T>();
216-
syclexp::image_descriptor imgDesc(sycl::range<3>(width, height, depth),
216+
// Bindless images use X,Y,Z order where X is dimension 0, Y is dimension 1,
217+
// Z is dimension 2. This differs from SYCL 2020 "fastest-incrementing"
218+
// convention. See sycl_ext_oneapi_bindless_images.asciidoc for details.
219+
syclexp::image_descriptor imgDesc(sycl::range<3>(depth, height, width),
217220
channels, syclType);
218221

219222
syclexp::image_mem_handle devHandle = syclexp::map_external_image_memory(
@@ -261,11 +264,14 @@ int runTest(
261264
h.depends_on(dependencyEvent);
262265

263266
sycl::accessor outAcc(checkBuf, h, sycl::write_only);
267+
// Range order matches image descriptor: (depth, height, width)
264268
h.parallel_for(
265-
sycl::range<3>(width, height, depth), [=](sycl::item<3> item) {
266-
int x = item.get_id(0);
269+
sycl::range<3>(depth, height, width), [=](sycl::item<3> item) {
270+
// Extract x from dimension 2, y from dimension 1, z from dimension
271+
// 0 to match bindless coordinate order
272+
int x = item.get_id(2);
267273
int y = item.get_id(1);
268-
int z = item.get_id(2);
274+
int z = item.get_id(0);
269275
size_t base = (z * width * height + y * width + x) * channels;
270276

271277
if (useSampled) {

‎sycl/test-e2e/bindless_images/dx12_interop/D3D12_sycl_interop_3D_write_unsampled.cpp‎

Lines changed: 10 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -160,7 +160,10 @@ int runTest(
160160
sycl::image_channel_type syclType = syclOverride.has_value()
161161
? syclOverride.value()
162162
: getSyclChannelType<T>();
163-
syclexp::image_descriptor imgDesc(sycl::range<3>(width, height, depth),
163+
// Bindless images use X,Y,Z order where X is dimension 0, Y is dimension 1,
164+
// Z is dimension 2. This differs from SYCL 2020 "fastest-incrementing"
165+
// convention. See sycl_ext_oneapi_bindless_images.asciidoc for details.
166+
syclexp::image_descriptor imgDesc(sycl::range<3>(depth, height, width),
164167
channels, syclType);
165168

166169
syclexp::image_mem_handle devHandle = syclexp::map_external_image_memory(
@@ -191,11 +194,14 @@ int runTest(
191194
if (useSemaphores)
192195
h.depends_on(dependencyEvent);
193196

197+
// Range order matches image descriptor: (depth, height, width)
194198
h.parallel_for(
195-
sycl::range<3>(width, height, depth), [=](sycl::item<3> item) {
196-
int x = item.get_id(0);
199+
sycl::range<3>(depth, height, width), [=](sycl::item<3> item) {
200+
// Extract x from dimension 2, y from dimension 1, z from dimension
201+
// 0 to match bindless coordinate order
202+
int x = item.get_id(2);
197203
int y = item.get_id(1);
198-
int z = item.get_id(2);
204+
int z = item.get_id(0);
199205

200206
size_t index = (z * width * height) + (y * width) + x;
201207
size_t totalPixels = width * height * depth;

0 commit comments

Comments
 (0)