Revert sum/product aggregation to always produce int64_t type

[SurajAralihalli]

Description

This pull request reverses the modifications made to the sum/product aggregation target type, ensuring it always produces int64. The changes implemented by PR 14679 which led to degraded performance when the aggregation column had an unsigned type, are reverted. Additional details can be found in the issue 14886.

Checklist

• I am familiar with the Contributing Guidelines.
• New or existing tests cover these changes.
• The documentation is up to date with these changes.

[ttnghia]

I still don't understand what is the reason that causes performance issue there? Why can't we fix that issue instead of reverting code like this?

[GregoryKimball]

Thank you @ttnghia for your message. I also would prefer to solve the root cause of the issue rather than revert the change. It appears that the degenerate performance is happening in a thrust::for_each_n call here in groupby.cu. I'm concerned that there isn't an obvious libcudf fix, and the root cause could be in CCCL. If the fix isn't a simple ~1-10 line change in libcudf... I think it's going to be too high risk for 24.02.

[ttnghia]

I'm fine to have this going into 24.02, but we should investigate further for it later. I suspect that it could be some thing messing with the device sum operator working on different operand types.

[karthikeyann]

@ttnghia The atomicAdd<T>( has no specialization for uint32_t and uint64_t. So it goes for CAS-while loop generic atomic implementation, which slows the groupby sum operation on uint column.

[ttnghia]

@ttnghia The atomicAdd<T>( has no specialization for uint32_t and uint64_t. So it goes for CAS-while loop generic atomic implementation, which slows the groupby sum operation on uint column.

Wait, from the cuda documentation I see that unsigned int/long are supported?
https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicadd

So there maybe something wrong with the groupby code that doesn't call native CUDA atomicAdd but instead calls to the generic CAS-while loop.

[GregoryKimball]

Thank you @karthikeyann and @ttnghia for this investigation. We would love your help making the groupby code work correctly with uint64 in 24.04. If this PR is ready to go, would you please request an admin merge in 24.02?

[bdice]

@karthikeyann's analysis is correct, from what I can tell. I think he was referring to the implementation in device_atomics.cuh, which has a limited implementation.

cudf/cpp/include/cudf/detail/utilities/device_atomics.cuh

Lines 160 to 162 in 5cc021a

	// specialized functions for operators
	// `atomicAdd` supports int32, float, double (signed int64 is not supported.)
	// `atomicMin`, `atomicMax` support int32_t, int64_t

That file only implements DeviceSum for float, double, int32_t, and int64_t. It seems reasonable to add uint32_t and uint64_t implementations, which could call the CUDA atomicAdd for their respective types. We would want to add static assertions like this to ensure that the type aliases like unsigned long long int are equivalent to uint64_t.

[davidwendt]

Maybe cuda::atomic_ref could be used instead of the device_atomics.cuh functions though I'm not sure if it is supported for unsigned types.
Reference: #13583
@PointKernel

[ttnghia]

From what I understand:

• device_atomics.cuh is to specialize for the types that are not natively supported by cuda atomicXXX functions. As such, in case of calling to atomicAdd for uint64_t, the compiler should call to the native atomicAdd function, not cudf::detail::atomicAdd function.
• For uint64_t , if cudf::detail::atomicAdd is called, which in turn calls to the generic CAS-while loop, there should be something wrong with the call chain.
• The caller is in cudf::detail:: namespace, and it just calls atomicAdd thus it is difficult to tell which atomicAdd is being executed.

[bdice]

device_atomics.cuh is to specialize for the types that are not natively supported by cuda atomicXXX functions. As such, in case of calling to atomicAdd for uint64_t, the compiler should call to the native atomicAdd function, not cudf::detail::atomicAdd function.

The overloads for float and double just pass through. We should do the same here for uint32_t and uint64_t, if my analysis is correct. We should always use cudf's internal implementation instead of the CUDA atomicAdd.

For uint64_t , if cudf::detail::atomicAdd is called, which in turn calls to the generic CAS-while loop, there should be something wrong with the call chain.

I'm tracing the call chain myself. I think it goes from aggregation::SUM, which calls atomicAdd and falls back to this generic case which calls this CAS loop.

The caller is in cudf::detail:: namespace, and it just calls atomicAdd thus it is difficult to tell which atomicAdd is being executed.

I'd be okay with adding an explicit namespace where we intend this to be called. I agree the name conflict is not ideal.

[PointKernel]

I'm not sure if it is supported for unsigned types.

cuda::atomic_ref works fine for 4-byte and 8-byte types but doesn't support bool or 2-byte types like int16_t. We can get rid of device_atomics.cuh once NVIDIA/cccl#1024 is resolved.

I'd be okay with adding an explicit namespace where we intend this to be called. I agree the name conflict is not ideal.

👍 explicit namespace or snake-case naming is probably the best temporary workaround before the cccl fix.

[GregoryKimball]

OK, thank you everyone for this discussion. If I understand the consensus solution correctly we should:

• add specializations for uint32_t and uint64_t as per the comment
• update namespace for atomicAdd in aggregation.cuh to specify the cudf::detail:: namespace
• extend the namespace updates to atomicMul atomicMin atomicMax in aggregation.cuh