 
             const mapped_type existing_value = hashtbl_values[current_index].second;
 
-            //hashtbl_values[current_index].second = x.second;
-            mapped_type prev_value = atomicExch(&(hashtbl_values[current_index].second), x.second);
-
-            printf("inserted new key: %d new value: %d old key: %d old value: %d \n", x.first, x.second, old_key, prev_value);
+            hashtbl_values[current_index].second = x.second;
+            //atomicExch(&(hashtbl_values[current_index].second), x.second);
+
+            printf("inserted new key: %d new value: %llu old key: %d old value: %llu \n", x.first, x.second, old_key, existing_value);
 
             insert_success = true;
           }

          // TODO: How to handle data types less than 32 bits?
          else if ( m_equal(x.first, old_key) ){

            mapped_type old_value = hashtbl_values[current_index].second;

            mapped_type expected = 0;

            mapped_type result_value = 0;
            printf("before update key: %d new_value: %d old_value: %d\n", x.first, x.second, old_value);

            // Attempt to perform the aggregation with existing_value and
            // store the result atomically
            // Guard against another thread's update to existing_value and
            // ensure that existing_value has been updated from its initial state
            // to ensure that the aggregation is valid
            while( expected != old_value || (old_value == m_unused_element) )
            {
              expected = old_value;

              result_value = op(x.second, old_value);
              printf("key: %d new value: %d old value: %d result: %d\n", x.first, x.second, old_value, result_value);

              old_value = atomicCAS(&(hashtbl_values[current_index].second), expected, result_value);
            }

            printf("updating key: %d with value: %d old_value: %d result value: %d \n", x.first, x.second, old_value, result_value);

            insert_success = true;
          }


        bool insert_success = false;

        const int idx = threadIdx.x + blockDim.x * blockIdx.x;
        if(idx < m_hashtbl_size){
          //printf("index: %d key: %d value: %llu\n", idx, hashtbl_values[idx].first, hashtbl_values[idx].second);
        }
        
        while (false == insert_success) {

#ifdef __CUDA_ARCH__

          mapped_type old_value = hashtbl_values[current_index].second;

          // Try and set the existing_key for the current hash bucket to insert_key
          //const key_type old_key = atomicCAS( existing_key, unused_key, insert_key);
