[GH-ISSUE #37] only uses the highest 14 bits

bsakaguchi commented

2026-10-26 21:10:42 +00:00

Owner

Originally created by @sisiwuxi on GitHub (Mar 4, 2025).
Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/37

3.5.2. Compute Hardware
Higher FP8 GEMM Accumulation Precision in Tensor Cores. In the current Tensor Core implementation of the NVIDIA Hopper architecture, FP8 GEMM (General Matrix Multiply) employs fixed-point accumulation, aligning the mantissa products by right-shifting based on the maximum exponent before addition. Our experiments reveal that it only uses the highest 14 bits of each mantissa product after sign-fill right shifting, and truncates bits exceeding this range. However, for example, to achieve precise FP32 results from the accumulation of 32 FP8×FP8 multiplications, at least 34-bit precision is required. Thus, we recommend that future chip designs increase accumulation precision in Tensor Cores to support full-precision accumulation, or select an appropriate accumulation bit-width according to the accuracy requirements of training and inference algorithms. This approach ensures that errors remain within acceptable bounds while maintaining computational efficiency.

请教 14bits 是怎样得到的？

实验步骤：

1、基于 tests/test_core.py 做修改得到 tests/test_core_simple.py，如附件

test_core_simple.py.txt

在 DeepGEMM目录执行 python3 tests/test_core_simple.py

修改部分：
m, k, n = 16, 256, 64
左支x，右支y 初始化为全 0，再分别赋值 x_[0,0] = 448, x_[0,1] = 7.5，y_[0,0] = 1, y_[0,1] = 1;
per_token_cast_to_fp8(x) 和 per_block_cast_to_fp8(y) 返回不做量化，scale全1的Tuple[torch.Tensor, torch.Tensor]；
实现调用 wgmma 指令的计算满足 a1 * b1 + a2 * b2 + 0 * 0 + ... + 0 * 0

2、实验结果发现调用 deep_gemm.gemm_fp8_fp8_bf16_nt(x_fp8, y_fp8, out) 后计算结果不满足预期，wgmma 指令 trunc 后保留的位数从5bit到8bit不等。

比如 448 * 1 + 7.5 * 1 + 0 + ... + 0 = 456 != 455.5

448 = (1.11)b * pow(2,8)
7.5 = (1.111)b * pow(2,2)
448 + 7.5 = (1.11)b * pow(2,8) + (0.00000 1111)b * pow(2,8) = (1.11000 1111)b * pow(2,8)
而
455.5 = (1.11000 1111)b * pow(2,8)
456 = (1.11001 0000)b * pow(2,8)
个人理解，仅保留了5bit的尾数

3、在 deep_gemm/include/deep_gemm/fp8_gemm.cuh/fp8_gemm_kernel末尾添加打印 printf("%f, ", (final_accum[0])); 得到 0.000000, 455.500000, 0.000000, 0.000000，观察到 448 + 7.5 此时的值为 455.5，满足预期。

printf("%f, ", (final_accum[0]));
修改为
printf("%.23f, ", (final_accum[0]));

Originally created by @sisiwuxi on GitHub (Mar 4, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/37 3.5.2. Compute Hardware Higher FP8 GEMM Accumulation Precision in Tensor Cores. In the current Tensor Core implementation of the NVIDIA Hopper architecture, FP8 GEMM (General Matrix Multiply) employs fixed-point accumulation, aligning the mantissa products by right-shifting based on the maximum exponent before addition. Our experiments reveal that it only uses the highest 14 bits of each mantissa product after sign-fill right shifting, and truncates bits exceeding this range. However, for example, to achieve precise FP32 results from the accumulation of 32 FP8×FP8 multiplications, at least 34-bit precision is required. Thus, we recommend that future chip designs increase accumulation precision in Tensor Cores to support full-precision accumulation, or select an appropriate accumulation bit-width according to the accuracy requirements of training and inference algorithms. This approach ensures that errors remain within acceptable bounds while maintaining computational efficiency. 请教 14bits 是怎样得到的？实验步骤： 1、基于 tests/test_core.py 做修改得到 tests/test_core_simple.py，如附件 [test_core_simple.py.txt](https://github.com/user-attachments/files/19070682/test_core_simple.py.txt) 在 DeepGEMM目录执行 python3 tests/test_core_simple.py 修改部分： m, k, n = 16, 256, 64 左支x，右支y 初始化为全 0，再分别赋值 x_[0,0] = 448, x_[0,1] = 7.5，y_[0,0] = 1, y_[0,1] = 1; per_token_cast_to_fp8(x) 和 per_block_cast_to_fp8(y) 返回不做量化，scale全1的Tuple[torch.Tensor, torch.Tensor]；实现调用 wgmma 指令的计算满足 a1 * b1 + a2 * b2 + 0 * 0 + ... + 0 * 0 2、实验结果发现调用 deep_gemm.gemm_fp8_fp8_bf16_nt(x_fp8, y_fp8, out) 后计算结果不满足预期，wgmma 指令 trunc 后保留的位数从5bit到8bit不等。比如 448 * 1 + 7.5 * 1 + 0 + ... + 0 = 456 != 455.5 448 = (1.11)b * pow(2,8) 7.5 = (1.111)b * pow(2,2) 448 + 7.5 = (1.11)b * pow(2,8) + (0.00000 1111)b * pow(2,8) = (1.11000 1111)b * pow(2,8) 而 455.5 = (1.11000 1111)b * pow(2,8) 456 = (1.11001 0000)b * pow(2,8) 个人理解，仅保留了5bit的尾数 3、在 deep_gemm/include/deep_gemm/fp8_gemm.cuh/fp8_gemm_kernel末尾添加打印 printf("%f, ", (final_accum[0])); 得到 0.000000, 455.500000, 0.000000, 0.000000，观察到 448 + 7.5 此时的值为 455.5，满足预期。 printf("%f, ", (final_accum[0])); 修改为 printf("%.23f, ", (final_accum[0]));

bsakaguchi closed this issue

2026-10-26 21:10:42 +00:00

bsakaguchi commented

2026-10-26 21:10:44 +00:00

Author

Owner

@jason-huang03 commented on GitHub (Mar 4, 2025):

@sisiwuxi Hi, you can refer to our work SageAttention2, which is the first paper (even earlier than DeepSeek V3) revealing the accumulator issue of Hopper/Ada fp8 tensor core. Actually we find that fp8 has 22 bits accumulator, rather than 14 bits as mentioned by deepseek. We conduct the experiment as follows:

We write a small kernel that performs one fp8 wgmma instruction (for example, m64n8k128), taking in A, B, C and performs D = A * B + C where A, B are in e4m3, C, D are in fp32.
We initialize A and B to be zero.
We varies the effective bits of C.
We find that as long as C has less or equal to 22 effective bits, then D = C. Else, the result D is equivalent to truncating the last 10 bits of C.

In this way, we conclude that the fp8 accumulator has 22 effective bits. In other word, the last 10 bits of accumulator are always truncated by fp8 tensor core operation.

@jason-huang03 commented on GitHub (Mar 4, 2025): @sisiwuxi Hi, you can refer to our work [SageAttention2](https://arxiv.org/abs/2411.10958), which is the first paper (even earlier than DeepSeek V3) revealing the accumulator issue of Hopper/Ada fp8 tensor core. Actually we find that fp8 has 22 bits accumulator, rather than 14 bits as mentioned by deepseek. We conduct the experiment as follows: 1. We write a small kernel that performs one fp8 wgmma instruction (for example, m64n8k128), taking in A, B, C and performs D = A * B + C where A, B are in e4m3, C, D are in fp32. 2. We initialize A and B to be zero. 3. We varies the effective bits of C. 4. We find that as long as C has less or equal to 22 effective bits, then D = C. Else, the result D is equivalent to truncating the last 10 bits of C. In this way, we conclude that the fp8 accumulator has 22 effective bits. In other word, the last 10 bits of accumulator are always truncated by fp8 tensor core operation.

bsakaguchi commented

2026-10-26 21:10:44 +00:00

Author

Owner

@soundOfDestiny commented on GitHub (Mar 4, 2025):

这里y的amax不是448，因此per_block_cast_to_fp8(y)返回的y的scale不会是1吧？

@soundOfDestiny commented on GitHub (Mar 4, 2025): > 左支x，右支y 初始化为全 0，再分别赋值 x_[0,0] = 448, x_[0,1] = 7.5，y_[0,0] = 1, y_[0,1] = 1; per_token_cast_to_fp8(x) 和 per_block_cast_to_fp8(y) 返回不做量化，scale全1的Tuple[torch.Tensor, torch.Tensor]；这里y的amax不是448，因此per_block_cast_to_fp8(y)返回的y的scale不会是1吧？

bsakaguchi commented

2026-10-26 21:10:45 +00:00

Author

Owner

@sisiwuxi commented on GitHub (Mar 5, 2025):

仅保留 14bit 复制到了
ok 16 * 1 + 0.029296875 * 1 + 0 + ... + 0, (16.029296875)d = (1. 0000 00000 1111)b * 2**(4)
err 32 * 1 + 0.029296875 * 1 + 0 + ... + 0, (32.029296875)d = (1. 00000 00000 1111)b * 2**(5)，但是final_accum[0]打印结果32.02734375

32.029296875 = (1. 00000 00000 1111)b * 2**(5)
32.02734375 = (1. 00000 00000 1110)b * 2**(5)
2. printf("\n [tid_0]: %.23f, %.23f \n", accum[0], final_accum[0]); print打印保留23bit精度

@sisiwuxi commented on GitHub (Mar 5, 2025): 1. 仅保留 14bit 复制到了 ok 16 * 1 + 0.029296875 * 1 + 0 + ... + 0, (16.029296875)d = (1. 0000 00000 1111)b * 2**(4) err 32 * 1 + 0.029296875 * 1 + 0 + ... + 0, (32.029296875)d = (1. 00000 00000 1111)b * 2**(5)，但是final_accum[0]打印结果32.02734375 32.029296875 = (1. 00000 00000 1111)b * 2**(5) 32.02734375 = (1. 00000 00000 1110)b * 2**(5) 2. printf("\n [tid_0]: %.23f, %.23f \n", accum[0], final_accum[0]); print打印保留23bit精度

bsakaguchi commented

2026-10-26 21:10:45 +00:00

Author

Owner

@soundOfDestiny commented on GitHub (Mar 5, 2025):

'14 bits' in DeepSeekV3 paper includes a sign bit and mantissa bits. 14bits + 8 exponent bits == 22 bits in total

@soundOfDestiny commented on GitHub (Mar 5, 2025): > [@sisiwuxi](https://github.com/sisiwuxi) Hi, you can refer to our work [SageAttention2](https://arxiv.org/abs/2411.10958), which is the first paper (even earlier than DeepSeek V3) revealing the accumulator issue of Hopper/Ada fp8 tensor core. Actually we find that fp8 has 22 bits accumulator, rather than 14 bits as mentioned by deepseek. We conduct the experiment as follows: > > 1. We write a small kernel that performs one fp8 wgmma instruction (for example, m64n8k128), taking in A, B, C and performs D = A * B + C where A, B are in e4m3, C, D are in fp32. > 2. We initialize A and B to be zero. > 3. We varies the effective bits of C. > 4. We find that as long as C has less or equal to 22 effective bits, then D = C. Else, the result D is equivalent to truncating the last 10 bits of C. > > In this way, we conclude that the fp8 accumulator has 22 effective bits. In other word, the last 10 bits of accumulator are always truncated by fp8 tensor core operation. '14 bits' in DeepSeekV3 paper includes a sign bit and mantissa bits. 14bits + 8 exponent bits == 22 bits in total

bsakaguchi commented

2026-10-26 21:10:45 +00:00

Author

Owner

@soundOfDestiny commented on GitHub (Mar 5, 2025):

这里y的amax不是448，因此per_block_cast_to_fp8(y)返回的y的scale不会是1吧？

@soundOfDestiny commented on GitHub (Mar 5, 2025): > 1. 仅保留 14bit 复制到了 > ok 16_1 + 0.029296875_1 + 0 + ... + 0, (16.029296875)d = (1. 0000 00000 1111)b * 2**(4) > err 32_1 + 0.029296875_1 + 0 + ... + 0, (32.029296875)d = (1. 00000 00000 1111)b * 2**(5)，但是final_accum[0]打印结果32.02734375 > > 32.029296875 = (1. 00000 00000 1111)b * 2**(5) 32.02734375 = (1. 00000 00000 1110)b * 2**(5) 2. printf("\n [tid_0]: %.23f, %.23f \n", accum[0], final_accum[0]); print打印保留23bit精度 > 左支x，右支y 初始化为全 0，再分别赋值 x_[0,0] = 448, x_[0,1] = 7.5，y_[0,0] = 1, y_[0,1] = 1; per_token_cast_to_fp8(x) 和 per_block_cast_to_fp8(y) 返回不做量化，scale全1的Tuple[torch.Tensor, torch.Tensor]；这里y的amax不是448，因此per_block_cast_to_fp8(y)返回的y的scale不会是1吧？

bsakaguchi commented

2026-10-26 21:10:45 +00:00

Author

Owner

@sisiwuxi commented on GitHub (Mar 5, 2025):

附件test_core_simple.py第26和第50行，左支右支的scale都手动修改为全1，避免wgmma之后cuda core乘scale的影响

@sisiwuxi commented on GitHub (Mar 5, 2025): > > 左支x，右支y 初始化为全 0，再分别赋值 x_[0,0] = 448, x_[0,1] = 7.5，y_[0,0] = 1, y_[0,1] = 1; > > per_token_cast_to_fp8(x) 和 per_block_cast_to_fp8(y) 返回不做量化，scale全1的Tuple[torch.Tensor, torch.Tensor]； > > 这里y的amax不是448，因此per_block_cast_to_fp8(y)返回的y的scale不会是1吧？附件test_core_simple.py第26和第50行，左支右支的scale都手动修改为全1，避免wgmma之后cuda core乘scale的影响

bsakaguchi commented

2026-10-26 21:10:46 +00:00

Author

Owner

@sisiwuxi commented on GitHub (Mar 5, 2025):

感谢回复，和实验结果相同，多谢

@sisiwuxi commented on GitHub (Mar 5, 2025): > > [@sisiwuxi](https://github.com/sisiwuxi) Hi, you can refer to our work [SageAttention2](https://arxiv.org/abs/2411.10958), which is the first paper (even earlier than DeepSeek V3) revealing the accumulator issue of Hopper/Ada fp8 tensor core. Actually we find that fp8 has 22 bits accumulator, rather than 14 bits as mentioned by deepseek. We conduct the experiment as follows: > > > > 1. We write a small kernel that performs one fp8 wgmma instruction (for example, m64n8k128), taking in A, B, C and performs D = A * B + C where A, B are in e4m3, C, D are in fp32. > > 2. We initialize A and B to be zero. > > 3. We varies the effective bits of C. > > 4. We find that as long as C has less or equal to 22 effective bits, then D = C. Else, the result D is equivalent to truncating the last 10 bits of C. > > > > In this way, we conclude that the fp8 accumulator has 22 effective bits. In other word, the last 10 bits of accumulator are always truncated by fp8 tensor core operation. > > '14 bits' in DeepSeekV3 paper includes a sign bit and mantissa bits. 14bits + 8 exponent bits == 22 bits in total 感谢回复，和实验结果相同，多谢

bsakaguchi commented

2026-10-26 21:10:46 +00:00

Author

Owner

@sisiwuxi commented on GitHub (Mar 5, 2025):

感谢回复，和实验结果相同，多谢

@sisiwuxi commented on GitHub (Mar 5, 2025): > [@sisiwuxi](https://github.com/sisiwuxi) Hi, you can refer to our work [SageAttention2](https://arxiv.org/abs/2411.10958), which is the first paper (even earlier than DeepSeek V3) revealing the accumulator issue of Hopper/Ada fp8 tensor core. Actually we find that fp8 has 22 bits accumulator, rather than 14 bits as mentioned by deepseek. We conduct the experiment as follows: > > 1. We write a small kernel that performs one fp8 wgmma instruction (for example, m64n8k128), taking in A, B, C and performs D = A * B + C where A, B are in e4m3, C, D are in fp32. > 2. We initialize A and B to be zero. > 3. We varies the effective bits of C. > 4. We find that as long as C has less or equal to 22 effective bits, then D = C. Else, the result D is equivalent to truncating the last 10 bits of C. > > In this way, we conclude that the fp8 accumulator has 22 effective bits. In other word, the last 10 bits of accumulator are always truncated by fp8 tensor core operation. 感谢回复，和实验结果相同，多谢

bsakaguchi commented

2026-10-26 21:10:47 +00:00

Author

Owner

@soundOfDestiny commented on GitHub (Mar 5, 2025):

455.5到456是结果转换成bf16的时候四舍五入了

@soundOfDestiny commented on GitHub (Mar 5, 2025): > 3.5.2. Compute Hardware Higher FP8 GEMM Accumulation Precision in Tensor Cores. In the current Tensor Core implementation of the NVIDIA Hopper architecture, FP8 GEMM (General Matrix Multiply) employs fixed-point accumulation, aligning the mantissa products by right-shifting based on the maximum exponent before addition. Our experiments reveal that it only uses the highest 14 bits of each mantissa product after sign-fill right shifting, and truncates bits exceeding this range. However, for example, to achieve precise FP32 results from the accumulation of 32 FP8×FP8 multiplications, at least 34-bit precision is required. Thus, we recommend that future chip designs increase accumulation precision in Tensor Cores to support full-precision accumulation, or select an appropriate accumulation bit-width according to the accuracy requirements of training and inference algorithms. This approach ensures that errors remain within acceptable bounds while maintaining computational efficiency. > > 请教 14bits 是怎样得到的？ > > 实验步骤： > > 1、基于 tests/test_core.py 做修改得到 tests/test_core_simple.py，如附件 > > [test_core_simple.py.txt](https://github.com/user-attachments/files/19070682/test_core_simple.py.txt) > > 在 DeepGEMM目录执行 python3 tests/test_core_simple.py > > 修改部分： m, k, n = 16, 256, 64 左支x，右支y 初始化为全 0，再分别赋值 x_[0,0] = 448, x_[0,1] = 7.5，y_[0,0] = 1, y_[0,1] = 1; per_token_cast_to_fp8(x) 和 per_block_cast_to_fp8(y) 返回不做量化，scale全1的Tuple[torch.Tensor, torch.Tensor]；实现调用 wgmma 指令的计算满足 a1 * b1 + a2 * b2 + 0 * 0 + ... + 0 * 0 > > 2、实验结果发现调用 deep_gemm.gemm_fp8_fp8_bf16_nt(x_fp8, y_fp8, out) 后计算结果不满足预期，wgmma 指令 trunc 后保留的位数从5bit到8bit不等。 > > 比如 448 * 1 + 7.5 * 1 + 0 + ... + 0 = 456 != 455.5 > > 448 = (1.11)b * pow(2,8) 7.5 = (1.111)b * pow(2,2) 448 + 7.5 = (1.11)b * pow(2,8) + (0.00000 1111)b * pow(2,8) = (1.11000 1111)b * pow(2,8) 而 455.5 = (1.11000 1111)b * pow(2,8) 456 = (1.11001 0000)b * pow(2,8) 个人理解，仅保留了5bit的尾数 > > 3、在 deep_gemm/include/deep_gemm/fp8_gemm.cuh/fp8_gemm_kernel末尾添加打印 printf("%f, ", (final_accum[0])); 得到 0.000000, 455.500000, 0.000000, 0.000000，观察到 448 + 7.5 此时的值为 455.5，满足预期。 > > printf("%f, ", (final_accum[0])); 修改为 printf("%.23f, ", (final_accum[0])); 455.5到456是结果转换成bf16的时候四舍五入了

bsakaguchi commented

2026-10-26 21:10:47 +00:00

Author

Owner

@seanxcwang commented on GitHub (Jun 9, 2025):

I followed this and wrote an MMA example running on a Hopper architecture GPU, but I didn't observe the result where C != D. Did I misunderstand something?

__global__ void tensor_core_example_16x8x32(float *D) {
    
    uint32_t a[4] = {0x0, 0x0, 0x0, 0x0};
    uint32_t b[2] = {0x0, 0x0};

    uint32_t t = 0b01000000010000000011000000010111;
    float t_f = *((float*)&t);

    float c[4] = {t_f, t_f, t_f, t_f};
    float d[4];
    
    asm(
        "mma.sync.aligned.m16n8k32.row.col.f32.e4m3.e4m3.f32 "
        "{%0,%1,%2,%3}, {%4,%5,%6,%7}, {%8,%9}, {%10,%11,%12,%13};\n"
        : "=f"(d[0]), "=f"(d[1]), "=f"(d[2]), "=f"(d[3])
        :
        "r"(a[0]), "r"(a[1]), "r"(a[2]), "r"(a[3]),
        "r"(b[0]), "r"(b[1]),
        "f"(c[0]), "f"(c[1]), "f"(c[2]), "f"(c[3])
    );

    if (threadIdx.x == 0) {
        printf("t_f = %f\n", t_f);
        printf("c[0] = %f, c[1] = %f, c[2] = %f, c[3] = %f\n", c[0], c[1], c[2], c[3]);
        printf("d[0] = %f, d[1] = %f, d[2] = %f, d[3] = %f\n", d[0], d[1], d[2], d[3]);
    }

    D[threadIdx.x * 4 + 0] = d[0];
    D[threadIdx.x * 4 + 1] = d[1];
    D[threadIdx.x * 4 + 2] = d[2];
    D[threadIdx.x * 4 + 3] = d[3];
}

mma_example_fp8.cu.txt

@seanxcwang commented on GitHub (Jun 9, 2025): > [@sisiwuxi](https://github.com/sisiwuxi) Hi, you can refer to our work [SageAttention2](https://arxiv.org/abs/2411.10958), which is the first paper (even earlier than DeepSeek V3) revealing the accumulator issue of Hopper/Ada fp8 tensor core. Actually we find that fp8 has 22 bits accumulator, rather than 14 bits as mentioned by deepseek. We conduct the experiment as follows: > > 1. We write a small kernel that performs one fp8 wgmma instruction (for example, m64n8k128), taking in A, B, C and performs D = A * B + C where A, B are in e4m3, C, D are in fp32. > 2. We initialize A and B to be zero. > 3. We varies the effective bits of C. > 4. We find that as long as C has less or equal to 22 effective bits, then D = C. Else, the result D is equivalent to truncating the last 10 bits of C. > > In this way, we conclude that the fp8 accumulator has 22 effective bits. In other word, the last 10 bits of accumulator are always truncated by fp8 tensor core operation. I followed this and wrote an MMA example running on a Hopper architecture GPU, but I didn't observe the result where C != D. Did I misunderstand something? ``` __global__ void tensor_core_example_16x8x32(float *D) { uint32_t a[4] = {0x0, 0x0, 0x0, 0x0}; uint32_t b[2] = {0x0, 0x0}; uint32_t t = 0b01000000010000000011000000010111; float t_f = *((float*)&t); float c[4] = {t_f, t_f, t_f, t_f}; float d[4]; asm( "mma.sync.aligned.m16n8k32.row.col.f32.e4m3.e4m3.f32 " "{%0,%1,%2,%3}, {%4,%5,%6,%7}, {%8,%9}, {%10,%11,%12,%13};\n" : "=f"(d[0]), "=f"(d[1]), "=f"(d[2]), "=f"(d[3]) : "r"(a[0]), "r"(a[1]), "r"(a[2]), "r"(a[3]), "r"(b[0]), "r"(b[1]), "f"(c[0]), "f"(c[1]), "f"(c[2]), "f"(c[3]) ); if (threadIdx.x == 0) { printf("t_f = %f\n", t_f); printf("c[0] = %f, c[1] = %f, c[2] = %f, c[3] = %f\n", c[0], c[1], c[2], c[3]); printf("d[0] = %f, d[1] = %f, d[2] = %f, d[3] = %f\n", d[0], d[1], d[2], d[3]); } D[threadIdx.x * 4 + 0] = d[0]; D[threadIdx.x * 4 + 1] = d[1]; D[threadIdx.x * 4 + 2] = d[2]; D[threadIdx.x * 4 + 3] = d[3]; } ``` [mma_example_fp8.cu.txt](https://github.com/user-attachments/files/20654312/mma_example_fp8.cu.txt)

bsakaguchi commented

2026-10-26 21:10:48 +00:00

Author

Owner

@jason-huang03 commented on GitHub (Jun 9, 2025):

@sisiwuxi on hopper, the mma fp8 instruction will fall back to fp16 (hence super slow). Instead, you should use wgmma on hopper to reproduce the problem.

@jason-huang03 commented on GitHub (Jun 9, 2025): > > [@sisiwuxi](https://github.com/sisiwuxi) Hi, you can refer to our work [SageAttention2](https://arxiv.org/abs/2411.10958), which is the first paper (even earlier than DeepSeek V3) revealing the accumulator issue of Hopper/Ada fp8 tensor core. Actually we find that fp8 has 22 bits accumulator, rather than 14 bits as mentioned by deepseek. We conduct the experiment as follows: > > > > 1. We write a small kernel that performs one fp8 wgmma instruction (for example, m64n8k128), taking in A, B, C and performs D = A * B + C where A, B are in e4m3, C, D are in fp32. > > 2. We initialize A and B to be zero. > > 3. We varies the effective bits of C. > > 4. We find that as long as C has less or equal to 22 effective bits, then D = C. Else, the result D is equivalent to truncating the last 10 bits of C. > > > > In this way, we conclude that the fp8 accumulator has 22 effective bits. In other word, the last 10 bits of accumulator are always truncated by fp8 tensor core operation. > > I followed this and wrote an MMA example running on a Hopper architecture GPU, but I didn't observe the result where C != D. Did I misunderstand something? > > ``` > __global__ void tensor_core_example_16x8x32(float *D) { > > uint32_t a[4] = {0x0, 0x0, 0x0, 0x0}; > uint32_t b[2] = {0x0, 0x0}; > > uint32_t t = 0b01000000010000000011000000010111; > float t_f = *((float*)&t); > > float c[4] = {t_f, t_f, t_f, t_f}; > float d[4]; > > asm( > "mma.sync.aligned.m16n8k32.row.col.f32.e4m3.e4m3.f32 " > "{%0,%1,%2,%3}, {%4,%5,%6,%7}, {%8,%9}, {%10,%11,%12,%13};\n" > : "=f"(d[0]), "=f"(d[1]), "=f"(d[2]), "=f"(d[3]) > : > "r"(a[0]), "r"(a[1]), "r"(a[2]), "r"(a[3]), > "r"(b[0]), "r"(b[1]), > "f"(c[0]), "f"(c[1]), "f"(c[2]), "f"(c[3]) > ); > > if (threadIdx.x == 0) { > printf("t_f = %f\n", t_f); > printf("c[0] = %f, c[1] = %f, c[2] = %f, c[3] = %f\n", c[0], c[1], c[2], c[3]); > printf("d[0] = %f, d[1] = %f, d[2] = %f, d[3] = %f\n", d[0], d[1], d[2], d[3]); > } > > D[threadIdx.x * 4 + 0] = d[0]; > D[threadIdx.x * 4 + 1] = d[1]; > D[threadIdx.x * 4 + 2] = d[2]; > D[threadIdx.x * 4 + 3] = d[3]; > } > ``` > > [mma_example_fp8.cu.txt](https://github.com/user-attachments/files/20654312/mma_example_fp8.cu.txt) @sisiwuxi on hopper, the mma fp8 instruction will fall back to fp16 (hence super slow). Instead, you should use wgmma on hopper to reproduce the problem.

bsakaguchi commented

2026-10-26 21:10:48 +00:00

Author

Owner

@seanxcwang commented on GitHub (Jun 10, 2025):

thx for the reply, wgmma produce the result where C != D

__global__ void tensor_core_wgmma_64x16x32() {
    uint32_t t = 0b01000000010000000011000001111111;
    float t_f = *((float*)&t);
    float d[8] = {t_f, t_f, t_f, t_f, t_f, t_f, t_f, t_f};

    __shared__ __align__(1024) uint8_t smem_buffer_a[4096];
    __shared__ __align__(1024) uint8_t smem_buffer_b[4096];

    if (threadIdx.x == 0) {
        for (int i = 0; i < 4096; i++) {
            smem_buffer_a[i] = 0;
            smem_buffer_b[i] = 0;
        }
    }

    __syncthreads();

    if (threadIdx.x == 0) {
        printf("\nhopper wgmma result!!!\n");
        printf("d[0] = %f, d[1] = %f, d[2] = %f, d[3] = %f\n", d[0], d[1], d[2], d[3]);
    }

    auto desc_a = make_smem_desc(smem_buffer_a, 1);
    auto desc_b = make_smem_desc(smem_buffer_b, 1);

    uint32_t scale_d = 1;

    cute::warpgroup_arrive();

    asm volatile("{\n"
                    ".reg .pred p;\n"
                    "setp.ne.b32 p, %10, 0;\n"
                    "wgmma.mma_async.sync.aligned.m64n16k32.f32.e4m3.e4m3"
                    "{%0,   %1,   %2,   %3,   %4,   %5,   %6,   %7},"
                    " %8,"
                    " %9,"
                    " p   , 1,    1;\n"
                    "}\n"
                    : "+f"(d[0]), "+f"(d[1]), "+f"(d[2]), "+f"(d[3]), "+f"(d[4]), "+f"(d[5]), "+f"(d[6]), "+f"(d[7])
                    : "l"(desc_a), "l"(desc_b), "r"(int32_t(scale_d)));

    cute::warpgroup_commit_batch();
    cute::warpgroup_wait<0>();

    if (threadIdx.x == 0) {
        printf("d[0] = %f, d[1] = %f, d[2] = %f, d[3] = %f\n", d[0], d[1], d[2], d[3]);
    }
}

// hopper wgmma result!!!
// d[0] = 3.002960, d[1] = 3.002960, d[2] = 3.002960, d[3] = 3.002960
// d[0] = 3.002930, d[1] = 3.002930, d[2] = 3.002930, d[3] = 3.002930

@seanxcwang commented on GitHub (Jun 10, 2025): > > > [@sisiwuxi](https://github.com/sisiwuxi) Hi, you can refer to our work [SageAttention2](https://arxiv.org/abs/2411.10958), which is the first paper (even earlier than DeepSeek V3) revealing the accumulator issue of Hopper/Ada fp8 tensor core. Actually we find that fp8 has 22 bits accumulator, rather than 14 bits as mentioned by deepseek. We conduct the experiment as follows: > > > > > > 1. We write a small kernel that performs one fp8 wgmma instruction (for example, m64n8k128), taking in A, B, C and performs D = A * B + C where A, B are in e4m3, C, D are in fp32. > > > 2. We initialize A and B to be zero. > > > 3. We varies the effective bits of C. > > > 4. We find that as long as C has less or equal to 22 effective bits, then D = C. Else, the result D is equivalent to truncating the last 10 bits of C. > > > > > > In this way, we conclude that the fp8 accumulator has 22 effective bits. In other word, the last 10 bits of accumulator are always truncated by fp8 tensor core operation. > > > > > > I followed this and wrote an MMA example running on a Hopper architecture GPU, but I didn't observe the result where C != D. Did I misunderstand something? > > ``` > > __global__ void tensor_core_example_16x8x32(float *D) { > > > > uint32_t a[4] = {0x0, 0x0, 0x0, 0x0}; > > uint32_t b[2] = {0x0, 0x0}; > > > > uint32_t t = 0b01000000010000000011000000010111; > > float t_f = *((float*)&t); > > > > float c[4] = {t_f, t_f, t_f, t_f}; > > float d[4]; > > > > asm( > > "mma.sync.aligned.m16n8k32.row.col.f32.e4m3.e4m3.f32 " > > "{%0,%1,%2,%3}, {%4,%5,%6,%7}, {%8,%9}, {%10,%11,%12,%13};\n" > > : "=f"(d[0]), "=f"(d[1]), "=f"(d[2]), "=f"(d[3]) > > : > > "r"(a[0]), "r"(a[1]), "r"(a[2]), "r"(a[3]), > > "r"(b[0]), "r"(b[1]), > > "f"(c[0]), "f"(c[1]), "f"(c[2]), "f"(c[3]) > > ); > > > > if (threadIdx.x == 0) { > > printf("t_f = %f\n", t_f); > > printf("c[0] = %f, c[1] = %f, c[2] = %f, c[3] = %f\n", c[0], c[1], c[2], c[3]); > > printf("d[0] = %f, d[1] = %f, d[2] = %f, d[3] = %f\n", d[0], d[1], d[2], d[3]); > > } > > > > D[threadIdx.x * 4 + 0] = d[0]; > > D[threadIdx.x * 4 + 1] = d[1]; > > D[threadIdx.x * 4 + 2] = d[2]; > > D[threadIdx.x * 4 + 3] = d[3]; > > } > > ``` > > > > > > > > > > > > > > > > > > > > > > > > [mma_example_fp8.cu.txt](https://github.com/user-attachments/files/20654312/mma_example_fp8.cu.txt) > > [@sisiwuxi](https://github.com/sisiwuxi) on hopper, the mma fp8 instruction will fall back to fp16 (hence super slow). Instead, you should use wgmma on hopper to reproduce the problem. thx for the reply, wgmma produce the result where C != D ``` __global__ void tensor_core_wgmma_64x16x32() { uint32_t t = 0b01000000010000000011000001111111; float t_f = *((float*)&t); float d[8] = {t_f, t_f, t_f, t_f, t_f, t_f, t_f, t_f}; __shared__ __align__(1024) uint8_t smem_buffer_a[4096]; __shared__ __align__(1024) uint8_t smem_buffer_b[4096]; if (threadIdx.x == 0) { for (int i = 0; i < 4096; i++) { smem_buffer_a[i] = 0; smem_buffer_b[i] = 0; } } __syncthreads(); if (threadIdx.x == 0) { printf("\nhopper wgmma result!!!\n"); printf("d[0] = %f, d[1] = %f, d[2] = %f, d[3] = %f\n", d[0], d[1], d[2], d[3]); } auto desc_a = make_smem_desc(smem_buffer_a, 1); auto desc_b = make_smem_desc(smem_buffer_b, 1); uint32_t scale_d = 1; cute::warpgroup_arrive(); asm volatile("{\n" ".reg .pred p;\n" "setp.ne.b32 p, %10, 0;\n" "wgmma.mma_async.sync.aligned.m64n16k32.f32.e4m3.e4m3" "{%0, %1, %2, %3, %4, %5, %6, %7}," " %8," " %9," " p , 1, 1;\n" "}\n" : "+f"(d[0]), "+f"(d[1]), "+f"(d[2]), "+f"(d[3]), "+f"(d[4]), "+f"(d[5]), "+f"(d[6]), "+f"(d[7]) : "l"(desc_a), "l"(desc_b), "r"(int32_t(scale_d))); cute::warpgroup_commit_batch(); cute::warpgroup_wait<0>(); if (threadIdx.x == 0) { printf("d[0] = %f, d[1] = %f, d[2] = %f, d[3] = %f\n", d[0], d[1], d[2], d[3]); } } // hopper wgmma result!!! // d[0] = 3.002960, d[1] = 3.002960, d[2] = 3.002960, d[3] = 3.002960 // d[0] = 3.002930, d[1] = 3.002930, d[2] = 3.002930, d[3] = 3.002930 ```

Rows
Columns

[GH-ISSUE #37] only uses the highest 14 bits #24