   // Create a separate table of the columns to be hashed
   std::unique_ptr< const gdf_table<size_type> > table_to_hash {new gdf_table<size_type>(num_cols_to_hash, 
                                                                                         gdf_columns_to_hash.data())};
+
+  gdf_error gdf_status{GDF_SUCCESS};
   switch(hash)
   {
     case GDF_HASH_MURMUR3:
       {
-        hash_partition_gdf_table<MurmurHash3_32>(*input_table, 
-                                                 *table_to_hash,
-                                                 num_partitions,
-                                                 partition_offsets,
-                                                 *output_table);
+        gdf_status = hash_partition_gdf_table<MurmurHash3_32>(*input_table, 
+                                                              *table_to_hash,
+                                                              num_partitions,
+                                                              partition_offsets,
+                                                              *output_table);
         break;
       }
     default:
-      return GDF_INVALID_HASH_FUNCTION;
-  }
-
-  return GDF_SUCCESS;
-}
-
+      gdf_status = GDF_INVALID_HASH_FUNCTION;
+  }
+
+  return gdf_status;
+}
+

  }
}

template <typename size_type>
__global__
void compute_output_locations( size_type const * const __restrict__ row_partition_numbers,
                               const size_type num_rows,
                               const size_type num_partitions,
                               size_type * __restrict__ partition_offsets,
                               size_type * __restrict__ output_locations)
{

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

                                          row_partition_numbers,
                                          partition_sizes);

  CUDA_TRY(cudaGetLastError());

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

  CUDA_TRY(cudaGetLastError());

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


  // Compute the output location for each row in-place based on it's 
  // partition number such that each partition will be contiguous in memory
  thrust::transform(row_partition_numbers,
                    row_partition_numbers + num_rows,
                    row_partition_numbers,
                    compute_row_output_location<size_type>(partition_sizes));


  CUDA_TRY(cudaGetLastError());

  CUDA_TRY(cudaFree(row_partition_numbers));
  CUDA_TRY(cudaFree(partition_sizes));

  return GDF_SUCCESS;
}

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