 
   found = this->the_map->find(0);
   EXPECT_EQ(11, found->second);
-
-}
-
+}
 
 
 TYPED_TEST(GroupByTest, AggregationTestDeviceAllSame)
 {
-  using value_type = typename TypeParam::value_type;
-  using pair_type = typename GroupByTest<TypeParam>::pair_type;
-  using op_type = typename GroupByTest<TypeParam>::op_type;
-
-  this->create_input(1, 1<<20);
-
-  const dim3 grid_size ((this->size + this->THREAD_BLOCK_SIZE -1) / this->THREAD_BLOCK_SIZE,1,1);
-  const dim3 block_size (this->THREAD_BLOCK_SIZE, 1, 1);
-
-  value_type * d_group = thrust::raw_pointer_cast(this->d_groupby_column.data());
-  value_type * d_agg = thrust::raw_pointer_cast(this->d_aggregation_column.data());
-
-  cudaDeviceSynchronize();
-  build_aggregation_table<<<grid_size, block_size>>>((this->the_map).get(), d_group, d_agg, this->size, op_type());
-  cudaDeviceSynchronize(); 
-
-  this->check_answer();
-
+  auto input = this->create_input(1, this->THREAD_BLOCK_SIZE);
+  this->build_aggregation_table_device(input);
+  this->verify_aggregation_table();
 }
 
 TYPED_TEST(GroupByTest, AggregationTestDeviceAllUnique)
 {
-  using value_type = typename TypeParam::value_type;
-  using pair_type = typename GroupByTest<TypeParam>::pair_type;
-  using op_type = typename GroupByTest<TypeParam>::op_type;
-  
-
-  this->create_input(1<<16, 1);
-
-  const dim3 grid_size ((this->size + this->THREAD_BLOCK_SIZE -1) / this->THREAD_BLOCK_SIZE,1,1);
-  const dim3 block_size (this->THREAD_BLOCK_SIZE, 1, 1);
-
-  value_type * d_group = thrust::raw_pointer_cast(this->d_groupby_column.data());
-  value_type * d_agg = thrust::raw_pointer_cast(this->d_aggregation_column.data());
-
-  cudaDeviceSynchronize();
-  build_aggregation_table<<<grid_size, block_size>>>((this->the_map).get(), d_group, d_agg, this->size, op_type());
-  cudaDeviceSynchronize(); 
-
-  this->check_answer();
+  auto input = this->create_input(1<<16, 1);
+  this->build_aggregation_table_device(input);
+  this->verify_aggregation_table();
 }
 
 TYPED_TEST(GroupByTest, AggregationTestDeviceWarpSame)
 {
-  using value_type = typename TypeParam::value_type;
-  using pair_type = typename GroupByTest<TypeParam>::pair_type;
-  using op_type = typename GroupByTest<TypeParam>::op_type;
-
-  this->create_input(1<<15, 32);
-
-  const dim3 grid_size ((this->size + this->THREAD_BLOCK_SIZE -1) / this->THREAD_BLOCK_SIZE,1,1);
-  const dim3 block_size (this->THREAD_BLOCK_SIZE, 1, 1);
-
-  value_type * d_group = thrust::raw_pointer_cast(this->d_groupby_column.data());
-  value_type * d_agg = thrust::raw_pointer_cast(this->d_aggregation_column.data());
-
-  cudaDeviceSynchronize();
-  build_aggregation_table<<<grid_size, block_size>>>((this->the_map).get(), d_group, d_agg, this->size, op_type());
-  cudaDeviceSynchronize(); 
-
-  this->check_answer();
+  auto input = this->create_input(1<<15, 32);
+  this->build_aggregation_table_device(input);
+  this->verify_aggregation_table();
 }
 
 TYPED_TEST(GroupByTest, AggregationTestDeviceBlockSame)
 {
-  using value_type = typename TypeParam::value_type;
-  using pair_type = typename GroupByTest<TypeParam>::pair_type;
-  using op_type = typename GroupByTest<TypeParam>::op_type;
-
-  this->create_input(1<<12, this->THREAD_BLOCK_SIZE);
-
-  const dim3 grid_size ((this->size + this->THREAD_BLOCK_SIZE -1) / this->THREAD_BLOCK_SIZE,1,1);
-  const dim3 block_size (this->THREAD_BLOCK_SIZE, 1, 1);
-
-  value_type * d_group = thrust::raw_pointer_cast(this->d_groupby_column.data());
-  value_type * d_agg = thrust::raw_pointer_cast(this->d_aggregation_column.data());
-
-  cudaDeviceSynchronize();
-  build_aggregation_table<<<grid_size, block_size>>>((this->the_map).get(), d_group, d_agg, this->size, op_type());
-  cudaDeviceSynchronize(); 
-
-  this->check_answer();
+  auto input = this->create_input(1<<12, this->THREAD_BLOCK_SIZE);
+  this->build_aggregation_table_device(input);
+  this->verify_aggregation_table();
 }
 
 

// and you can't have a device lambda inside a private member
// function
template<typename value_type>
  struct max_op
  {
    __host__ __device__
    value_type operator()(value_type a, value_type b)
    {
      return (a > b? a : b);
    }
  };

template<typename value_type>
  struct min_op
  {
    __host__ __device__
    value_type operator()(value_type a, value_type b)
    {
      return (a < b? a : b);
    }
  };


// A new instance of this class will be created for each *TEST(GroupByTest, ...)
// Put all repeated stuff for each test here

// to nest multiple types inside of the KeyValueTypes struct above
// KeyValueTypes<type1, type2> implies key_type = type1, value_type = type2
// This list is the types across which Google Test will run our tests
typedef ::testing::Types< KeyValueTypes<int, int, max_op>,
                          KeyValueTypes<int, float, max_op>,
                          KeyValueTypes<int, double, max_op>,
                          KeyValueTypes<int, long long int, max_op>,
                          KeyValueTypes<int, unsigned long long int, max_op>,
                          KeyValueTypes<unsigned long long int, int, max_op>,
                          KeyValueTypes<unsigned long long int, float, max_op>,
                          KeyValueTypes<unsigned long long int, double, max_op>,
                          KeyValueTypes<unsigned long long int, long long int, max_op>,
                          KeyValueTypes<unsigned long long int, unsigned long long int, max_op>,
                          KeyValueTypes<int, int, min_op>,
                          KeyValueTypes<int, float, min_op>,
                          KeyValueTypes<int, double, min_op>,
                          KeyValueTypes<int, long long int, min_op>,
                          KeyValueTypes<int, unsigned long long int, min_op>,
                          KeyValueTypes<unsigned long long int, int, min_op>,
                          KeyValueTypes<unsigned long long int, float, min_op>,
                          KeyValueTypes<unsigned long long int, double, min_op>,
                          KeyValueTypes<unsigned long long int, long long int, min_op>,
                          KeyValueTypes<unsigned long long int, unsigned long long int, min_op>
                          > Implementations;

TYPED_TEST_CASE(GroupByTest, Implementations);

TYPED_TEST(GroupByTest, AggregationTestHost)
{
  using key_type = typename TypeParam::key_type;
  using value_type = typename TypeParam::value_type;

