[BUG] [CUTLASS] Integer-Gemm gives incorrect results

Which component has the problem?

CUTLASS C++

Bug Report

Describe the bug
When using Cutlass for u8 to s32 Gemm-operations, the results seem to be wrong after a certain point, specifically when the accumulator is large enough to not fit into 24 bits anymore. This is most likely due to a conversion to float in the Default-Epilogue-Definition in gemm/device/default_gemm_configuration.h.

Steps/Code to reproduce bug
mwe.c

#include "cutlass/cutlass.h"
#include "cutlass/gemm/device/gemm.h"
#include <cuda_runtime.h>

__global__ void prepareDot(uint8_t *A, uint8_t *B, int upTo, int K){
    for(int i = 0; i < K; i++){
        if(i < upTo){
            A[i] = 255;
            B[i] = 255;
        }else{
            A[i] = 0;
            B[i] = 0;
        }
    }
}
int main(){
    int M = 16;
    int N = 16;
    int K = 16;

    uint8_t *A = new uint8_t[M * K];    //RowMajor
    uint8_t *B = new uint8_t[K * N];    //ColMajor
    int32_t *C = new int32_t[M * N];
    for(int i = 0; i < M * K; i++)
        A[i] = 0;
    for(int i = 0; i < K * N; i++)
        B[i] = 0;
    for(int i = 0; i < M * N; i++)
        C[i] = 0;
    
    //Init first row of A and first column of B
    for(int i = 0; i < K; i++){
        A[i] = 255;
        B[i] = 255;
    }
    //Init first entry of C to a large value
    C[0] = 16646400;

    uint8_t *A_dev, *B_dev;
    int32_t *C_dev;
    cudaMalloc(&A_dev, M * K * sizeof(uint8_t));
    cudaMalloc(&B_dev, K * N * sizeof(uint8_t));
    cudaMalloc(&C_dev, M * N * sizeof(int32_t));

    cudaMemcpy(A_dev, A, M * K * sizeof(uint8_t), cudaMemcpyHostToDevice);
    cudaMemcpy(B_dev, B, K * N * sizeof(uint8_t), cudaMemcpyHostToDevice);
    cudaMemcpy(C_dev, C, M * N * sizeof(int32_t), cudaMemcpyHostToDevice);

    using CutlassGemm_TensorOp = cutlass::gemm::device::Gemm<
        uint8_t,
        cutlass::layout::RowMajor,
        uint8_t,
        cutlass::layout::ColumnMajor,
        int32_t,
        cutlass::layout::RowMajor,
        int32_t,
        cutlass::arch::OpClassTensorOp,
        cutlass::arch::Sm80>;
    CutlassGemm_TensorOp gemm_operator_tensorop;
    const int lda = M;
    const int ldb = K;
    const int ldc = M;
    int alpha = 1;
    int beta = 1;
    cudaDeviceSynchronize();
    typename CutlassGemm_TensorOp::Arguments args({M,N,K},
        {A_dev, lda},
        {B_dev, ldb},
        {C_dev, ldc},
        {C_dev, ldc},
        {alpha, beta});

    int32_t *result_gpu;
    result_gpu = new int32_t;
    for(int i = 1; i < K; i++){
        //CPU
        int32_t result_cpu = C[0];
        for(int j = 0; j < i; j++){
            result_cpu += A[j] * B[j];
        }
        //CUTLASS, GPU
        prepareDot<<<1, 1>>>(A_dev, B_dev, i, K);
        cudaDeviceSynchronize();
        gemm_operator_tensorop(args);
        cudaMemcpy(result_gpu, C_dev, sizeof(int32_t), cudaMemcpyDeviceToHost);
        printf("i = %d, cpu = %d, gpu = %d\n", i, result_cpu, *result_gpu);
        //Re-init C 
        cudaMemcpy(C_dev, C, M * N * sizeof(int32_t), cudaMemcpyHostToDevice);
    }
}

This essentially computes a dot-product with the accumulator being initialized to a large value. After a couple of additional accumulations, the results between CPU and GPU differ by exactly one (as the last one in the bit-pattern of the accumulator is likely lost due to a conversion to float). The result of this looks like this:

Expected behavior
The results between cpu and gpu should match. When modifying gemm/device/default_gemm_configuration.h
by setting the compute-type in the epilogue-definition for the correct template-specialization to int32_t instead of float, the results are computed correctly, see e.g.

(This is probably an issue for the other specializations, but I only tested the one.)

Environment details (please complete the following information):
Tested on a HPC-system with A100, tested with cutlass-4.2.1 as well as 3.5.1. From looking at the above mentioned file, this might still be an issue in the more recent versions as well.

Additional context
Add any other context about the problem here.

Start with gemm/device/default_gemm_configuration.h and build the linked mwe.c against the reported CUTLASS versions on an A100 or comparable setup. Compare the CPU and GPU results as the accumulator grows, then verify the affected GEMM template specializations and their expected integer results.