reduction fails due to accuracy.

[ingridli]

Hi experts,

When working on reduction problems, I often fail due to precision issues on A100 and H100. Later, I directly submitted several codes from the solution section, but I still encountered the same problem. The specific manifestation is as follows. Could someone please help me figure it out?

I tried the answer from "SteadfastTactician088" and "EnchantingSavant288".
Test failed! Here are the inputs:
input = [-1.0558, 6.6916, 6.4315, 1.4499, -5.5541, ..., 8.1268, 8.7411, -5.1739, -1.3938, 8.3171]
N = 50000000
Mismatch in 'output'
Expected: [-1883.689453125]
Got: [-1883.7373046875]
Max abs diff: 0.0478515625
Performance run failed correctness check

Thanks,
Ingrid.

[bagasjiwanta]

I came across this issue a few days ago too. Using precision='ieee' for tl.dot operation did it for me (I use triton).

Since A100s introduced TF32, have you tried disabling TF32 matmul? I believe in c++ (assuming you're using CUDA) you can do this.

[HoiM]

I am not the contributor of this project. But can you show the code? Maybe there does exist numerical issues in your code.

[ingridli]

I am not the contributor of this project. But can you show the code? Maybe there does exist numerical issues in your code.

Hi,
Thanks for your help.
I reproduce this abnormal behavios.
Again I choose the code from "Solutions" section, which I think it should be the correct answer. The username is "SteadfastTactician088" and his code is as follows. No offence. Just randomly pick one code.
`#include <cuda_runtime.h>
#define threadPertask 8
global void cuda_solve(const float* input, float* arr, int N){
shared float smem[256];
int idx = threadIdx.x + threadPertaskblockDim.xblockIdx.x;
int tid = threadIdx.x;
int laneid = tid%32;
int warp_id = tid/32;

float acc = 0.0;
for(int i = 0;i<threadPertask;i++){
    int gmem_idx = idx + i*256;
    if(gmem_idx<N){
        acc += input[gmem_idx];
    }
}
smem[tid] = acc;


__syncthreads();

float tmp = smem[tid];
for(int i = 32/2; i>0;i/=2){
    tmp +=__shfl_down_sync(0xffffffff,tmp,i);
}
if(laneid==0){
    smem[warp_id] = tmp;
}
__syncthreads();
if(warp_id == 0){
    tmp = laneid<8? smem[laneid] :0;
}
__syncthreads();

for(int i=8/2;i>0;i/=2){
    tmp += __shfl_down_sync(0xffffffff,tmp,i); 
}
if(tid == 0){
    atomicAdd(arr,tmp);
}


return;

}

// input, output are device pointers
extern "C" void solve(const float* input, float* output, int N) {
int threadPerblock = 256;
int blockPergrid = (N+threadPertaskthreadPerblock-1)/(threadPertaskthreadPerblock);
cuda_solve<<<blockPergrid,threadPerblock>>>(input,output,N);
}`

But I still don't pass the submission as the following results.
Test failed! Here are the inputs: input = [-8.1341, -4.7439, 0.8717, 6.6786, -8.0403, ..., -8.6186, -6.7947, 9.4231, 3.3224, 9.8672] N = 50000000 Mismatch in 'output'

I submit the code directly on the website.

BR,
Ingrid.

[ingridli]

I came across this issue a few days ago too. Using precision='ieee' for tl.dot operation did it for me (I use triton).

Since A100s introduced TF32, have you tried disabling TF32 matmul? I believe in c++ (assuming you're using CUDA) you can do this.

Hi,
Thanks for your help and sorry for the late reply. I will try it later.
BR,
Ingrid.

[HoiM]

I am not the contributor of this project. But can you show the code? Maybe there does exist numerical issues in your code.

Hi, Thanks for your help. I reproduce this abnormal behavios. Again I choose the code from "Solutions" section, which I think it should be the correct answer. The username is "SteadfastTactician088" and his code is as follows. No offence. Just randomly pick one code. `#include <cuda_runtime.h> #define threadPertask 8 global void cuda_solve(const float* input, float* arr, int N){ shared float smem[256]; int idx = threadIdx.x + threadPertask_blockDim.x_blockIdx.x; int tid = threadIdx.x; int laneid = tid%32; int warp_id = tid/32;

float acc = 0.0;
for(int i = 0;i<threadPertask;i++){
    int gmem_idx = idx + i*256;
    if(gmem_idx<N){
        acc += input[gmem_idx];
    }
}
smem[tid] = acc;


__syncthreads();

float tmp = smem[tid];
for(int i = 32/2; i>0;i/=2){
    tmp +=__shfl_down_sync(0xffffffff,tmp,i);
}
if(laneid==0){
    smem[warp_id] = tmp;
}
__syncthreads();
if(warp_id == 0){
    tmp = laneid<8? smem[laneid] :0;
}
__syncthreads();

for(int i=8/2;i>0;i/=2){
    tmp += __shfl_down_sync(0xffffffff,tmp,i); 
}
if(tid == 0){
    atomicAdd(arr,tmp);
}


return;

}

// input, output are device pointers extern "C" void solve(const float* input, float* output, int N) { int threadPerblock = 256; int blockPergrid = (N+threadPertask_threadPerblock-1)/(threadPertask_threadPerblock); cuda_solve<<<blockPergrid,threadPerblock>>>(input,output,N); }`

But I still don't pass the submission as the following results. Test failed! Here are the inputs: input = [-8.1341, -4.7439, 0.8717, 6.6786, -8.0403, ..., -8.6186, -6.7947, 9.4231, 3.3224, 9.8672] N = 50000000 Mismatch in 'output'

I submit the code directly on the website.

BR, Ingrid.

I also encountered this numerical issue on this "N = 50000000" test case. Here are my solutions:

1. Use Kahan sum instead of "+". When the number is large after iterative summation, adding a small number can result in data loss.
2. Use double precision.

Either of the two solutions can avoid the numerical problem.

[kunal-mansukhani]

I believe our test case is too large on this problem. We will reduce it and rerun all submissions.