[SYCLomatic][PTX] Support migration of PTX instruction add.f16x2 and sub.f16x2 (#2727)

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

Author: tomflinda

clang/lib/DPCT/RulesAsm/AsmMigration.cpp

@@ -1043,10 +1043,12 @@ class SYCLGen : public SYCLGenBase {
           BI->getKind() != InlineAsmBuiltinType::u32 &&
           BI->getKind() != InlineAsmBuiltinType::s64 &&
           BI->getKind() != InlineAsmBuiltinType::u64 &&
+          BI->getKind() != InlineAsmBuiltinType::f16x2 &&
           BI->getKind() != InlineAsmBuiltinType::s16x2 &&
           BI->getKind() != InlineAsmBuiltinType::u16x2)
         return false;
       isVec = BI->getKind() == InlineAsmBuiltinType::s16x2 ||
+              BI->getKind() == InlineAsmBuiltinType::f16x2 ||
               BI->getKind() == InlineAsmBuiltinType::u16x2;
     } else {
       return false;
@@ -1088,12 +1090,21 @@ class SYCLGen : public SYCLGenBase {
         OS() << MapNames::getClNamespace()
              << llvm::formatv("sub_sat({0}, {1})", Op[0], Op[1]);
     } else {
-      if (Inst->is(asmtok::op_add))
-        OS() << llvm::formatv("{0} + {1}", Op[0], Op[1]);
-      else
-        OS() << llvm::formatv("{0} - {1}", Op[0], Op[1]);
+      if (const auto *BI = dyn_cast<InlineAsmBuiltinType>(Inst->getType(0))) {
+        std::string operatorStr = Inst->is(asmtok::op_add) ? "+" : "-";
+
+        if (BI->getKind() == InlineAsmBuiltinType::f16x2) {
+          std::string FormatTemp =
+              "(((sycl::vec<int, 1>({0})).as<sycl::vec<sycl::half, 2>>() {1} "
+              "(sycl::vec<int, 1>({2})).as<sycl::vec<sycl::half, "
+              "2>>()).as<sycl::vec<int, 1>>()).x();";
+          OS() << llvm::formatv(FormatTemp.c_str(), Op[0], operatorStr, Op[1]);
+
+        } else {
+          OS() << llvm::formatv("{0} {1} {2}", Op[0], operatorStr, Op[1]);
+        }
+      }
     }
-
     endstmt();
     return SYCLGenSuccess();
   }

clang/test/dpct/asm/add.cu

@@ -59,4 +59,15 @@ __global__ void add() {
   asm("add.u16x2 %0, {1, 1}, %1;" : "=r"(u16x2) : "r"(ua));
 }
 
+// CHECK: inline uint32_t add(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 add(uint32_t a, uint32_t b) {
+  uint32_t c;
+  asm volatile("add.f16x2 %0, %1, %2;\n" : "=r"(c) : "r"(a), "r"(b));
+  return c;
+}
+
 // clang-format off

clang/test/dpct/asm/sub.cu

@@ -59,4 +59,15 @@ __global__ void sub() {
   asm("sub.u16x2 %0, {1, 1}, %1;" : "=r"(u16x2) : "r"(ua));
 }
 
+// CHECK: inline uint32_t sub(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 sub(uint32_t a, uint32_t b) {
+  uint32_t c;
+  asm volatile("sub.f16x2 %0, %1, %2;\n" : "=r"(c) : "r"(a), "r"(b));
+  return c;
+}
+
 // clang-format on