From 18967b072febfca6038d14971646e5080e011f7b Mon Sep 17 00:00:00 2001
From: Olivier Lapicque <olapicque@nvidia.com>
Date: Mon, 10 Feb 2020 12:07:52 -0800
Subject: [PATCH] Cosmetic Item #3: Remove warp size dependency

---
 cpp/src/io/orc/stats_enc.cu | 11 +++++------
 1 file changed, 5 insertions(+), 6 deletions(-)

diff --git a/cpp/src/io/orc/stats_enc.cu b/cpp/src/io/orc/stats_enc.cu
index 615f16124fd..5500f9c22c4 100644
--- a/cpp/src/io/orc/stats_enc.cu
+++ b/cpp/src/io/orc/stats_enc.cu
@@ -42,17 +42,16 @@ gpu_init_statistics_groups(statistics_group *groups, const stats_column_desc *co
   uint32_t t = threadIdx.x & 0x1f;
   volatile statistics_group *group = &group_g[threadIdx.x >> 5];
   if (chunk_id < num_rowgroups) {
-    if (!t) {
+    if (t == 0) {
       uint32_t num_rows = cols[col_id].num_rows;
       group->col = &cols[col_id];
       group->start_row = chunk_id * row_index_stride;
       group->num_rows = min(num_rows - min(chunk_id * row_index_stride, num_rows), row_index_stride);
-      __threadfence_block();
-    }
-    SYNCWARP();
-    if (t < sizeof(statistics_group) / sizeof(uint32_t)) {
-      reinterpret_cast<uint32_t *>(&groups[col_id * num_rowgroups + chunk_id])[t] = reinterpret_cast<volatile uint32_t *>(group)[t];
     }
+  }
+  __syncthreads();
+  if (chunk_id < num_rowgroups && t < sizeof(statistics_group) / sizeof(uint32_t)) {
+    reinterpret_cast<uint32_t *>(&groups[col_id * num_rowgroups + chunk_id])[t] = reinterpret_cast<volatile uint32_t *>(group)[t];
   } 
 }