[SYCL] Build kernel only for device 0 when using sycl::get_kernel_bundle on the platform with multiple devices

fengyuan14 commented 3 months ago

Describe the bug

Simple description. After calling get_kernel_bundle(const context& ctxt, const std::vector<kernel_id>& kernelIds);, the kernel only is built on device 0, even the kernel is compatible on device 1. Launching kernel on device 1 will raise runtime error PI_ERROR_INVALID_QUEUE.

The issue was filed internally. We have got the root cause. We synced with the SYCL compiler community, and they suggested using another API as an alternative solution to fix the issue temporarily. In parallel, the SYCL compiler community will fix the issue. The purpose of the issue here is to show Intel's transparency in PyTorch community.

Problematic spot: https://github.com/intel/llvm/blob/sycl/sycl/source/detail/program_manager/program_manager.cpp#L2508-L2513

To reproduce

#include <sycl/sycl.hpp>

class KernelName;

int main() {
  sycl::platform platform;
  auto devices = platform.get_devices();
  std::cout << "Init devices, count: " << devices.size() << std::endl;

  auto dev0 = devices[0];
  auto dev1 = devices[1];
  auto ctx = sycl::context({dev0, dev1});
  auto q1 = sycl::queue(ctx, dev1);

  sycl::kernel_id kid = sycl::get_kernel_id<KernelName>();
    sycl::kernel_bundle KernelBundleExecutable =
        sycl::get_kernel_bundle<sycl::bundle_state::executable>(ctx, {kid});
  q1.submit([=](sycl::handler &cgh) {
    cgh.use_kernel_bundle(KernelBundleExecutable);
    cgh.single_task<class KernelName>([=]() {
    });
  });
  q1.wait();

Environment

Compiler: pytorch-gpu-dev-0.5, Intel(R) oneAPI DPC+/C+ Compiler 2024.1.3 (2024.1.3.20240604)

UMD: PyTorch Prerequisites for Intel® GPUs

ii libigc-dev 1.0.16900.24-914~22.04 amd64 Intel graphics co> ii libigc-tools 1.0.16900.24-914~22.04 amd64 Intel graphics co> ii libigc1 1.0.16900.24-914~22.04 amd64 Intel graphics co> ii libigdfcl-dev 1.0.16900.24-914~22.04 amd64 Intel graphics co> ii libigdfcl1 1.0.16900.24-914~22.04 amd64 Intel graphics co> ii intel-level-zero-gpu 1.3.29735.27-914~22.04 amd64 Intel(R) Graphics>

Additional context

No response

fengyuan14 commented 3 months ago

https://github.com/intel/torch-xpu-ops/issues/745

againull commented 1 month ago

Addressed in https://github.com/intel/llvm/pull/15546

intel / llvm