Inquiry: How does Triton handle precision differences caused by summation order in tl.sum?

### Summary We are evaluating Triton's numerical precision for reduction operations across different platforms . We observed that `tl.sum` can produce results that differ from PyTorch's `torch.sum`, and based on our analysis, these discrepancies appear to stem from **different summation orders** (e.g., tree reduction vs. sequential accumulation). We would like to understand the Triton community's perspective on this phenomenon and how such cases are typically handled. --- ### Environment | Component | Version / Detail | |-----------|------------------| | Triton | 3.0.0 | | PyTorch | 2.4.0 | | CUDA | 12.4 | | NVIDIA Driver | 550.54.15 | | GPU | NVIDIA A100-SXM4-80GB | | OS | Linux | --- ### Context We wrote a suite of Triton kernels that perform `tl.sum` over dimension 0 for tensors of various ranks (1D through 8D). Each kernel loads the entire input tensor into a single block using `tl.arange` and then calls `tl.sum(x, 0)`. The reference is PyTorch's `torch.sum` on the same data. **Kernel pattern (example for 2D):** ```python @triton.jit def triton_sum_2D_dim0(in_ptr0, out_ptr0, L: tl.constexpr, M: tl.constexpr): lblk_idx = tl.arange(0, L) mblk_idx = tl.arange(0, M) idx = mblk_idx[None, :] + lblk_idx[:, None] * M x = tl.load(in_ptr0 + idx) ret = tl.sum(x, 0) tl.store(out_ptr0 + mblk_idx, ret) ``` **Data types tested:** `int8`, `int16`, `int32`, `int64`, `float16`, `float32`, `bfloat16`, `bool` **Shapes tested:** A wide range from 1D to 8D, including edge cases where the reduction dimension is size 1 or large (e.g., 1024). --- ### Observation We found that a number of test cases show precision mismatches between Triton and PyTorch. The pattern suggests the root cause is **summation order** rather than algorithmic bugs: 1. **`float16` and `bfloat16`** are the most affected, which is expected given their limited precision — different accumulation orders amplify rounding errors. 2. **`float32`** in higher-dimensional tensors (5D–8D) also shows occasional mismatches of a few ULPs. 3. Even **integer types** (`int32`, `int64`) show differences in some configurations. --- ### Questions for the Community We are not necessarily asking for a "fix" — we understand that summation order is often an implementation detail that can vary across backends and frameworks. Instead, we would like to understand: 1. **How does the Triton project view precision differences caused by summation order?** Are these considered expected behavior, or is bit-exactness with PyTorch a design goal for reduction ops? 2. **What reduction algorithm does `tl.sum` use internally?** For example, is it a tree-based reduction, warp shuffle, or sequential accumulation? Knowing this would help us reason about the deviations we observe. 3. **What is the general policy for handling such precision reports?** Should users adjust their validation tolerances to account for summation order differences, or does the team aim to align with PyTorch's reduction order where possible? 4. **Are there any plans or existing mechanisms** to give users more control over reduction precision (e.g., higher-precision accumulators, or a way to specify reduction order)? --- ```