Removed item_ct1 in favor of free functions

Author: TejaX-Alaghari

clang/lib/DPCT/RulesAsm/AsmMigration.cpp

@@ -1344,16 +1344,13 @@ class SYCLGen : public SYCLGenBase {
     if (emitStmt(Src)) {
       return SYCLGenError();
     }
-    OS() << ", ";
     for (unsigned Inst = 0; Inst != VE->getNumElements(); ++Inst) {
       if (isa<InlineAsmDiscardExpr>(VE->getElement(Inst)))
         continue;
-      OS() << "&";
+      OS() << ", &";
       if (emitStmt(VE->getElement(Inst)))
         return SYCLGenError();
-      OS() << ", ";
     }
-    OS() << DpctGlobalInfo::getItem(GAS);
     if (Inst->hasAttr(InstAttr::trans))
       OS() << ", true";
     OS() << ");";

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

@@ -2061,13 +2061,12 @@ class joint_matrix {
 /// \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
-/// \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, 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();
+template <typename T>
+void ldmatrix(uintptr_t addr, T *m, bool trans = false, unsigned mat = 0) {
+  auto sg = sycl::ext::oneapi::this_work_item::get_sub_group();
+  int lane = sg.get_local_linear_id();
 
   int lane_group8_row = lane / 8;
   int lane_group8_col = lane % 8;
@@ -2079,8 +2078,8 @@ void ldmatrix(uintptr_t addr, T *m, const ItemT &item, bool trans = false,
       src_lane += 1;
 
     // Broadcast the address from the source lane
-    auto recv_addr_uintp = dpct::select_from_sub_group(
-        item.get_sub_group(), addr, mat * 8 + src_lane);
+    auto recv_addr_uintp =
+        dpct::select_from_sub_group(sg, 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);
@@ -2092,10 +2091,10 @@ void ldmatrix(uintptr_t addr, T *m, const ItemT &item, bool trans = false,
     int src_lane = (lane % 4) * 2;
 
     // Broadcast the address from the source lane
-    auto recv_addr_uintp_1 = dpct::select_from_sub_group(
-        item.get_sub_group(), addr, mat * 8 + src_lane);
-    auto recv_addr_uintp_2 = dpct::select_from_sub_group(
-        item.get_sub_group(), addr, mat * 8 + src_lane + 1);
+    auto recv_addr_uintp_1 =
+        dpct::select_from_sub_group(sg, addr, mat * 8 + src_lane);
+    auto recv_addr_uintp_2 =
+        dpct::select_from_sub_group(sg, 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 +2117,13 @@ void ldmatrix(uintptr_t addr, T *m, const ItemT &item, bool trans = false,
 /// \param [in] addr The address of the matrix in local memory
 /// \param [in] m1 The private memory to store data of 1st matrix
 /// \param [in] m2 The private memory to store data of 2nd matrix
-/// \param [in] item The sycl::nd_item index space class
 /// \param [in] trans Indicates whether the matrix to be loaded transposed
-template <typename T, typename ItemT>
-void ldmatrix(uintptr_t addr, T *m1, T *m2, const ItemT &item,
-              bool trans = false) {
+template <typename T>
+void ldmatrix(uintptr_t addr, T *m1, T *m2, bool trans = false) {
   // Load 1st matrix
-  ldmatrix(addr, m1, item, trans, 0);
+  ldmatrix(addr, m1, trans, 0);
   // Load 2nd matrix
-  ldmatrix(addr, m2, item, trans, 1);
+  ldmatrix(addr, m2, trans, 1);
 }
 
 /// Loads 4 8x8 b16 matrix from local memory to private memory (32-bits per wi)
@@ -2137,19 +2134,17 @@ void ldmatrix(uintptr_t addr, T *m1, T *m2, const ItemT &item,
 /// \param [in] m2 The private memory to store data of 2nd matrix
 /// \param [in] m3 The private memory to store data of 3rd matrix
 /// \param [in] m4 The private memory to store data of 4th matrix
-/// \param [in] item The sycl::nd_item index space class
 /// \param [in] trans Indicates whether the matrix to be loaded transposed
-template <typename T, typename ItemT>
-void ldmatrix(uintptr_t addr, T *m1, T *m2, T *m3, T *m4, const ItemT &item,
-              bool trans = false) {
+template <typename T>
+void ldmatrix(uintptr_t addr, T *m1, T *m2, T *m3, T *m4, bool trans = false) {
   // Load 1st matrix
-  ldmatrix(addr, m1, item, trans, 0);
+  ldmatrix(addr, m1, trans, 0);
   // Load 2nd matrix
-  ldmatrix(addr, m2, item, trans, 1);
+  ldmatrix(addr, m2, trans, 1);
   // Load 3rd matrix
-  ldmatrix(addr, m3, item, trans, 2);
+  ldmatrix(addr, m3, trans, 2);
   // Load 4th matrix
-  ldmatrix(addr, m4, item, trans, 3);
+  ldmatrix(addr, m4, trans, 3);
 }
 
 } // namespace matrix

clang/test/dpct/asm/ldmatrix.cu

@@ -22,7 +22,7 @@ __device__ void load_matrix_x1(void *sh_r_addr, int *r) {
   // CHECK: auto addr = sh_r_addr;
   uint32_t addr = static_cast<uint32_t>(__cvta_generic_to_shared(sh_r_addr));
 
-  // CHECK: dpct::experimental::matrix::ldmatrix((uintptr_t)addr, &r[0], item_ct1);
+  // CHECK: dpct::experimental::matrix::ldmatrix((uintptr_t)addr, &r[0]);
   asm volatile("ldmatrix.sync.aligned.m8n8.x1.shared.b16 {%0}, [%1];\n"
                 : "=r"(r[0])
                 : "r"(addr));
@@ -32,7 +32,7 @@ __device__ void load_matrix_x2(void *sh_r_addr, int *r) {
   // CHECK: auto addr = sh_r_addr;
   uint32_t addr = static_cast<uint32_t>(__cvta_generic_to_shared(sh_r_addr));
 
-  // CHECK: dpct::experimental::matrix::ldmatrix((uintptr_t)addr, &r[0], &r[1], item_ct1);
+  // CHECK: dpct::experimental::matrix::ldmatrix((uintptr_t)addr, &r[0], &r[1]);
   asm volatile("ldmatrix.sync.aligned.m8n8.x2.shared.b16 {%0, %1}, [%2];\n"
                 : "=r"(r[0]), "=r"(r[1])
                 : "r"(addr));
@@ -42,7 +42,7 @@ __device__ void load_matrix_x4(void *sh_r_addr, int *r) {
   // CHECK: auto addr = sh_r_addr;
   uint32_t addr = static_cast<uint32_t>(__cvta_generic_to_shared(sh_r_addr));
 
-  // CHECK: dpct::experimental::matrix::ldmatrix((uintptr_t)addr, &r[0], &r[1], &r[2], &r[3], item_ct1);
+  // CHECK: dpct::experimental::matrix::ldmatrix((uintptr_t)addr, &r[0], &r[1], &r[2], &r[3]);
   asm volatile("ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%0, %1, %2, %3}, [%4];\n"
                 : "=r"(r[0]), "=r"(r[1]), "=r"(r[2]), "=r"(r[3])
                 : "r"(addr));
@@ -52,7 +52,7 @@ __device__ void load_matrix_x1_trans(void *sh_r_addr, int *r) {
   // CHECK: auto addr = sh_r_addr;
   uint32_t addr = static_cast<uint32_t>(__cvta_generic_to_shared(sh_r_addr));
 
-  // CHECK: dpct::experimental::matrix::ldmatrix((uintptr_t)addr, &r[0], item_ct1, true);
+  // CHECK: dpct::experimental::matrix::ldmatrix((uintptr_t)addr, &r[0], true);
   asm volatile("ldmatrix.sync.aligned.m8n8.x1.trans.shared.b16 {%0}, [%1];\n"
                 : "=r"(r[0])
                 : "r"(addr));
@@ -62,7 +62,7 @@ __device__ void load_matrix_x2_trans(void *sh_r_addr, int *r) {
   // CHECK: auto addr = sh_r_addr;
   uint32_t addr = static_cast<uint32_t>(__cvta_generic_to_shared(sh_r_addr));
 
-  // CHECK: dpct::experimental::matrix::ldmatrix((uintptr_t)addr, &r[0], &r[1], item_ct1, true);
+  // CHECK: dpct::experimental::matrix::ldmatrix((uintptr_t)addr, &r[0], &r[1], true);
   asm volatile("ldmatrix.sync.aligned.m8n8.x2.trans.shared.b16 {%0, %1}, [%2];\n"
                 : "=r"(r[0]), "=r"(r[1])
                 : "r"(addr));
@@ -72,7 +72,7 @@ __device__ void load_matrix_x4_trans(void *sh_r_addr, int *r) {
   // CHECK: auto addr = sh_r_addr;
   uint32_t addr = static_cast<uint32_t>(__cvta_generic_to_shared(sh_r_addr));
 
-  // CHECK: dpct::experimental::matrix::ldmatrix((uintptr_t)addr, &r[0], &r[1], &r[2], &r[3], item_ct1, true);
+  // CHECK: dpct::experimental::matrix::ldmatrix((uintptr_t)addr, &r[0], &r[1], &r[2], &r[3], true);
   asm volatile("ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%0, %1, %2, %3}, [%4];\n"
                 : "=r"(r[0]), "=r"(r[1]), "=r"(r[2]), "=r"(r[3])
                 : "r"(addr));