 
   // Get each row's partition number, and get it's output location by 
   // incrementing block's offset counter for that partition number
+  // and store the row's output location in-place
   while( row_number < num_rows )
   {
+    // Get partition number of this row
     const size_type partition_number = row_partition_numbers[row_number];
 
+    // Get output location based on partition number by incrementing the corresponding
+    // partition offset for this block
     const size_type row_output_location = atomicAdd(&(shared_partition_offsets[partition_number]), size_type(1));
 
+    // Store the row's output location in-place
     row_partition_numbers[row_number] = row_output_location;
 
     row_number += blockDim.x * gridDim.x;
   }
 }
-
-
-
-/* --------------------------------------------------------------------------*/
-/** 
- * @Synopsis  Functor used to map a row's partition number to it's final output location
- */
-/* ----------------------------------------------------------------------------*/
-template <typename size_type>
-struct compute_row_output_location
-{
-  compute_row_output_location(size_type * const offsets) 
-    : partition_offsets{offsets}{};
-
-  __device__
-  size_type operator()(size_type partition_number)
-  {
-    const size_type output_location = atomicAdd(&(partition_offsets[partition_number]), size_type(1));
-    return output_location;
-  }
-
-  size_type * const __restrict__ partition_offsets;
-};
-
-
 
 /* --------------------------------------------------------------------------*/
 /** 


  // Compute the hash value for each row, store it to the array of hash values
  // and compute the partition to which the hash value belongs and increment
  // the counter for that partition
  while( row_number < num_rows)
  {
    // See here why template disambiguator is required: 

  CUDA_TRY( cudaMalloc(&row_partition_numbers, num_rows * sizeof(hash_value_type)) );

  // Array to hold the size of each partition computed by each block
  //  i.e., { {block0 partition0 size, block1 partition0 size, ...}, 
  //          {block0 partition1 size, block1 partition1 size, ...},
  //          ...
  //          {block0 partition(num_partitions-1) size, block1 partition(num_partitions -1) size, ...} }
  size_type * block_partition_sizes{nullptr};
  CUDA_TRY(cudaMalloc(&block_partition_sizes, (grid_size * num_partitions) * sizeof(size_type)));

  // Holds the total number of rows in each partition
  size_type * global_partition_sizes{nullptr};
  CUDA_TRY( cudaMalloc(&global_partition_sizes, num_partitions * sizeof(size_type)) );
  CUDA_TRY( cudaMemsetAsync(global_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
