From fff3a8f5f22e0eb6ba11de2837f3613a6edf0b06 Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Fri, 4 Sep 2026 02:12:30 +0200 Subject: [PATCH 1/4] Fix kernel range check for interop kernels --- sycl/source/detail/scheduler/commands.cpp | 16 +++++++++++----- 1 file changed, 11 insertions(+), 5 deletions(-) diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index cd5759c26339..952b17306b8d 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -2919,12 +2919,18 @@ void enqueueImpKernel( NDRDesc, static_cast(std::numeric_limits::max())); if (isRangeGreaterThanIntMax) { uint32_t IdQueryRangeProp = 0; - // Get device image of kernel and retrieve the id queries range property. - if (MSyclKernel != nullptr && !MSyclKernel->isInteropOrSourceBased()) { - DeviceImageImpl = &MSyclKernel->getDeviceImage(); - IdQueryRangeProp = - DeviceImageImpl->get_bin_image_ref()->getIdQueriesRangeProperties(); + if (MSyclKernel != nullptr) { + // Interop and source-based kernels do not have device images and they + // are compiled with assumption that id/range will fit in size_t. + if (MSyclKernel->isInteropOrSourceBased()) { + IdQueryRangeProp = 2; // size_t range + } + else { + DeviceImageImpl = &MSyclKernel->getDeviceImage(); + IdQueryRangeProp = + DeviceImageImpl->get_bin_image_ref()->getIdQueriesRangeProperties(); + } } else if (DeviceImageImpl != nullptr) { IdQueryRangeProp = DeviceImageImpl->get_bin_image_ref()->getIdQueriesRangeProperties(); From 599632d98da0e5aa876e715795053df1d2a25f41 Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Fri, 4 Sep 2026 02:43:24 +0200 Subject: [PATCH 2/4] Add e2e test --- sycl/source/detail/scheduler/commands.cpp | 3 +- .../Regression/interop_kernel_large_range.cpp | 61 +++++++++++++++++++ 2 files changed, 62 insertions(+), 2 deletions(-) create mode 100644 sycl/test-e2e/Regression/interop_kernel_large_range.cpp diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index 952b17306b8d..57aa30008150 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -2925,8 +2925,7 @@ void enqueueImpKernel( // are compiled with assumption that id/range will fit in size_t. if (MSyclKernel->isInteropOrSourceBased()) { IdQueryRangeProp = 2; // size_t range - } - else { + } else { DeviceImageImpl = &MSyclKernel->getDeviceImage(); IdQueryRangeProp = DeviceImageImpl->get_bin_image_ref()->getIdQueriesRangeProperties(); diff --git a/sycl/test-e2e/Regression/interop_kernel_large_range.cpp b/sycl/test-e2e/Regression/interop_kernel_large_range.cpp new file mode 100644 index 000000000000..f5d3320c154a --- /dev/null +++ b/sycl/test-e2e/Regression/interop_kernel_large_range.cpp @@ -0,0 +1,61 @@ +// REQUIRES: opencl, opencl_icd, gpu + +// RUN: %{build} -o %t.out %opencl_lib +// RUN: %{run} %t.out + +// Interop kernels have no device image and always use size_t +// id/range semantics, so the launch must succeed. + +#include +#include +#include +#include +#include + +#include + +using namespace sycl; + +const char KernelSource[] = "__kernel void big_range(__global uchar *out) {\ + if (get_global_id(0) == 0) out[0] = 1;\ + }"; +const size_t KernelSourceSize = sizeof(KernelSource); +const char *Sources[1] = {KernelSource}; + +int main() { + queue Q{}; + context Ctx = Q.get_context(); + device Dev = Q.get_device(); + + cl_int Err; + cl_program Prog = clCreateProgramWithSource( + get_native(Ctx), 1, Sources, &KernelSourceSize, &Err); + assert(Err == CL_SUCCESS); + cl_device_id CLDev = get_native(Dev); + Err = clBuildProgram(Prog, 1, &CLDev, nullptr, nullptr, nullptr); + assert(Err == CL_SUCCESS); + cl_kernel CLKernel = clCreateKernel(Prog, "big_range", &Err); + assert(Err == CL_SUCCESS); + + kernel Kernel = make_kernel(CLKernel, Ctx); + + uint8_t *Out = malloc_shared(1, Q); + + auto Launch = [&](size_t Global, size_t Local) { + *Out = 0; + Q.submit([&](handler &CGH) { + CGH.set_arg(0, Out); + CGH.parallel_for(nd_range<1>{Global, Local}, Kernel); + }).wait_and_throw(); + assert(*Out == 1 && "Kernel did not run"); + }; + + constexpr size_t IntMax = std::numeric_limits::max(); + Launch(1024, 16); // sanity: below INT_MAX + Launch(IntMax + 1ull, 16); // the regression: above INT_MAX + + free(Out, Q); + clReleaseKernel(CLKernel); + clReleaseProgram(Prog); + return 0; +} From 5280a1cce79f08ff1eccc259bbd6291654446339 Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Fri, 4 Sep 2026 03:06:53 +0200 Subject: [PATCH 3/4] Only skip the id-range check for kernels without SYCL metadata isInteropOrSourceBased() is also true for kernel_compiler images built from SYCL source. Those do have a device image carrying the idQueriesRange property and do truncate id queries to int by default, so skipping the check for them turns a clear exception into silently wrong ids. hasSYCLMetadata() excludes exactly the interop and non-SYCL-source kernels, which have no property to read. --- sycl/source/detail/scheduler/commands.cpp | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index 57aa30008150..fe4f3b9c83e4 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -2921,9 +2921,12 @@ void enqueueImpKernel( uint32_t IdQueryRangeProp = 0; // Get device image of kernel and retrieve the id queries range property. if (MSyclKernel != nullptr) { - // Interop and source-based kernels do not have device images and they - // are compiled with assumption that id/range will fit in size_t. - if (MSyclKernel->isInteropOrSourceBased()) { + // Interop kernels and kernels built from a non-SYCL source language + // (OpenCL C, SPIR-V) carry no SYCL metadata, so there is no id queries + // range property to read; their id queries are size_t by definition. + // Kernels built from SYCL source do have a device image with the + // property, so they must still be checked. + if (!MSyclKernel->hasSYCLMetadata()) { IdQueryRangeProp = 2; // size_t range } else { DeviceImageImpl = &MSyclKernel->getDeviceImage(); From c024c72d2a35ea5278df141e4f9cd7581bcb64d9 Mon Sep 17 00:00:00 2001 From: Udit Kumar Agarwal Date: Thu, 3 Sep 2026 18:26:31 -0700 Subject: [PATCH 4/4] Potential fix for pull request finding Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> --- sycl/test-e2e/Regression/interop_kernel_large_range.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/Regression/interop_kernel_large_range.cpp b/sycl/test-e2e/Regression/interop_kernel_large_range.cpp index f5d3320c154a..3c3fd73ca061 100644 --- a/sycl/test-e2e/Regression/interop_kernel_large_range.cpp +++ b/sycl/test-e2e/Regression/interop_kernel_large_range.cpp @@ -1,4 +1,4 @@ -// REQUIRES: opencl, opencl_icd, gpu +// REQUIRES: opencl, opencl_icd, gpu, aspect-usm_shared_allocations // RUN: %{build} -o %t.out %opencl_lib // RUN: %{run} %t.out