[Bug Report] CPU and NPU produce different BFP8_B values from the same BF16 tensor.

[namhyeong-kim]

Describe the bug
The CPU and NPU produce different BFP8_B values from the same BF16 tensor.

To Reproduce
Steps to reproduce the behavior:

1. go to branch namhyeong/.... Link to branch
2. run tests/ttnn/unit_tests/test_bfp8_bf16_conversion.py Link to test file
   1. test_typecast_bf16_to_bfp8_b fails, while test_typecast_bfp8_b_to_bf16 passes.

Expected behavior
The CPU and NPU should produce the same BFP8_B values.

Screenshots

Please complete the following environment information:

• OS: [Ubuntu 20.04]

Additional context
No additional context

[rdjogoTT]

I have started investigating and should have an update by tomorrow.

[rdjogoTT]

I've managed to trace the data through the HW path it takes, and it looks like the issue occurs in the Packer. I will be investigating this further

[rdjogoTT]

The issue comes from how the packer handles bfp8_b datums, with rounding and shifting with rounding again. I'm looking into best course of action now.

Can you please provide some context to why the values need to match exactly? Is this just a bug report or is the accuracy needed somewhere specific? One potential fix would come at the expense of perf due to not using full packer bandwidth, so it would help to know the application.

[namhyeong-kim]

1. We will use BFP8_B for training large language models. I can't fully anticipate the model-level side effects caused by conversion inconsistencies between the CPU and NPU. Could you implement an option for precise conversion, even if it slightly reduces packer bandwidth?
2. CPU converts correctly, while the NPU converts approximately. Is that correct?

[rdjogoTT]

Yes, the CPU is producing the more accurate values in this case. Ok I will look into how that can be done.

[jvasilje]

@rdjogoTT any updates on this P0 bug?

[rdjogoTT]

I have pushed a branch origin/rd/bfp8b_pack with the workaround we found. It involves changing the pack src format for bfp8b to fp16b instead, which would cut packer bandwidth in half for WHB0 (but shouldn't affect BH).

This raises the question of how to handle this tradeoff. The way to decide would be to try training with bfp8b and see whether or not the accuracy is needed/tradeoff is worth it.

[namhyeong-kim]

It looks that I cannot apply the accurate conversion in program-object-level. Could you make an option that toggles this feature?

[rdjogoTT]

Ok, I will add an option ASAP.

[rdjogoTT]

#14822 contains the new flag bfp8_pack_precise which allows you to toggle the feature. Setting it to true in the ComputeConfig will result in the test above passing. I will work to merge this change quickly.

In the case of typecast, this is how it can be set to True:

auto eltwise_unary_kernel_group_1_id = tt::tt_metal::CreateKernel(
        program,
        "tt_metal/kernels/compute/eltwise_sfpu.cpp",
        core_group_1,
        tt::tt_metal::ComputeConfig{
            .math_fidelity = MathFidelity::HiFi4,
            .fp32_dest_acc_en = args.fp32_dest_acc_en,
            .unpack_to_dest_mode = unpack_to_dest_mode,
            .bfp8_pack_precise = true, // <----------- add this line
            .math_approx_mode = math_approx_mode,
            .compile_args = compute_kernel_args_group_1,
            .defines = unary_defines});

[namhyeong-kim]

Thank you! Could I check the feature in various input?
The test test_typecast_bf16_to_bfp8_b only demonstrates the existence of the bug; it doesn't have sufficient coverage.

[rdjogoTT]

The feature should enable fp16b or fp32 dest to be packed out accurately. Yes, please test it with various inputs.

[namhyeong-kim]

What is enabling fp16b?

[rdjogoTT]

Sorry I should have clarified - typecast works by unpacking the fp16b to Dest and using the packer to convert to bfp8b. Similar with fp32 to bfp8b typecast. This bug was caused by that packer conversion, so this feature should enable accurate conversion now.

[rdjogoTT]

Hi @namhyeong-kim, can you please share your tests with me when you make them. I would like to include them with my PR: #14822 before merging to main.

[namhyeong-kim]

Ok. I will share it

[namhyeong-kim]

I pushed the tests to my branch that is rebased on your branch.
Every test in test_typecast_bf16_to_bfp8_b_various_input passes.
But when insert_nan is true in test_typecast_bf16_to_bfp8_b_with_inf_nan, the test fails.
cpu makes nan but npu makes -inf.

idx=(tensor(0), tensor(1)) cpu=nan npu=-inf
idx=(tensor(0), tensor(4)) cpu=nan npu=-inf
idx=(tensor(0), tensor(18)) cpu=nan npu=-inf
idx=(tensor(0), tensor(29)) cpu=nan npu=-inf
idx=(tensor(0), tensor(35)) cpu=nan npu=-inf

Could you check these failed tests?

[rdjogoTT]

I've found that the issue is again to do with the packer, will post further updates soon.

[rdjogoTT]

Tensix team has confirmed that we have had issues in the past as well with these special values. Unfortunately nothing can be done from the HW/LLK side to remedy this.

LLK recommendation is to make the checks for special values aware that NaN in Bfp8_b may be converted to -Inf. If there is a case that you must know if there is a NaN, maybe check before converting to Bfp8_b?

[namhyeong-kim]

I understood the limitation. I will check before conversion to bfp8_b as your recommendation if nan should be handled.
Lastly please double check the tests that I provided have enough coverage.

[jvasilje]

@namhyeong-kim can we close this P0 issue?

[rdjogoTT]

@jvasilje I am working on adding the tests and merging the PR today.

[zzigler-tt]

@rdjogoTT Please advise when this is fully complete so we can close out this issue, thanks.

[rdjogoTT]

This has been resolved with PR: #14822. Can now be closed