[SYCL] Fix build log for L0 plugin (#2479)

The build log was not preserved for failed compile and link
operations in the L0 plugin, and this prevented the build log
information from being included in the SYCL "compile_program_error"
exception.  Fix this and add a test.

Author: gmlueck

sycl/plugins/level_zero/pi_level_zero.cpp

@@ -2231,16 +2231,22 @@ pi_result piProgramLink(pi_context Context, pi_uint32 NumDevices,
       ZeHandles.push_back(Input->ZeModule);
     }
 
-    // Link all the modules together.  If this fails (or if we catch an
-    // exception below), we need to release the reference counts on the input
-    // modules, delete any copies, etc.
+    // Link all the modules together.
     ze_module_build_log_handle_t ZeBuildLog;
-    ZE_CALL(zeModuleDynamicLinkMock(ZeHandles.size(), ZeHandles.data(),
-                                    &ZeBuildLog));
+    ze_result_t ZeResult = ZE_CALL_NOCHECK(zeModuleDynamicLinkMock(
+        ZeHandles.size(), ZeHandles.data(), &ZeBuildLog));
 
     // Construct a new program object to represent the linked executable.  This
-    // new object holds a reference to all the input programs.
-    *RetProgram = new _pi_program(Context, std::move(Inputs), ZeBuildLog);
+    // new object holds a reference to all the input programs.  Note that we
+    // create this program object even if the link fails with "link failure"
+    // because we need the new program object to hold the buid log (which has
+    // the description of the failure).
+    if (ZeResult == ZE_RESULT_SUCCESS ||
+        ZeResult == ZE_RESULT_ERROR_MODULE_LINK_FAILURE) {
+      *RetProgram = new _pi_program(Context, std::move(Inputs), ZeBuildLog);
+    }
+    if (ZeResult != ZE_RESULT_SUCCESS)
+      return mapError(ZeResult);
   } catch (const std::bad_alloc &) {
     return PI_OUT_OF_HOST_MEMORY;
   } catch (...) {
@@ -2356,9 +2362,8 @@ static pi_result compileOrBuild(pi_program Program, pi_uint32 NumDevices,
   ze_device_handle_t ZeDevice = DeviceList[0]->ZeDevice;
   ze_context_handle_t ZeContext = Program->Context->ZeContext;
   ze_module_handle_t ZeModule;
-  ze_module_build_log_handle_t ZeBuildLog;
   ZE_CALL(zeModuleCreate(ZeContext, ZeDevice, &ZeModuleDesc, &ZeModule,
-                         &ZeBuildLog));
+                         &Program->ZeBuildLog));
 
   // Check if this module imports any symbols, which we need to know if we
   // end up linking this module later.  See comments in piProgramLink() for
@@ -2371,7 +2376,6 @@ static pi_result compileOrBuild(pi_program Program, pi_uint32 NumDevices,
   // The caller must set the State to Object or Exe as appropriate.
   Program->Code.reset();
   Program->ZeModule = ZeModule;
-  Program->ZeBuildLog = ZeBuildLog;
   return PI_SUCCESS;
 }
 

sycl/test/kernel-and-program/build-log.cpp

@@ -0,0 +1,51 @@
+// XFAIL: cuda
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple  %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+//==--- build-log.cpp - Test log message from faild build ----------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===--------------------------------------------------------------===//
+
+#include <CL/sycl.hpp>
+
+SYCL_EXTERNAL
+void symbol_that_does_not_exist();
+
+void test() {
+  cl::sycl::queue Queue;
+
+  // Submitting this kernel should result in a compile_program_error exception
+  // with a message indicating that "symbol_that_does_not_exist" is undefined.
+  auto Kernel = []() {
+#ifdef __SYCL_DEVICE_ONLY__
+    symbol_that_does_not_exist();
+#endif
+  };
+
+  std::string Msg;
+  int Result;
+
+  try {
+    Queue.submit([&](cl::sycl::handler &CGH) {
+      CGH.single_task<class SingleTask>(Kernel);
+    });
+    assert(false && "There must be compilation error");
+  } catch (const cl::sycl::compile_program_error &e) {
+    std::string Msg(e.what());
+    assert(Msg.find("symbol_that_does_not_exist") != std::string::npos);
+  } catch (...) {
+    assert(false && "There must be cl::sycl::compile_program_error");
+  }
+}
+
+int main() {
+  test();
+
+  return 0;
+}