Updated logic

Author: TejaX-Alaghari

clang/lib/DPCT/RulesAsm/AsmMigration.cpp

@@ -556,16 +556,18 @@ bool SYCLGenBase::emitVectorType(const InlineAsmVectorType *T) {
     return SYCLGenError();
   OS() << ", ";
   switch (T->getKind()) {
-  case InlineAsmVectorType::v2:
   case InlineAsmVectorType::x1:
+    OS() << 1;
+    break;
+  case InlineAsmVectorType::v2:
+  case InlineAsmVectorType::x2:
     OS() << 2;
     break;
   case InlineAsmVectorType::v4:
-  case InlineAsmVectorType::x2:
+  case InlineAsmVectorType::x4:
     OS() << 4;
     break;
   case InlineAsmVectorType::v8:
-  case InlineAsmVectorType::x4:
     OS() << 8;
     break;
   }
@@ -1322,14 +1324,18 @@ class SYCLGen : public SYCLGenBase {
       return SYCLGenError();
     }
     OS() << ", ";
-    const auto *VE = dyn_cast<InlineAsmVectorExpr>(Inst->getOutputOperand());
-    for (unsigned Inst = 0, E = VE->getNumElements(); Inst != E; ++Inst) {
-      if (isa<InlineAsmDiscardExpr>(VE->getElement(Inst)))
-        continue;
-      OS() << "&";
-      if (emitStmt(VE->getElement(Inst)))
-        return SYCLGenError();
-      OS() << ", ";
+    if (const auto *VE =
+            dyn_cast<InlineAsmVectorExpr>(Inst->getOutputOperand())) {
+      for (unsigned Inst = 0; Inst != VE->getNumElements(); ++Inst) {
+        if (isa<InlineAsmDiscardExpr>(VE->getElement(Inst)))
+          continue;
+        OS() << "&";
+        if (emitStmt(VE->getElement(Inst)))
+          return SYCLGenError();
+        OS() << ", ";
+      }
+    } else {
+      return SYCLGenError();
     }
     OS() << DpctGlobalInfo::getItem(GAS);
     if (Inst->hasAttr(InstAttr::trans))

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

@@ -2061,13 +2061,13 @@ class joint_matrix {
 /// \tparam [in] T The type of result variable
 /// \param [in] addr The address of the matrix in shared memory
 /// \param [in] m The local memory to store the matrix
-/// \param [in] item_ct1 The sycl::nd_item object
+/// \param [in] item The sycl::nd_item index space class
 /// \param [in] trans Indicates whether the matrix to be loaded transposed
 /// \param [in] mat The matrix index to be loaded
-template <typename T>
-void ldmatrix(uintptr_t addr, T *m, const sycl::nd_item<3> &item_ct1,
-              bool trans = false, unsigned mat = 0) {
-  int lane = item_ct1.get_local_id(2) % 32;
+template <typename T, typename ItemT>
+void ldmatrix(uintptr_t addr, T *m, const ItemT &item, bool trans = false,
+              unsigned mat = 0) {
+  int lane = item.get_sub_group().get_local_linear_id();
 
   int lane_group8_row = lane / 8;
   int lane_group8_col = lane % 8;
@@ -2080,7 +2080,7 @@ void ldmatrix(uintptr_t addr, T *m, const sycl::nd_item<3> &item_ct1,
 
     // 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);
+        item.get_sub_group(), addr, mat * 8 + src_lane);
 
     // Cast the received address from uintptr_t to the type of 'm'
     auto recv_addr = reinterpret_cast<T *>(recv_addr_uintp);
@@ -2093,9 +2093,9 @@ void ldmatrix(uintptr_t addr, T *m, const sycl::nd_item<3> &item_ct1,
 
     // Broadcast the address from the source lane
     auto recv_addr_uintp_1 = dpct::select_from_sub_group(
-        item_ct1.get_sub_group(), addr, mat * 8 + src_lane);
+        item.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);
+        item.get_sub_group(), addr, mat * 8 + src_lane + 1);
 
     // Cast the received address from uintptr_t to 'half *'
     auto recv_addr_1 = reinterpret_cast<sycl::half *>(recv_addr_uintp_1);
@@ -2118,15 +2118,15 @@ void ldmatrix(uintptr_t addr, T *m, const sycl::nd_item<3> &item_ct1,
 /// \param [in] addr The address of the matrix in shared memory
 /// \param [in] m1 The local memory to store data of 1st matrix
 /// \param [in] m2 The local memory to store data of 2nd matrix
-/// \param [in] item_ct1 The sycl::nd_item object
+/// \param [in] item The sycl::nd_item index space class
 /// \param [in] trans Indicates whether the matrix to be loaded transposed
-template <typename T>
-void ldmatrix(uintptr_t addr, T *m1, T *m2, const sycl::nd_item<3> &item_ct1,
+template <typename T, typename ItemT>
+void ldmatrix(uintptr_t addr, T *m1, T *m2, const ItemT &item,
               bool trans = false) {
   // Load 1st matrix
-  ldmatrix(addr, m1, item_ct1, trans, 0);
+  ldmatrix(addr, m1, item, trans, 0);
   // Load 2nd matrix
-  ldmatrix(addr, m2, item_ct1, trans, 1);
+  ldmatrix(addr, m2, item, trans, 1);
 }
 
 /// Loads 4 8x8 b16 matrix from shared memory to local memory (32-bits per wi)
@@ -2137,19 +2137,19 @@ void ldmatrix(uintptr_t addr, T *m1, T *m2, const sycl::nd_item<3> &item_ct1,
 /// \param [in] m2 The local memory to store data of 2nd matrix
 /// \param [in] m3 The local memory to store data of 3rd matrix
 /// \param [in] m4 The local memory to store data of 4th matrix
-/// \param [in] item_ct1 The sycl::nd_item object
+/// \param [in] item The sycl::nd_item index space class
 /// \param [in] trans Indicates whether the matrix to be loaded transposed
-template <typename T>
-void ldmatrix(uintptr_t addr, T *m1, T *m2, T *m3, T *m4,
-              const sycl::nd_item<3> &item_ct1, bool trans = false) {
+template <typename T, typename ItemT>
+void ldmatrix(uintptr_t addr, T *m1, T *m2, T *m3, T *m4, const ItemT &item,
+              bool trans = false) {
   // Load 1st matrix
-  ldmatrix(addr, m1, item_ct1, trans, 0);
+  ldmatrix(addr, m1, item, trans, 0);
   // Load 2nd matrix
-  ldmatrix(addr, m2, item_ct1, trans, 1);
+  ldmatrix(addr, m2, item, trans, 1);
   // Load 3rd matrix
-  ldmatrix(addr, m3, item_ct1, trans, 2);
+  ldmatrix(addr, m3, item, trans, 2);
   // Load 4th matrix
-  ldmatrix(addr, m4, item_ct1, trans, 3);
+  ldmatrix(addr, m4, item, trans, 3);
 }
 
 } // namespace matrix