patching grid reduction to be volta-safe

38bada23 · Jie · 438f6f9f · 38bada23
Commit 38bada23 authored Jan 17, 2019 by Jie
Hide whitespace changes
Inline Side-by-side

Showing with 2 additions and 0 deletions

csrc/welford.cu csrc/welford.cu +2 -0

No files found.
--- a/csrc/welford.cu
+++ b/csrc/welford.cu
@@ -508,6 +508,7 @@ welford_kernel_c_last(
    }

    __threadfence();
+    __syncthreads(); // ensuring writes to staging_ is visible to all blocks

    __shared__ bool is_last_block_done;
    // mark block done
@@ -724,6 +725,7 @@ __global__ void reduce_bn_c_last_kernel(
    }

    __threadfence();
+    __syncthreads(); // ensuring writes to staging_ is visible to all blocks

    __shared__ bool is_last_block_done;
    // mark block done