[SYCLomatic][PTX] Support migration of PTX instruction neg.f16x2 (#2775)

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

Author: tomflinda

clang/lib/DPCT/RulesAsm/AsmMigration.cpp

@@ -1575,6 +1575,7 @@ class SYCLGen : public SYCLGenBase {
 
     if (Type->getKind() != InlineAsmBuiltinType::s16 &&
         Type->getKind() != InlineAsmBuiltinType::s32 &&
+        Type->getKind() != InlineAsmBuiltinType::f16x2 &&
         Type->getKind() != InlineAsmBuiltinType::s64) {
       return SYCLGenError();
     }
@@ -1594,7 +1595,14 @@ class SYCLGen : public SYCLGenBase {
 
     if (Inst->is(asmtok::op_abs))
       OS() << MapNames::getClNamespace() << "abs(" << Op << ")";
-    else
+    else if (Inst->is(asmtok::op_neg) &&
+             Type->getKind() == InlineAsmBuiltinType::f16x2) {
+      std::string FormatTemp =
+          "(-sycl::vec<int, 1>({0}).as<sycl::vec<sycl::half, "
+          "2>>()).as<sycl::vec<int, 1>>().x()";
+
+      OS() << llvm::formatv(FormatTemp.c_str(), Op);
+    } else
       OS() << "-" << Op;
 
     endstmt();

clang/test/dpct/asm/neg.cu

@@ -6,6 +6,7 @@
 
 // clang-format off
 #include <cuda_runtime.h>
+#include <cuda_fp16.h>
 #include <cstdint>
 
 __global__ void neg() {
@@ -24,4 +25,17 @@ __global__ void neg() {
   asm("neg.s64 %0, %1;" : "=r"(i64) : "r"(x));
 }
 
+// CHECK: inline void negate_half2(sycl::half2 *addr) {
+// CHECK-NEXT:    unsigned reg[2];
+// CHECK-NEXT:    reg[0] = *reinterpret_cast<unsigned int*>(addr);
+// CHECK-NEXT:    reg[0] = (-sycl::vec<int, 1>(reg[0]).as<sycl::vec<sycl::half, 2>>()).as<sycl::vec<int, 1>>().x();
+// CHECK-NEXT:    *reinterpret_cast<unsigned int*>(addr) = reg[0];
+// CHECK-NEXT:}
+__device__ inline void negate_half2(__half2 *addr) {
+    unsigned reg[2];
+    reg[0] = *reinterpret_cast<unsigned int*>(addr);
+    asm volatile("neg.f16x2 %0, %0;" : "+r"(reg[0]));
+    *reinterpret_cast<unsigned int*>(addr) = reg[0];
+}
+
 // clang-format on