[SYCLomatic] Fix the migration of __hmul2 and __hsub2 when --use-dpcpp-extensions=intel_device_math is specified (#2728)

Signed-off-by: Jiang, Zhiwei <zhiwei.jiang@intel.com>

Author: zhiweij1

clang/lib/DPCT/RulesLang/Math/RewriterHalf2ArithmeticFunctions.cpp

@@ -261,22 +261,30 @@ RewriterMap dpct::createHalf2ArithmeticFunctionsRewriterMap() {
                               makeCallArgCreatorWithCall(2)),
                            LITERAL("{0.f, 0.f}"), LITERAL("{1.f, 1.f}"))))))
       // __hmul2
-      MATH_API_REWRITER_DEVICE(
-          "__hmul2",
-          MATH_API_DEVICE_NODES(
-              EMPTY_FACTORY_ENTRY("__hmul2"),
-              MATH_API_SPECIFIC_ELSE_EMU(
-                  CheckArgType(0, "__half2"),
+      MATH_API_REWRITER_DEVICE_OVERLOAD(
+          CheckArgType(0, "__half2"),
+          MATH_API_REWRITERS_V2(
+              "__hmul2",
+              MATH_API_REWRITER_PAIR(
+                  math::Tag::math_libdevice,
                   HEADER_INSERT_FACTORY(
                       HeaderType::HT_SYCL_Math,
                       CALL_FACTORY_ENTRY("__hmul2",
                                          CALL(MapNames::getClNamespace() +
                                                   "ext::intel::math::hmul2",
                                               ARG(0), ARG(1))))),
-              EMPTY_FACTORY_ENTRY("__hmul2"),
-              BINARY_OP_FACTORY_ENTRY("__hmul2", BinaryOperatorKind::BO_Mul,
-                                      makeCallArgCreatorWithCall(0),
-                                      makeCallArgCreatorWithCall(1))))
+              MATH_API_REWRITER_PAIR(
+                  math::Tag::emulation,
+                  BINARY_OP_FACTORY_ENTRY("__hmul2", BinaryOperatorKind::BO_Mul,
+                                          makeCallArgCreatorWithCall(0),
+                                          makeCallArgCreatorWithCall(1)))),
+          MATH_API_REWRITERS_V2(
+              "__hmul2",
+              MATH_API_REWRITER_PAIR(
+                  math::Tag::emulation,
+                  BINARY_OP_FACTORY_ENTRY("__hmul2", BinaryOperatorKind::BO_Mul,
+                                          makeCallArgCreatorWithCall(0),
+                                          makeCallArgCreatorWithCall(1)))))
       // __hmul2_rn
       MATH_API_REWRITER_DEVICE(
           "__hmul2_rn",
@@ -343,22 +351,30 @@ RewriterMap dpct::createHalf2ArithmeticFunctionsRewriterMap() {
               UNARY_OP_FACTORY_ENTRY("__hneg2", UnaryOperatorKind::UO_Minus,
                                      makeCallArgCreatorWithCall(0))))
       // __hsub2
-      MATH_API_REWRITER_DEVICE(
-          "__hsub2",
-          MATH_API_DEVICE_NODES(
-              EMPTY_FACTORY_ENTRY("__hsub2"),
-              MATH_API_SPECIFIC_ELSE_EMU(
-                  CheckArgType(0, "__half2"),
+      MATH_API_REWRITER_DEVICE_OVERLOAD(
+          CheckArgType(0, "__half2"),
+          MATH_API_REWRITERS_V2(
+              "__hsub2",
+              MATH_API_REWRITER_PAIR(
+                  math::Tag::math_libdevice,
                   HEADER_INSERT_FACTORY(
                       HeaderType::HT_SYCL_Math,
                       CALL_FACTORY_ENTRY("__hsub2",
                                          CALL(MapNames::getClNamespace() +
                                                   "ext::intel::math::hsub2",
                                               ARG(0), ARG(1))))),
-              EMPTY_FACTORY_ENTRY("__hsub2"),
-              BINARY_OP_FACTORY_ENTRY("__hsub2", BinaryOperatorKind::BO_Sub,
-                                      makeCallArgCreatorWithCall(0),
-                                      makeCallArgCreatorWithCall(1))))
+              MATH_API_REWRITER_PAIR(
+                  math::Tag::emulation,
+                  BINARY_OP_FACTORY_ENTRY("__hsub2", BinaryOperatorKind::BO_Sub,
+                                          makeCallArgCreatorWithCall(0),
+                                          makeCallArgCreatorWithCall(1)))),
+          MATH_API_REWRITERS_V2(
+              "__hsub2",
+              MATH_API_REWRITER_PAIR(
+                  math::Tag::emulation,
+                  BINARY_OP_FACTORY_ENTRY("__hsub2", BinaryOperatorKind::BO_Sub,
+                                          makeCallArgCreatorWithCall(0),
+                                          makeCallArgCreatorWithCall(1)))))
       // __hsub2_rn
       MATH_API_REWRITER_DEVICE(
           "__hsub2_rn",

clang/test/dpct/math/bfloat16/bfloat16_ext.cu

@@ -11,6 +11,10 @@ __global__ void kernelFuncBfloat162Arithmetic() {
   __nv_bfloat162 bf162, bf162_1, bf162_2;
   // CHECK: bf162 = bf162_1 / bf162_2;
   bf162 = __h2div(bf162_1, bf162_2);
+  // CHECK: bf162 = bf162_1 * bf162_2;
+  bf162 = __hmul2(bf162_1, bf162_2);
+  // CHECK: bf162 = bf162_1 - bf162_2;
+  bf162 = __hsub2(bf162_1, bf162_2);
 }
 
 __global__ void kernelFuncBfloat16Comparison() {