Added test case for xn

Author: TejaX-Alaghari

clang/lib/DPCT/RulesAsm/AsmMigration.cpp

@@ -1312,6 +1312,23 @@ class SYCLGen : public SYCLGenBase {
     if (Inst->getNumInputOperands() != 1)
       return SYCLGenError();
 
+    const InlineAsmVectorExpr *VE;
+    if (VE = dyn_cast<InlineAsmVectorExpr>(Inst->getOutputOperand())) {
+      auto numOutputOperands = VE->getNumElements();
+      if (Inst->hasAttr(InstAttr::x1)) {
+        if (numOutputOperands != 1)
+          return SYCLGenError();
+      } else if (Inst->hasAttr(InstAttr::x2)) {
+        if (numOutputOperands != 2)
+          return SYCLGenError();
+      } else if (Inst->hasAttr(InstAttr::x4)) {
+        if (numOutputOperands != 4)
+          return SYCLGenError();
+      }
+    } else {
+      return SYCLGenError();
+    }
+
     llvm::SaveAndRestore<const InlineAsmInstruction *> Store(CurrInst);
     CurrInst = Inst;
     const auto *Src =
@@ -1324,18 +1341,13 @@ class SYCLGen : public SYCLGenBase {
       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();
+    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))

clang/test/dpct/asm/ldmatrix.cu

@@ -97,4 +97,21 @@ int main () {
   return 0;
 }
 
+__device__ void test_xn(uint32_t addr, int *r) {
+  // CHECK: DPCT1053:{{.*}}: Migration of device assembly code is not supported.
+  asm volatile("ldmatrix.sync.aligned.m8n8.x1.shared.b16 {%0, %1}, [%2];\n"
+                : "=r"(r[0]), "=r"(r[1])
+                : "r"(addr));
+
+  // CHECK: DPCT1053:{{.*}}: Migration of device assembly code is not supported.
+  asm volatile("ldmatrix.sync.aligned.m8n8.x2.shared.b16 {%0}, [%0];\n"
+                :
+                : "r"(addr));
+
+  // CHECK: DPCT1053:{{.*}}: Migration of device assembly code is not supported.
+  asm volatile("ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%0, %1, %2}, [%3];\n"
+                : "=r"(r[0]), "=r"(r[1]), "=r"(r[2])
+                : "r"(addr));
+}
+
 // clang-format on