Revert "Roll back to original usage of sycl::get_kernel_bundle (#1935)" (#2026)

To solve #2025.
This reverts commit 83a1555 to quickly
bypass the performance regression caused by usage of
sycl::get_kernel_bundle .

Author: CuiYifeng

src/comm/DeviceProperties.h

@@ -16,8 +16,13 @@ static int64_t syclMaxWorkGroupSize(
   auto dev = q.get_device();
 
   auto kid = ::sycl::get_kernel_id<KernelClass>();
-  auto kbundle =
-      ::sycl::get_kernel_bundle<::sycl::bundle_state::executable>(ctx, {kid});
+  // The kernel won't be built for devices except for the first device.
+  // Launching kernel on devices except for the first device will raise
+  // runtime error. Here is an alternative as a temporary solution to
+  // provide an extra hint to SYCL runtime.
+  // https://github.com/intel/llvm/issues/15127
+  auto kbundle = ::sycl::get_kernel_bundle<::sycl::bundle_state::executable>(
+      ctx, {dev}, {kid});
 
   ::sycl::kernel k = kbundle.get_kernel(kid);
   return k.get_info<::sycl::info::kernel_device_specific::work_group_size>(dev);