search

intel / torch-xpu-ops

Apache License 2.0
30 stars 21 forks source link

Performance: Host overhead: Severe host overhead in sycl::get_kernel_bundle. #1016

Open fengyuan14 opened 4 weeks ago

fengyuan14 commented 4 weeks ago

🐛 Describe the bug

We are using kernel specific max work group size to avoid platform compatibility issue. The routine is,

  auto kid = ::sycl::get_kernel_id<KernelClass>();
  auto kbundle = ::sycl::get_kernel_bundle<::sycl::bundle_state::executable>(
      ctx, {dev}, {kid});
  sycl::kernel k = kbundle.get_kernel(kid);
  int max_work_group_size =  k.get_info<::sycl::info::kernel_device_specific::work_group_size>(dev);

sycl::get_kernel_bundles gets severe host overhead. The data is as below, image

Impacts: All kernels in torch-xpu-ops launched with kernel specific max work group are impacted.

  1. 40us overhead is not acceptable for some single batch inference cases, since latency of kernels might be less than 10us.
  2. CUDA runtime usually spends ~6us for a kernel launch.

https://github.com/intel/llvm/issues/15824

Versions