[SYCLomatic][PTX] Support migration of PTX instruction cvt.rn.f16x2.f32 (#2735)

Signed-off-by: chenwei.sun <chenwei.sun@intel.com>

Author: tomflinda

clang/lib/DPCT/RulesAsm/AsmMigration.cpp

@@ -2428,7 +2428,44 @@ class SYCLGen : public SYCLGenBase {
     }
   }
 
+  bool HandleCvtVec(const InlineAsmInstruction *Inst) {
+
+    if (emitStmt(Inst->getOutputOperand()))
+      return SYCLGenError();
+    std::string Op;
+    if (tryEmitStmt(Op, Inst->getInputOperand(0)))
+      return SYCLGenError();
+    OS() << " = ";
+    std::string FormatTemp =
+        "(sycl::ushort2(sycl::vec<float, 1>({0}).convert<sycl::half, "
+        "sycl::rounding_mode::rte>().as<sycl::vec<uint16_t, 1>>().x(),"
+        "sycl::vec<float, 1>({1}).convert<sycl::half, "
+        "sycl::rounding_mode::rte>().as<sycl::vec<uint16_t, 1>>().x()))"
+        ".as<sycl::vec<int, 1>>().x()";
+
+    std::string InputOp[2];
+    for (unsigned I = 0; I < Inst->getNumInputOperands(); ++I) {
+      if (tryEmitStmt(InputOp[I], Inst->getInputOperand(I)))
+        return SYCLGenError();
+      if (Inst->hasAttr(InstAttr::sat))
+        InputOp[I] = Cast(Inst->getType(0), Inst->getInputOperand(I)->getType(),
+                          InputOp[I]);
+    }
+
+    OS() << llvm::formatv(FormatTemp.c_str(), InputOp[1], InputOp[0]);
+
+    endstmt();
+    return SYCLGenSuccess();
+  }
+
   bool handle_cvt(const InlineAsmInstruction *Inst) override {
+
+    if (Inst->getNumInputOperands() == 2 && Inst->getNumTypes() == 2 &&
+        isa<InlineAsmBuiltinType>(Inst->getType(0)) &&
+        isa<InlineAsmBuiltinType>(Inst->getType(1))) {
+      return HandleCvtVec(Inst);
+    }
+
     if (Inst->getNumInputOperands() != 1 || Inst->getNumTypes() != 2 ||
         !isa<InlineAsmBuiltinType>(Inst->getType(0)) ||
         !isa<InlineAsmBuiltinType>(Inst->getType(1)))

clang/test/dpct/asm/cvt.cu

@@ -400,4 +400,21 @@ __global__ void cvt() {
 }
 #endif
 
+// CHECK: inline uint32_t float2_to_half2(sycl::float2 f) {
+// CHECK-NEXT:   union {
+// CHECK-NEXT:     uint32_t u32;
+// CHECK-NEXT:     uint16_t u16[2];
+// CHECK-NEXT:   } tmp;
+// CHECK-NEXT:   tmp.u32 = (sycl::ushort2(sycl::vec<float, 1>(f.x()).convert<sycl::half, sycl::rounding_mode::rte>().as<sycl::vec<uint16_t, 1>>().x(),sycl::vec<float, 1>(f.y()).convert<sycl::half, sycl::rounding_mode::rte>().as<sycl::vec<uint16_t, 1>>().x())).as<sycl::vec<int, 1>>().x();
+// CHECK-NEXT:   return tmp.u32;
+// CHECK-NEXT: }
+inline __device__ uint32_t float2_to_half2(float2 f) {
+  union {
+    uint32_t u32;
+    uint16_t u16[2];
+  } tmp;
+  asm volatile("cvt.rn.f16x2.f32 %0, %1, %2;\n" : "=r"(tmp.u32) : "f"(f.y), "f"(f.x));
+  return tmp.u32;
+}
+
 // clang-format on