     }
   }
 
-  unsigned int extract_groupby_result_device()
+  size_t extract_groupby_result_device()
   {
 
     const dim3 grid_size ((this->hash_table_size + this->THREAD_BLOCK_SIZE - 1) / this->THREAD_BLOCK_SIZE, 1, 1);

{
  const int num_keys = 1<<12;
  const int num_values_per_key = 1;
  auto input = this->create_input(num_keys, num_values_per_key, 1000, 1000);
  auto expected_values = this->compute_reference_solution();
  this->build_aggregation_table_device(input);
  this->verify_aggregation_table(expected_values);
  unsigned int result_size = this->extract_groupby_result_device();
  this->verify_groupby_result(expected_values);

  // The size of the result should be equal to the number of unique keys
  auto begin = this->d_groupby_column.begin();
  auto end = this->d_groupby_column.end();
  thrust::sort(begin, end);
  unsigned int unique_count = thrust::unique(begin, end) - begin;
  EXPECT_EQ(unique_count, result_size);
}

TYPED_TEST(GroupByTest, AggregationTestDeviceWarpSame)
{
  const int num_keys = 1<<12;
  const int num_values_per_key = 32;
  auto input = this->create_input(num_keys, num_values_per_key, 1000, 1000);
  auto expected_values = this->compute_reference_solution();
  this->build_aggregation_table_device(input);
  this->verify_aggregation_table(expected_values);
  unsigned int result_size = this->extract_groupby_result_device();
  this->verify_groupby_result(expected_values);

  // The size of the result should be equal to the number of unique keys
  auto begin = this->d_groupby_column.begin();
  auto end = this->d_groupby_column.end();
  thrust::sort(begin, end);
  unsigned int unique_count = thrust::unique(begin, end) - begin;
  EXPECT_EQ(unique_count, result_size);
}

TYPED_TEST(GroupByTest, AggregationTestDeviceBlockSame)
{
  const int num_keys = 1<<8;
  const int num_values_per_key = this->THREAD_BLOCK_SIZE;
  auto input = this->create_input(num_keys, num_values_per_key, 1000, 1000);
  auto expected_values = this->compute_reference_solution();
  this->build_aggregation_table_device(input);
  this->verify_aggregation_table(expected_values);
  unsigned int result_size = this->extract_groupby_result_device();
  this->verify_groupby_result(expected_values);

  // The size of the result should be equal to the number of unique keys
  auto begin = this->d_groupby_column.begin();
  auto end = this->d_groupby_column.end();
  thrust::sort(begin, end);
  unsigned int unique_count = thrust::unique(begin, end) - begin;
  EXPECT_EQ(unique_count, result_size);
}

TYPED_TEST(GroupByTest, GroupBy)
{
  const int num_keys = 1<<12;
  const int num_values_per_key = 1;
  auto input = this->create_input(num_keys, num_values_per_key, 1000, 1000);
  auto expected_values = this->compute_reference_solution();

  const unsigned int result_size = this->groupby(input.first, input.second);
  this->verify_groupby_result(expected_values);

  // The size of the result should be equal to the number of unique keys
  auto begin = this->d_groupby_column.begin();
  auto end = this->d_groupby_column.end();
  thrust::sort(begin, end);
  unsigned int unique_count = thrust::unique(begin, end) - begin;
  EXPECT_EQ(unique_count, result_size);
}



{
  const int num_keys = 1;
  const int num_values_per_key = 1<<12;

  auto input = this->create_input(num_keys, num_values_per_key);
  auto expected_values = this->compute_reference_solution();

  this->build_aggregation_table_device(input);
  this->verify_aggregation_table(expected_values);

  size_t computed_result_size = this->extract_groupby_result_device();
  this->verify_groupby_result(computed_result_size, expected_values);

}

// TODO Update the create_input function to ensure all keys are actually unique
