Removed item_ct1 in favor of free functions

Author: TejaX-Alaghari

clang/lib/DPCT/RulesAsm/AsmMigration.cpp

@@ -1344,15 +1344,13 @@ class SYCLGen : public SYCLGenBase {
     if (emitStmt(Dst)) {
       return SYCLGenError();
     }
-    OS() << ", ";
     for (unsigned Inst = 0; Inst != VE->getNumElements(); ++Inst) {
       if (isa<InlineAsmDiscardExpr>(VE->getElement(Inst)))
         continue;
+      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 matrix elements
 /// \param [in] addr The address of the matrix in local memory
 /// \param [in] m The private memory containing data of matrix
-/// \param [in] item The sycl::nd_item index space class
 /// \param [in] trans Indicates whether the matrix to be stored transposed
 /// \param [in] mat The matrix index to be stored
-template <typename T, typename ItemT>
-void stmatrix(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 stmatrix(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 stmatrix(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 stmatrix(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);
@@ -2117,15 +2116,13 @@ void stmatrix(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 containing data of 1st matrix
 /// \param [in] m2 The private memory containing data of 2nd matrix
-/// \param [in] item The sycl::nd_item index space class
 /// \param [in] trans Indicates whether the matrix to be stored transposed
-template <typename T, typename ItemT>
-void stmatrix(uintptr_t addr, T m1, T m2, const ItemT &item,
-              bool trans = false) {
+template <typename T>
+void stmatrix(uintptr_t addr, T m1, T m2, bool trans = false) {
   // Store 1st matrix
-  stmatrix(addr, m1, item, trans, 0);
+  stmatrix(addr, m1, trans, 0);
   // Store 2nd matrix
-  stmatrix(addr, m2, item, trans, 1);
+  stmatrix(addr, m2, trans, 1);
 }
 
 /// Stores 4 8x8 b16 matrix from private memory to local memory (32-bits per wi)
@@ -2136,19 +2133,17 @@ void stmatrix(uintptr_t addr, T m1, T m2, const ItemT &item,
 /// \param [in] m2 The private memory containing data of 2nd matrix
 /// \param [in] m3 The private memory containing data of 3rd matrix
 /// \param [in] m4 The private memory containing data of 4th matrix
-/// \param [in] item The sycl::nd_item index space class
 /// \param [in] trans Indicates whether the matrix to be stored transposed
-template <typename T, typename ItemT>
-void stmatrix(uintptr_t addr, T m1, T m2, T m3, T m4, const ItemT &item,
-              bool trans = false) {
+template <typename T>
+void stmatrix(uintptr_t addr, T m1, T m2, T m3, T m4, bool trans = false) {
   // Store 1st matrix
-  stmatrix(addr, m1, item, trans, 0);
+  stmatrix(addr, m1, trans, 0);
   // Store 2nd matrix
-  stmatrix(addr, m2, item, trans, 1);
+  stmatrix(addr, m2, trans, 1);
   // Store 3rd matrix
-  stmatrix(addr, m3, item, trans, 2);
+  stmatrix(addr, m3, trans, 2);
   // Store 4th matrix
-  stmatrix(addr, m4, item, trans, 3);
+  stmatrix(addr, m4, trans, 3);
 }
 
 } // namespace matrix

clang/test/dpct/asm/stmatrix.cu

@@ -22,7 +22,7 @@ __device__ void store_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::stmatrix((uintptr_t)addr, r[0], item_ct1);
+  // CHECK: dpct::experimental::matrix::stmatrix((uintptr_t)addr, r[0]);
   asm volatile("stmatrix.sync.aligned.m8n8.x1.shared.b16 [%0], {%1};\n"
                 : 
                 : "r"(addr), "r"(r[0]));
@@ -32,7 +32,7 @@ __device__ void store_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::stmatrix((uintptr_t)addr, r[0], r[1], item_ct1);
+  // CHECK: dpct::experimental::matrix::stmatrix((uintptr_t)addr, r[0], r[1]);
   asm volatile("stmatrix.sync.aligned.m8n8.x2.shared.b16 [%0], {%1, %2};\n"
                 : 
                 : "r"(addr), "r"(r[0]), "r"(r[1]));
@@ -42,7 +42,7 @@ __device__ void store_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::stmatrix((uintptr_t)addr, r[0], r[1], r[2], r[3], item_ct1);
+  // CHECK: dpct::experimental::matrix::stmatrix((uintptr_t)addr, r[0], r[1], r[2], r[3]);
   asm volatile("stmatrix.sync.aligned.m8n8.x4.shared.b16 [%0], {%1, %2, %3, %4};\n"
                 : 
                 : "r"(addr), "r"(r[0]), "r"(r[1]), "r"(r[2]), "r"(r[3]));
@@ -52,7 +52,7 @@ __device__ void store_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::stmatrix((uintptr_t)addr, r[0], item_ct1, true);
+  // CHECK: dpct::experimental::matrix::stmatrix((uintptr_t)addr, r[0], true);
   asm volatile("stmatrix.sync.aligned.m8n8.x1.trans.shared.b16 [%0], {%1};\n"
                 : 
                 : "r"(addr), "r"(r[0]));
@@ -62,7 +62,7 @@ __device__ void store_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::stmatrix((uintptr_t)addr, r[0], r[1], item_ct1, true);
+  // CHECK: dpct::experimental::matrix::stmatrix((uintptr_t)addr, r[0], r[1], true);
   asm volatile("stmatrix.sync.aligned.m8n8.x2.trans.shared.b16 [%0], {%1, %2};\n"
                 : 
                 : "r"(addr), "r"(r[0]), "r"(r[1]));
@@ -72,7 +72,7 @@ __device__ void store_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::stmatrix((uintptr_t)addr, r[0], r[1], r[2], r[3], item_ct1, true);
+  // CHECK: dpct::experimental::matrix::stmatrix((uintptr_t)addr, r[0], r[1], r[2], r[3], true);
   asm volatile("stmatrix.sync.aligned.m8n8.x4.trans.shared.b16 [%0], {%1, %2, %3, %4};\n"
                 : 
                 : "r"(addr), "r"(r[0]), "r"(r[1]), "r"(r[2]), "r"(r[3]));