From 7e3b2b0f48be1698395069bc109b0d52f76dc898 Mon Sep 17 00:00:00 2001 From: Duncan Brawley Date: Mon, 21 Jul 2025 13:41:56 +0100 Subject: [PATCH 1/4] [SYCL][BINDLESS][L0] Add support for usm max image width and height queries Added support for usm image max image width and height queries and updated test to cover this --- .../bindless_images/image_reqs_get_info.cpp | 36 +++++++++++++------ .../source/adapters/level_zero/device.cpp | 24 +++++++++++-- 2 files changed, 48 insertions(+), 12 deletions(-) diff --git a/sycl/test-e2e/bindless_images/image_reqs_get_info.cpp b/sycl/test-e2e/bindless_images/image_reqs_get_info.cpp index 758abf49ef36b..8c9f735f25ef3 100644 --- a/sycl/test-e2e/bindless_images/image_reqs_get_info.cpp +++ b/sycl/test-e2e/bindless_images/image_reqs_get_info.cpp @@ -1,7 +1,8 @@ // REQUIRES: aspect-ext_oneapi_bindless_images -// UNSUPPORTED: level_zero -// UNSUPPORTED-INTENDED: The feature is not implemented in the Level Zero stack. +// These features are only partly implemented in the Level Zero stack. +// Only max_image_linear_width and max_image_linear_height are supported in the +// Level Zero stack. // https://github.com/intel/llvm/issues/17663 // RUN: %{build} -o %t.out @@ -26,20 +27,35 @@ int main() { // These can be different depending on the device so we cannot test that the // values are correct // But we should at least see that the query itself works - auto pitchAlign = dev.get_info< - sycl::ext::oneapi::experimental::info::device::image_row_pitch_align>(); - auto maxPitch = dev.get_info(); - auto maxWidth = dev.get_info(); - auto maxheight = dev.get_info(); + + sycl::backend backend = dev.get_backend(); + + // Level Zero does not currently support these queries. Only CUDA does. + if (backend == sycl::backend::ext_oneapi_cuda) { + auto pitchAlign = dev.get_info(); + auto maxPitch = dev.get_info(); + } + + if (backend == sycl::backend::ext_oneapi_cuda || + backend == sycl::backend::ext_oneapi_level_zero) { + auto maxWidth = dev.get_info(); + auto maxheight = dev.get_info(); + } #ifdef VERBOSE_PRINT + if (backend == sycl::backend::ext_oneapi_cuda) { std::cout << "image_row_pitch_align: " << pitchAlign << "\nmax_image_linear_row_pitch: " << maxPitch << "\nmax_image_linear_width: " << maxWidth << "\nmax_image_linear_height: " << maxheight << "\n"; + } else if (backend == sycl::backend::ext_oneapi_level_zero) { + std::cout << "\nmax_image_linear_width: " << maxWidth + << "\nmax_image_linear_height: " << maxheight << "\n"; + } #endif } catch (sycl::exception e) { diff --git a/unified-runtime/source/adapters/level_zero/device.cpp b/unified-runtime/source/adapters/level_zero/device.cpp index 6392b3802a199..a5c02a3d80aa6 100644 --- a/unified-runtime/source/adapters/level_zero/device.cpp +++ b/unified-runtime/source/adapters/level_zero/device.cpp @@ -1143,9 +1143,29 @@ ur_result_t urDeviceGetInfo( return ReturnValue(Device->Platform->ZeBindlessImagesExtensionSupported && Device->ZeDeviceImageProperties->maxImageDims2D > 0); } + case UR_DEVICE_INFO_MAX_IMAGE_LINEAR_WIDTH_EXP: { + ze_device_image_properties_t imageProps = {}; + imageProps.stype = ZE_STRUCTURE_TYPE_DEVICE_IMAGE_PROPERTIES; + ze_device_pitched_alloc_exp_properties_t imageAllocProps = {}; + imageAllocProps.stype = ZE_STRUCTURE_TYPE_PITCHED_ALLOC_DEVICE_EXP_PROPERTIES; + imageProps.pNext = (void *)&imageAllocProps; + + ZE_CALL_NOCHECK(zeDeviceGetImageProperties, (ZeDevice, &imageProps)); + + return ReturnValue(imageAllocProps.maxImageLinearWidth); + } + case UR_DEVICE_INFO_MAX_IMAGE_LINEAR_HEIGHT_EXP: { + ze_device_image_properties_t imageProps = {}; + imageProps.stype = ZE_STRUCTURE_TYPE_DEVICE_IMAGE_PROPERTIES; + ze_device_pitched_alloc_exp_properties_t imageAllocProps = {}; + imageAllocProps.stype = ZE_STRUCTURE_TYPE_PITCHED_ALLOC_DEVICE_EXP_PROPERTIES; + imageProps.pNext = (void *)&imageAllocProps; + + ZE_CALL_NOCHECK(zeDeviceGetImageProperties, (ZeDevice, &imageProps)); + + return ReturnValue(imageAllocProps.maxImageLinearHeight); + } case UR_DEVICE_INFO_IMAGE_PITCH_ALIGN_EXP: - case UR_DEVICE_INFO_MAX_IMAGE_LINEAR_WIDTH_EXP: - case UR_DEVICE_INFO_MAX_IMAGE_LINEAR_HEIGHT_EXP: case UR_DEVICE_INFO_MAX_IMAGE_LINEAR_PITCH_EXP: UR_LOG(ERR, "Unsupported ParamName in urGetDeviceInfo"); UR_LOG(ERR, "ParamName=%{}(0x{})", ParamName, logger::toHex(ParamName)); From fee699947778a3476ea2fe7adca5f1c81537b8ab Mon Sep 17 00:00:00 2001 From: Duncan Brawley Date: Mon, 21 Jul 2025 13:54:32 +0100 Subject: [PATCH 2/4] format code --- .../test-e2e/bindless_images/image_reqs_get_info.cpp | 12 ++++++------ .../source/adapters/level_zero/device.cpp | 6 ++++-- 2 files changed, 10 insertions(+), 8 deletions(-) diff --git a/sycl/test-e2e/bindless_images/image_reqs_get_info.cpp b/sycl/test-e2e/bindless_images/image_reqs_get_info.cpp index 8c9f735f25ef3..361b7b788d83d 100644 --- a/sycl/test-e2e/bindless_images/image_reqs_get_info.cpp +++ b/sycl/test-e2e/bindless_images/image_reqs_get_info.cpp @@ -48,13 +48,13 @@ int main() { #ifdef VERBOSE_PRINT if (backend == sycl::backend::ext_oneapi_cuda) { - std::cout << "image_row_pitch_align: " << pitchAlign - << "\nmax_image_linear_row_pitch: " << maxPitch - << "\nmax_image_linear_width: " << maxWidth - << "\nmax_image_linear_height: " << maxheight << "\n"; + std::cout << "image_row_pitch_align: " << pitchAlign + << "\nmax_image_linear_row_pitch: " << maxPitch + << "\nmax_image_linear_width: " << maxWidth + << "\nmax_image_linear_height: " << maxheight << "\n"; } else if (backend == sycl::backend::ext_oneapi_level_zero) { - std::cout << "\nmax_image_linear_width: " << maxWidth - << "\nmax_image_linear_height: " << maxheight << "\n"; + std::cout << "\nmax_image_linear_width: " << maxWidth + << "\nmax_image_linear_height: " << maxheight << "\n"; } #endif diff --git a/unified-runtime/source/adapters/level_zero/device.cpp b/unified-runtime/source/adapters/level_zero/device.cpp index a5c02a3d80aa6..d51a2f1bf4b11 100644 --- a/unified-runtime/source/adapters/level_zero/device.cpp +++ b/unified-runtime/source/adapters/level_zero/device.cpp @@ -1147,7 +1147,8 @@ ur_result_t urDeviceGetInfo( ze_device_image_properties_t imageProps = {}; imageProps.stype = ZE_STRUCTURE_TYPE_DEVICE_IMAGE_PROPERTIES; ze_device_pitched_alloc_exp_properties_t imageAllocProps = {}; - imageAllocProps.stype = ZE_STRUCTURE_TYPE_PITCHED_ALLOC_DEVICE_EXP_PROPERTIES; + imageAllocProps.stype = + ZE_STRUCTURE_TYPE_PITCHED_ALLOC_DEVICE_EXP_PROPERTIES; imageProps.pNext = (void *)&imageAllocProps; ZE_CALL_NOCHECK(zeDeviceGetImageProperties, (ZeDevice, &imageProps)); @@ -1158,7 +1159,8 @@ ur_result_t urDeviceGetInfo( ze_device_image_properties_t imageProps = {}; imageProps.stype = ZE_STRUCTURE_TYPE_DEVICE_IMAGE_PROPERTIES; ze_device_pitched_alloc_exp_properties_t imageAllocProps = {}; - imageAllocProps.stype = ZE_STRUCTURE_TYPE_PITCHED_ALLOC_DEVICE_EXP_PROPERTIES; + imageAllocProps.stype = + ZE_STRUCTURE_TYPE_PITCHED_ALLOC_DEVICE_EXP_PROPERTIES; imageProps.pNext = (void *)&imageAllocProps; ZE_CALL_NOCHECK(zeDeviceGetImageProperties, (ZeDevice, &imageProps)); From db60a094fa5dbbe76922be072fa6abdae87304d2 Mon Sep 17 00:00:00 2001 From: Duncan Brawley Date: Mon, 21 Jul 2025 18:00:28 +0100 Subject: [PATCH 3/4] Fix variables not being available causing compiler error when verbose print is used in test --- .../bindless_images/image_reqs_get_info.cpp | 13 +++++++++---- 1 file changed, 9 insertions(+), 4 deletions(-) diff --git a/sycl/test-e2e/bindless_images/image_reqs_get_info.cpp b/sycl/test-e2e/bindless_images/image_reqs_get_info.cpp index 361b7b788d83d..a272bae030322 100644 --- a/sycl/test-e2e/bindless_images/image_reqs_get_info.cpp +++ b/sycl/test-e2e/bindless_images/image_reqs_get_info.cpp @@ -30,19 +30,24 @@ int main() { sycl::backend backend = dev.get_backend(); + size_t pitchAlign = 0; + size_t maxPitch = 0; + size_t maxWidth = 0; + size_t maxheight = 0; + // Level Zero does not currently support these queries. Only CUDA does. if (backend == sycl::backend::ext_oneapi_cuda) { - auto pitchAlign = dev.get_info(); - auto maxPitch = dev.get_info(); } if (backend == sycl::backend::ext_oneapi_cuda || backend == sycl::backend::ext_oneapi_level_zero) { - auto maxWidth = dev.get_info(); - auto maxheight = dev.get_info(); } From 0f225957fcad0c30ac43ebbea6dbc3ae9ca50ae9 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Peter=20=C5=BDu=C5=BEek?= Date: Tue, 22 Jul 2025 16:18:29 +0100 Subject: [PATCH 4/4] Fix formatting --- .../bindless_images/image_reqs_get_info.cpp | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/sycl/test-e2e/bindless_images/image_reqs_get_info.cpp b/sycl/test-e2e/bindless_images/image_reqs_get_info.cpp index a272bae030322..a6d548cdc902a 100644 --- a/sycl/test-e2e/bindless_images/image_reqs_get_info.cpp +++ b/sycl/test-e2e/bindless_images/image_reqs_get_info.cpp @@ -37,18 +37,18 @@ int main() { // Level Zero does not currently support these queries. Only CUDA does. if (backend == sycl::backend::ext_oneapi_cuda) { - pitchAlign = dev.get_info(); - maxPitch = dev.get_info(); + pitchAlign = dev.get_info(); + maxPitch = dev.get_info(); } if (backend == sycl::backend::ext_oneapi_cuda || backend == sycl::backend::ext_oneapi_level_zero) { - maxWidth = dev.get_info(); - maxheight = dev.get_info(); + maxWidth = dev.get_info(); + maxheight = dev.get_info(); } #ifdef VERBOSE_PRINT