From 2985a5280fcfac1bd006c41f757a3a1576309fdb Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Mon, 10 Feb 2020 15:12:27 +0300 Subject: [PATCH 01/10] [SYCL] Store original message and code of build result Signed-off-by: Sergey Kanaev --- .../CL/sycl/detail/kernel_program_cache.hpp | 6 ++++++ .../detail/program_manager/program_manager.cpp | 17 +++++++++++++++-- 2 files changed, 21 insertions(+), 2 deletions(-) diff --git a/sycl/include/CL/sycl/detail/kernel_program_cache.hpp b/sycl/include/CL/sycl/detail/kernel_program_cache.hpp index 470650d84f824..527956b492cf7 100644 --- a/sycl/include/CL/sycl/detail/kernel_program_cache.hpp +++ b/sycl/include/CL/sycl/detail/kernel_program_cache.hpp @@ -25,6 +25,11 @@ namespace detail { class context_impl; class KernelProgramCache { public: + struct BuildResultT { + std::string Msg; + cl_int Code; + }; + /// Denotes pointer to some entity with its state. /// 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. @@ -33,6 +38,7 @@ class KernelProgramCache { struct EntityWithState { std::atomic Ptr; std::atomic State; + std::unique_ptr BuildResult; EntityWithState(T* P, int S) : Ptr{P}, State{S} diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 72971d16a70ad..fcd4fac0adae2 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -121,11 +121,16 @@ waitUntilBuilt(KernelProgramCache &Cache, return State == BS_Done || State == BS_Failed; }); + if (WithBuildState->BuildResult.get()) { + using BuildResult = KernelProgramCache::BuildResultT; + const BuildResult &Res = *WithBuildState->BuildResult.get(); + throw ExceptionT(Res.Msg, Res.Code); + } + RetT *Result = WithBuildState->Ptr.load(); if (!Result) - throw ExceptionT("The other thread tried to build the program/kernel but " - "did not succeed."); + throw ExceptionT("Build of the program/kernel did not succeed previously."); return Result; } @@ -190,6 +195,14 @@ RetT *getOrBuild(KernelProgramCache &KPCache, const KeyT &CacheKey, KPCache.notifyAllBuild(); return Desired; + } catch (const exception &Ex) { + using BuildResultT = KernelProgramCache::BuildResultT; + WithState->BuildResult.reset(new BuildResultT{Ex.what(), Ex.get_cl_code()}); + WithState->State.store(BS_Failed); + + KPCache.notifyAllBuild(); + + std::rethrow_exception(std::current_exception()); } catch (...) { WithState->State.store(BS_Failed); From eb8f4f5e8d31feeabd5964729967bf5eb6ef6efd Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Tue, 11 Feb 2020 14:55:51 +0300 Subject: [PATCH 02/10] [SYCL] Fix review comments Signed-off-by: Sergey Kanaev --- sycl/include/CL/sycl/detail/kernel_program_cache.hpp | 8 ++++---- .../source/detail/program_manager/program_manager.cpp | 11 +++++------ 2 files changed, 9 insertions(+), 10 deletions(-) diff --git a/sycl/include/CL/sycl/detail/kernel_program_cache.hpp b/sycl/include/CL/sycl/detail/kernel_program_cache.hpp index 527956b492cf7..2a48bcc664b29 100644 --- a/sycl/include/CL/sycl/detail/kernel_program_cache.hpp +++ b/sycl/include/CL/sycl/detail/kernel_program_cache.hpp @@ -35,25 +35,25 @@ class KernelProgramCache { /// State of the entity is provided by the user of cache instance. /// Currently there is only a single user - ProgramManager class. template - struct EntityWithState { + struct EntityWithBuildResult { std::atomic Ptr; std::atomic State; std::unique_ptr BuildResult; - EntityWithState(T* P, int S) + EntityWithBuildResult(T* P, int S) : Ptr{P}, State{S} {} }; using PiProgramT = std::remove_pointer::type; using PiProgramPtrT = std::atomic; - using ProgramWithBuildStateT = EntityWithState; + using ProgramWithBuildStateT = EntityWithBuildResult; using ProgramCacheT = std::map; using ContextPtr = context_impl *; using PiKernelT = std::remove_pointer::type; using PiKernelPtrT = std::atomic; - using KernelWithBuildStateT = EntityWithState; + using KernelWithBuildStateT = EntityWithBuildResult; using KernelByNameT = std::map; using KernelCacheT = std::map; diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index fcd4fac0adae2..43b7daf8e5a6a 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -110,9 +110,9 @@ DeviceImage &ProgramManager::getDeviceImage(OSModuleHandle M, } template -RetT * -waitUntilBuilt(KernelProgramCache &Cache, - KernelProgramCache::EntityWithState *WithBuildState) { +RetT *waitUntilBuilt( + KernelProgramCache &Cache, + KernelProgramCache::EntityWithBuildResult *WithBuildState) { // any thread which will find nullptr in cache will wait until the pointer // is not null anymore Cache.waitUntilBuilt([WithBuildState]() { @@ -129,8 +129,7 @@ waitUntilBuilt(KernelProgramCache &Cache, RetT *Result = WithBuildState->Ptr.load(); - if (!Result) - throw ExceptionT("Build of the program/kernel did not succeed previously."); + assert(Result && "An exception should have been thrown"); return Result; } @@ -157,7 +156,7 @@ template *WithState; + KernelProgramCache::EntityWithBuildResult *WithState; { auto LockedCache = Acquire(KPCache); From 96e92a8ef978f4478c2de90e7ad0804293188dcc Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Tue, 11 Feb 2020 14:58:02 +0300 Subject: [PATCH 03/10] [SYCL] Add test Signed-off-by: Sergey Kanaev --- .../kernel-and-program/cache-build-result.cpp | 61 +++++++++++++++++++ 1 file changed, 61 insertions(+) create mode 100644 sycl/test/kernel-and-program/cache-build-result.cpp diff --git a/sycl/test/kernel-and-program/cache-build-result.cpp b/sycl/test/kernel-and-program/cache-build-result.cpp new file mode 100644 index 0000000000000..366e7f850076e --- /dev/null +++ b/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 +// FIXME do not use internal methods in tests. +#include + +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(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; +} From 54958894fc17ae2cb29cad255609d5315cbeddca Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Wed, 12 Feb 2020 16:48:06 +0300 Subject: [PATCH 04/10] [SYCL] Fix review comment Signed-off-by: Sergey Kanaev --- .../CL/sycl/detail/kernel_program_cache.hpp | 17 +++++---- .../program_manager/program_manager.cpp | 38 +++++++++---------- 2 files changed, 28 insertions(+), 27 deletions(-) diff --git a/sycl/include/CL/sycl/detail/kernel_program_cache.hpp b/sycl/include/CL/sycl/detail/kernel_program_cache.hpp index 2a48bcc664b29..e49484774767a 100644 --- a/sycl/include/CL/sycl/detail/kernel_program_cache.hpp +++ b/sycl/include/CL/sycl/detail/kernel_program_cache.hpp @@ -25,35 +25,36 @@ namespace detail { class context_impl; class KernelProgramCache { public: - struct BuildResultT { + struct BuildError { std::string Msg; cl_int Code; + bool FilledIn; }; - /// Denotes pointer to some entity with its state. + /// 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 - struct EntityWithBuildResult { + struct BuildResult { std::atomic Ptr; std::atomic State; - std::unique_ptr BuildResult; + BuildError Error; - EntityWithBuildResult(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::type; using PiProgramPtrT = std::atomic; - using ProgramWithBuildStateT = EntityWithBuildResult; + using ProgramWithBuildStateT = BuildResult; using ProgramCacheT = std::map; using ContextPtr = context_impl *; using PiKernelT = std::remove_pointer::type; using PiKernelPtrT = std::atomic; - using KernelWithBuildStateT = EntityWithBuildResult; + using KernelWithBuildStateT = BuildResult; using KernelByNameT = std::map; using KernelCacheT = std::map; diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 43b7daf8e5a6a..c669668b93ae1 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -110,24 +110,22 @@ DeviceImage &ProgramManager::getDeviceImage(OSModuleHandle M, } template -RetT *waitUntilBuilt( - KernelProgramCache &Cache, - KernelProgramCache::EntityWithBuildResult *WithBuildState) { +RetT *waitUntilBuilt(KernelProgramCache &Cache, + KernelProgramCache::BuildResult *BuildResult) { // 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; }); - if (WithBuildState->BuildResult.get()) { - using BuildResult = KernelProgramCache::BuildResultT; - const BuildResult &Res = *WithBuildState->BuildResult.get(); - throw ExceptionT(Res.Msg, Res.Code); + if (BuildResult->Error.FilledIn) { + const KernelProgramCache::BuildError &Error = BuildResult->Error; + throw ExceptionT(Error.Msg, Error.Code); } - RetT *Result = WithBuildState->Ptr.load(); + RetT *Result = BuildResult->Ptr.load(); assert(Result && "An exception should have been thrown"); @@ -156,7 +154,7 @@ template *WithState; + KernelProgramCache::BuildResult *BuildResult; { auto LockedCache = Acquire(KPCache); @@ -166,13 +164,13 @@ 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(KPCache, WithState); + return waitUntilBuilt(KPCache, BuildResult); } // only the building thread will run this, and only once. @@ -182,28 +180,30 @@ RetT *getOrBuild(KernelProgramCache &KPCache, const KeyT &CacheKey, #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); #endif - WithState->State.store(BS_Done); + BuildResult->State.store(BS_Done); KPCache.notifyAllBuild(); return Desired; } catch (const exception &Ex) { - using BuildResultT = KernelProgramCache::BuildResultT; - WithState->BuildResult.reset(new BuildResultT{Ex.what(), Ex.get_cl_code()}); - WithState->State.store(BS_Failed); + 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(); From 0c19b88ac29178d09977863b5eb2345ba7b2bfd7 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Thu, 13 Feb 2020 14:51:14 +0300 Subject: [PATCH 05/10] [SYCL] Fix review comments. Signed-off-by: Sergey Kanaev --- .../CL/sycl/detail/kernel_program_cache.hpp | 8 ++++- .../program_manager/program_manager.cpp | 32 ++++++++++++++++--- 2 files changed, 34 insertions(+), 6 deletions(-) diff --git a/sycl/include/CL/sycl/detail/kernel_program_cache.hpp b/sycl/include/CL/sycl/detail/kernel_program_cache.hpp index e49484774767a..cbfb85b33cf8e 100644 --- a/sycl/include/CL/sycl/detail/kernel_program_cache.hpp +++ b/sycl/include/CL/sycl/detail/kernel_program_cache.hpp @@ -25,9 +25,15 @@ namespace detail { class context_impl; class KernelProgramCache { public: + /// Denotes build error data. The data is filled in from cl::sycl::exception + /// class instance. struct BuildError { std::string Msg; - cl_int Code; + 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; }; diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index c669668b93ae1..b2c0fa57c5524 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -111,7 +111,8 @@ DeviceImage &ProgramManager::getDeviceImage(OSModuleHandle M, template RetT *waitUntilBuilt(KernelProgramCache &Cache, - KernelProgramCache::BuildResult *BuildResult) { + KernelProgramCache::BuildResult *BuildResult, + bool &TryAgain) { // any thread which will find nullptr in cache will wait until the pointer // is not null anymore Cache.waitUntilBuilt([BuildResult]() { @@ -127,7 +128,9 @@ RetT *waitUntilBuilt(KernelProgramCache &Cache, RetT *Result = BuildResult->Ptr.load(); - assert(Result && "An exception should have been thrown"); + // 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; } @@ -170,10 +173,29 @@ RetT *getOrBuild(KernelProgramCache &KPCache, const KeyT &CacheKey, // no insertion took place, thus some other thread has already inserted smth // in the cache if (!InsertionTookPlace) { - return waitUntilBuilt(KPCache, BuildResult); + bool TryAgain = false; + + for (;;) { + RetT *Result = waitUntilBuilt(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(); @@ -184,7 +206,7 @@ RetT *getOrBuild(KernelProgramCache &KPCache, const KeyT &CacheKey, // 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 BuildResult->State.store(BS_Done); From 5e64424d074143b78e618f8c41fd7d0b831ddef7 Mon Sep 17 00:00:00 2001 From: s-kanaev <57672082+s-kanaev@users.noreply.github.com> Date: Thu, 13 Feb 2020 15:52:36 +0300 Subject: [PATCH 06/10] Fix review comment Signed-off-by: Sergey Kanaev --- sycl/include/CL/sycl/detail/kernel_program_cache.hpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/sycl/include/CL/sycl/detail/kernel_program_cache.hpp b/sycl/include/CL/sycl/detail/kernel_program_cache.hpp index cbfb85b33cf8e..b2ae9e0c80870 100644 --- a/sycl/include/CL/sycl/detail/kernel_program_cache.hpp +++ b/sycl/include/CL/sycl/detail/kernel_program_cache.hpp @@ -32,8 +32,9 @@ class KernelProgramCache { 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 + /// due to the possibility of error code being equal to zero even in case + // if build is failed and cl::sycl::exception is thrown + bool FilledIn; }; From be30bc32e74462f7372b9d65b7938be59c26e5db Mon Sep 17 00:00:00 2001 From: s-kanaev <57672082+s-kanaev@users.noreply.github.com> Date: Thu, 13 Feb 2020 16:27:26 +0300 Subject: [PATCH 07/10] Update sycl/include/CL/sycl/detail/kernel_program_cache.hpp Signed-off-by: Sergey Kanaev Co-Authored-By: Mikhail Lychkov <51128024+mlychkov@users.noreply.github.com> --- sycl/include/CL/sycl/detail/kernel_program_cache.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/CL/sycl/detail/kernel_program_cache.hpp b/sycl/include/CL/sycl/detail/kernel_program_cache.hpp index b2ae9e0c80870..5e01bc1757826 100644 --- a/sycl/include/CL/sycl/detail/kernel_program_cache.hpp +++ b/sycl/include/CL/sycl/detail/kernel_program_cache.hpp @@ -33,7 +33,7 @@ class KernelProgramCache { /// Equals to true if the Msg and Code are initialized. This flag is added /// due to the possibility of error code being equal to zero even in case - // if build is failed and cl::sycl::exception is thrown + /// if build is failed and cl::sycl::exception is thrown. bool FilledIn; }; From 0a4dbad7e3dd52ad4def0f83e323096c80f8ea0c Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Fri, 14 Feb 2020 12:28:46 +0300 Subject: [PATCH 08/10] [SYCL] Remove unwanted empty line Signed-off-by: Sergey Kanaev --- sycl/include/CL/sycl/detail/kernel_program_cache.hpp | 1 - 1 file changed, 1 deletion(-) diff --git a/sycl/include/CL/sycl/detail/kernel_program_cache.hpp b/sycl/include/CL/sycl/detail/kernel_program_cache.hpp index 5e01bc1757826..b0de93e6cf616 100644 --- a/sycl/include/CL/sycl/detail/kernel_program_cache.hpp +++ b/sycl/include/CL/sycl/detail/kernel_program_cache.hpp @@ -34,7 +34,6 @@ class KernelProgramCache { /// Equals to true if the Msg and Code are initialized. This flag is added /// due to the possibility of error code being equal to zero even in case /// if build is failed and cl::sycl::exception is thrown. - bool FilledIn; }; From 44fa8c5e112f3d20a85acb33527a69d9f78d2e64 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Fri, 14 Feb 2020 18:36:03 +0300 Subject: [PATCH 09/10] [SYCL] Fix review comments Signed-off-by: Sergey Kanaev --- .../detail/program_manager/program_manager.cpp | 3 +-- .../test/kernel-and-program/cache-build-result.cpp | 14 -------------- 2 files changed, 1 insertion(+), 16 deletions(-) diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index b2c0fa57c5524..3c9dd77411f3f 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -173,9 +173,8 @@ RetT *getOrBuild(KernelProgramCache &KPCache, const KeyT &CacheKey, // no insertion took place, thus some other thread has already inserted smth // in the cache if (!InsertionTookPlace) { - bool TryAgain = false; - for (;;) { + bool TryAgain = false; RetT *Result = waitUntilBuilt(KPCache, BuildResult, TryAgain); if (TryAgain) { diff --git a/sycl/test/kernel-and-program/cache-build-result.cpp b/sycl/test/kernel-and-program/cache-build-result.cpp index 366e7f850076e..adf2bf2706d61 100644 --- a/sycl/test/kernel-and-program/cache-build-result.cpp +++ b/sycl/test/kernel-and-program/cache-build-result.cpp @@ -2,20 +2,6 @@ // RUN: %CPU_RUN_PLACEHOLDER %t.out #include -// FIXME do not use internal methods in tests. -#include - -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(); From 4e7bcb8d30f25cfa3785fc6a9ccae695365ef931 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Mon, 17 Feb 2020 17:13:03 +0300 Subject: [PATCH 10/10] [SYCL] Fix review comment Signed-off-by: Sergey Kanaev --- .../program_manager/program_manager.cpp | 30 +++++++------------ 1 file changed, 10 insertions(+), 20 deletions(-) diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 3c9dd77411f3f..5a4ef59b40051 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -111,8 +111,7 @@ DeviceImage &ProgramManager::getDeviceImage(OSModuleHandle M, template RetT *waitUntilBuilt(KernelProgramCache &Cache, - KernelProgramCache::BuildResult *BuildResult, - bool &TryAgain) { + KernelProgramCache::BuildResult *BuildResult) { // any thread which will find nullptr in cache will wait until the pointer // is not null anymore Cache.waitUntilBuilt([BuildResult]() { @@ -128,10 +127,6 @@ RetT *waitUntilBuilt(KernelProgramCache &Cache, RetT *Result = BuildResult->Ptr.load(); - // 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; } @@ -174,23 +169,18 @@ RetT *getOrBuild(KernelProgramCache &KPCache, const KeyT &CacheKey, // in the cache if (!InsertionTookPlace) { for (;;) { - bool TryAgain = false; - RetT *Result = waitUntilBuilt(KPCache, BuildResult, TryAgain); + RetT *Result = waitUntilBuilt(KPCache, BuildResult); - 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 (Result) + return Result; - if (BuildResult->State.compare_exchange_strong(Expected, Desired)) { - // this thread is the building thread now - break; - } + // 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; - continue; - } else // no need to try once more - return Result; + if (BuildResult->State.compare_exchange_strong(Expected, Desired)) + break; // this thread is the building thread now } }