remove unnecessary if-statement.

ooreilly · Jan 12, 2021 · 4ace6a8 · 4ace6a8
1 parent 84c8f89
commit 4ace6a8
Showing 1 changed file with 7 additions and 10 deletions.
diff --git 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<kernel>(&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();
         }