diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index cd5759c263397..fe4f3b9c83e4d 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -2919,12 +2919,20 @@ 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 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(); + IdQueryRangeProp = + DeviceImageImpl->get_bin_image_ref()->getIdQueriesRangeProperties(); + } } else if (DeviceImageImpl != nullptr) { 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 0000000000000..3c3fd73ca0615 --- /dev/null +++ b/sycl/test-e2e/Regression/interop_kernel_large_range.cpp @@ -0,0 +1,61 @@ +// REQUIRES: opencl, opencl_icd, gpu, aspect-usm_shared_allocations + +// 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; +}