   {
     return GDF_SUCCESS;
   }
+
+  // TODO Check if the num_rows is > MAX_ROWS (MAX_INT)
 
   // check that the columns data are not null, have matching types,
   // and the same number of rows

#include "hashmap/hash_functions.cuh"
#include "int_fastdiv.h"

constexpr int HASH_KERNEL_BLOCK_SIZE = 256;
constexpr int HASH_KERNEL_ROWS_PER_THREAD = 1;

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

    // a partitioning operator on the hash value. Also computes the number of
    // rows in each partition
    compute_row_partition_numbers<hash_function>
    <<<grid_size, BLOCK_SIZE, num_partitions * sizeof(size_type)>>>(table_to_hash, 
                                                                    num_rows,
                                                                    num_partitions,
                                                                    partitioner_type(num_partitions),
                                                                    row_partition_numbers,
                                                                    block_partition_sizes,
                                                                    partition_sizes);
  }


  CUDA_CHECK_LAST();

  // Compute exclusive scan of all blocks' partition sizes in-place to determine 
  // the starting point for each blocks portion of each partition in the output
  size_type * scanned_block_partition_sizes{block_partition_sizes};
  thrust::exclusive_scan(thrust::cuda::par,
                         block_partition_sizes, 
                         block_partition_sizes + (grid_size * num_partitions), 
                         scanned_block_partition_sizes);
  CUDA_CHECK_LAST();


  // Compute exclusive scan of size of each partition to determine offset location
  // of each partition in final output. This can be done independently on a separate stream
  cudaStream_t s1{};
  cudaStreamCreate(&s1);
  size_type * scanned_partition_sizes{partition_sizes};
  thrust::exclusive_scan(thrust::cuda::par.on(s1),
                         partition_sizes, 
                         partition_sizes + num_partitions,
                         scanned_partition_sizes);
  CUDA_CHECK_LAST();

