[Bug Report] Incorrect Result in `mul_tiles` Calculation

[hschoi4448]

Describe the bug While debugging the op, I discovered that the result of mul_tiles is incorrect.

    // copy mask
    cb_wait_front(cb_mask_w, onetile);

    tile_regs_acquire();
    copy_tile_init();
    const uint32_t mask_dst = 8; // maybe this is problem, but why?
    copy_tile(cb_mask_w, first_tile, mask_dst);
    tile_regs_commit();

    tile_regs_wait();
    tile_regs_release();

    // multiplication
    cb_wait_front(cb_x, onetile);
    UNPACK(DPRINT << "ZZZZZZZZZZZZZZZ cb_x \n" << TSLICE(cb_x, 0, SliceRange{ .h0 = 0, .h1 = 4, .hs = 1, .w0 = 0, .w1 = 4, .ws = 1 }) << "\n";)

    cb_reserve_back(cb_xmm, onetile);

    tile_regs_acquire();
    mul_tiles_init(cb_x, cb_x);
    mul_tiles(cb_x, cb_x, 0, 0, dst0);
    tile_regs_commit();

    tile_regs_wait();
    pack_tile(dst0, cb_xmm);
    tile_regs_release();

    cb_push_back(cb_xmm, onetile);

    // print value
    cb_wait_front(cb_xmm, onetile);
    UNPACK(DPRINT << "ZZZZZZZZZZZZZZZ cb_xmm2 \n" << TSLICE(cb_xmm, 0, SliceRange{ .h0 = 0, .h1 = 4, .hs = 1, .w0 = 0, .w1 = 4, .ws = 1 }) << "\n";)
    cb_pop_front(cb_xmm, onetile);

compute code: https://github.com/tenstorrent/tt-metal/blob/hyungsuk/debug_mul_tiles/ttnn/cpp/ttnn/deprecated/tt_dnn/op_library/moreh_layernorm/kernels/moreh_layernorm_small_kernel.cpp#L41

The test scenario is as follows:

1. The input tensor for the test is a tiled input tensor of size [1,1,1,1] filled with zeros and padded with 'nan'.
2. The reader kernel generates 1 tile size arbitrary data.
3. In the compute kernel, arbitrary data is copied to destination register address 8.
4. In the compute kernel, multiplies the input tensor x input tensor
5. Upon checking the multiplied value, it is 0 x 0, but it contain 1.

ZZZZZZZZZZZZZZZ cb_x 
0 nan nan nan
nan nan nan nan
nan nan nan nan
nan nan nan nan

ZZZZZZZZZZZZZZZ2 cb_xmm2 
1 0 0 0  <--- 0 x 0 should be 0, but it's showing as 1.
0 0 0 0
0 0 0 0
0 0 0 0

This is output.

My guess is that the issue occurred because there are only 8 available destination registers, but a copy was made to the 9th register.

1. However, there is currently no way to set the number of available destination registers. Some documents indicate that there are 16 destination registers.
2. it seems wrong that a mistake in the copy operation affects mul_tiles

To Reproduce Steps to reproduce the behavior:

1. Go to https://github.com/tenstorrent/tt-metal/tree/hyungsuk/debug_mul_tiles
2. run pytest ./tests/tt_eager/python_api_testing/unit_testing/misc/test_moreh_groupnorm.py

Please complete the following environment information:

• OS: [e.g. Ubuntu 20.04]
• Version of software (eg. commit 6c9512c01b6aa275e86bd04bc57afa01e5523a3a)
• Arch: wormhole_b0

[rdjogoTT]

@hschoi4448 your intuition about what is causing the problem is correct, the issue is indeed due to the fact that you are copying a tile to the 9th tile spot of the Dest register. I will try to clarify what is going wrong:

The destination register (Dest) is one big register that can fit a total of 16 tiles (assuming 16bit datums). In most cases we want to use double-buffering for Dest to improve performance, so this reduces the number of usable tiles to 8 at a time. The way double-buffering works is that tiles 0-7 are used by math, and then we switch to using tiles 8-15 while 0-7 are being packed out. Then we switch back to 0-7 while packing 8-15, etc. The tile_regs_commit(); and tile_regs_release(); calls are what switch which half of Dest is in use by the Math/Packer.

This is important because by copying a tile to index 8 of Dest, you are actually putting it into index "0" of the following operation, which causes the corruption you are seeing in this test. If you change const uint32_t mask_dst = 8; to be any other value other than 8, you would no longer see the data corruption, because only index "0" of Dest is used by mul_tiles. This is why when you are using double-buffering for dest, you can't use more than 8 tiles of dest for your operation without causing issues for the next op.

*Note that this is all for 16bit data, and if you are using 32bit data then all the number get divided by 2 (8 total tiles in Dest with 4 usable at one time if using double-buffering).

Let me know if you have further questions or need more explanation

[hschoi4448]

Thank you for explaining the cause of the issue. However, I didn't want to know the cause of the problem; I created the issue because I believe that errors like this can negatively affect the development environment.

Here are the reasons why I think this is an issue:

1. Accessing an unusable register doesn't explicitly trigger an error, but silently corrupts the value.

   • Moreover, the point where the value gets corrupted is not where the mistake is made, but in the subsequent operation, leading to more time spent on debugging.

2. According to the current copy_tile documentation, it states that there are 16 dst registers.

   • The information about double buffering and the fact that the number of usable registers is halved in 32-bit is also being communicated through conversations between individuals, rather than through documented sources that everyone can access. Thank you for the explanation, but other developers will not be aware of this conversation and will likely encounter the same issues.

It seems that some form of improvement regarding this problem is necessary.

@rdjogoTT @razorback3

[rdjogoTT]

I agree with your point regarding the need for better clarity regarding this issue, I'll defer the decision regarding the priority of this issue to @ttmtrajkovic.

[ttmtrajkovic]

hello @hschoi4448, @razorback3,

Thank you for flagging the problem, it's been basically baked into our programming model from day 1, since all acquire/release functions always operate on halfs only because of the parallel processing of packer and math thread and double-buffering scheme on dest registers (pack thread operates on one half while math thread operates on other). Having asserts for out of bounds tile id is something we'd definitely want to have, but I would argue the priority of this given that it has been part of the model for a long time and it shouldn't be gating progress.

Could you please downgrade to P2 bug that we will address once all P0/P1 bugs are cleared out. Let me know what you think.

Milos