[SYCLomatic][PTX] Support migration of PTX instruction mul.f16x2 (#2771)

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

Author: tomflinda

clang/lib/DPCT/RulesAsm/AsmMigration.cpp

@@ -1363,6 +1363,7 @@ class SYCLGen : public SYCLGenBase {
            Type->getKind() == InlineAsmBuiltinType::s32 ||
            Type->getKind() == InlineAsmBuiltinType::u32 ||
            Type->getKind() == InlineAsmBuiltinType::s64 ||
+           Type->getKind() == InlineAsmBuiltinType::f16x2 ||
            Type->getKind() == InlineAsmBuiltinType::u64;
   }
 
@@ -1419,6 +1420,13 @@ class SYCLGen : public SYCLGenBase {
       OS() << Cast(GetWiderTypeAsString(Type), Op[0]) << " * "
            << Cast(GetWiderTypeAsString(Type), Op[1]);
       // mul.lo
+    } else if (Type->getKind() == InlineAsmBuiltinType::f16x2) {
+      std::string FormatTemp =
+          "((sycl::vec<int, 1>({0})).as<sycl::vec<sycl::half, 2>>() * "
+          "(sycl::vec<int, 1>({1})).as<sycl::vec<sycl::half, "
+          "2>>()).as<sycl::vec<int, 1>>().x()";
+
+      OS() << llvm::formatv(FormatTemp.c_str(), Op[0], Op[1]);
     } else {
       // Need to add a new help function.
       // OS() << Op[0] << " * " << Op[1];

clang/test/dpct/asm/mul.cu

@@ -54,4 +54,14 @@ __global__ void mul() {
   asm("mul.wide.u64 %0, %1, %2;" : "=r"(u64) : "r"(x), "r"(y));
 }
 
+// CHECK: inline uint32_t mul_f16x2(uint32_t a, uint32_t b) {
+// CHECK-NEXT:    uint32_t c;
+// CHECK-NEXT:    c = ((sycl::vec<int, 1>(a)).as<sycl::vec<sycl::half, 2>>() * (sycl::vec<int, 1>(b)).as<sycl::vec<sycl::half, 2>>()).as<sycl::vec<int, 1>>().x();
+// CHECK-NEXT:    return c;
+// CHECK-NEXT: }
+inline __device__ uint32_t mul_f16x2(uint32_t a, uint32_t b) {
+    uint32_t c;
+    asm volatile("mul.f16x2 %0, %1, %2;\n" : "=r"(c) : "r"(a), "r"(b));
+    return c;
+}
 // clang-format on