[SYCL] Preserve original message and code of kernel/program build result

sycl/include/CL/sycl/detail/kernel_program_cache.hpp

@@ -25,29 +25,42 @@ namespace detail {
 class context_impl;
 class KernelProgramCache {
 public:
-  /// Denotes pointer to some entity with its state.
+  /// Denotes build error data. The data is filled in from cl::sycl::exception
+  /// class instance.
+  struct BuildError {
+    std::string Msg;
+    pi_int32 Code;
+
+    /// Equals to true if the Msg and Code are initialized. This flag is added
+    /// due to possibility of Code equal to zero even if there is a
+    /// cl::sycl::exception thrown
+    bool FilledIn;
+  };
+
+  /// Denotes pointer to some entity with its general state and build error.
   /// The pointer is not null if and only if the entity is usable.
   /// State of the entity is provided by the user of cache instance.
   /// Currently there is only a single user - ProgramManager class.
   template<typename T>
-  struct EntityWithState {
+  struct BuildResult {
     std::atomic<T *> Ptr;
     std::atomic<int> State;
+    BuildError Error;
 
-    EntityWithState(T* P, int S)
-      : Ptr{P}, State{S}
+    BuildResult(T* P, int S)
+      : Ptr{P}, State{S}, Error{"", 0, false}
     {}
   };
 
   using PiProgramT = std::remove_pointer<RT::PiProgram>::type;
   using PiProgramPtrT = std::atomic<PiProgramT *>;
-  using ProgramWithBuildStateT = EntityWithState<PiProgramT>;
+  using ProgramWithBuildStateT = BuildResult<PiProgramT>;
   using ProgramCacheT = std::map<OSModuleHandle, ProgramWithBuildStateT>;
   using ContextPtr = context_impl *;
 
   using PiKernelT = std::remove_pointer<RT::PiKernel>::type;
   using PiKernelPtrT = std::atomic<PiKernelT *>;
-  using KernelWithBuildStateT = EntityWithState<PiKernelT>;
+  using KernelWithBuildStateT = BuildResult<PiKernelT>;
   using KernelByNameT = std::map<string_class, KernelWithBuildStateT>;
   using KernelCacheT = std::map<RT::PiProgram, KernelByNameT>;
 

sycl/source/detail/program_manager/program_manager.cpp

@@ -110,22 +110,27 @@ DeviceImage &ProgramManager::getDeviceImage(OSModuleHandle M,
 }
 
 template <typename ExceptionT, typename RetT>
-RetT *
-waitUntilBuilt(KernelProgramCache &Cache,
-               KernelProgramCache::EntityWithState<RetT> *WithBuildState) {
+RetT *waitUntilBuilt(KernelProgramCache &Cache,
+                     KernelProgramCache::BuildResult<RetT> *BuildResult,
+                     bool &TryAgain) {
   // any thread which will find nullptr in cache will wait until the pointer
   // is not null anymore
-  Cache.waitUntilBuilt([WithBuildState]() {
-    int State = WithBuildState->State.load();
+  Cache.waitUntilBuilt([BuildResult]() {
+    int State = BuildResult->State.load();
 
     return State == BS_Done || State == BS_Failed;
   });
 
-  RetT *Result = WithBuildState->Ptr.load();
+  if (BuildResult->Error.FilledIn) {
+    const KernelProgramCache::BuildError &Error = BuildResult->Error;
+    throw ExceptionT(Error.Msg, Error.Code);
+  }
+
+  RetT *Result = BuildResult->Ptr.load();
 
-  if (!Result)
-    throw ExceptionT("The other thread tried to build the program/kernel but "
-                     "did not succeed.");
+  // if the result is still null then there was no SYCL exception and we may try
+  // to build kernel/program once more to generate the original exception
+  TryAgain = !Result;
 
   return Result;
 }
@@ -152,7 +157,7 @@ template <typename RetT, typename ExceptionT, typename KeyT, typename AcquireFT,
 RetT *getOrBuild(KernelProgramCache &KPCache, const KeyT &CacheKey,
                  AcquireFT &&Acquire, GetCacheFT &&GetCache, BuildFT &&Build) {
   bool InsertionTookPlace;
-  KernelProgramCache::EntityWithState<RetT> *WithState;
+  KernelProgramCache::BuildResult<RetT> *BuildResult;
 
   {
     auto LockedCache = Acquire(KPCache);
@@ -162,36 +167,65 @@ RetT *getOrBuild(KernelProgramCache &KPCache, const KeyT &CacheKey,
                       std::forward_as_tuple(nullptr, BS_InProgress));
 
     InsertionTookPlace = Inserted.second;
-    WithState = &Inserted.first->second;
+    BuildResult = &Inserted.first->second;
   }
 
   // no insertion took place, thus some other thread has already inserted smth
   // in the cache
   if (!InsertionTookPlace) {
-    return waitUntilBuilt<ExceptionT>(KPCache, WithState);
+    bool TryAgain = false;
+
+    for (;;) {
+      RetT *Result = waitUntilBuilt<ExceptionT>(KPCache, BuildResult, TryAgain);
+
+      if (TryAgain) {
+        // Previous build is failed. There was no SYCL exception though.
+        // We might try to build once more.
+        int Expected = BS_Failed;
+        int Desired = BS_InProgress;
+
+        if (BuildResult->State.compare_exchange_strong(Expected, Desired)) {
+          // this thread is the building thread now
+          break;
+        }
+
+        continue;
+      } else // no need to try once more
+        return Result;
+    }
   }
 
-  // only the building thread will run this, and only once.
+  // only the building thread will run this
   try {
     RetT *Desired = Build();
 
 #ifndef NDEBUG
     RetT *Expected = nullptr;
 
-    if (!WithState->Ptr.compare_exchange_strong(Expected, Desired))
+    if (!BuildResult->Ptr.compare_exchange_strong(Expected, Desired))
       // We've got a funny story here
       assert(false && "We've build an entity that is already have been built.");
 #else
-    WithState->Ptr.store(Desired);
+    BuildResult->Ptr.store(Desired);
 #endif
 
-    WithState->State.store(BS_Done);
+    BuildResult->State.store(BS_Done);
 
     KPCache.notifyAllBuild();
 
     return Desired;
+  } catch (const exception &Ex) {
+    BuildResult->Error.Msg = Ex.what();
+    BuildResult->Error.Code = Ex.get_cl_code();
+    BuildResult->Error.FilledIn = true;
+
+    BuildResult->State.store(BS_Failed);
+
+    KPCache.notifyAllBuild();
+
+    std::rethrow_exception(std::current_exception());
   } catch (...) {
-    WithState->State.store(BS_Failed);
+    BuildResult->State.store(BS_Failed);
 
     KPCache.notifyAllBuild();
 

sycl/test/kernel-and-program/cache-build-result.cpp

@@ -0,0 +1,61 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+
+#include <CL/sycl.hpp>
+// FIXME do not use internal methods in tests.
+#include <CL/sycl/detail/program_impl.hpp>
+
+namespace RT = cl::sycl::RT;
+namespace detail = cl::sycl::detail;
+namespace pi = detail::pi;
+
+using ProgramCacheT = detail::KernelProgramCache::ProgramCacheT;
+using KernelCacheT = detail::KernelProgramCache::KernelCacheT;
+
+class Functor {
+public:
+  void operator()(cl::sycl::item<1> Item) { (void)Item; }
+};
+
+SYCL_EXTERNAL
+void undefined();
+
+void test() {
+  cl::sycl::queue Queue;
+
+  auto Kernel = []() {
+#ifdef __SYCL_DEVICE_ONLY__
+    undefined();
+#endif
+  };
+
+  std::string Msg;
+  int Result;
+
+  for (int Idx = 0; Idx < 2; ++Idx) {
+    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) {
+      fprintf(stderr, "Exception: %s, %d\n", e.what(), e.get_cl_code());
+      if (Idx == 0) {
+        Msg = e.what();
+        Result = e.get_cl_code();
+      } else {
+        // Exception constantly adds info on its error code in the message
+        assert(Msg.find_first_of(e.what()) == 0 && "Exception text differs");
+        assert(Result == e.get_cl_code() && "Exception code differs");
+      }
+    } catch (...) {
+      assert(false && "There must be cl::sycl::compile_program_error");
+    }
+  }
+}
+
+int main() {
+  test();
+
+  return 0;
+}