finegrained-fp8: [bugfix] add epsilon to act_quant scale to avoid NaNs.#680
finegrained-fp8: [bugfix] add epsilon to act_quant scale to avoid NaNs.#680wz-ml wants to merge 2 commits into
Conversation
…s are 0 have a scale factor of 0 (yielding 0/0 = NaN)
|
Hi @wz-ml, thanks for your interest in contributing! This project requires that pull request authors are vouched, and you are not in the list of vouched users. This PR will be closed automatically. See https://github.com/huggingface/kernels-community/blob/main/CONTRIBUTING.md for more details. |
|
Hi @wz-ml, thanks for your interest in contributing! This project requires that pull request authors are vouched, and you are not in the list of vouched users. This PR will be closed automatically. See https://github.com/huggingface/kernels-community/blob/main/CONTRIBUTING.md for more details. |
|
@IlyasMoutawwakil could you do a review? |
| offs = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE) | ||
| x = tl.load(x_ptr + offs).to(tl.float32) | ||
| s = tl.max(tl.abs(x)) / 448.0 # float8_e4m3fn max | ||
| s = tl.maximum(s, 1e-12) |
There was a problem hiding this comment.
yes this actually needed, thanks for the fix.
|
it seems i can't approve the pr 🥲 |
|
@IlyasMoutawwakil this looks good to you? |
|
@sayakpaul yes i couldn't approve it before |
4 participants
A quick single-line change to fix this issue I was running into with this FP8 kernel:
When an X block is all zeros, the scale factor's also zero, which causes the kernel to return 0/0 = NaN.