[GH-ISSUE #32] Use f16 accumulator instead of f32 #21

New issue

Closed

opened 2026-10-26 21:10:39 +00:00 by bsakaguchi · 2 comments

bsakaguchi commented

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

Owner

Originally created by @i-zhen on GitHub (Mar 3, 2025).
Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/32

Hi, thanks a lot for your great contributions.

I just wonder that why a full-precision F32 accumulator is used instead of F16. Such as following assembly code:

wgmma.mma_async.sync.aligned.m64n16k32.f16.e4m3.e4m3

Can we also achieve an acceptable precision with f16 accumulator?

Originally created by @i-zhen on GitHub (Mar 3, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/32 Hi, thanks a lot for your great contributions. I just wonder that why a full-precision F32 accumulator is used instead of F16. Such as following assembly code: ``` wgmma.mma_async.sync.aligned.m64n16k32.f16.e4m3.e4m3 ``` Can we also achieve an acceptable precision with f16 accumulator?

bsakaguchi closed this issue

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

bsakaguchi commented

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

Author

Owner

@LyricZhao commented on GitHub (Mar 3, 2025):

Even FP32 accumulator is not enough (that's why we use two-level accumulation: tensor core (low-precision FP32 accum, hardward "bugs") + CUDA core (full FP32 accum)).

"Not enough" means bad for training.

However, for inference, I guess it is related to different model settings: one-level tensor core FP32 accum (works for DeepSeek-V2 series), or one-level tensor core FP16 (not tested on any model, you can try by yourself).

@LyricZhao commented on GitHub (Mar 3, 2025): Even FP32 accumulator is not enough (that's why we use two-level accumulation: tensor core (low-precision FP32 accum, hardward "bugs") + CUDA core (full FP32 accum)). "Not enough" means bad for training. However, for inference, I guess it is related to different model settings: one-level tensor core FP32 accum (works for DeepSeek-V2 series), or one-level tensor core FP16 (not tested on any model, you can try by yourself).

bsakaguchi commented

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

Author

Owner

@i-zhen commented on GitHub (Mar 4, 2025):

Even FP32 accumulator is not enough (that's why we use two-level accumulation: tensor core (low-precision FP32 accum, hardward "bugs") + CUDA core (full FP32 accum)).

"Not enough" means bad for training.

However, for inference, I guess it is related to different model settings: one-level tensor core FP32 accum (works for DeepSeek-V2 series), or one-level tensor core FP16 (not tested on any model, you can try by yourself).

Thank you for your kind reply~

@i-zhen commented on GitHub (Mar 4, 2025): > Even FP32 accumulator is not enough (that's why we use two-level accumulation: tensor core (low-precision FP32 accum, hardward "bugs") + CUDA core (full FP32 accum)). > > "Not enough" means bad for training. > > However, for inference, I guess it is related to different model settings: one-level tensor core FP32 accum (works for DeepSeek-V2 series), or one-level tensor core FP16 (not tested on any model, you can try by yourself). Thank you for your kind reply~

bsakaguchi referenced this issue

2026-10-26 21:14:52 +00:00

[PR #21] [CLOSED] pytest Integration #207