Only run __syncthreads when there is data to preload

7ddeb944 · Paul · 2783c649 · 7ddeb944
Commit 7ddeb944 authored Jun 30, 2022 by Paul
Hide whitespace changes
Inline Side-by-side

Showing with 2 additions and 1 deletion

src/targets/gpu/kernels/include/migraphx/kernels/preload.hpp src/targets/gpu/kernels/include/migraphx/kernels/preload.hpp +2 -1

No files found.
--- a/src/targets/gpu/kernels/include/migraphx/kernels/preload.hpp
+++ b/src/targets/gpu/kernels/include/migraphx/kernels/preload.hpp
@@ -186,7 +186,8 @@ __device__ auto auto_preload(index idx)
 {
    return make_transform([=](auto f, auto... xs) {
        auto invoke = [=](auto... ys) {
-            __syncthreads();
+            if constexpr((Bs or ...))
+                __syncthreads();
            f(ys...);
        };
        join(invoke, preload_copy<Bs>(idx, xs)...);