[DeviceSanitizer] Print log when build/link fail (#17521)

yingcong-wu · web-flow · commit fae8cc2630b5 · 2025-04-02T10:04:57.000+02:00
Print logs when ur build/link fail. The main purpose for this patch is
for OpenMP program rather than SYCL program since SYCL runtime has
already do a good job printing those logs.

Also, this PR fix a bug in L0 adapter where the error log does not set
when `checkUnresolvedSymbols()` fails in `urProgramLinkExp()`. This bug
would prevent the actual ze_log that contains the failure
message(missing symbols) being added to the `UrProgram`.
diff --git a/sycl/test-e2e/AddressSanitizer/common/print-build-log.cpp b/sycl/test-e2e/AddressSanitizer/common/print-build-log.cpp
@@ -0,0 +1,37 @@
+// This test is adapted from sycl/test-e2e/KernelAndProgram/build-log.cpp
+// REQUIRES: linux, cpu || (gpu && level_zero)
+// RUN: %{build} %device_asan_flags -DGPU -o %t_gpu.out
+// RUN: %{build} %device_asan_flags -o %t.out
+// RUN: %{run} not --crash %if gpu %{ %t_gpu.out %} %else %{ %t.out %} 2>&1 | FileCheck %s
+
+#include <iostream>
+#include <sycl/detail/core.hpp>
+SYCL_EXTERNAL
+void symbol_that_does_not_exist();
+
+void test() {
+  sycl::queue Queue;
+
+  // Submitting this kernel should result in an exception with error code
+  // `sycl::errc::build` and a message indicating
+  // "PI_ERROR_BUILD_PROGRAM_FAILURE".
+  auto Kernel = []() {
+#ifdef __SYCL_DEVICE_ONLY__
+#ifdef GPU
+    asm volatile("undefined\n");
+#else  // GPU
+    symbol_that_does_not_exist();
+#endif // GPU
+#endif // __SYCL_DEVICE_ONLY__
+  };
+
+  Queue.submit(
+      [&](sycl::handler &CGH) { CGH.single_task<class SingleTask>(Kernel); });
+}
+
+// CHECK: <SANITIZER>[ERROR]: Printing build log for program
+
+int main() {
+  test();
+  return 0;
+}
diff --git a/unified-runtime/source/adapters/level_zero/program.cpp b/unified-runtime/source/adapters/level_zero/program.cpp
@@ -490,15 +490,16 @@ ur_result_t urProgramLinkExp(
       // because the ZeBuildLog tells which symbols are unresolved.
       if (ZeResult == ZE_RESULT_SUCCESS) {
         ZeResult = checkUnresolvedSymbols(ZeModule, &ZeBuildLog);
-        if (ZeResult != ZE_RESULT_SUCCESS) {
-          return ze2urResult(ZeResult);
-        }
+        UrResult = ze2urResult(ZeResult);
       }
       UrProgram->setZeModule(ZeDevice, ZeModule);
       UrProgram->setBuildLog(ZeDevice, ZeBuildLog);
       UrProgram->setState(ZeDevice, (UrResult == UR_RESULT_SUCCESS)
                                         ? ur_program_handle_t_::Exe
                                         : ur_program_handle_t_::Invalid);
+      if (ZeResult != ZE_RESULT_SUCCESS) {
+        return UrResult;
+      }
     }
   } catch (const std::bad_alloc &) {
     return UR_RESULT_ERROR_OUT_OF_HOST_MEMORY;
diff --git a/unified-runtime/source/loader/layers/sanitizer/asan/asan_ddi.cpp b/unified-runtime/source/loader/layers/sanitizer/asan/asan_ddi.cpp
@@ -322,7 +322,12 @@ __urdlllocal ur_result_t UR_APICALL urProgramBuild(
 
   getContext()->logger.debug("==== urProgramBuild");
 
-  UR_CALL(pfnProgramBuild(hContext, hProgram, pOptions));
+  auto UrRes = pfnProgramBuild(hContext, hProgram, pOptions);
+  if (UrRes != UR_RESULT_SUCCESS) {
+    auto Devices = GetDevices(hContext);
+    PrintUrBuildLog(hProgram, Devices.data(), Devices.size());
+    return UrRes;
+  }
 
   UR_CALL(getAsanInterceptor()->registerProgram(hProgram));
 
@@ -348,7 +353,12 @@ __urdlllocal ur_result_t UR_APICALL urProgramBuildExp(
 
   getContext()->logger.debug("==== urProgramBuildExp");
 
-  UR_CALL(pfnBuildExp(hProgram, numDevices, phDevices, pOptions));
+  auto UrRes = pfnBuildExp(hProgram, numDevices, phDevices, pOptions);
+  if (UrRes != UR_RESULT_SUCCESS) {
+    PrintUrBuildLog(hProgram, phDevices, numDevices);
+    return UrRes;
+  }
+
   UR_CALL(getAsanInterceptor()->registerProgram(hProgram));
 
   return UR_RESULT_SUCCESS;
@@ -375,7 +385,12 @@ __urdlllocal ur_result_t UR_APICALL urProgramLink(
 
   getContext()->logger.debug("==== urProgramLink");
 
-  UR_CALL(pfnProgramLink(hContext, count, phPrograms, pOptions, phProgram));
+  auto UrRes = pfnProgramLink(hContext, count, phPrograms, pOptions, phProgram);
+  if (UrRes != UR_RESULT_SUCCESS) {
+    auto Devices = GetDevices(hContext);
+    PrintUrBuildLog(*phProgram, Devices.data(), Devices.size());
+    return UrRes;
+  }
 
   UR_CALL(getAsanInterceptor()->insertProgram(*phProgram));
   UR_CALL(getAsanInterceptor()->registerProgram(*phProgram));
@@ -408,8 +423,12 @@ ur_result_t UR_APICALL urProgramLinkExp(
 
   getContext()->logger.debug("==== urProgramLinkExp");
 
-  UR_CALL(pfnProgramLinkExp(hContext, numDevices, phDevices, count, phPrograms,
-                            pOptions, phProgram));
+  auto UrRes = pfnProgramLinkExp(hContext, numDevices, phDevices, count,
+                                 phPrograms, pOptions, phProgram);
+  if (UrRes != UR_RESULT_SUCCESS) {
+    PrintUrBuildLog(*phProgram, phDevices, numDevices);
+    return UrRes;
+  }
 
   UR_CALL(getAsanInterceptor()->insertProgram(*phProgram));
   UR_CALL(getAsanInterceptor()->registerProgram(*phProgram));
diff --git a/unified-runtime/source/loader/layers/sanitizer/msan/msan_ddi.cpp b/unified-runtime/source/loader/layers/sanitizer/msan/msan_ddi.cpp
@@ -265,7 +265,12 @@ ur_result_t urProgramBuild(
 
   getContext()->logger.debug("==== urProgramBuild");
 
-  UR_CALL(pfnProgramBuild(hContext, hProgram, pOptions));
+  auto UrRes = pfnProgramBuild(hContext, hProgram, pOptions);
+  if (UrRes != UR_RESULT_SUCCESS) {
+    auto Devices = GetDevices(hContext);
+    PrintUrBuildLog(hProgram, Devices.data(), Devices.size());
+    return UrRes;
+  }
 
   UR_CALL(getMsanInterceptor()->registerProgram(hProgram));
 
@@ -287,7 +292,12 @@ ur_result_t urProgramBuildExp(
 
   getContext()->logger.debug("==== urProgramBuildExp");
 
-  UR_CALL(pfnBuildExp(hProgram, numDevices, phDevices, pOptions));
+  auto UrRes = pfnBuildExp(hProgram, numDevices, phDevices, pOptions);
+  if (UrRes != UR_RESULT_SUCCESS) {
+    PrintUrBuildLog(hProgram, phDevices, numDevices);
+    return UrRes;
+  }
+
   UR_CALL(getMsanInterceptor()->registerProgram(hProgram));
 
   return UR_RESULT_SUCCESS;
@@ -310,7 +320,12 @@ ur_result_t urProgramLink(
 
   getContext()->logger.debug("==== urProgramLink");
 
-  UR_CALL(pfnProgramLink(hContext, count, phPrograms, pOptions, phProgram));
+  auto UrRes = pfnProgramLink(hContext, count, phPrograms, pOptions, phProgram);
+  if (UrRes != UR_RESULT_SUCCESS) {
+    auto Devices = GetDevices(hContext);
+    PrintUrBuildLog(*phProgram, Devices.data(), Devices.size());
+    return UrRes;
+  }
 
   UR_CALL(getMsanInterceptor()->insertProgram(*phProgram));
   UR_CALL(getMsanInterceptor()->registerProgram(*phProgram));
@@ -339,8 +354,12 @@ ur_result_t urProgramLinkExp(
 
   getContext()->logger.debug("==== urProgramLinkExp");
 
-  UR_CALL(pfnProgramLinkExp(hContext, numDevices, phDevices, count, phPrograms,
-                            pOptions, phProgram));
+  auto UrRes = pfnProgramLinkExp(hContext, numDevices, phDevices, count,
+                                 phPrograms, pOptions, phProgram);
+  if (UrRes != UR_RESULT_SUCCESS) {
+    PrintUrBuildLog(*phProgram, phDevices, numDevices);
+    return UrRes;
+  }
 
   UR_CALL(getMsanInterceptor()->insertProgram(*phProgram));
   UR_CALL(getMsanInterceptor()->registerProgram(*phProgram));
diff --git a/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_utils.cpp b/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_utils.cpp
@@ -256,4 +256,36 @@ ur_result_t EnqueueUSMBlockingSet(ur_queue_handle_t Queue, void *Ptr,
       Queue, Ptr, 1, &Value, Size, NumEvents, EventWaitList, OutEvent);
 }
 
+void PrintUrBuildLog(ur_program_handle_t hProgram,
+                     ur_device_handle_t *phDevices, size_t numDevices) {
+  getContext()->logger.error("Printing build log for program {}",
+                             (void *)hProgram);
+  for (size_t i = 0; i < numDevices; i++) {
+    std::vector<char> LogBuf;
+    size_t LogSize = 0;
+    auto hDevice = phDevices[i];
+
+    auto UrRes = getContext()->urDdiTable.Program.pfnGetBuildInfo(
+        hProgram, hDevice, UR_PROGRAM_BUILD_INFO_LOG, 0, nullptr, &LogSize);
+    if (UrRes != UR_RESULT_SUCCESS) {
+      getContext()->logger.error("For device {}: failed to get build log size.",
+                                 (void *)hDevice);
+      continue;
+    }
+
+    LogBuf.resize(LogSize);
+    UrRes = getContext()->urDdiTable.Program.pfnGetBuildInfo(
+        hProgram, hDevice, UR_PROGRAM_BUILD_INFO_LOG, LogSize, LogBuf.data(),
+        nullptr);
+    if (UrRes != UR_RESULT_SUCCESS) {
+      getContext()->logger.error("For device {}: failed to get build log.",
+                                 (void *)hDevice);
+      continue;
+    }
+
+    getContext()->logger.error("For device {}:\n{}", (void *)hDevice,
+                               LogBuf.data());
+  }
+}
+
 } // namespace ur_sanitizer_layer
diff --git a/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_utils.hpp b/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_utils.hpp
@@ -65,4 +65,7 @@ EnqueueUSMBlockingSet(ur_queue_handle_t Queue, void *Ptr, char Value,
                       const ur_event_handle_t *EventWaitList = nullptr,
                       ur_event_handle_t *OutEvent = nullptr);
 
+void PrintUrBuildLog(ur_program_handle_t hProgram,
+                     ur_device_handle_t *phDevices, size_t numDevices);
+
 } // namespace ur_sanitizer_layer
