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/lib/DPCT/RulesAsm/Parser/AsmNodes.h

@@ -340,8 +340,6 @@ class InlineAsmInstruction : public InlineAsmStmt {
   /// therest are input operands.
   SmallVector<InlineAsmExpr *, 4> InputOps;
 
-  SmallVector<InlineAsmExpr *, 4> OutputOps;
-
 public:
   InlineAsmInstruction(InlineAsmIdentifierInfo *Op,
                        SmallVector<AsmStateSpace, 4> AsmStateSpaces,

clang/test/dpct/asm/ldmatrix.cu

@@ -2,7 +2,7 @@
 // UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2
 // RUN: dpct --format-range=none -out-root %T/ldmatrix %s --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only
 // RUN: FileCheck %s --match-full-lines --input-file %T/ldmatrix/ldmatrix.dp.cpp
-// RUN: %if build_lit %{icpx -c -fsycl %T/ldmatrix/ldmatrix.dp.cpp -o %T/ldmatrix/ldmatrix.dp.o %}
+// RUN: %if build_lit %{icpx -c -DNO_BUILD_TEST -fsycl %T/ldmatrix/ldmatrix.dp.cpp -o %T/ldmatrix/ldmatrix.dp.o %}
 
 // clang-format off
 #include <cuda_runtime.h>
@@ -97,4 +97,23 @@ int main () {
   return 0;
 }
 
+#ifndef NO_BUILD_TEST
+__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));
+}
+#endif // NO_BUILD_TEST
+
 // clang-format on