PR feedback

Author: Jakub Chlanda

sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp

@@ -376,7 +376,6 @@ static void setupTool(ClangTool &Tool, const std::string &DPCPPRoot,
 static void setGPUTarget(BinaryFormat Format,
                          SmallVector<std::string> &CommandLine) {
   auto [CPU, _] = Translator::getTargetCPUAndFeatureAttrs(nullptr, "", Format);
-  CommandLine.push_back("-fsycl");
   if (Format == BinaryFormat::PTX) {
     CommandLine.push_back("-fsycl-targets=nvptx64-nvidia-cuda");
     CommandLine.push_back("-Xsycl-target-backend");
@@ -603,9 +602,9 @@ Error jit_compiler::linkDeviceLibraries(llvm::Module &Module,
     return createStringError("Could not determine list of device libraries: %s",
                              BuildLog.c_str());
   }
-  const bool IsGPUTarget =
+  const bool IsCudaHIP =
       Format == BinaryFormat::PTX || Format == BinaryFormat::AMDGCN;
-  if (IsGPUTarget) {
+  if (IsCudaHIP) {
     // Based on the OS and the format decide on the version of libspirv.
     // NOTE: this will be problematic if cross-compiling between OSes.
     std::string Libclc{"clc/"};
@@ -638,7 +637,7 @@ Error jit_compiler::linkDeviceLibraries(llvm::Module &Module,
   }
 
   // For GPU targets we need to link against vendor provided libdevice.
-  if (IsGPUTarget) {
+  if (IsCudaHIP) {
     Triple T{Module.getTargetTriple()};
     Driver D{(Twine(DPCPPRoot) + "/bin/clang++").str(), T.getTriple(), Diags};
     auto [CPU, _] =
@@ -680,10 +679,9 @@ Error jit_compiler::linkDeviceLibraries(llvm::Module &Module,
     }
 #endif
     for (auto &LibDeviceFile : LibDeviceFiles) {
-      auto Res = LinkInLib(LibDeviceFile);
       // llvm::Error converts to false on success.
-      if (Res) {
-        return Res;
+      if (auto Error = LinkInLib(LibDeviceFile)) {
+        return Error;
       }
     }
   }

sycl/source/detail/device_image_impl.hpp

@@ -705,8 +705,8 @@ class device_image_impl {
       return ::jit_compiler::BinaryFormat::AMDGCN;
     default:
       throw sycl::exception(
-          sycl::make_error_code(sycl::errc::feature_not_supported),
-          "Backend unsupported by kernel fusion");
+          sycl::make_error_code(sycl::errc::invalid),
+          "Backend does not support kernel_compiler extension");
     }
   }
 

sycl/test-e2e/KernelCompiler/sycl_imf.cpp

@@ -1,4 +1,4 @@
-//==--- sycl.cpp --- kernel_compiler extension imf tests -------------------==//
+//==--- sycl_imf.cpp --- kernel_compiler extension imf tests ---------------==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -37,6 +37,7 @@ void imf_kernel(float *ptr) {
 
 int main() {
 #ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER
+  sycl::queue q;
   namespace syclex = sycl::ext::oneapi::experimental;
   using source_kb = sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source>;
   using exe_kb = sycl::kernel_bundle<sycl::bundle_state::executable>;