[Level Zero] sycl::parallel_for with ranges larger than INT_MAX deadlocks or aborts

[masterleinad]

Describe the bug
Running

#include <iostream>
#include <CL/sycl.hpp>

int main(int, char**) {
   cl::sycl::default_selector device_selector;
   cl::sycl::queue queue(device_selector);
   std::cout << "Running on "
             << queue.get_device().get_info<cl::sycl::info::device::name>()
             << "\n";
   size_t N = INT_MAX; //breaks for CUDA
   // size_t N = 5000000000; // breaks for Intel
   sycl::range<1> range(N+1);
   auto parallel_for_event = queue.submit([&](sycl::handler& cgh) {
     cgh.parallel_for(range, [=](sycl::item<1> /*item*/) {});
   });

   return 0;
}

deadlocks on CUDA devices or gives

C++ exception with description "PI backend failed. PI backend returns: -54 (CL_INVALID_WORK_GROUP_SIZE) -54 (CL_INVALID_WORK_GROUP_SIZE)" thrown in the test body.

on Intel GPUs when compiled and run via

clang++ -fsycl -fsycl-unnamed-lambda -fno-sycl-id-queries-fit-in-int -fsycl-targets=nvptx64-nvidia-cuda-sycldevice && ./a.out

resp.

clang++ -fsycl -fsycl-unnamed-lambda -fno-sycl-id-queries-fit-in-int dummy.cc && ./a.out

Environment:

• OS: Linux
• Target device and vendor: Intel GPU, NVIDIA GPU
• DPC++ version: nightly release 20210621

[zjin-lcf]

Running on Intel(R) UHD Graphics P630 [0x3e96]
terminate called after throwing an instance of 'cl::sycl::runtime_error'
what(): Provided range is out of integer limits. Pass `-fno-sycl-id-queries-fit-in-int' to disable range check. -30 (CL_INVALID_VALUE)

Is it right that the real issue is sycl::range should not be limited to the range of an integer ? Thanks.

[bader]

Is it right that the real issue is sycl::range should not be limited to the range of an integer ? Thanks.

Yes. It's done for performance reasons and can be relaxed with -fno-sycl-id-queries-fit-in-int flag if needed.

[masterleinad]

As said in the initial post, I was using -fno-sycl-id-queries-fit-in-in already.

[masterleinad]

Enabling

size_t N = 5000000000lu; // breaks for Intel

the test still fails with

Running on Intel(R) Graphics [0x020a]
terminate called after throwing an instance of 'cl::sycl::runtime_error'
  what():  PI backend failed. PI backend returns: -54 (CL_INVALID_WORK_GROUP_SIZE) -54 (CL_INVALID_WORK_GROUP_SIZE)
Aborted

on Intel GPUs with a nightly build from 10/25.

[AerialMantis]

Now that #5095 is merged this should address the problem for the CUDA backend, so I will remove the CUDA label.

@bader I believe the remaining issue here is with the OpenCL/Level Zero backend.