No test is available for in/int8 and out/int32 #7

Activity

1. rwvo commented on Sep 27, 2023

   

   rwvo

   on Sep 27, 2023

   Contributor

   More actions

   Mapping matrix entries to registers for the MFMA instructions can be highly confusing. Fortunately, Joe Greathouse made a tool that helps tremendously: https://github.com/RadeonOpenCompute/amd_matrix_instruction_calculator.

   For the instruction you are looking at, for instance (showing partial output) :

   $ ./matrix_calculator.py --architecture CDNA2 -i v_mfma_i32_16x16x16i8 --matrix-layout -A
   Architecture: CDNA2
   Instruction: V_MFMA_I32_16X16X16I8
   +--------+------------+-------------+--------------+--------------+
   |   lane | v0.[7:0]   | v0.[15:8]   | v0.[23:16]   | v0.[31:24]   |
   +========+============+=============+==============+==============+
   |      0 | A[0][0]    | A[0][1]     | A[0][2]      | A[0][3]      |
   +--------+------------+-------------+--------------+--------------+
   |      1 | A[1][0]    | A[1][1]     | A[1][2]      | A[1][3]      |
   +--------+------------+-------------+--------------+--------------+
   |      2 | A[2][0]    | A[2][1]     | A[2][2]      | A[2][3]      |
   +--------+------------+-------------+--------------+--------------+
   |      3 | A[3][0]    | A[3][1]     | A[3][2]      | A[3][3]      |
   +--------+------------+-------------+--------------+--------------+
   |      4 | A[4][0]    | A[4][1]     | A[4][2]      | A[4][3]      |
   +--------+------------+-------------+--------------+--------------+
   |      5 | A[5][0]    | A[5][1]     | A[5][2]      | A[5][3]      |
   +--------+------------+-------------+--------------+--------------+
   |      6 | A[6][0]    | A[6][1]     | A[6][2]      | A[6][3]      |
   +--------+------------+-------------+--------------+--------------+
   |      7 | A[7][0]    | A[7][1]     | A[7][2]      | A[7][3]      |
   +--------+------------+-------------+--------------+--------------+
   |      8 | A[8][0]    | A[8][1]     | A[8][2]      | A[8][3]      |
   +--------+------------+-------------+--------------+--------------+
   

   The mapping for matrices B, C and D can be queried similarly.

   ❤️React with ❤️1Reacted by Gina Sitaraman
2. mmoadeli commented on Sep 27, 2023

   

   mmoadeli

   on Sep 27, 2023

   Last edited by mmoadeli

   Author

   More actions

   @rwvo Thanks so much for reply.

   I already was aware of the amd matrix calculator. To me the above instruction is similar to __builtin_amdgcn_mfma_f32_16x16x16f16, used in sample.
   A difference to me seems to be a and b types.
   a and b in this sample are of below type
   using float16x4 = __attribute__((__vector_size__(4 * sizeof(float16_t)))) float16_t;
   while they are int32_t for __builtin_amdgcn_mfma_i32_16x16x16i8 case. That's why I have populated an int32_t type with 4 int8_t values.
   Would you be able to advise how should I prepare a and b for above instruction please?
   Ideally, I'd be grateful to have revised sample for __builtin_amdgcn_mfma_i32_16x16x16i8.

   Thanks,
3. rwvo commented on Sep 27, 2023

   

   rwvo

   on Sep 27, 2023

   Contributor

   More actions

   In the line of code that you quoted, the left-shift over 8 * (3-i) bits seems suspicious to me. Wouldn't that put A[0][0] in the higher-order (left-most) byte of a for thread (0,0), while it's supposed to go into the lower-order (right-most) byte?

   In any case: one of my colleagues wrote a working example for __builtin_amdgcn_mfma_i32_16x16x16i8. It should appear on the blog soon, but I attach it here for your reference. It uses arrays of int8_t values, and then casts to int32_t in the intrinsics call. Attaching as *.txt because github complains about not supporting *.cpp.

   mfma_i32_16x16x16i8.txt

   👍React with 👍1Reacted by mmoadeli
4. mmoadeli commented on Sep 27, 2023

   

   mmoadeli

   on Sep 27, 2023

   Author

   More actions

   Many many thanks @rwvo. I really appreciated your quick support.

   👍React with 👍1Reacted by René van Oostrum
5. amd

   locked and limited conversation to collaborators on Sep 27, 2023
6. jychang48

   converted this issue into a discussion #8 on Sep 27, 2023