 /* ----------------------------------------------------------------------------*/
 template < template <typename> class hash_function,
            typename size_type>
-void hash_partition_gdf_table(gdf_table<size_type> const & input_table,
-                              gdf_table<size_type> const & table_to_hash,
-                              const size_type num_partitions,
-                              size_type * partition_offsets,
-                              gdf_table<size_type> & partitioned_output)
-{
-
-  
+gdf_error hash_partition_gdf_table(gdf_table<size_type> const & input_table,
+                                   gdf_table<size_type> const & table_to_hash,
+                                   const size_type num_partitions,
+                                   size_type * partition_offsets,
+                                   gdf_table<size_type> & partitioned_output)
+{
+
   // Determines how the mapping between hash value and partition number is computed
   using partitioner_type = modulo_partitioner<hash_value_type, size_type, size_type>;
 
-  // Compute the hash value for all the rows in the table to hash
-  //table_row_hasher<hash_function, size_type> the_hasher(table_to_hash);
-  //thrust::tabulate(row_partition_numbers.begin(), 
-  //                 row_partition_numbers.end(), 
-  //                 the_hasher);
   const size_type num_rows = table_to_hash.get_column_length();
 
-
+  // Allocate array to hold which partition each row belongs to
   size_type * row_partition_numbers{nullptr};
-  cudaMalloc(&row_partition_numbers, num_rows * sizeof(hash_value_type));
-
+  CUDA_TRY( cudaMalloc(&row_partition_numbers, num_rows * sizeof(hash_value_type)) );
+
+  // Array to hold the size of each partition
   size_type * partition_sizes{nullptr};
-  cudaMalloc(&partition_sizes, num_partitions * sizeof(size_type));
-  cudaMemsetAsync(partition_sizes, 0, num_partitions * sizeof(size_type));
+  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;
-
 
   // 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

                                          row_partition_numbers,
                                          partition_sizes);

  // Compute exclusive scan of the partition sizes in-place to determine 
  // the starting point for each partition in the output
  thrust::exclusive_scan(partition_sizes, 
                         partition_sizes + num_partitions, 
                         partition_sizes);

  // Copy the result of the exlusive scan to the output offsets array
  // to indicate the starting point for each partition in the output
  cudaMemcpyAsync(partition_offsets, 
                  partition_sizes, 
                  num_partitions * sizeof(size_type),
                  cudaMemcpyDeviceToHost);

  size_type * row_output_locations{nullptr};
  cudaMalloc(&row_output_locations, num_rows * sizeof(size_type));

  compute_output_locations<<<grid_size, HASH_KERNEL_BLOCK_SIZE>>>(row_partition_numbers,
                                                                  num_rows,
                                                                  num_partitions,
                                                                  partition_sizes,
                                                                  row_output_locations);

  cudaFree(row_partition_numbers);
  cudaFree(partition_sizes);
  cudaFree(row_output_locations);
}

/* --------------------------------------------------------------------------*/


    const size_type output_location = atomicAdd(&partition_offsets[row_partition_number], size_type(1));

    row_partition_numbers[row_number] = output_location;

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