 
     row_partition_numbers[row_number] = partition_number;
 
-    atomicAdd(&(partition_sizes[partition_number]), size_type(1));
+    atomicAdd(&(shared_partition_sizes[partition_number]), size_type(1));
 
     row_number += blockDim.x * gridDim.x;
+  }
+
+  __syncthreads();
+
+  // Flush shared memory histogram to global memory
+  i = threadIdx.x;
+  while(i < num_partitions)
+  {
+    size_type old = atomicAdd(&(partition_sizes[i]), shared_partition_sizes[i]);
+    i += blockDim.x;
   }
 }
 

    // Determines how the mapping between hash value and partition number is computed
    using partitioner_type = bitwise_partitioner<hash_value_type, size_type, size_type>;

    compute_row_partition_numbers<hash_function>
    <<<grid_size, HASH_KERNEL_BLOCK_SIZE>>>(table_to_hash, 
                                            num_rows,
                                            partitioner_type(static_cast<int>(num_partitions)),
                                            row_partition_numbers,
                                            partition_sizes);

  }
  else

    // Determines how the mapping between hash value and partition number is computed
    using partitioner_type = modulo_partitioner<hash_value_type, size_type, size_type>;

    // Computes which partition each row belongs to by hashing the row and performing
    // a partitioning operator on the hash value. Also computes the number of
    // rows in each partition
    compute_row_partition_numbers<hash_function>
    <<<grid_size, HASH_KERNEL_BLOCK_SIZE, num_partitions * sizeof(size_type)>>>(table_to_hash, 
                                                                                num_rows,
                                                                                num_partitions,
                                                                                partitioner_type(num_partitions),
                                                                                row_partition_numbers,
                                                                                partition_sizes);
  }


  CUDA_CHECK_LAST();

