diff --git a/clang/lib/DPCT/APINames.inc b/clang/lib/DPCT/APINames.inc
index 6567500ffec7..3e25e7b1a6dd 100644
--- a/clang/lib/DPCT/APINames.inc
+++ b/clang/lib/DPCT/APINames.inc
@@ -624,7 +624,7 @@ ENTRY(h2floor, h2floor, true, NO_FLAG, P4, "Successful")
 ENTRY(h2log, h2log, true, NO_FLAG, P4, "Successful")
 ENTRY(h2log10, h2log10, true, NO_FLAG, P4, "Successful")
 ENTRY(h2log2, h2log2, true, NO_FLAG, P4, "Successful")
-ENTRY(h2rcp, h2rcp, false, NO_FLAG, P4, "no direct mapping yet")
+ENTRY(h2rcp, h2rcp, true, NO_FLAG, P4, "Successful")
 ENTRY(h2rint, h2rint, true, NO_FLAG, P4, "Successful")
 ENTRY(h2rsqrt, h2rsqrt, true, NO_FLAG, P4, "Successful")
 ENTRY(h2sin, h2sin, true, NO_FLAG, P4, "Successful")
diff --git a/clang/lib/DPCT/APINamesMath.inc b/clang/lib/DPCT/APINamesMath.inc
index 66f9daa490d8..5fc732b4c9cf 100644
--- a/clang/lib/DPCT/APINamesMath.inc
+++ b/clang/lib/DPCT/APINamesMath.inc
@@ -16,37 +16,38 @@ ENTRY_RENAMED("__hisinf", MapNames::getClNamespace(false, true) + "isinf")
 ENTRY_RENAMED("__hisnan", MapNames::getClNamespace(false, true) + "isnan")
 
 // Half Math Functions
-ENTRY_RENAMED("hceil", MapNames::getClNamespace(false, true) + "ceil")
-ENTRY_RENAMED("hcos", MapNames::getClNamespace(false, true) + "cos")
-ENTRY_RENAMED("hexp", MapNames::getClNamespace(false, true) + "exp")
-ENTRY_RENAMED("hexp10", MapNames::getClNamespace(false, true) + "exp10")
-ENTRY_RENAMED("hexp2", MapNames::getClNamespace(false, true) + "exp2")
-ENTRY_RENAMED("hfloor", MapNames::getClNamespace(false, true) + "floor")
-ENTRY_RENAMED("hlog", MapNames::getClNamespace(false, true) + "log")
-ENTRY_RENAMED("hlog10", MapNames::getClNamespace(false, true) + "log10")
-ENTRY_RENAMED("hlog2", MapNames::getClNamespace(false, true) + "log2")
-ENTRY_RENAMED("hrcp", MapNames::getClNamespace(false, true) + "half_precision::recip<float>")
-ENTRY_RENAMED("hrint", MapNames::getClNamespace(false, true) + "rint")
-ENTRY_RENAMED("hrsqrt", MapNames::getClNamespace(false, true) + "rsqrt")
-ENTRY_RENAMED("hsin", MapNames::getClNamespace(false, true) + "sin")
-ENTRY_RENAMED("hsqrt", MapNames::getClNamespace(false, true) + "sqrt")
-ENTRY_RENAMED("htrunc", MapNames::getClNamespace(false, true) + "trunc")
+ENTRY_REWRITE("hceil")
+ENTRY_REWRITE("hcos")
+ENTRY_REWRITE("hexp")
+ENTRY_REWRITE("hexp10")
+ENTRY_REWRITE("hexp2")
+ENTRY_REWRITE("hfloor")
+ENTRY_REWRITE("hlog")
+ENTRY_REWRITE("hlog10")
+ENTRY_REWRITE("hlog2")
+ENTRY_REWRITE("hrcp")
+ENTRY_REWRITE("hrint")
+ENTRY_REWRITE("hrsqrt")
+ENTRY_REWRITE("hsin")
+ENTRY_REWRITE("hsqrt")
+ENTRY_REWRITE("htrunc")
 
 // Half2 Math Functions
-ENTRY_RENAMED("h2ceil", MapNames::getClNamespace(false, true) + "ceil")
-ENTRY_RENAMED("h2cos", MapNames::getClNamespace(false, true) + "cos")
-ENTRY_RENAMED("h2exp", MapNames::getClNamespace(false, true) + "exp")
-ENTRY_RENAMED("h2exp10", MapNames::getClNamespace(false, true) + "exp10")
-ENTRY_RENAMED("h2exp2", MapNames::getClNamespace(false, true) + "exp2")
-ENTRY_RENAMED("h2floor", MapNames::getClNamespace(false, true) + "floor")
-ENTRY_RENAMED("h2log", MapNames::getClNamespace(false, true) + "log")
-ENTRY_RENAMED("h2log10", MapNames::getClNamespace(false, true) + "log10")
-ENTRY_RENAMED("h2log2", MapNames::getClNamespace(false, true) + "log2")
-ENTRY_RENAMED("h2rint", MapNames::getClNamespace(false, true) + "rint")
-ENTRY_RENAMED("h2rsqrt", MapNames::getClNamespace(false, true) + "rsqrt")
-ENTRY_RENAMED("h2sin", MapNames::getClNamespace(false, true) + "sin")
-ENTRY_RENAMED("h2sqrt", MapNames::getClNamespace(false, true) + "sqrt")
-ENTRY_RENAMED("h2trunc", MapNames::getClNamespace(false, true) + "trunc")
+ENTRY_REWRITE("h2ceil")
+ENTRY_REWRITE("h2cos")
+ENTRY_REWRITE("h2exp")
+ENTRY_REWRITE("h2exp10")
+ENTRY_REWRITE("h2exp2")
+ENTRY_REWRITE("h2floor")
+ENTRY_REWRITE("h2log")
+ENTRY_REWRITE("h2log10")
+ENTRY_REWRITE("h2log2")
+ENTRY_REWRITE("h2rcp")
+ENTRY_REWRITE("h2rint")
+ENTRY_REWRITE("h2rsqrt")
+ENTRY_REWRITE("h2sin")
+ENTRY_REWRITE("h2sqrt")
+ENTRY_REWRITE("h2trunc")
 
 // Single Precision Mathematical Functions
 ENTRY_RENAMED_SINGLE("acosf", MapNames::getClNamespace(false, true) + "acos")
@@ -663,9 +664,6 @@ ENTRY_REWRITE("__hne2_mask")
 ENTRY_REWRITE("__hneu2")
 ENTRY_REWRITE("__hneu2_mask")
 
-// Half2 Math Functions
-ENTRY_UNSUPPORTED("h2rcp")
-
 // Single Precision Mathematical Functions
 ENTRY_REWRITE("cyl_bessel_i0f")
 ENTRY_REWRITE("cyl_bessel_i1f")
diff --git a/clang/lib/DPCT/APINamesMathRewrite.inc b/clang/lib/DPCT/APINamesMathRewrite.inc
index 3cb3c2b6b706..b0fe3eb0623f 100644
--- a/clang/lib/DPCT/APINamesMathRewrite.inc
+++ b/clang/lib/DPCT/APINamesMathRewrite.inc
@@ -1701,6 +1701,828 @@ MATH_API_REWRITER_DEVICE(
             Diagnostics::MATH_EMULATION_EXPRESSION, std::string("__stwt"),
             std::string("'='"))))
 
+// Half Math Functions
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half"),
+    MATH_API_REWRITER_DEVICE(
+        "hceil",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("hceil"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("hceil", CALL(MapNames::getClNamespace() +
+                                                     "ext::intel::math::ceil",
+                                                 ARG(0)))),
+            EMPTY_FACTORY_ENTRY("hceil"),
+            CALL_FACTORY_ENTRY(
+                "hceil",
+                CALL(MapNames::getClNamespace(false, true) + "ceil", ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "hceil",
+        CALL_FACTORY_ENTRY("hceil", CALL(MapNames::getClNamespace(false, true) +
+                                             "ext::oneapi::experimental::ceil",
+                                         ARG(0))),
+        CALL_FACTORY_ENTRY("hceil",
+                           CALL(MapNames::getClNamespace(false, true) + "ceil",
+                                CALL("float", ARG(0))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half"),
+    MATH_API_REWRITER_DEVICE(
+        "hcos",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("hcos"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("hcos", CALL(MapNames::getClNamespace() +
+                                                    "ext::intel::math::cos",
+                                                ARG(0)))),
+            EMPTY_FACTORY_ENTRY("hcos"),
+            CALL_FACTORY_ENTRY(
+                "hcos",
+                CALL(MapNames::getClNamespace(false, true) + "cos", ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "hcos",
+        CALL_FACTORY_ENTRY("hcos", CALL(MapNames::getClNamespace(false, true) +
+                                            "ext::oneapi::experimental::cos",
+                                        ARG(0))),
+        CALL_FACTORY_ENTRY("hcos",
+                           CALL(MapNames::getClNamespace(false, true) + "cos",
+                                CALL("float", ARG(0))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half"),
+    MATH_API_REWRITER_DEVICE(
+        "hexp",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("hexp"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("hexp", CALL(MapNames::getClNamespace() +
+                                                    "ext::intel::math::exp",
+                                                ARG(0)))),
+            EMPTY_FACTORY_ENTRY("hexp"),
+            CALL_FACTORY_ENTRY(
+                "hexp",
+                CALL(MapNames::getClNamespace(false, true) + "exp", ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "hexp",
+        CALL_FACTORY_ENTRY("hexp", CALL(MapNames::getClNamespace(false, true) +
+                                            "ext::oneapi::experimental::exp",
+                                        ARG(0))),
+        CALL_FACTORY_ENTRY("hexp",
+                           CALL(MapNames::getClNamespace(false, true) + "exp",
+                                CALL("float", ARG(0))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half"),
+    MATH_API_REWRITER_DEVICE(
+        "hexp10",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("hexp10"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("hexp10", CALL(MapNames::getClNamespace() +
+                                                      "ext::intel::math::exp10",
+                                                  ARG(0)))),
+            EMPTY_FACTORY_ENTRY("hexp10"),
+            CALL_FACTORY_ENTRY(
+                "hexp10", CALL(MapNames::getClNamespace(false, true) + "exp10",
+                               ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "hexp10",
+        CALL_FACTORY_ENTRY("hexp10",
+                           CALL(MapNames::getClNamespace(false, true) +
+                                    "ext::oneapi::experimental::exp10",
+                                ARG(0))),
+        CALL_FACTORY_ENTRY("hexp10",
+                           CALL(MapNames::getClNamespace(false, true) + "exp10",
+                                CALL("float", ARG(0))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half"),
+    MATH_API_REWRITER_DEVICE(
+        "hexp2",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("hexp2"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("hexp2", CALL(MapNames::getClNamespace() +
+                                                     "ext::intel::math::exp2",
+                                                 ARG(0)))),
+            EMPTY_FACTORY_ENTRY("hexp2"),
+            CALL_FACTORY_ENTRY(
+                "hexp2",
+                CALL(MapNames::getClNamespace(false, true) + "exp2", ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "hexp2",
+        CALL_FACTORY_ENTRY("hexp2", CALL(MapNames::getClNamespace(false, true) +
+                                             "ext::oneapi::experimental::exp2",
+                                         ARG(0))),
+        CALL_FACTORY_ENTRY("hexp2",
+                           CALL(MapNames::getClNamespace(false, true) + "exp2",
+                                CALL("float", ARG(0))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half"),
+    MATH_API_REWRITER_DEVICE(
+        "hfloor",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("hfloor"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("hfloor", CALL(MapNames::getClNamespace() +
+                                                      "ext::intel::math::floor",
+                                                  ARG(0)))),
+            EMPTY_FACTORY_ENTRY("hfloor"),
+            CALL_FACTORY_ENTRY(
+                "hfloor", CALL(MapNames::getClNamespace(false, true) + "floor",
+                               ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "hfloor",
+        CALL_FACTORY_ENTRY("hfloor",
+                           CALL(MapNames::getClNamespace(false, true) +
+                                    "ext::oneapi::experimental::floor",
+                                ARG(0))),
+        CALL_FACTORY_ENTRY("hfloor",
+                           CALL(MapNames::getClNamespace(false, true) + "floor",
+                                CALL("float", ARG(0))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half"),
+    MATH_API_REWRITER_DEVICE(
+        "hlog",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("hlog"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("hlog", CALL(MapNames::getClNamespace() +
+                                                    "ext::intel::math::log",
+                                                ARG(0)))),
+            EMPTY_FACTORY_ENTRY("hlog"),
+            CALL_FACTORY_ENTRY(
+                "hlog",
+                CALL(MapNames::getClNamespace(false, true) + "log", ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "hlog",
+        CALL_FACTORY_ENTRY("hlog", CALL(MapNames::getClNamespace(false, true) +
+                                            "ext::oneapi::experimental::log",
+                                        ARG(0))),
+        CALL_FACTORY_ENTRY("hlog",
+                           CALL(MapNames::getClNamespace(false, true) + "log",
+                                CALL("float", ARG(0))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half"),
+    MATH_API_REWRITER_DEVICE(
+        "hlog10",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("hlog10"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("hlog10", CALL(MapNames::getClNamespace() +
+                                                      "ext::intel::math::log10",
+                                                  ARG(0)))),
+            EMPTY_FACTORY_ENTRY("hlog10"),
+            CALL_FACTORY_ENTRY(
+                "hlog10", CALL(MapNames::getClNamespace(false, true) + "log10",
+                               ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "hlog10",
+        CALL_FACTORY_ENTRY("hlog10",
+                           CALL(MapNames::getClNamespace(false, true) +
+                                    "ext::oneapi::experimental::log10",
+                                ARG(0))),
+        CALL_FACTORY_ENTRY("hlog10",
+                           CALL(MapNames::getClNamespace(false, true) + "log10",
+                                CALL("float", ARG(0))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half"),
+    MATH_API_REWRITER_DEVICE(
+        "hlog2",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("hlog2"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("hlog2", CALL(MapNames::getClNamespace() +
+                                                     "ext::intel::math::log2",
+                                                 ARG(0)))),
+            EMPTY_FACTORY_ENTRY("hlog2"),
+            CALL_FACTORY_ENTRY(
+                "hlog2",
+                CALL(MapNames::getClNamespace(false, true) + "log2", ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "hlog2",
+        CALL_FACTORY_ENTRY("hlog2", CALL(MapNames::getClNamespace(false, true) +
+                                             "ext::oneapi::experimental::log2",
+                                         ARG(0))),
+        CALL_FACTORY_ENTRY("hlog2",
+                           CALL(MapNames::getClNamespace(false, true) + "log2",
+                                CALL("float", ARG(0))))))
+
+MATH_API_REWRITER_DEVICE(
+    "hrcp",
+    MATH_API_DEVICE_NODES(
+        EMPTY_FACTORY_ENTRY("hrcp"),
+        MATH_API_SPECIFIC_ELSE_EMU(
+            CheckArgType(0, "__half"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("hrcp", CALL(MapNames::getClNamespace() +
+                                                    "ext::intel::math::inv",
+                                                ARG(0))))),
+        EMPTY_FACTORY_ENTRY("hrcp"),
+        CALL_FACTORY_ENTRY("hrcp", CALL(MapNames::getClNamespace(false, true) +
+                                            "half_precision::recip",
+                                        CALL("float", ARG(0))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half"),
+    MATH_API_REWRITER_DEVICE(
+        "hrint",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("hrint"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("hrint", CALL(MapNames::getClNamespace() +
+                                                     "ext::intel::math::rint",
+                                                 ARG(0)))),
+            EMPTY_FACTORY_ENTRY("hrint"),
+            CALL_FACTORY_ENTRY(
+                "hrint",
+                CALL(MapNames::getClNamespace(false, true) + "rint", ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "hrint",
+        CALL_FACTORY_ENTRY("hrint", CALL(MapNames::getClNamespace(false, true) +
+                                             "ext::oneapi::experimental::rint",
+                                         ARG(0))),
+        CALL_FACTORY_ENTRY("hrint",
+                           CALL(MapNames::getClNamespace(false, true) + "rint",
+                                CALL("float", ARG(0))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half"),
+    MATH_API_REWRITER_DEVICE(
+        "hrsqrt",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("hrsqrt"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("hrsqrt", CALL(MapNames::getClNamespace() +
+                                                      "ext::intel::math::rsqrt",
+                                                  ARG(0)))),
+            EMPTY_FACTORY_ENTRY("hrsqrt"),
+            CALL_FACTORY_ENTRY(
+                "hrsqrt", CALL(MapNames::getClNamespace(false, true) + "rsqrt",
+                               ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "hrsqrt",
+        CALL_FACTORY_ENTRY("hrsqrt",
+                           CALL(MapNames::getClNamespace(false, true) +
+                                    "ext::oneapi::experimental::rsqrt",
+                                ARG(0))),
+        CALL_FACTORY_ENTRY("hrsqrt",
+                           CALL(MapNames::getClNamespace(false, true) + "rsqrt",
+                                CALL("float", ARG(0))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half"),
+    MATH_API_REWRITER_DEVICE(
+        "hsin",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("hsin"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("hsin", CALL(MapNames::getClNamespace() +
+                                                    "ext::intel::math::sin",
+                                                ARG(0)))),
+            EMPTY_FACTORY_ENTRY("hsin"),
+            CALL_FACTORY_ENTRY(
+                "hsin",
+                CALL(MapNames::getClNamespace(false, true) + "sin", ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "hsin",
+        CALL_FACTORY_ENTRY("hsin", CALL(MapNames::getClNamespace(false, true) +
+                                            "ext::oneapi::experimental::sin",
+                                        ARG(0))),
+        CALL_FACTORY_ENTRY("hsin",
+                           CALL(MapNames::getClNamespace(false, true) + "sin",
+                                CALL("float", ARG(0))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half"),
+    MATH_API_REWRITER_DEVICE(
+        "hsqrt",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("hsqrt"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("hsqrt", CALL(MapNames::getClNamespace() +
+                                                     "ext::intel::math::sqrt",
+                                                 ARG(0)))),
+            EMPTY_FACTORY_ENTRY("hsqrt"),
+            CALL_FACTORY_ENTRY(
+                "hsqrt",
+                CALL(MapNames::getClNamespace(false, true) + "sqrt", ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "hsqrt",
+        CALL_FACTORY_ENTRY("hsqrt", CALL(MapNames::getClNamespace(false, true) +
+                                             "ext::oneapi::experimental::sqrt",
+                                         ARG(0))),
+        CALL_FACTORY_ENTRY("hsqrt",
+                           CALL(MapNames::getClNamespace(false, true) + "sqrt",
+                                CALL("float", ARG(0))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half"),
+    MATH_API_REWRITER_DEVICE(
+        "htrunc",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("htrunc"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("htrunc", CALL(MapNames::getClNamespace() +
+                                                      "ext::intel::math::trunc",
+                                                  ARG(0)))),
+            EMPTY_FACTORY_ENTRY("htrunc"),
+            CALL_FACTORY_ENTRY(
+                "htrunc", CALL(MapNames::getClNamespace(false, true) + "trunc",
+                               ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "htrunc",
+        CALL_FACTORY_ENTRY("htrunc",
+                           CALL(MapNames::getClNamespace(false, true) +
+                                    "ext::oneapi::experimental::trunc",
+                                ARG(0))),
+        CALL_FACTORY_ENTRY("htrunc",
+                           CALL(MapNames::getClNamespace(false, true) + "trunc",
+                                CALL("float", ARG(0))))))
+
+// Half2 Math Functions
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half2"),
+    MATH_API_REWRITER_DEVICE(
+        "h2ceil",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("h2ceil"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("h2ceil", CALL(MapNames::getClNamespace() +
+                                                      "ext::intel::math::ceil",
+                                                  ARG(0)))),
+            EMPTY_FACTORY_ENTRY("h2ceil"),
+            CALL_FACTORY_ENTRY(
+                "h2ceil",
+                CALL(MapNames::getClNamespace(false, true) + "ceil", ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "h2ceil",
+        CALL_FACTORY_ENTRY("h2ceil",
+                           CALL(MapNames::getClNamespace(false, true) +
+                                    "ext::oneapi::experimental::ceil",
+                                ARG(0))),
+        CALL_FACTORY_ENTRY(
+            "h2ceil",
+            CALL(MapNames::getClNamespace() + "marray<" +
+                     MapNames::getClNamespace() + "ext::oneapi::bfloat16, 2>",
+                 CALL(MapNames::getClNamespace(false, true) + "ceil",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("0")))),
+                 CALL(MapNames::getClNamespace(false, true) + "ceil",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("1"))))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half2"),
+    MATH_API_REWRITER_DEVICE(
+        "h2cos",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("h2cos"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("h2cos", CALL(MapNames::getClNamespace() +
+                                                     "ext::intel::math::cos",
+                                                 ARG(0)))),
+            EMPTY_FACTORY_ENTRY("h2cos"),
+            CALL_FACTORY_ENTRY(
+                "h2cos",
+                CALL(MapNames::getClNamespace(false, true) + "cos", ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "h2cos",
+        CALL_FACTORY_ENTRY("h2cos", CALL(MapNames::getClNamespace(false, true) +
+                                             "ext::oneapi::experimental::cos",
+                                         ARG(0))),
+        CALL_FACTORY_ENTRY(
+            "h2cos",
+            CALL(MapNames::getClNamespace() + "marray<" +
+                     MapNames::getClNamespace() + "ext::oneapi::bfloat16, 2>",
+                 CALL(MapNames::getClNamespace(false, true) + "cos",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("0")))),
+                 CALL(MapNames::getClNamespace(false, true) + "cos",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("1"))))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half2"),
+    MATH_API_REWRITER_DEVICE(
+        "h2exp",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("h2exp"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("h2exp", CALL(MapNames::getClNamespace() +
+                                                     "ext::intel::math::exp",
+                                                 ARG(0)))),
+            EMPTY_FACTORY_ENTRY("h2exp"),
+            CALL_FACTORY_ENTRY(
+                "h2exp",
+                CALL(MapNames::getClNamespace(false, true) + "exp", ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "h2exp",
+        CALL_FACTORY_ENTRY("h2exp", CALL(MapNames::getClNamespace(false, true) +
+                                             "ext::oneapi::experimental::exp",
+                                         ARG(0))),
+        CALL_FACTORY_ENTRY(
+            "h2exp",
+            CALL(MapNames::getClNamespace() + "marray<" +
+                     MapNames::getClNamespace() + "ext::oneapi::bfloat16, 2>",
+                 CALL(MapNames::getClNamespace(false, true) + "exp",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("0")))),
+                 CALL(MapNames::getClNamespace(false, true) + "exp",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("1"))))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half2"),
+    MATH_API_REWRITER_DEVICE(
+        "h2exp10",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("h2exp10"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("h2exp10",
+                                   CALL(MapNames::getClNamespace() +
+                                            "ext::intel::math::exp10",
+                                        ARG(0)))),
+            EMPTY_FACTORY_ENTRY("h2exp10"),
+            CALL_FACTORY_ENTRY(
+                "h2exp10", CALL(MapNames::getClNamespace(false, true) + "exp10",
+                                ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "h2exp10",
+        CALL_FACTORY_ENTRY("h2exp10",
+                           CALL(MapNames::getClNamespace(false, true) +
+                                    "ext::oneapi::experimental::exp10",
+                                ARG(0))),
+        CALL_FACTORY_ENTRY(
+            "h2exp10",
+            CALL(MapNames::getClNamespace() + "marray<" +
+                     MapNames::getClNamespace() + "ext::oneapi::bfloat16, 2>",
+                 CALL(MapNames::getClNamespace(false, true) + "exp10",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("0")))),
+                 CALL(MapNames::getClNamespace(false, true) + "exp10",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("1"))))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half2"),
+    MATH_API_REWRITER_DEVICE(
+        "h2exp2",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("h2exp2"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("h2exp2", CALL(MapNames::getClNamespace() +
+                                                      "ext::intel::math::exp2",
+                                                  ARG(0)))),
+            EMPTY_FACTORY_ENTRY("h2exp2"),
+            CALL_FACTORY_ENTRY(
+                "h2exp2",
+                CALL(MapNames::getClNamespace(false, true) + "exp2", ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "h2exp2",
+        CALL_FACTORY_ENTRY("h2exp2",
+                           CALL(MapNames::getClNamespace(false, true) +
+                                    "ext::oneapi::experimental::exp2",
+                                ARG(0))),
+        CALL_FACTORY_ENTRY(
+            "h2exp2",
+            CALL(MapNames::getClNamespace() + "marray<" +
+                     MapNames::getClNamespace() + "ext::oneapi::bfloat16, 2>",
+                 CALL(MapNames::getClNamespace(false, true) + "exp2",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("0")))),
+                 CALL(MapNames::getClNamespace(false, true) + "exp2",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("1"))))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half2"),
+    MATH_API_REWRITER_DEVICE(
+        "h2floor",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("h2floor"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("h2floor",
+                                   CALL(MapNames::getClNamespace() +
+                                            "ext::intel::math::floor",
+                                        ARG(0)))),
+            EMPTY_FACTORY_ENTRY("h2floor"),
+            CALL_FACTORY_ENTRY(
+                "h2floor", CALL(MapNames::getClNamespace(false, true) + "floor",
+                                ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "h2floor",
+        CALL_FACTORY_ENTRY("h2floor",
+                           CALL(MapNames::getClNamespace(false, true) +
+                                    "ext::oneapi::experimental::floor",
+                                ARG(0))),
+        CALL_FACTORY_ENTRY(
+            "h2floor",
+            CALL(MapNames::getClNamespace() + "marray<" +
+                     MapNames::getClNamespace() + "ext::oneapi::bfloat16, 2>",
+                 CALL(MapNames::getClNamespace(false, true) + "floor",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("0")))),
+                 CALL(MapNames::getClNamespace(false, true) + "floor",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("1"))))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half2"),
+    MATH_API_REWRITER_DEVICE(
+        "h2log",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("h2log"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("h2log", CALL(MapNames::getClNamespace() +
+                                                     "ext::intel::math::log",
+                                                 ARG(0)))),
+            EMPTY_FACTORY_ENTRY("h2log"),
+            CALL_FACTORY_ENTRY(
+                "h2log",
+                CALL(MapNames::getClNamespace(false, true) + "log", ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "h2log",
+        CALL_FACTORY_ENTRY("h2log", CALL(MapNames::getClNamespace(false, true) +
+                                             "ext::oneapi::experimental::log",
+                                         ARG(0))),
+        CALL_FACTORY_ENTRY(
+            "h2log",
+            CALL(MapNames::getClNamespace() + "marray<" +
+                     MapNames::getClNamespace() + "ext::oneapi::bfloat16, 2>",
+                 CALL(MapNames::getClNamespace(false, true) + "log",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("0")))),
+                 CALL(MapNames::getClNamespace(false, true) + "log",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("1"))))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half2"),
+    MATH_API_REWRITER_DEVICE(
+        "h2log10",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("h2log10"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("h2log10",
+                                   CALL(MapNames::getClNamespace() +
+                                            "ext::intel::math::log10",
+                                        ARG(0)))),
+            EMPTY_FACTORY_ENTRY("h2log10"),
+            CALL_FACTORY_ENTRY(
+                "h2log10", CALL(MapNames::getClNamespace(false, true) + "log10",
+                                ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "h2log10",
+        CALL_FACTORY_ENTRY("h2log10",
+                           CALL(MapNames::getClNamespace(false, true) +
+                                    "ext::oneapi::experimental::log10",
+                                ARG(0))),
+        CALL_FACTORY_ENTRY(
+            "h2log10",
+            CALL(MapNames::getClNamespace() + "marray<" +
+                     MapNames::getClNamespace() + "ext::oneapi::bfloat16, 2>",
+                 CALL(MapNames::getClNamespace(false, true) + "log10",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("0")))),
+                 CALL(MapNames::getClNamespace(false, true) + "log10",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("1"))))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half2"),
+    MATH_API_REWRITER_DEVICE(
+        "h2log2",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("h2log2"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("h2log2", CALL(MapNames::getClNamespace() +
+                                                      "ext::intel::math::log2",
+                                                  ARG(0)))),
+            EMPTY_FACTORY_ENTRY("h2log2"),
+            CALL_FACTORY_ENTRY(
+                "h2log2",
+                CALL(MapNames::getClNamespace(false, true) + "log2", ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "h2log2",
+        CALL_FACTORY_ENTRY("h2log2",
+                           CALL(MapNames::getClNamespace(false, true) +
+                                    "ext::oneapi::experimental::log2",
+                                ARG(0))),
+        CALL_FACTORY_ENTRY(
+            "h2log2",
+            CALL(MapNames::getClNamespace() + "marray<" +
+                     MapNames::getClNamespace() + "ext::oneapi::bfloat16, 2>",
+                 CALL(MapNames::getClNamespace(false, true) + "log2",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("0")))),
+                 CALL(MapNames::getClNamespace(false, true) + "log2",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("1"))))))))
+
+MATH_API_REWRITER_DEVICE(
+    "h2rcp",
+    MATH_API_DEVICE_NODES(
+        EMPTY_FACTORY_ENTRY("h2rcp"),
+        MATH_API_SPECIFIC_ELSE_EMU(
+            CheckArgType(0, "__half2"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("h2rcp", CALL(MapNames::getClNamespace() +
+                                                     "ext::intel::math::inv",
+                                                 ARG(0))))),
+        EMPTY_FACTORY_ENTRY("h2rcp"),
+        CONDITIONAL_FACTORY_ENTRY(
+            CheckArgType(0, "__half2"),
+            CALL_FACTORY_ENTRY(
+                "h2rcp",
+                CALL(MapNames::getClNamespace() + "half2",
+                     CALL(MapNames::getClNamespace(false, true) +
+                              "half_precision::recip",
+                          CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("0")))),
+                     CALL(MapNames::getClNamespace(false, true) +
+                              "half_precision::recip",
+                          CALL("float",
+                               ARRAY_SUBSCRIPT(ARG(0), LITERAL("1")))))),
+            CONDITIONAL_FACTORY_ENTRY(
+                math::UseBFloat16,
+                CALL_FACTORY_ENTRY(
+                    "h2rcp",
+                    CALL(MapNames::getClNamespace() + "marray<" +
+                             MapNames::getClNamespace() +
+                             "ext::oneapi::bfloat16, 2>",
+                         CALL(MapNames::getClNamespace(false, true) +
+                                  "half_precision::recip",
+                              CALL("float",
+                                   ARRAY_SUBSCRIPT(ARG(0), LITERAL("0")))),
+                         CALL(MapNames::getClNamespace(false, true) +
+                                  "half_precision::recip",
+                              CALL("float",
+                                   ARRAY_SUBSCRIPT(ARG(0), LITERAL("1")))))),
+                UNSUPPORT_FACTORY_ENTRY("h2rcp", Diagnostics::API_NOT_MIGRATED,
+                                        ARG("h2rcp"))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half2"),
+    MATH_API_REWRITER_DEVICE(
+        "h2rint",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("h2rint"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("h2rint", CALL(MapNames::getClNamespace() +
+                                                      "ext::intel::math::rint",
+                                                  ARG(0)))),
+            EMPTY_FACTORY_ENTRY("h2rint"),
+            CALL_FACTORY_ENTRY(
+                "h2rint",
+                CALL(MapNames::getClNamespace(false, true) + "rint", ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "h2rint",
+        CALL_FACTORY_ENTRY("h2rint",
+                           CALL(MapNames::getClNamespace(false, true) +
+                                    "ext::oneapi::experimental::rint",
+                                ARG(0))),
+        CALL_FACTORY_ENTRY(
+            "h2rint",
+            CALL(MapNames::getClNamespace() + "marray<" +
+                     MapNames::getClNamespace() + "ext::oneapi::bfloat16, 2>",
+                 CALL(MapNames::getClNamespace(false, true) + "rint",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("0")))),
+                 CALL(MapNames::getClNamespace(false, true) + "rint",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("1"))))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half2"),
+    MATH_API_REWRITER_DEVICE(
+        "h2rsqrt",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("h2rsqrt"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("h2rsqrt",
+                                   CALL(MapNames::getClNamespace() +
+                                            "ext::intel::math::rsqrt",
+                                        ARG(0)))),
+            EMPTY_FACTORY_ENTRY("h2rsqrt"),
+            CALL_FACTORY_ENTRY(
+                "h2rsqrt", CALL(MapNames::getClNamespace(false, true) + "rsqrt",
+                                ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "h2rsqrt",
+        CALL_FACTORY_ENTRY("h2rsqrt",
+                           CALL(MapNames::getClNamespace(false, true) +
+                                    "ext::oneapi::experimental::rsqrt",
+                                ARG(0))),
+        CALL_FACTORY_ENTRY(
+            "h2rsqrt",
+            CALL(MapNames::getClNamespace() + "marray<" +
+                     MapNames::getClNamespace() + "ext::oneapi::bfloat16, 2>",
+                 CALL(MapNames::getClNamespace(false, true) + "rsqrt",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("0")))),
+                 CALL(MapNames::getClNamespace(false, true) + "rsqrt",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("1"))))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half2"),
+    MATH_API_REWRITER_DEVICE(
+        "h2sin",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("h2sin"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("h2sin", CALL(MapNames::getClNamespace() +
+                                                     "ext::intel::math::sin",
+                                                 ARG(0)))),
+            EMPTY_FACTORY_ENTRY("h2sin"),
+            CALL_FACTORY_ENTRY(
+                "h2sin",
+                CALL(MapNames::getClNamespace(false, true) + "sin", ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "h2sin",
+        CALL_FACTORY_ENTRY("h2sin", CALL(MapNames::getClNamespace(false, true) +
+                                             "ext::oneapi::experimental::sin",
+                                         ARG(0))),
+        CALL_FACTORY_ENTRY(
+            "h2sin",
+            CALL(MapNames::getClNamespace() + "marray<" +
+                     MapNames::getClNamespace() + "ext::oneapi::bfloat16, 2>",
+                 CALL(MapNames::getClNamespace(false, true) + "sin",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("0")))),
+                 CALL(MapNames::getClNamespace(false, true) + "sin",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("1"))))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half2"),
+    MATH_API_REWRITER_DEVICE(
+        "h2sqrt",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("h2sqrt"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("h2sqrt", CALL(MapNames::getClNamespace() +
+                                                      "ext::intel::math::sqrt",
+                                                  ARG(0)))),
+            EMPTY_FACTORY_ENTRY("h2sqrt"),
+            CALL_FACTORY_ENTRY(
+                "h2sqrt",
+                CALL(MapNames::getClNamespace(false, true) + "sqrt", ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "h2sqrt",
+        CALL_FACTORY_ENTRY("h2sqrt",
+                           CALL(MapNames::getClNamespace(false, true) +
+                                    "ext::oneapi::experimental::sqrt",
+                                ARG(0))),
+        CALL_FACTORY_ENTRY(
+            "h2sqrt",
+            CALL(MapNames::getClNamespace() + "marray<" +
+                     MapNames::getClNamespace() + "ext::oneapi::bfloat16, 2>",
+                 CALL(MapNames::getClNamespace(false, true) + "sqrt",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("0")))),
+                 CALL(MapNames::getClNamespace(false, true) + "sqrt",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("1"))))))))
+
+MATH_API_REWRITER_DEVICE_OVERLOAD(
+    CheckArgType(0, "__half2"),
+    MATH_API_REWRITER_DEVICE(
+        "h2trunc",
+        MATH_API_DEVICE_NODES(
+            EMPTY_FACTORY_ENTRY("h2trunc"),
+            HEADER_INSERT_FACTORY(
+                HeaderType::HT_SYCL_Math,
+                CALL_FACTORY_ENTRY("h2trunc",
+                                   CALL(MapNames::getClNamespace() +
+                                            "ext::intel::math::trunc",
+                                        ARG(0)))),
+            EMPTY_FACTORY_ENTRY("h2trunc"),
+            CALL_FACTORY_ENTRY(
+                "h2trunc", CALL(MapNames::getClNamespace(false, true) + "trunc",
+                                ARG(0))))),
+    MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16(
+        "h2trunc",
+        CALL_FACTORY_ENTRY("h2trunc",
+                           CALL(MapNames::getClNamespace(false, true) +
+                                    "ext::oneapi::experimental::trunc",
+                                ARG(0))),
+        CALL_FACTORY_ENTRY(
+            "h2trunc",
+            CALL(MapNames::getClNamespace() + "marray<" +
+                     MapNames::getClNamespace() + "ext::oneapi::bfloat16, 2>",
+                 CALL(MapNames::getClNamespace(false, true) + "trunc",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("0")))),
+                 CALL(MapNames::getClNamespace(false, true) + "trunc",
+                      CALL("float", ARRAY_SUBSCRIPT(ARG(0), LITERAL("1"))))))))
+
 // Bfloat16 Precision Conversion and Data Movement
 CALL_FACTORY_ENTRY("__bfloat1622float2",
                    CALL(MapNames::getClNamespace() + "float2",
diff --git a/clang/test/dpct/math/bfloat16/bfloat16.cu b/clang/test/dpct/math/bfloat16/bfloat16.cu
index fc3030f7688a..579116114c94 100644
--- a/clang/test/dpct/math/bfloat16/bfloat16.cu
+++ b/clang/test/dpct/math/bfloat16/bfloat16.cu
@@ -296,4 +296,74 @@ void test_conversions() {
   bf16 = __float2bfloat16(f);
 }
 
+__global__ void kernelFuncBfloat16Math() {
+  // CHECK: sycl::ext::oneapi::bfloat16 bf16, bf16_1;
+  __nv_bfloat16 bf16, bf16_1;
+  // CHECK: bf16_1 = sycl::ceil(float(bf16));
+  bf16_1 = hceil(bf16);
+  // CHECK: bf16_1 = sycl::cos(float(bf16));
+  bf16_1 = hcos(bf16);
+  // CHECK: bf16_1 = sycl::exp(float(bf16));
+  bf16_1 = hexp(bf16);
+  // CHECK: bf16_1 = sycl::exp10(float(bf16));
+  bf16_1 = hexp10(bf16);
+  // CHECK: bf16_1 = sycl::exp2(float(bf16));
+  bf16_1 = hexp2(bf16);
+  // CHECK: bf16_1 = sycl::floor(float(bf16));
+  bf16_1 = hfloor(bf16);
+  // CHECK: bf16_1 = sycl::log(float(bf16));
+  bf16_1 = hlog(bf16);
+  // CHECK: bf16_1 = sycl::log10(float(bf16));
+  bf16_1 = hlog10(bf16);
+  // CHECK: bf16_1 = sycl::log2(float(bf16));
+  bf16_1 = hlog2(bf16);
+  // CHECK: bf16_1 = sycl::half_precision::recip(float(bf16));
+  bf16_1 = hrcp(bf16);
+  // CHECK: bf16_1 = sycl::rint(float(bf16));
+  bf16_1 = hrint(bf16);
+  // CHECK: bf16_1 = sycl::rsqrt(float(bf16));
+  bf16_1 = hrsqrt(bf16);
+  // CHECK: bf16_1 = sycl::sin(float(bf16));
+  bf16_1 = hsin(bf16);
+  // CHECK: bf16_1 = sycl::sqrt(float(bf16));
+  bf16_1 = hsqrt(bf16);
+  // CHECK: bf16_1 = sycl::trunc(float(bf16));
+  bf16_1 = htrunc(bf16);
+}
+
+__global__ void kernelFuncBfloat162Math() {
+  // CHECK: sycl::marray<sycl::ext::oneapi::bfloat16, 2> bf162, bf162_1;
+  __nv_bfloat162 bf162, bf162_1;
+  // CHECK: bf162_1 = sycl::marray<sycl::ext::oneapi::bfloat16, 2>(sycl::ceil(float(bf162[0])), sycl::ceil(float(bf162[1])));
+  bf162_1 = h2ceil(bf162);
+  // CHECK: bf162_1 = sycl::marray<sycl::ext::oneapi::bfloat16, 2>(sycl::cos(float(bf162[0])), sycl::cos(float(bf162[1])));
+  bf162_1 = h2cos(bf162);
+  // CHECK: bf162_1 = sycl::marray<sycl::ext::oneapi::bfloat16, 2>(sycl::exp(float(bf162[0])), sycl::exp(float(bf162[1])));
+  bf162_1 = h2exp(bf162);
+  // CHECK: bf162_1 = sycl::marray<sycl::ext::oneapi::bfloat16, 2>(sycl::exp10(float(bf162[0])), sycl::exp10(float(bf162[1])));
+  bf162_1 = h2exp10(bf162);
+  // CHECK: bf162_1 = sycl::marray<sycl::ext::oneapi::bfloat16, 2>(sycl::exp2(float(bf162[0])), sycl::exp2(float(bf162[1])));
+  bf162_1 = h2exp2(bf162);
+  // CHECK: bf162_1 = sycl::marray<sycl::ext::oneapi::bfloat16, 2>(sycl::floor(float(bf162[0])), sycl::floor(float(bf162[1])));
+  bf162_1 = h2floor(bf162);
+  // CHECK: bf162_1 = sycl::marray<sycl::ext::oneapi::bfloat16, 2>(sycl::log(float(bf162[0])), sycl::log(float(bf162[1])));
+  bf162_1 = h2log(bf162);
+  // CHECK: bf162_1 = sycl::marray<sycl::ext::oneapi::bfloat16, 2>(sycl::log10(float(bf162[0])), sycl::log10(float(bf162[1])));
+  bf162_1 = h2log10(bf162);
+  // CHECK: bf162_1 = sycl::marray<sycl::ext::oneapi::bfloat16, 2>(sycl::log2(float(bf162[0])), sycl::log2(float(bf162[1])));
+  bf162_1 = h2log2(bf162);
+  // CHECK: bf162_1 = sycl::marray<sycl::ext::oneapi::bfloat16, 2>(sycl::half_precision::recip(float(bf162[0])), sycl::half_precision::recip(float(bf162[1])));
+  bf162_1 = h2rcp(bf162);
+  // CHECK: bf162_1 = sycl::marray<sycl::ext::oneapi::bfloat16, 2>(sycl::rint(float(bf162[0])), sycl::rint(float(bf162[1])));
+  bf162_1 = h2rint(bf162);
+  // CHECK: bf162_1 = sycl::marray<sycl::ext::oneapi::bfloat16, 2>(sycl::rsqrt(float(bf162[0])), sycl::rsqrt(float(bf162[1])));
+  bf162_1 = h2rsqrt(bf162);
+  // CHECK: bf162_1 = sycl::marray<sycl::ext::oneapi::bfloat16, 2>(sycl::sin(float(bf162[0])), sycl::sin(float(bf162[1])));
+  bf162_1 = h2sin(bf162);
+  // CHECK: bf162_1 = sycl::marray<sycl::ext::oneapi::bfloat16, 2>(sycl::sqrt(float(bf162[0])), sycl::sqrt(float(bf162[1])));
+  bf162_1 = h2sqrt(bf162);
+  // CHECK: bf162_1 = sycl::marray<sycl::ext::oneapi::bfloat16, 2>(sycl::trunc(float(bf162[0])), sycl::trunc(float(bf162[1])));
+  bf162_1 = h2trunc(bf162);
+}
+
 int main() { return 0; }
diff --git a/clang/test/dpct/math/bfloat16/bfloat16_experimental.cu b/clang/test/dpct/math/bfloat16/bfloat16_experimental.cu
index 589d0f1b69b1..ddaf2801260e 100644
--- a/clang/test/dpct/math/bfloat16/bfloat16_experimental.cu
+++ b/clang/test/dpct/math/bfloat16/bfloat16_experimental.cu
@@ -77,4 +77,74 @@ __global__ void kernelFuncBfloat162Arithmetic() {
   bf162 = __hsub2_sat(bf162_1, bf162_2);
 }
 
+__global__ void kernelFuncBfloat16Math() {
+  // CHECK: sycl::ext::oneapi::bfloat16 bf16, bf16_1;
+  __nv_bfloat16 bf16, bf16_1;
+  // CHECK: bf16_1 = sycl::ext::oneapi::experimental::ceil(bf16);
+  bf16_1 = hceil(bf16);
+  // CHECK: bf16_1 = sycl::ext::oneapi::experimental::cos(bf16);
+  bf16_1 = hcos(bf16);
+  // CHECK: bf16_1 = sycl::ext::oneapi::experimental::exp(bf16);
+  bf16_1 = hexp(bf16);
+  // CHECK: bf16_1 = sycl::ext::oneapi::experimental::exp10(bf16);
+  bf16_1 = hexp10(bf16);
+  // CHECK: bf16_1 = sycl::ext::oneapi::experimental::exp2(bf16);
+  bf16_1 = hexp2(bf16);
+  // CHECK: bf16_1 = sycl::ext::oneapi::experimental::floor(bf16);
+  bf16_1 = hfloor(bf16);
+  // CHECK: bf16_1 = sycl::ext::oneapi::experimental::log(bf16);
+  bf16_1 = hlog(bf16);
+  // CHECK: bf16_1 = sycl::ext::oneapi::experimental::log10(bf16);
+  bf16_1 = hlog10(bf16);
+  // CHECK: bf16_1 = sycl::ext::oneapi::experimental::log2(bf16);
+  bf16_1 = hlog2(bf16);
+  // CHECK: bf16_1 = sycl::half_precision::recip(float(bf16));
+  bf16_1 = hrcp(bf16);
+  // CHECK: bf16_1 = sycl::ext::oneapi::experimental::rint(bf16);
+  bf16_1 = hrint(bf16);
+  // CHECK: bf16_1 = sycl::ext::oneapi::experimental::rsqrt(bf16);
+  bf16_1 = hrsqrt(bf16);
+  // CHECK: bf16_1 = sycl::ext::oneapi::experimental::sin(bf16);
+  bf16_1 = hsin(bf16);
+  // CHECK: bf16_1 = sycl::ext::oneapi::experimental::sqrt(bf16);
+  bf16_1 = hsqrt(bf16);
+  // CHECK: bf16_1 = sycl::ext::oneapi::experimental::trunc(bf16);
+  bf16_1 = htrunc(bf16);
+}
+
+__global__ void kernelFuncBfloat162Math() {
+  // CHECK: sycl::marray<sycl::ext::oneapi::bfloat16, 2> bf162, bf162_1;
+  __nv_bfloat162 bf162, bf162_1;
+  // CHECK: bf162_1 = sycl::ext::oneapi::experimental::ceil(bf162);
+  bf162_1 = h2ceil(bf162);
+  // CHECK: bf162_1 = sycl::ext::oneapi::experimental::cos(bf162);
+  bf162_1 = h2cos(bf162);
+  // CHECK: bf162_1 = sycl::ext::oneapi::experimental::exp(bf162);
+  bf162_1 = h2exp(bf162);
+  // CHECK: bf162_1 = sycl::ext::oneapi::experimental::exp10(bf162);
+  bf162_1 = h2exp10(bf162);
+  // CHECK: bf162_1 = sycl::ext::oneapi::experimental::exp2(bf162);
+  bf162_1 = h2exp2(bf162);
+  // CHECK: bf162_1 = sycl::ext::oneapi::experimental::floor(bf162);
+  bf162_1 = h2floor(bf162);
+  // CHECK: bf162_1 = sycl::ext::oneapi::experimental::log(bf162);
+  bf162_1 = h2log(bf162);
+  // CHECK: bf162_1 = sycl::ext::oneapi::experimental::log10(bf162);
+  bf162_1 = h2log10(bf162);
+  // CHECK: bf162_1 = sycl::ext::oneapi::experimental::log2(bf162);
+  bf162_1 = h2log2(bf162);
+  // CHECK: bf162_1 = sycl::marray<sycl::ext::oneapi::bfloat16, 2>(sycl::half_precision::recip(float(bf162[0])), sycl::half_precision::recip(float(bf162[1])));
+  bf162_1 = h2rcp(bf162);
+  // CHECK: bf162_1 = sycl::ext::oneapi::experimental::rint(bf162);
+  bf162_1 = h2rint(bf162);
+  // CHECK: bf162_1 = sycl::ext::oneapi::experimental::rsqrt(bf162);
+  bf162_1 = h2rsqrt(bf162);
+  // CHECK: bf162_1 = sycl::ext::oneapi::experimental::sin(bf162);
+  bf162_1 = h2sin(bf162);
+  // CHECK: bf162_1 = sycl::ext::oneapi::experimental::sqrt(bf162);
+  bf162_1 = h2sqrt(bf162);
+  // CHECK: bf162_1 = sycl::ext::oneapi::experimental::trunc(bf162);
+  bf162_1 = h2trunc(bf162);
+}
+
 int main() { return 0; }
diff --git a/clang/test/dpct/math/cuda-math-extension.cu b/clang/test/dpct/math/cuda-math-extension.cu
index a127fbb3047d..529263a782c0 100644
--- a/clang/test/dpct/math/cuda-math-extension.cu
+++ b/clang/test/dpct/math/cuda-math-extension.cu
@@ -161,6 +161,39 @@ __global__ void kernelFuncHalf() {
   b = __hne(h, h_1);
   // CHECK: b = sycl::ext::intel::math::hneu(h, h_1);
   b = __hneu(h, h_1);
+
+  // Half Math Functions
+
+  // CHECK: h_2 = sycl::ext::intel::math::ceil(h);
+  h_2 = hceil(h);
+  // CHECK: h_2 = sycl::ext::intel::math::cos(h);
+  h_2 = hcos(h);
+  // CHECK: h_2 = sycl::ext::intel::math::exp(h);
+  h_2 = hexp(h);
+  // CHECK: h_2 = sycl::ext::intel::math::exp10(h);
+  h_2 = hexp10(h);
+  // CHECK: h_2 = sycl::ext::intel::math::exp2(h);
+  h_2 = hexp2(h);
+  // CHECK: h_2 = sycl::ext::intel::math::floor(h);
+  h_2 = hfloor(h);
+  // CHECK: h_2 = sycl::ext::intel::math::log(h);
+  h_2 = hlog(h);
+  // CHECK: h_2 = sycl::ext::intel::math::log10(h);
+  h_2 = hlog10(h);
+  // CHECK: h_2 = sycl::ext::intel::math::log2(h);
+  h_2 = hlog2(h);
+  // CHECK: h_2 = sycl::ext::intel::math::inv(h);
+  h_2 = hrcp(h);
+  // CHECK: h_2 = sycl::ext::intel::math::rint(h);
+  h_2 = hrint(h);
+  // CHECK: h_2 = sycl::ext::intel::math::rsqrt(h);
+  h_2 = hrsqrt(h);
+  // CHECK: h_2 = sycl::ext::intel::math::sin(h);
+  h_2 = hsin(h);
+  // CHECK: h_2 = sycl::ext::intel::math::sqrt(h);
+  h_2 = hsqrt(h);
+  // CHECK: h_2 = sycl::ext::intel::math::trunc(h);
+  h_2 = htrunc(h);
 }
 
 __global__ void kernelFuncHalf2() {
@@ -241,6 +274,39 @@ __global__ void kernelFuncHalf2() {
   h2_2 = __hne2(h2, h2_1);
   // CHECK: sycl::ext::intel::math::hneu2(h2, h2);
   __hneu2(h2, h2);
+
+  // Half2 Math Functions
+
+  // CHECK: h2_2 = sycl::ext::intel::math::ceil(h2);
+  h2_2 = h2ceil(h2);
+  // CHECK: h2_2 = sycl::ext::intel::math::cos(h2);
+  h2_2 = h2cos(h2);
+  // CHECK: h2_2 = sycl::ext::intel::math::exp(h2);
+  h2_2 = h2exp(h2);
+  // CHECK: h2_2 = sycl::ext::intel::math::exp10(h2);
+  h2_2 = h2exp10(h2);
+  // CHECK: h2_2 = sycl::ext::intel::math::exp2(h2);
+  h2_2 = h2exp2(h2);
+  // CHECK: h2_2 = sycl::ext::intel::math::floor(h2);
+  h2_2 = h2floor(h2);
+  // CHECK: h2_2 = sycl::ext::intel::math::log(h2);
+  h2_2 = h2log(h2);
+  // CHECK: h2_2 = sycl::ext::intel::math::log10(h2);
+  h2_2 = h2log10(h2);
+  // CHECK: h2_2 = sycl::ext::intel::math::log2(h2);
+  h2_2 = h2log2(h2);
+  // CHECK: h2_2 = sycl::ext::intel::math::inv(h2);
+  h2_2 = h2rcp(h2);
+  // CHECK: h2_2 = sycl::ext::intel::math::rint(h2);
+  h2_2 = h2rint(h2);
+  // CHECK: h2_2 = sycl::ext::intel::math::rsqrt(h2);
+  h2_2 = h2rsqrt(h2);
+  // CHECK: h2_2 = sycl::ext::intel::math::sin(h2);
+  h2_2 = h2sin(h2);
+  // CHECK: h2_2 = sycl::ext::intel::math::sqrt(h2);
+  h2_2 = h2sqrt(h2);
+  // CHECK: h2_2 = sycl::ext::intel::math::trunc(h2);
+  h2_2 = h2trunc(h2);
 }
 
 __global__ void kernelFuncInt() {
diff --git a/clang/test/dpct/math/cuda-math-intrinsics.cu b/clang/test/dpct/math/cuda-math-intrinsics.cu
index acbba3af44d9..49ae92da6eb6 100644
--- a/clang/test/dpct/math/cuda-math-intrinsics.cu
+++ b/clang/test/dpct/math/cuda-math-intrinsics.cu
@@ -207,7 +207,7 @@ __global__ void kernelFuncHalf(double *deviceArrayDouble) {
   h_2 = hlog10(h);
   // CHECK: h_2 = sycl::log2(h);
   h_2 = hlog2(h);
-  // CHECK: h_2 = sycl::half_precision::recip<float>(h);
+  // CHECK: h_2 = sycl::half_precision::recip(float(h));
   h_2 = hrcp(h);
   // CHECK: h_2 = sycl::rint(h);
   h_2 = hrint(h);
@@ -240,9 +240,7 @@ __global__ void kernelFuncHalf(double *deviceArrayDouble) {
   h2_2 = h2log10(h2);
   // CHECK: h2_2 = sycl::log2(h2);
   h2_2 = h2log2(h2);
-  // CHECK: /*
-  // CHECK-NEXT: DPCT1007:{{[0-9]+}}: Migration of h2rcp is not supported.
-  // CHECK-NEXT: */
+  // CHECK: h2_2 = sycl::half2(sycl::half_precision::recip(float(h2[0])), sycl::half_precision::recip(float(h2[1])));
   h2_2 = h2rcp(h2);
   // CHECK: h2_2 = sycl::rint(h2);
   h2_2 = h2rint(h2);