Why does double type get rounded when adding in a static_map

[avithemad]

I was trying to write tpch query (q3 in particular). Here we need to groupby and aggregate the revenue column which is in double format.
So I wrote a kernel something like this

template <typename Map, typename AggMap>
__global__ void probe(Map order_hash,
                      Map customer_hash,
                      int32_t* l_shipdate,
                      int64_t* l_orderkey,
                      int64_t* o_custkey,
                      int32_t* o_orderdate,
                      int64_t* o_shippriority,
                      double* l_extendedprice,
                      double* l_discount,
                      size_t l_size, size_t o_size,
                      double* result, AggMap agg_map) {
  int tid = (threadIdx.x + blockIdx.x*blockDim.x) / TILE_SIZE;
  if (tid >= l_size) return;

  // filter line item by shipdate
  if (l_shipdate[tid] <= 9204) return; 
  auto order_idx = order_hash.find(l_orderkey[tid]);
  if (order_idx == order_hash.end()) return;
  auto cust_idx = customer_hash.find(o_custkey[order_idx->second]);
  if (cust_idx == customer_hash.end()) return;
  // prepare the packed key
  int64_t aggkey = l_orderkey[tid];
  aggkey |= (o_orderdate[order_idx->second] << 32);
  double res = l_extendedprice[tid]*(1.0f - l_discount[tid]);
  // perform the aggregation
  auto [slot, is_new_key]= agg_map.insert_and_find(cuco::pair{aggkey, res});
  if (!is_new_key) {
    auto ref = cuda::atomic_ref<typename AggMap::mapped_type, cuda::thread_scope_device>{slot->second};
    ref.fetch_add(res, cuda::memory_order_relaxed);
  }
}

The results I see are rounded

406181
405839
390324
384538
378673
378377
375154
373133
371407
367371

the expected results should look something like this

406181.0111	
405838.6989	
390324.0610	
384537.9359	
378673.0558	
378376.7952	
375153.9215	
373133.3094	
371407.4595	
367371.1452

Is this device dependent? I am running on RTX 2050 with cuda version 12.4.

[PointKernel]

@avithemad Is it possible to get a minimal repro in https://godbolt.org/z/1GvxoKbP3?

The cudf groupby and aggregate operations also utilize cuco hash tables, but we haven't observed any of the rounding errors mentioned here. Since cuco hash tables store data exactly as provided, the issue is more likely with the aggregation. You might want to examine the following details more closely:

  double res = l_extendedprice[tid]*(1.0f - l_discount[tid]);

Since 1.0f is a float, you might want to use 1. or 1.0 to prevent narrow conversions. Also, could it be an issue with how the output is being printed, for example, by converting the value to an integer before printing?

[avithemad]

Apologies, you are right, there was an issue in the way I was printing the result.
I was not aware of setting precision in the std::cout, therefore it was rounding up, I set it up with std::cout << std::setprecision(10);
and Now I get the correct results.