Report build log in exception message

Author: sommerlukas

dpctl/_backend.pxd

@@ -458,9 +458,11 @@ cdef extern from "syclinterface/dpctl_sycl_kernel_bundle_interface.h":
     cdef struct DPCTLBuildOptionList
     cdef struct DPCTLKernelNameList
     cdef struct DPCTLVirtualHeaderList
+    cdef struct DPCTLKernelBuildLog
     ctypedef DPCTLBuildOptionList* DPCTLBuildOptionListRef
     ctypedef DPCTLKernelNameList* DPCTLKernelNameListRef
     ctypedef DPCTLVirtualHeaderList* DPCTLVirtualHeaderListRef
+    ctypedef DPCTLKernelBuildLog* DPCTLKernelBuildLogRef
 
     cdef DPCTLBuildOptionListRef DPCTLBuildOptionList_Create()
     cdef void DPCTLBuildOptionList_Delete(DPCTLBuildOptionListRef Ref)
@@ -478,13 +480,18 @@ cdef extern from "syclinterface/dpctl_sycl_kernel_bundle_interface.h":
                                             const char *Name,
                                             const char *Content)
 
+    cdef DPCTLKernelBuildLogRef DPCTLKernelBuildLog_Create()
+    cdef void DPCTLKernelBuildLog_Delete(DPCTLKernelBuildLogRef Ref)
+    cdef const char *DPCTLKernelBuildLog_Get(DPCTLKernelBuildLogRef)
+
     cdef DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource(
         const DPCTLSyclContextRef Ctx,
         const DPCTLSyclDeviceRef Dev,
         const char *Source,
         DPCTLVirtualHeaderListRef Headers,
         DPCTLKernelNameListRef Names,
-        DPCTLBuildOptionListRef BuildOptions)
+        DPCTLBuildOptionListRef BuildOptions,
+        DPCTLKernelBuildLogRef BuildLog)
 
     cdef DPCTLSyclKernelRef DPCTLKernelBundle_GetSyclKernel(
                                                 DPCTLSyclKernelBundleRef KBRef,

dpctl/program/_program.pyx

@@ -42,6 +42,10 @@ from dpctl._backend cimport (  # noqa: E211, E402;
     DPCTLKernel_GetPreferredWorkGroupSizeMultiple,
     DPCTLKernel_GetPrivateMemSize,
     DPCTLKernel_GetWorkGroupSize,
+    DPCTLKernelBuildLog_Create,
+    DPCTLKernelBuildLog_Delete,
+    DPCTLKernelBuildLog_Get,
+    DPCTLKernelBuildLogRef,
     DPCTLKernelBundle_Copy,
     DPCTLKernelBundle_CreateFromOCLSource,
     DPCTLKernelBundle_CreateFromSpirv,
@@ -383,7 +387,8 @@ cpdef create_program_from_sycl_source(SyclQueue q, unicode source,
 
         Raises:
             SyclProgramCompilationError
-                If a SYCL kernel bundle could not be created.
+                If a SYCL kernel bundle could not be created. The exception
+                message contains the build log for more details.
     """
     cdef DPCTLSyclKernelBundleRef KBref
     cdef DPCTLSyclContextRef CRef = q.get_sycl_context().get_context_ref()
@@ -397,6 +402,7 @@ cpdef create_program_from_sycl_source(SyclQueue q, unicode source,
     cdef const char* sName
     cdef bytes bContent
     cdef const char* sContent
+    cdef const char* buildLogContent
     for opt in copts:
         if not isinstance(opt, unicode):
             DPCTLBuildOptionList_Delete(BuildOpts)
@@ -430,19 +436,26 @@ cpdef create_program_from_sycl_source(SyclQueue q, unicode source,
         sContent = <const char*>bContent
         DPCTLVirtualHeaderList_Append(VirtualHeaders, sName, sContent)
 
+    cdef DPCTLKernelBuildLogRef BuildLog
+    BuildLog = DPCTLKernelBuildLog_Create()
+
     KBref = DPCTLKernelBundle_CreateFromSYCLSource(CRef, DRef, Src,
                                                    VirtualHeaders, KernelNames,
-                                                   BuildOpts)
+                                                   BuildOpts, BuildLog)
 
     if KBref is NULL:
+        buildLogContent = DPCTLKernelBuildLog_Get(BuildLog)
+        buildLogStr = str(buildLogContent, "utf-8")
         DPCTLBuildOptionList_Delete(BuildOpts)
         DPCTLKernelNameList_Delete(KernelNames)
         DPCTLVirtualHeaderList_Delete(VirtualHeaders)
-        raise SyclProgramCompilationError()
+        DPCTLKernelBuildLog_Delete(BuildLog)
+        raise SyclProgramCompilationError(buildLogStr)
 
     DPCTLBuildOptionList_Delete(BuildOpts)
     DPCTLKernelNameList_Delete(KernelNames)
     DPCTLVirtualHeaderList_Delete(VirtualHeaders)
+    DPCTLKernelBuildLog_Delete(BuildLog)
 
     return SyclProgram._create(KBref, True)
 

dpctl/tests/test_sycl_program.py

@@ -268,9 +268,9 @@ def test_create_program_from_invalid_src_ocl():
 
 def test_create_program_from_sycl_source():
     try:
-        q = dpctl.SyclQueue("level_zero")
+        q = dpctl.SyclQueue("opencl")
     except dpctl.SyclQueueCreationError:
-        pytest.skip("No Level-zero queue is available")
+        pytest.skip("No OpenCL queue is available")
 
     if not q.get_sycl_device().can_compile("sycl"):
         pytest.skip("SYCL source compilation not supported")
@@ -374,3 +374,40 @@ def test_create_program_from_sycl_source():
         assert type(cmsgsz) is int
 
     _check_cpython_api_SyclProgram_GetKernelBundleRef(prog)
+
+
+def test_create_program_from_invalid_src_sycl():
+    try:
+        q = dpctl.SyclQueue("opencl")
+    except dpctl.SyclQueueCreationError:
+        pytest.skip("No OpenCL queue is available")
+
+    if not q.get_sycl_device().can_compile("sycl"):
+        pytest.skip("SYCL source compilation not supported")
+
+    sycl_source = """
+    #include <sycl/sycl.hpp>
+
+    namespace syclext = sycl::ext::oneapi::experimental;
+
+    extern "C" SYCL_EXTERNAL
+    SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclext::nd_range_kernel<1>))
+    void vector_add(int* in1, int* in2, int* out){
+        sycl::nd_item<1> item =
+                        sycl::ext::oneapi::this_work_item::get_nd_item<1>();
+        size_t globalID = item.get_global_linear_id()
+        out[globalID] = in1[globalID] + in2[globalID];
+    }
+    """
+    try:
+        _ = dpctl.program.create_program_from_sycl_source(
+            q,
+            sycl_source,
+            headers=[],
+            registered_names=[],
+            copts=[],
+        )
+        assert False
+    except dpctl_prog.SyclProgramCompilationError as prog_error:
+        print(str(prog_error))
+        assert "error: expected ';' at end of declaration" in str(prog_error)

libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h

@@ -132,6 +132,7 @@ DPCTLKernelBundle_Copy(__dpctl_keep const DPCTLSyclKernelBundleRef KBRef);
 typedef struct DPCTLBuildOptionList *DPCTLBuildOptionListRef;
 typedef struct DPCTLKernelNameList *DPCTLKernelNameListRef;
 typedef struct DPCTLVirtualHeaderList *DPCTLVirtualHeaderListRef;
+typedef struct DPCTLKernelBuildLog *DPCTLKernelBuildLogRef;
 
 /*!
  * @brief Create an empty list of build options.
@@ -145,7 +146,7 @@ __dpctl_give DPCTLBuildOptionListRef DPCTLBuildOptionList_Create();
 /*!
  * @brief Frees the DPCTLBuildOptionListRef pointer.
  *
- * @param   KBRef           Opaque pointer to a list of build options
+ * @param   Ref           Opaque pointer to a list of build options
  * @ingroup KernelBundleInterface
  */
 DPCTL_API void
@@ -173,7 +174,7 @@ __dpctl_give DPCTLKernelNameListRef DPCTLKernelNameList_Create();
 /*!
  * @brief Frees the DPCTLKernelNameListRef pointer.
  *
- * @param   KBRef           Opaque pointer to a list of kernels to register
+ * @param   Ref           Opaque pointer to a list of kernels to register
  * @ingroup KernelBundleInterface
  */
 DPCTL_API void
@@ -200,7 +201,7 @@ __dpctl_give DPCTLVirtualHeaderListRef DPCTLVirtualHeaderList_Create();
 /*!
  * @brief Frees the DPCTLVirtualHeaderListRef pointer.
  *
- * @param   KBRef           Opaque pointer to a list of virtual headers
+ * @param   Ref           Opaque pointer to a list of virtual headers
  * @ingroup KernelBundleInterface
  */
 DPCTL_API void
@@ -218,6 +219,33 @@ void DPCTLVirtualHeaderList_Append(__dpctl_keep DPCTLVirtualHeaderListRef Ref,
                                    __dpctl_keep const char *Name,
                                    __dpctl_keep const char *Content);
 
+/*!
+ * @brief Create an empty kernel build log.
+ *
+ * @return Opaque pointer to the kernel build log.
+ * @ingroup KernelBundleInterface
+ */
+DPCTL_API __dpctl_give DPCTLKernelBuildLogRef DPCTLKernelBuildLog_Create();
+
+/*!
+ * @brief Frees the DPCTLKernelBuildLogRef pointer.
+ *
+ * @param   Ref           Opaque pointer to a kernel build log.
+ * @ingroup KernelBundleInterface
+ */
+DPCTL_API
+void DPCTLKernelBuildLog_Delete(__dpctl_take DPCTLKernelBuildLogRef Ref);
+
+/*!
+ * @brief Get the content of the build log.
+ *
+ * @param Ref   Opaque pointer to the kernel build log.
+ * @return      Content of the build log
+ * @ingroup     KernelBundleInterface
+ */
+DPCTL_API const char *
+DPCTLKernelBuildLog_Get(__dpctl_keep DPCTLKernelBuildLogRef);
+
 /*!
  * @brief Create a SYCL kernel bundle from an SYCL kernel source string.
  *
@@ -238,7 +266,8 @@ __dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource(
     __dpctl_keep const char *Source,
     __dpctl_keep DPCTLVirtualHeaderListRef Headers,
     __dpctl_keep DPCTLKernelNameListRef Names,
-    __dpctl_keep DPCTLBuildOptionListRef BuildOptions);
+    __dpctl_keep DPCTLBuildOptionListRef BuildOptions,
+    __dpctl_keep DPCTLKernelBuildLogRef BuildLog);
 
 /*!
  * @brief Returns the SyclKernel with given name from the program compiled from

libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp

@@ -833,6 +833,27 @@ void DPCTLVirtualHeaderList_Append(__dpctl_keep DPCTLVirtualHeaderListRef Ref,
     reinterpret_cast<virtual_header_list_t *>(Ref)->push_back(Header);
 }
 
+using kernel_build_log_t = std::string;
+
+__dpctl_give DPCTLKernelBuildLogRef DPCTLKernelBuildLog_Create()
+{
+    auto BuildLog =
+        std::unique_ptr<kernel_build_log_t>(new kernel_build_log_t(""));
+    auto *RetVal = reinterpret_cast<DPCTLKernelBuildLogRef>(BuildLog.get());
+    BuildLog.release();
+    return RetVal;
+}
+
+void DPCTLKernelBuildLog_Delete(__dpctl_take DPCTLKernelBuildLogRef Ref)
+{
+    delete reinterpret_cast<kernel_build_log_t *>(Ref);
+}
+
+const char *DPCTLKernelBuildLog_Get(__dpctl_keep DPCTLKernelBuildLogRef Ref)
+{
+    return reinterpret_cast<kernel_build_log_t *>(Ref)->data();
+}
+
 namespace syclex = sycl::ext::oneapi::experimental;
 
 #if defined(SYCL_EXT_ONEAPI_KERNEL_COMPILER) &&                                \
@@ -880,7 +901,8 @@ __dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource(
     __dpctl_keep const char *Source,
     __dpctl_keep DPCTLVirtualHeaderListRef Headers,
     __dpctl_keep DPCTLKernelNameListRef Names,
-    __dpctl_keep DPCTLBuildOptionListRef BuildOptions)
+    __dpctl_keep DPCTLBuildOptionListRef BuildOptions,
+    __dpctl_keep DPCTLKernelBuildLogRef BuildLog)
 {
 #if (SUPPORTS_SYCL_COMPILATION > 0)
     context *SyclCtx = unwrap<context>(Ctx);
@@ -941,7 +963,8 @@ __dpctl_give DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSYCLSource(
         return wrap<kernel_bundle<bundle_state::executable>>(
             ResultBundle.release());
     } catch (const std::exception &e) {
-        error_handler(e, __FILE__, __func__, __LINE__);
+        auto *RawBuildLog = reinterpret_cast<kernel_build_log_t *>(BuildLog);
+        *RawBuildLog = e.what();
         return nullptr;
     }
 #else

libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp

@@ -335,11 +335,13 @@ struct TestSYCLKernelBundleFromSource
             DPCTLVirtualHeaderListRef VHRef = DPCTLVirtualHeaderList_Create();
             DPCTLVirtualHeaderList_Append(VHRef, Header1Name, header1_content);
             DPCTLVirtualHeaderList_Append(VHRef, Header2Name, header2_content);
+            DPCTLKernelBuildLogRef KBLRef = DPCTLKernelBuildLog_Create();
             KBRef = DPCTLKernelBundle_CreateFromSYCLSource(
-                CRef, DRef, sycl_source, VHRef, KNRef, BORef);
+                CRef, DRef, sycl_source, VHRef, KNRef, BORef, KBLRef);
             DPCTLVirtualHeaderList_Delete(VHRef);
             DPCTLKernelNameList_Delete(KNRef);
             DPCTLBuildOptionList_Delete(BORef);
+            DPCTLKernelBuildLog_Delete(KBLRef);
         }
     }