From 714774daa16beab499f1744d95edecaa51ca7c44 Mon Sep 17 00:00:00 2001
From: Greg Lueck <gregory.m.lueck@intel.com>
Date: Tue, 15 Sep 2020 05:08:40 -0700
Subject: [PATCH] [SYCL] Fix build log for L0 plugin

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.
---
 sycl/plugins/level_zero/pi_level_zero.cpp  | 24 +++++-----
 sycl/test/kernel-and-program/build-log.cpp | 51 ++++++++++++++++++++++
 2 files changed, 65 insertions(+), 10 deletions(-)
 create mode 100644 sycl/test/kernel-and-program/build-log.cpp

diff --git a/sycl/plugins/level_zero/pi_level_zero.cpp b/sycl/plugins/level_zero/pi_level_zero.cpp
index ae141b3c1f78e..a9998b6abe11a 100644
--- a/sycl/plugins/level_zero/pi_level_zero.cpp
+++ b/sycl/plugins/level_zero/pi_level_zero.cpp
@@ -2187,16 +2187,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 (...) {
@@ -2310,9 +2316,8 @@ static pi_result compileOrBuild(pi_program Program, pi_uint32 NumDevices,
   ze_device_handle_t ZeDevice = Program->Context->Device->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
@@ -2325,7 +2330,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;
 }
 
diff --git a/sycl/test/kernel-and-program/build-log.cpp b/sycl/test/kernel-and-program/build-log.cpp
new file mode 100644
index 0000000000000..cc3c3bedbe28b
--- /dev/null
+++ b/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;
+}