 
   const int size;
 
-
-  MapTest(const int hash_table_size = 100)
+  const int THREAD_BLOCK_SIZE{256};
+
+  std::vector<key_type> keys;
+  std::vector<value_type> values;
+  thrust::device_vector<key_type> d_keys;
+  thrust::device_vector<value_type> d_values;
+
+  std::unordered_map<key_type, value_type> expected_values;
+
+  MapTest(const int hash_table_size = 10000)
     : size(hash_table_size), the_map(new map_type(hash_table_size))
   {
+  }
+
+  int create_input(const int num_unique_keys, const int num_values_per_key, const int max_key = RAND_MAX, const int max_value = RAND_MAX)
+  {
+
+
+    const int TOTAL_PAIRS = num_unique_keys * num_values_per_key;
+
+    this->the_map.reset(new map_type(2*TOTAL_PAIRS));
+
+    keys.reserve(TOTAL_PAIRS);
+    values.reserve(TOTAL_PAIRS);
+
+    // Always use the same seed so the random sequence is the same each time
+    std::srand(0);
+
+    for(int i = 0; i < num_unique_keys; ++i )
+    {
+      // Create random key
+      key_type current_key = std::rand() % max_key;
+
+      // Don't use unused_key
+      while(current_key == this->unused_key)
+      {
+        current_key = std::rand();
+      }
+
+      // For the current key, generate random values
+      for(int j = 0; j < num_values_per_key; ++j)
+      {
+        value_type current_value = std::rand() % max_value;
+
+        // Don't use unused_value
+        while(current_value == this->unused_value)
+        {
+          current_value = std::rand();
+        }
+
+        // Store current key and value
+        keys.push_back(current_key);
+        values.push_back(current_value);
+
+        // Use a STL map to keep track of the max value for each key
+        auto found = expected_values.find(current_key);
+
+        // Key doesn't exist yet, insert it
+        if(found == expected_values.end())
+        {
+          expected_values.insert(std::make_pair(current_key,current_value));
+        }
+        // Key exists, update the value with the max
+        else
+        {
+          max_op<value_type> op;
+          value_type new_value = op(found->second, current_value);
+          found->second = new_value;
+        }
+      }
+    }
+
+    d_keys = keys;
+    d_values = values;
+    return TOTAL_PAIRS;
+  }
+
+  void check_answer(){
+
+    for(auto const &k : this->expected_values)
+    {
+      key_type test_key = k.first;
+
+      value_type expected_value = k.second;
+
+      auto found = this->the_map->find(test_key);
+
+      ASSERT_NE(this->the_map->end(), found);
+
+      value_type test_value = found->second;
+
+      EXPECT_EQ(expected_value, test_value) << "Key is: " << test_key;
+    }
   }
 
   ~MapTest(){


}

// Have to use a functor instead of a device lambda because
// you can't create a device lambda inside of a Google Test
// because the macro expands into a private member function
// 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);
  }
};


TYPED_TEST(MapTest, MaxAggregationTestDevice)

  using value_type = typename TypeParam::value_type;
  using size_type = typename MapTest<TypeParam>::map_type::size_type;

  const size_type input_size = this->create_input(512, 256*256);

  key_type *k = thrust::raw_pointer_cast(this->d_keys.data());
  value_type *v = thrust::raw_pointer_cast(this->d_values.data());

  const dim3 grid_size ((input_size + this->THREAD_BLOCK_SIZE -1) / this->THREAD_BLOCK_SIZE,1,1);
  const dim3 block_size (this->THREAD_BLOCK_SIZE, 1, 1);

  std::cout << "Input Size: " << input_size << " Grid Size: " << grid_size.x << " Block Size: " << block_size.x << std::endl;

  cudaDeviceSynchronize();
  build_table<<<grid_size, block_size>>>((this->the_map).get(), k, v, input_size, max_op<value_type>());
  cudaDeviceSynchronize(); 

  this->check_answer();

}

