From 4ace6a824139a7b1a7bff971727ce49d7d85fe81 Mon Sep 17 00:00:00 2001 From: Ossian O'Reilly Date: Mon, 11 Jan 2021 22:45:56 -0800 Subject: [PATCH] remove unnecessary if-statement. --- opt_32_7.cuh | 17 +++++++---------- 1 file changed, 7 insertions(+), 10 deletions(-) diff --git a/opt_32_7.cuh b/opt_32_7.cuh index 2332856..c0858b1 100644 --- a/opt_32_7.cuh +++ b/opt_32_7.cuh @@ -103,18 +103,15 @@ __global__ void opt7wl79_32x32x32(float *in) { // Apply wavelet transform line by line in the z-direction opt7ds79_compute(&smem[idx + snxy * idy], snx); - __syncthreads(); + __syncwarp(); - // Write all (x,z) planes back to global memory - if (batch_y * block_y + idy < 32) { - // Process an entire 32 x 32 plane - for (int tile_z = 0; tile_z < 32 ; ++tile_z) { - size_t sptr = idx + snx * tile_z + snxy * idy; - size_t gptr = idx + 1024 * tile_z + 32 * idy; - in[batch_y * planes * 32 + gptr + block_idx] = smem[sptr]; - } + // Write all (x,z) planes back to global memory + for (int tile_z = 0; tile_z < 32; ++tile_z) { + size_t sptr = idx + snx * tile_z + snxy * idy; + size_t gptr = idx + 1024 * tile_z + 32 * idy; + in[batch_y * planes * 32 + gptr + block_idx] = smem[sptr]; } - + __syncthreads(); }