Updated trans logic

Author: TejaX-Alaghari

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

@@ -2057,9 +2057,11 @@ void ldmatrix(uintptr_t addr, T *m,
   int group = lane / 8;
   int sub = lane % 8;
   int src_base = group * 2;
-  int src_lane = (sub / 4) ? (src_base + 1) : src_base;
 
   if (!trans) {
+    // calculate the source lane
+    int src_lane = (sub / 4) ? (src_base + 1) : src_base;
+
     // Broadcast the address from the source lane
     auto recv_addr_uintp = dpct::select_from_sub_group(
         item_ct1.get_sub_group(), addr, mat * 8 + src_lane);
@@ -2073,16 +2075,21 @@ void ldmatrix(uintptr_t addr, T *m,
     uint16_t bits1 = sycl::bit_cast<unsigned short, sycl::half>(val1);
     *m = ((uint32_t)bits1 << 16) | bits0;
   } else {
+    // calculate the source lane
+    int src_lane = (lane % 4) * 2;
+
     // Broadcast the address from the source lane:
-    auto recv_addr_uintp = dpct::select_from_sub_group(
-        item_ct1.get_sub_group(), addr, mat * 8);
-    auto recv_addr = reinterpret_cast<sycl::half *>(recv_addr_uintp);
-    recv_addr += src_lane;
+    auto recv_addr_uintp_1 = dpct::select_from_sub_group(
+        item_ct1.get_sub_group(), addr, mat * 8 + src_lane);
+    auto recv_addr_uintp_2 = dpct::select_from_sub_group(
+        item_ct1.get_sub_group(), addr, mat * 8 + src_lane + 1);
+    auto recv_addr_1 = reinterpret_cast<sycl::half *>(recv_addr_uintp_1);
+    auto recv_addr_2 = reinterpret_cast<sycl::half *>(recv_addr_uintp_2);
 
     // Transposed load
-    int index = (lane % 4) * 8 * 2;
-    sycl::half val0 = recv_addr[index];
-    sycl::half val1 = recv_addr[index + 8];
+    int index = (lane / 4);
+    sycl::half val0 = recv_addr_1[index];
+    sycl::half val1 = recv_addr_2[index];
     uint16_t bits0 = sycl::bit_cast<unsigned short, sycl::half>(val0);
     uint16_t bits1 = sycl::bit_cast<unsigned short, sycl::half>(val1);
     *m = ((uint32_t)bits1 << 16) | bits0;