 
   const size_type num_rows = table_to_hash.get_column_length();
 
+  constexpr int rows_per_block = BLOCK_SIZE * ROWS_PER_THREAD;
+  const size_type grid_size = (num_rows + rows_per_block - 1) / rows_per_block;
+
   // Allocate array to hold which partition each row belongs to
   size_type * row_partition_numbers{nullptr};
   CUDA_TRY( cudaMalloc(&row_partition_numbers, num_rows * sizeof(hash_value_type)) );
 
-  // Array to hold the size of each partition
+  // Array to hold the size of each partition computed by each block
+  size_type * block_partition_sizes{nullptr};
+  CUDA_TRY(cudaMalloc(&block_partition_sizes, (grid_size * num_partitions) * sizeof(size_type)));
+
+  // Holds the total size of each partition
   size_type * partition_sizes{nullptr};
-  CUDA_TRY(cudaMalloc(&partition_sizes, num_partitions * sizeof(size_type)));
-  CUDA_TRY(cudaMemsetAsync(partition_sizes, 0, num_partitions * sizeof(size_type)));
-
-  constexpr int rows_per_block = HASH_KERNEL_BLOCK_SIZE * HASH_KERNEL_ROWS_PER_THREAD;
-  const int grid_size = (num_rows + rows_per_block - 1) / rows_per_block;
-
-
-
-
+  CUDA_TRY( cudaMalloc(&partition_sizes, num_partitions * sizeof(size_type)) );
+  CUDA_TRY( cudaMemsetAsync(partition_sizes, 0, num_partitions * sizeof(size_type)) );
 
   // If the number of partitions is a power of two, we can compute the partition 
   // number of each row more efficiently with bitwise operations

    // 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);

  }
  else

  __syncthreads();

  // Flush shared memory histogram to global memory
  partition_number = threadIdx.x;
  while(partition_number < num_partitions)
  {
    const size_type block_partition_size = shared_partition_sizes[partition_number];

    // Update global size of each partition
    atomicAdd(&global_partition_sizes[partition_number], block_partition_size);

    // Record the size of this partition in this block
    const size_type write_location = partition_number * gridDim.x + blockIdx.x;
    block_partition_sizes[write_location] = block_partition_size;
    partition_number += blockDim.x;
  }
}

template <typename size_type>
__global__ 
void compute_row_output_locations(size_type * row_partition_numbers, 
                                  const size_type num_rows,
                                  const size_type num_partitions,
                                  size_type * block_partition_offsets)
{
  // Shared array that holds the offset of this blocks partitions in 
  // global memory
  extern __shared__ size_type shared_partition_offsets[];

  // Initialize array of this blocks offsets from global array
  size_type partition_number= threadIdx.x;
  while(partition_number < num_partitions)
  {
    shared_partition_offsets[partition_number] = block_partition_offsets[partition_number * gridDim.x + blockIdx.x];
    partition_number += blockDim.x;
  }
  __syncthreads();

  size_type row_number = threadIdx.x + blockIdx.x * blockDim.x;

  // Get each row's partition number, and get it's output location by 
  // incrementing block's offset counter for that partition number
  while( row_number < num_rows )
  {
    const size_type partition_number = row_partition_numbers[row_number];

    const size_type row_output_location = atomicAdd(&(shared_partition_offsets[partition_number]), size_type(1));

    row_partition_numbers[row_number] = row_output_location;

    row_number += blockDim.x * gridDim.x;
  }
}

