Added more info about functionality

Author: TejaX-Alaghari

clang/runtime/dpct-rt/include/dpct/math.hpp

@@ -2056,8 +2056,24 @@ class joint_matrix {
   const size_t num_elements;
 };
 
-/// Loads 1 8x8 b16 matrix from local memory to private memory (32-bits per wi)
-/// Requires the sub-group size of kernel calling this function to be 32
+/// Loads 1 8x8 b16 matrix from private memory to local memory per sub-group.
+/// Requires the sub-group size of kernel calling this function to be 32.
+/// Each of the first 8 work items contain the starting address of their
+/// respective matrix row.
+/// Each of the 32 work items load 32-bits (2 packed 16-bit data) for a total
+/// of 128 bytes.
+/// Row Major: Each row of the matrix is loaded by a group of 4 work items
+/// r0: t0 t1 t2 t3
+/// r1: t4 t5 t6 t7
+/// ...
+/// r7: t24 t25 t26 t27
+/// r7: t28 t29 t30 t31
+/// Col Major: Each col of the matrix is loadedd by a group of 4 work items
+/// r0: t0 t4 t8 ... t28
+/// r1: t0 t4 t8 ... t28
+/// ...
+/// r6: t3 t7 t11 ... t31
+/// r7: t3 t7 t11 ... t31
 /// \tparam [in] T The type of result variable
 /// \param [in] addr The address of the matrix in local memory
 /// \param [in] m The private memory to store the matrix
@@ -2111,8 +2127,12 @@ void ldmatrix(uintptr_t addr, T *m, bool trans = false, unsigned mat = 0) {
   }
 }
 
-/// Loads 2 8x8 b16 matrix from local memory to private memory (32-bits per wi)
-/// Requires the sub-group size of kernel calling this function to be 32
+/// Loads 2 8x8 b16 matrix from private memory to local memory per sub-group.
+/// Requires the sub-group size of kernel calling this function to be 32.
+/// Each of the first 16 work items contain the starting address of their
+/// respective matrix row.
+/// Each of the 32 work items load 64-bits (32-bit per matrix) for a total
+/// of 256 bytes.
 /// \tparam [in] T The type of result variable
 /// \param [in] addr The address of the matrix in local memory
 /// \param [in] m1 The private memory to store data of 1st matrix
@@ -2126,8 +2146,12 @@ void ldmatrix(uintptr_t addr, T *m1, T *m2, bool trans = false) {
   ldmatrix(addr, m2, trans, 1);
 }
 
-/// Loads 4 8x8 b16 matrix from local memory to private memory (32-bits per wi)
-/// Requires the sub-group size of kernel calling this function to be 32
+/// Loads 4 8x8 b16 matrix from private memory to local memory per sub-group.
+/// Requires the sub-group size of kernel calling this function to be 32.
+/// Each of the 32 work items contain the starting address of their
+/// respective matrix row.
+/// Each of the 32 work items load 128-bits (32-bit per matrix) for a total
+/// of 512 bytes.
 /// \tparam [in] T The type of result variable
 /// \param [in] addr The address of the matrix in local memory
 /// \param [in] m1 The private memory to store data of 1st matrix