diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index e8630740a3ee6..578c9080a3e58 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -11834,9 +11834,6 @@ def err_sycl_restrict : Error<
 def err_sycl_external_global : Error<
   "invalid reference to 'device_global' variable; external 'device_global'"
   " variable must be marked with SYCL_EXTERNAL macro">;
-def err_sycl_external_no_rdc : Error<
-  "invalid %select{declaration|definition}0 of SYCL_EXTERNAL function in non-relocatable "
-   "device code mode">;
 def warn_sycl_kernel_too_big_args : Warning<
   "size of kernel arguments (%0 bytes) may exceed the supported maximum "
   "of %1 bytes on some devices">, InGroup<SyclStrict>, ShowInSystemHeader;
diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp
index 91db4555b44e3..7e8fff4c254c5 100644
--- a/clang/lib/Frontend/InitPreprocessor.cpp
+++ b/clang/lib/Frontend/InitPreprocessor.cpp
@@ -1297,7 +1297,8 @@ static void InitializePredefinedMacros(const TargetInfo &TI,
   // SYCL device compiler which doesn't produce host binary.
   if (LangOpts.SYCLIsDevice) {
     Builder.defineMacro("__SYCL_DEVICE_ONLY__");
-    Builder.defineMacro("SYCL_EXTERNAL", "__attribute__((sycl_device))");
+    if (LangOpts.GPURelocatableDeviceCode)
+      Builder.defineMacro("SYCL_EXTERNAL", "__attribute__((sycl_device))");
 
     const llvm::Triple &DeviceTriple = TI.getTriple();
     const llvm::Triple::SubArchType DeviceSubArch = DeviceTriple.getSubArch();
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 11c2951086f9a..6de1aa761c6d1 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -10294,15 +10294,6 @@ Sema::ActOnFunctionDeclarator(Scope *S, Declarator &D, DeclContext *DC,
     }
   }
 
-  if (getLangOpts().SYCLIsDevice && !getLangOpts().GPURelocatableDeviceCode &&
-      NewFD->hasAttr<SYCLDeviceAttr>() &&
-      !getSourceManager().isInSystemHeader(NewFD->getLocation())) {
-    Diag(NewFD->getLocation(), diag::err_sycl_external_no_rdc)
-        << (D.getFunctionDefinitionKind() ==
-            clang::FunctionDefinitionKind::Definition);
-    NewFD->setInvalidDecl();
-  }
-
   if (!getLangOpts().CPlusPlus) {
     // Perform semantic checking on the function declaration.
     if (!NewFD->isInvalidDecl() && NewFD->isMain())
diff --git a/clang/test/SemaSYCL/sycl-no-rdc.cpp b/clang/test/SemaSYCL/sycl-no-rdc.cpp
index b3466ee8fa396..8d2a4a6413c46 100644
--- a/clang/test/SemaSYCL/sycl-no-rdc.cpp
+++ b/clang/test/SemaSYCL/sycl-no-rdc.cpp
@@ -1,13 +1,10 @@
 // RUN: %clang_cc1 -fsycl-is-device -verify -fsyntax-only -fno-gpu-rdc -internal-isystem %S/Inputs %s
 
-// Check that declarations of SYCL_EXTERNAL functions throw an error if -fno-gpu-rdc is passed
+// Check that uses of SYCL_EXTERNAL throw an error if -fno-gpu-rdc is passed
 #include "sycl.hpp"
 
-// expected-error@+1{{invalid declaration of SYCL_EXTERNAL function in non-relocatable device code mode}}
-SYCL_EXTERNAL void syclExternalDecl();
-
-// expected-error@+1{{invalid definition of SYCL_EXTERNAL function in non-relocatable device code mode}}
-SYCL_EXTERNAL void syclExternalDefn() {}
+// expected-error@+1{{unknown type name 'SYCL_EXTERNAL'}}
+SYCL_EXTERNAL void syclExternal() {}
 
 using namespace sycl;
 queue q;
diff --git a/sycl/include/CL/__spirv/spirv_ops.hpp b/sycl/include/CL/__spirv/spirv_ops.hpp
index aa07bda9f4282..0f841b1707a16 100644
--- a/sycl/include/CL/__spirv/spirv_ops.hpp
+++ b/sycl/include/CL/__spirv/spirv_ops.hpp
@@ -27,7 +27,7 @@
 template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
           __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
           __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *
+extern __DPCPP_SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *
 __spirv_JointMatrixLoadINTEL(T *Ptr, std::size_t Stride,
                              __spv::MatrixLayout Layout = L,
                              __spv::Scope::Flag Sc = S, int MemOperand = 0);
@@ -35,7 +35,7 @@ __spirv_JointMatrixLoadINTEL(T *Ptr, std::size_t Stride,
 template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
           __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
           __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern SYCL_EXTERNAL void __spirv_JointMatrixStoreINTEL(
+extern __DPCPP_SYCL_EXTERNAL void __spirv_JointMatrixStoreINTEL(
     T *Ptr, __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *Object,
     std::size_t Stride, __spv::MatrixLayout Layout = L,
     __spv::Scope::Flag Sc = S, int MemOperand = 0);
@@ -46,12 +46,13 @@ template <typename T1, typename T2, std::size_t M, std::size_t K, std::size_t N,
           __spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
           __spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
           __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T2, M, N, LC, S, UC> *
-__spirv_JointMatrixMadINTEL(
-    __spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S, UA> *A,
-    __spv::__spirv_JointMatrixINTEL<T1, K, N, LB, S, UB> *B,
-    __spv::__spirv_JointMatrixINTEL<T2, M, N, LC, S, UC> *C,
-    __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
+extern __DPCPP_SYCL_EXTERNAL
+    __spv::__spirv_JointMatrixINTEL<T2, M, N, LC, S, UC> *
+    __spirv_JointMatrixMadINTEL(
+        __spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S, UA> *A,
+        __spv::__spirv_JointMatrixINTEL<T1, K, N, LB, S, UB> *B,
+        __spv::__spirv_JointMatrixINTEL<T2, M, N, LC, S, UC> *C,
+        __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
 
 template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
           std::size_t N, __spv::MatrixUse UA, __spv::MatrixUse UB,
@@ -60,12 +61,13 @@ template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
           __spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
           __spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
           __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *
-__spirv_JointMatrixUUMadINTEL(
-    __spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S, UA> *A,
-    __spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S, UB> *B,
-    __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *C,
-    __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
+extern __DPCPP_SYCL_EXTERNAL
+    __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *
+    __spirv_JointMatrixUUMadINTEL(
+        __spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S, UA> *A,
+        __spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S, UB> *B,
+        __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *C,
+        __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
 
 template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
           std::size_t N, __spv::MatrixUse UA, __spv::MatrixUse UB,
@@ -74,12 +76,13 @@ template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
           __spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
           __spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
           __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *
-__spirv_JointMatrixUSMadINTEL(
-    __spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S, UA> *A,
-    __spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S, UB> *B,
-    __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *C,
-    __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
+extern __DPCPP_SYCL_EXTERNAL
+    __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *
+    __spirv_JointMatrixUSMadINTEL(
+        __spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S, UA> *A,
+        __spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S, UB> *B,
+        __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *C,
+        __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
 
 template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
           std::size_t N, __spv::MatrixUse UA, __spv::MatrixUse UB,
@@ -88,42 +91,43 @@ template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
           __spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
           __spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
           __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *
-__spirv_JointMatrixSUMadINTEL(
-    __spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S, UA> *A,
-    __spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S, UB> *B,
-    __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *C,
-    __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
+extern __DPCPP_SYCL_EXTERNAL
+    __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *
+    __spirv_JointMatrixSUMadINTEL(
+        __spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S, UA> *A,
+        __spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S, UB> *B,
+        __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *C,
+        __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
 
 template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
           __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
           __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *
+extern __DPCPP_SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *
 __spirv_CompositeConstruct(const T v);
 
 template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
           __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
           __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern SYCL_EXTERNAL size_t __spirv_JointMatrixWorkItemLengthINTEL(
+extern __DPCPP_SYCL_EXTERNAL size_t __spirv_JointMatrixWorkItemLengthINTEL(
     __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *);
 
 template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
           __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
           __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern SYCL_EXTERNAL T __spirv_VectorExtractDynamic(
+extern __DPCPP_SYCL_EXTERNAL T __spirv_VectorExtractDynamic(
     __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *, size_t i);
 
 template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
           __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
           __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *
+extern __DPCPP_SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *
 __spirv_VectorInsertDynamic(__spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *,
                             T val, size_t i);
 #else
 template <typename T, std::size_t R, std::size_t C,
           __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
           __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *
+extern __DPCPP_SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *
 __spirv_JointMatrixLoadINTEL(T *Ptr, std::size_t Stride,
                              __spv::MatrixLayout Layout = L,
                              __spv::Scope::Flag Sc = S, int MemOperand = 0);
@@ -131,7 +135,7 @@ __spirv_JointMatrixLoadINTEL(T *Ptr, std::size_t Stride,
 template <typename T, std::size_t R, std::size_t C,
           __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
           __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern SYCL_EXTERNAL void __spirv_JointMatrixStoreINTEL(
+extern __DPCPP_SYCL_EXTERNAL void __spirv_JointMatrixStoreINTEL(
     T *Ptr, __spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *Object,
     std::size_t Stride, __spv::MatrixLayout Layout = L,
     __spv::Scope::Flag Sc = S, int MemOperand = 0);
@@ -141,7 +145,7 @@ template <typename T1, typename T2, std::size_t M, std::size_t K, std::size_t N,
           __spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
           __spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
           __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T2, M, N, LC, S> *
+extern __DPCPP_SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T2, M, N, LC, S> *
 __spirv_JointMatrixMadINTEL(
     __spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S> *A,
     __spv::__spirv_JointMatrixINTEL<T1, K, N, LB, S> *B,
@@ -153,7 +157,7 @@ template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
           __spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
           __spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
           __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S> *
+extern __DPCPP_SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S> *
 __spirv_JointMatrixUUMadINTEL(
     __spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S> *A,
     __spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S> *B,
@@ -165,7 +169,7 @@ template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
           __spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
           __spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
           __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S> *
+extern __DPCPP_SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S> *
 __spirv_JointMatrixUSMadINTEL(
     __spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S> *A,
     __spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S> *B,
@@ -177,7 +181,7 @@ template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
           __spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
           __spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
           __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S> *
+extern __DPCPP_SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S> *
 __spirv_JointMatrixSUMadINTEL(
     __spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S> *A,
     __spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S> *B,
@@ -187,25 +191,25 @@ __spirv_JointMatrixSUMadINTEL(
 template <typename T, std::size_t R, std::size_t C,
           __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
           __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *
+extern __DPCPP_SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *
 __spirv_CompositeConstruct(const T v);
 
 template <typename T, std::size_t R, std::size_t C,
           __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
           __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern SYCL_EXTERNAL size_t __spirv_JointMatrixWorkItemLengthINTEL(
+extern __DPCPP_SYCL_EXTERNAL size_t __spirv_JointMatrixWorkItemLengthINTEL(
     __spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *);
 
 template <typename T, std::size_t R, std::size_t C,
           __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
           __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern SYCL_EXTERNAL T __spirv_VectorExtractDynamic(
+extern __DPCPP_SYCL_EXTERNAL T __spirv_VectorExtractDynamic(
     __spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *, size_t i);
 
 template <typename T, std::size_t R, std::size_t C,
           __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
           __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *
+extern __DPCPP_SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *
 __spirv_VectorInsertDynamic(__spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *,
                             T val, size_t i);
 #endif // SYCL_EXT_ONEAPI_MATRIX_VERSION
@@ -216,94 +220,94 @@ __spirv_VectorInsertDynamic(__spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *,
 #endif
 
 template <typename RetT, typename ImageT>
-extern SYCL_EXTERNAL RetT __spirv_ImageQueryFormat(ImageT);
+extern __DPCPP_SYCL_EXTERNAL RetT __spirv_ImageQueryFormat(ImageT);
 
 template <typename RetT, typename ImageT>
-extern SYCL_EXTERNAL RetT __spirv_ImageQueryOrder(ImageT);
+extern __DPCPP_SYCL_EXTERNAL RetT __spirv_ImageQueryOrder(ImageT);
 
 template <typename RetT, typename ImageT>
-extern SYCL_EXTERNAL RetT __spirv_ImageQuerySize(ImageT);
+extern __DPCPP_SYCL_EXTERNAL RetT __spirv_ImageQuerySize(ImageT);
 
 template <typename ImageT, typename CoordT, typename ValT>
-extern SYCL_EXTERNAL void __spirv_ImageWrite(ImageT, CoordT, ValT);
+extern __DPCPP_SYCL_EXTERNAL void __spirv_ImageWrite(ImageT, CoordT, ValT);
 
 template <class RetT, typename ImageT, typename TempArgT>
-extern SYCL_EXTERNAL RetT __spirv_ImageRead(ImageT, TempArgT);
+extern __DPCPP_SYCL_EXTERNAL RetT __spirv_ImageRead(ImageT, TempArgT);
 
 template <typename ImageT, typename SampledType>
-extern SYCL_EXTERNAL SampledType __spirv_SampledImage(ImageT, __ocl_sampler_t);
+extern __DPCPP_SYCL_EXTERNAL SampledType __spirv_SampledImage(ImageT,
+                                                              __ocl_sampler_t);
 
 template <typename SampledType, typename TempRetT, typename TempArgT>
-extern SYCL_EXTERNAL TempRetT __spirv_ImageSampleExplicitLod(SampledType,
-                                                             TempArgT, int,
-                                                             float);
+extern __DPCPP_SYCL_EXTERNAL TempRetT
+__spirv_ImageSampleExplicitLod(SampledType, TempArgT, int, float);
 
 #define __SYCL_OpGroupAsyncCopyGlobalToLocal __spirv_GroupAsyncCopy
 #define __SYCL_OpGroupAsyncCopyLocalToGlobal __spirv_GroupAsyncCopy
 
 // Atomic SPIR-V builtins
 #define __SPIRV_ATOMIC_LOAD(AS, Type)                                          \
-  extern SYCL_EXTERNAL Type __spirv_AtomicLoad(                                \
+  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicLoad(                        \
       AS const Type *P, __spv::Scope::Flag S,                                  \
       __spv::MemorySemanticsMask::Flag O);
 #define __SPIRV_ATOMIC_STORE(AS, Type)                                         \
-  extern SYCL_EXTERNAL void __spirv_AtomicStore(                               \
+  extern __DPCPP_SYCL_EXTERNAL void __spirv_AtomicStore(                       \
       AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O,    \
       Type V);
 #define __SPIRV_ATOMIC_EXCHANGE(AS, Type)                                      \
-  extern SYCL_EXTERNAL Type __spirv_AtomicExchange(                            \
+  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicExchange(                    \
       AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O,    \
       Type V);
 #define __SPIRV_ATOMIC_CMP_EXCHANGE(AS, Type)                                  \
-  extern SYCL_EXTERNAL Type __spirv_AtomicCompareExchange(                     \
+  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicCompareExchange(             \
       AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag E,    \
       __spv::MemorySemanticsMask::Flag U, Type V, Type C);
 #define __SPIRV_ATOMIC_IADD(AS, Type)                                          \
-  extern SYCL_EXTERNAL Type __spirv_AtomicIAdd(                                \
+  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicIAdd(                        \
       AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O,    \
       Type V);
 #define __SPIRV_ATOMIC_ISUB(AS, Type)                                          \
-  extern SYCL_EXTERNAL Type __spirv_AtomicISub(                                \
+  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicISub(                        \
       AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O,    \
       Type V);
 #define __SPIRV_ATOMIC_FADD(AS, Type)                                          \
-  extern SYCL_EXTERNAL Type __spirv_AtomicFAddEXT(                             \
+  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicFAddEXT(                     \
       AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O,    \
       Type V);
 #define __SPIRV_ATOMIC_SMIN(AS, Type)                                          \
-  extern SYCL_EXTERNAL Type __spirv_AtomicSMin(                                \
+  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicSMin(                        \
       AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O,    \
       Type V);
 #define __SPIRV_ATOMIC_UMIN(AS, Type)                                          \
-  extern SYCL_EXTERNAL Type __spirv_AtomicUMin(                                \
+  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicUMin(                        \
       AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O,    \
       Type V);
 #define __SPIRV_ATOMIC_FMIN(AS, Type)                                          \
-  extern SYCL_EXTERNAL Type __spirv_AtomicFMinEXT(                             \
+  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicFMinEXT(                     \
       AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O,    \
       Type V);
 #define __SPIRV_ATOMIC_SMAX(AS, Type)                                          \
-  extern SYCL_EXTERNAL Type __spirv_AtomicSMax(                                \
+  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicSMax(                        \
       AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O,    \
       Type V);
 #define __SPIRV_ATOMIC_UMAX(AS, Type)                                          \
-  extern SYCL_EXTERNAL Type __spirv_AtomicUMax(                                \
+  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicUMax(                        \
       AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O,    \
       Type V);
 #define __SPIRV_ATOMIC_FMAX(AS, Type)                                          \
-  extern SYCL_EXTERNAL Type __spirv_AtomicFMaxEXT(                             \
+  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicFMaxEXT(                     \
       AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O,    \
       Type V);
 #define __SPIRV_ATOMIC_AND(AS, Type)                                           \
-  extern SYCL_EXTERNAL Type __spirv_AtomicAnd(                                 \
+  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicAnd(                         \
       AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O,    \
       Type V);
 #define __SPIRV_ATOMIC_OR(AS, Type)                                            \
-  extern SYCL_EXTERNAL Type __spirv_AtomicOr(                                  \
+  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicOr(                          \
       AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O,    \
       Type V);
 #define __SPIRV_ATOMIC_XOR(AS, Type)                                           \
-  extern SYCL_EXTERNAL Type __spirv_AtomicXor(                                 \
+  extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicXor(                         \
       AS Type *P, __spv::Scope::Flag S, __spv::MemorySemanticsMask::Flag O,    \
       Type V);
 
@@ -496,101 +500,107 @@ __SYCL_GenericCastToPtrExplicit_ToPrivate(const volatile void *Ptr) noexcept {
 }
 
 template <typename dataT>
-__SYCL_CONVERGENT__ extern SYCL_EXTERNAL dataT
+__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
 __spirv_SubgroupShuffleINTEL(dataT Data, uint32_t InvocationId) noexcept;
 template <typename dataT>
-__SYCL_CONVERGENT__ extern SYCL_EXTERNAL dataT __spirv_SubgroupShuffleDownINTEL(
-    dataT Current, dataT Next, uint32_t Delta) noexcept;
+__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
+__spirv_SubgroupShuffleDownINTEL(dataT Current, dataT Next,
+                                 uint32_t Delta) noexcept;
 template <typename dataT>
-__SYCL_CONVERGENT__ extern SYCL_EXTERNAL dataT __spirv_SubgroupShuffleUpINTEL(
-    dataT Previous, dataT Current, uint32_t Delta) noexcept;
+__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
+__spirv_SubgroupShuffleUpINTEL(dataT Previous, dataT Current,
+                               uint32_t Delta) noexcept;
 template <typename dataT>
-__SYCL_CONVERGENT__ extern SYCL_EXTERNAL dataT
+__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
 __spirv_SubgroupShuffleXorINTEL(dataT Data, uint32_t Value) noexcept;
 
 template <typename dataT>
-__SYCL_CONVERGENT__ extern SYCL_EXTERNAL dataT __spirv_SubgroupBlockReadINTEL(
-    const __attribute__((opencl_global)) uint8_t *Ptr) noexcept;
+__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
+__spirv_SubgroupBlockReadINTEL(const __attribute__((opencl_global))
+                               uint8_t *Ptr) noexcept;
 
 template <typename dataT>
-__SYCL_CONVERGENT__ extern SYCL_EXTERNAL void
+__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL void
 __spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_global)) uint8_t *Ptr,
                                 dataT Data) noexcept;
 
 template <typename dataT>
-__SYCL_CONVERGENT__ extern SYCL_EXTERNAL dataT __spirv_SubgroupBlockReadINTEL(
-    const __attribute__((opencl_global)) uint16_t *Ptr) noexcept;
+__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
+__spirv_SubgroupBlockReadINTEL(const __attribute__((opencl_global))
+                               uint16_t *Ptr) noexcept;
 
 template <typename dataT>
-__SYCL_CONVERGENT__ extern SYCL_EXTERNAL void
+__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL void
 __spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_global)) uint16_t *Ptr,
                                 dataT Data) noexcept;
 
 template <typename dataT>
-__SYCL_CONVERGENT__ extern SYCL_EXTERNAL dataT __spirv_SubgroupBlockReadINTEL(
-    const __attribute__((opencl_global)) uint32_t *Ptr) noexcept;
+__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
+__spirv_SubgroupBlockReadINTEL(const __attribute__((opencl_global))
+                               uint32_t *Ptr) noexcept;
 
 template <typename dataT>
-__SYCL_CONVERGENT__ extern SYCL_EXTERNAL void
+__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL void
 __spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_global)) uint32_t *Ptr,
                                 dataT Data) noexcept;
 
 template <typename dataT>
-__SYCL_CONVERGENT__ extern SYCL_EXTERNAL dataT __spirv_SubgroupBlockReadINTEL(
-    const __attribute__((opencl_global)) uint64_t *Ptr) noexcept;
+__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
+__spirv_SubgroupBlockReadINTEL(const __attribute__((opencl_global))
+                               uint64_t *Ptr) noexcept;
 
 template <typename dataT>
-__SYCL_CONVERGENT__ extern SYCL_EXTERNAL void
+__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL void
 __spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_global)) uint64_t *Ptr,
                                 dataT Data) noexcept;
 template <int W, int rW>
-extern SYCL_EXTERNAL sycl::detail::ap_int<rW>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<rW>
 __spirv_FixedSqrtINTEL(sycl::detail::ap_int<W> a, bool S, int32_t I, int32_t rI,
                        int32_t Quantization = 0, int32_t Overflow = 0) noexcept;
 template <int W, int rW>
-extern SYCL_EXTERNAL sycl::detail::ap_int<rW>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<rW>
 __spirv_FixedRecipINTEL(sycl::detail::ap_int<W> a, bool S, int32_t I,
                         int32_t rI, int32_t Quantization = 0,
                         int32_t Overflow = 0) noexcept;
 template <int W, int rW>
-extern SYCL_EXTERNAL sycl::detail::ap_int<rW>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<rW>
 __spirv_FixedRsqrtINTEL(sycl::detail::ap_int<W> a, bool S, int32_t I,
                         int32_t rI, int32_t Quantization = 0,
                         int32_t Overflow = 0) noexcept;
 template <int W, int rW>
-extern SYCL_EXTERNAL sycl::detail::ap_int<rW>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<rW>
 __spirv_FixedSinINTEL(sycl::detail::ap_int<W> a, bool S, int32_t I, int32_t rI,
                       int32_t Quantization = 0, int32_t Overflow = 0) noexcept;
 template <int W, int rW>
-extern SYCL_EXTERNAL sycl::detail::ap_int<rW>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<rW>
 __spirv_FixedCosINTEL(sycl::detail::ap_int<W> a, bool S, int32_t I, int32_t rI,
                       int32_t Quantization = 0, int32_t Overflow = 0) noexcept;
 template <int W, int rW>
-extern SYCL_EXTERNAL sycl::detail::ap_int<2 * rW>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<2 * rW>
 __spirv_FixedSinCosINTEL(sycl::detail::ap_int<W> a, bool S, int32_t I,
                          int32_t rI, int32_t Quantization = 0,
                          int32_t Overflow = 0) noexcept;
 template <int W, int rW>
-extern SYCL_EXTERNAL sycl::detail::ap_int<rW>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<rW>
 __spirv_FixedSinPiINTEL(sycl::detail::ap_int<W> a, bool S, int32_t I,
                         int32_t rI, int32_t Quantization = 0,
                         int32_t Overflow = 0) noexcept;
 template <int W, int rW>
-extern SYCL_EXTERNAL sycl::detail::ap_int<rW>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<rW>
 __spirv_FixedCosPiINTEL(sycl::detail::ap_int<W> a, bool S, int32_t I,
                         int32_t rI, int32_t Quantization = 0,
                         int32_t Overflow = 0) noexcept;
 template <int W, int rW>
-extern SYCL_EXTERNAL sycl::detail::ap_int<2 * rW>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<2 * rW>
 __spirv_FixedSinCosPiINTEL(sycl::detail::ap_int<W> a, bool S, int32_t I,
                            int32_t rI, int32_t Quantization = 0,
                            int32_t Overflow = 0) noexcept;
 template <int W, int rW>
-extern SYCL_EXTERNAL sycl::detail::ap_int<rW>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<rW>
 __spirv_FixedLogINTEL(sycl::detail::ap_int<W> a, bool S, int32_t I, int32_t rI,
                       int32_t Quantization = 0, int32_t Overflow = 0) noexcept;
 template <int W, int rW>
-extern SYCL_EXTERNAL sycl::detail::ap_int<rW>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<rW>
 __spirv_FixedExpINTEL(sycl::detail::ap_int<W> a, bool S, int32_t I, int32_t rI,
                       int32_t Quantization = 0, int32_t Overflow = 0) noexcept;
 
@@ -598,14 +608,14 @@ __spirv_FixedExpINTEL(sycl::detail::ap_int<W> a, bool S, int32_t I, int32_t rI,
 // a floating point variable should be equal to sum of corresponding
 // exponent width E, mantissa width M and 1 for sign bit. I.e. WA = EA + MA + 1.
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatCastINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                 int32_t Mout, int32_t EnableSubnormals = 0,
                                 int32_t RoundingMode = 0,
                                 int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatCastFromIntINTEL(sycl::detail::ap_int<WA> A, int32_t Mout,
                                        bool FromSign = false,
                                        int32_t EnableSubnormals = 0,
@@ -613,7 +623,7 @@ __spirv_ArbitraryFloatCastFromIntINTEL(sycl::detail::ap_int<WA> A, int32_t Mout,
                                        int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatCastToIntINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                      bool ToSign = false,
                                      int32_t EnableSubnormals = 0,
@@ -621,79 +631,87 @@ __spirv_ArbitraryFloatCastToIntINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                      int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int WB, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout> __spirv_ArbitraryFloatAddINTEL(
-    sycl::detail::ap_int<WA> A, int32_t MA, sycl::detail::ap_int<WB> B,
-    int32_t MB, int32_t Mout, int32_t EnableSubnormals = 0,
-    int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept;
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+__spirv_ArbitraryFloatAddINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
+                               sycl::detail::ap_int<WB> B, int32_t MB,
+                               int32_t Mout, int32_t EnableSubnormals = 0,
+                               int32_t RoundingMode = 0,
+                               int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int WB, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout> __spirv_ArbitraryFloatSubINTEL(
-    sycl::detail::ap_int<WA> A, int32_t MA, sycl::detail::ap_int<WB> B,
-    int32_t MB, int32_t Mout, int32_t EnableSubnormals = 0,
-    int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept;
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+__spirv_ArbitraryFloatSubINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
+                               sycl::detail::ap_int<WB> B, int32_t MB,
+                               int32_t Mout, int32_t EnableSubnormals = 0,
+                               int32_t RoundingMode = 0,
+                               int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int WB, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout> __spirv_ArbitraryFloatMulINTEL(
-    sycl::detail::ap_int<WA> A, int32_t MA, sycl::detail::ap_int<WB> B,
-    int32_t MB, int32_t Mout, int32_t EnableSubnormals = 0,
-    int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept;
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+__spirv_ArbitraryFloatMulINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
+                               sycl::detail::ap_int<WB> B, int32_t MB,
+                               int32_t Mout, int32_t EnableSubnormals = 0,
+                               int32_t RoundingMode = 0,
+                               int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int WB, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout> __spirv_ArbitraryFloatDivINTEL(
-    sycl::detail::ap_int<WA> A, int32_t MA, sycl::detail::ap_int<WB> B,
-    int32_t MB, int32_t Mout, int32_t EnableSubnormals = 0,
-    int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept;
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+__spirv_ArbitraryFloatDivINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
+                               sycl::detail::ap_int<WB> B, int32_t MB,
+                               int32_t Mout, int32_t EnableSubnormals = 0,
+                               int32_t RoundingMode = 0,
+                               int32_t RoundingAccuracy = 0) noexcept;
 
 // Comparison built-ins don't use Subnormal Support, Rounding Mode and
 // Rounding Accuracy.
 template <int WA, int WB>
-extern SYCL_EXTERNAL bool
+extern __DPCPP_SYCL_EXTERNAL bool
 __spirv_ArbitraryFloatGTINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                               sycl::detail::ap_int<WB> B, int32_t MB) noexcept;
 
 template <int WA, int WB>
-extern SYCL_EXTERNAL bool
+extern __DPCPP_SYCL_EXTERNAL bool
 __spirv_ArbitraryFloatGEINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                               sycl::detail::ap_int<WB> B, int32_t MB) noexcept;
 
 template <int WA, int WB>
-extern SYCL_EXTERNAL bool
+extern __DPCPP_SYCL_EXTERNAL bool
 __spirv_ArbitraryFloatLTINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                               sycl::detail::ap_int<WB> B, int32_t MB) noexcept;
 
 template <int WA, int WB>
-extern SYCL_EXTERNAL bool
+extern __DPCPP_SYCL_EXTERNAL bool
 __spirv_ArbitraryFloatLEINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                               sycl::detail::ap_int<WB> B, int32_t MB) noexcept;
 
 template <int WA, int WB>
-extern SYCL_EXTERNAL bool
+extern __DPCPP_SYCL_EXTERNAL bool
 __spirv_ArbitraryFloatEQINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                               sycl::detail::ap_int<WB> B, int32_t MB) noexcept;
 
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatRecipINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                  int32_t Mout, int32_t EnableSubnormals = 0,
                                  int32_t RoundingMode = 0,
                                  int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatRSqrtINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                  int32_t Mout, int32_t EnableSubnormals = 0,
                                  int32_t RoundingMode = 0,
                                  int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatCbrtINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                 int32_t Mout, int32_t EnableSubnormals = 0,
                                 int32_t RoundingMode = 0,
                                 int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int WB, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatHypotINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                  sycl::detail::ap_int<WB> B, int32_t MB,
                                  int32_t Mout, int32_t EnableSubnormals = 0,
@@ -701,77 +719,77 @@ __spirv_ArbitraryFloatHypotINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                  int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatSqrtINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                 int32_t Mout, int32_t EnableSubnormals = 0,
                                 int32_t RoundingMode = 0,
                                 int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatLogINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                int32_t Mout, int32_t EnableSubnormals = 0,
                                int32_t RoundingMode = 0,
                                int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatLog2INTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                 int32_t Mout, int32_t EnableSubnormals = 0,
                                 int32_t RoundingMode = 0,
                                 int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatLog10INTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                  int32_t Mout, int32_t EnableSubnormals = 0,
                                  int32_t RoundingMode = 0,
                                  int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatLog1pINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                  int32_t Mout, int32_t EnableSubnormals = 0,
                                  int32_t RoundingMode = 0,
                                  int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatExpINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                int32_t Mout, int32_t EnableSubnormals = 0,
                                int32_t RoundingMode = 0,
                                int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatExp2INTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                 int32_t Mout, int32_t EnableSubnormals = 0,
                                 int32_t RoundingMode = 0,
                                 int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatExp10INTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                  int32_t Mout, int32_t EnableSubnormals = 0,
                                  int32_t RoundingMode = 0,
                                  int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatExpm1INTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                  int32_t Mout, int32_t EnableSubnormals = 0,
                                  int32_t RoundingMode = 0,
                                  int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatSinINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                int32_t Mout, int32_t EnableSubnormals = 0,
                                int32_t RoundingMode = 0,
                                int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatCosINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                int32_t Mout, int32_t EnableSubnormals = 0,
                                int32_t RoundingMode = 0,
@@ -780,21 +798,21 @@ __spirv_ArbitraryFloatCosINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
 // Result value contains both values of sine and cosine and so has the size of
 // 2 * Wout where Wout is equal to (1 + Eout + Mout).
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<2 * Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<2 * Wout>
 __spirv_ArbitraryFloatSinCosINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                   int32_t Mout, int32_t EnableSubnormals = 0,
                                   int32_t RoundingMode = 0,
                                   int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatSinPiINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                  int32_t Mout, int32_t EnableSubnormals = 0,
                                  int32_t RoundingMode = 0,
                                  int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatCosPiINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                  int32_t Mout, int32_t EnableSubnormals = 0,
                                  int32_t RoundingMode = 0,
@@ -803,56 +821,56 @@ __spirv_ArbitraryFloatCosPiINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
 // Result value contains both values of sine(A*pi) and cosine(A*pi) and so has
 // the size of 2 * Wout where Wout is equal to (1 + Eout + Mout).
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<2 * Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<2 * Wout>
 __spirv_ArbitraryFloatSinCosPiINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                     int32_t Mout, int32_t EnableSubnormals = 0,
                                     int32_t RoundingMode = 0,
                                     int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatASinINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                 int32_t Mout, int32_t EnableSubnormals = 0,
                                 int32_t RoundingMode = 0,
                                 int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatASinPiINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                   int32_t Mout, int32_t EnableSubnormals = 0,
                                   int32_t RoundingMode = 0,
                                   int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatACosINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                 int32_t Mout, int32_t EnableSubnormals = 0,
                                 int32_t RoundingMode = 0,
                                 int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatACosPiINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                   int32_t Mout, int32_t EnableSubnormals = 0,
                                   int32_t RoundingMode = 0,
                                   int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatATanINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                 int32_t Mout, int32_t EnableSubnormals = 0,
                                 int32_t RoundingMode = 0,
                                 int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatATanPiINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                   int32_t Mout, int32_t EnableSubnormals = 0,
                                   int32_t RoundingMode = 0,
                                   int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int WB, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatATan2INTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                  sycl::detail::ap_int<WB> B, int32_t MB,
                                  int32_t Mout, int32_t EnableSubnormals = 0,
@@ -860,153 +878,163 @@ __spirv_ArbitraryFloatATan2INTEL(sycl::detail::ap_int<WA> A, int32_t MA,
                                  int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int WB, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout> __spirv_ArbitraryFloatPowINTEL(
-    sycl::detail::ap_int<WA> A, int32_t MA, sycl::detail::ap_int<WB> B,
-    int32_t MB, int32_t Mout, int32_t EnableSubnormals = 0,
-    int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept;
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+__spirv_ArbitraryFloatPowINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
+                               sycl::detail::ap_int<WB> B, int32_t MB,
+                               int32_t Mout, int32_t EnableSubnormals = 0,
+                               int32_t RoundingMode = 0,
+                               int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int WB, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout> __spirv_ArbitraryFloatPowRINTEL(
-    sycl::detail::ap_int<WA> A, int32_t MA, sycl::detail::ap_int<WB> B,
-    int32_t MB, int32_t Mout, int32_t EnableSubnormals = 0,
-    int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept;
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+__spirv_ArbitraryFloatPowRINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
+                                sycl::detail::ap_int<WB> B, int32_t MB,
+                                int32_t Mout, int32_t EnableSubnormals = 0,
+                                int32_t RoundingMode = 0,
+                                int32_t RoundingAccuracy = 0) noexcept;
 
 // PowN built-in calculates `A^B` where `A` is arbitrary precision floating
 // point number and `B` is signed or unsigned arbitrary precision integer,
 // i.e. its width doesn't depend on sum of exponent and mantissa.
 template <int WA, int WB, int Wout>
-extern SYCL_EXTERNAL sycl::detail::ap_int<Wout> __spirv_ArbitraryFloatPowNINTEL(
-    sycl::detail::ap_int<WA> A, int32_t MA, sycl::detail::ap_int<WB> B,
-    bool SignOfB, int32_t Mout, int32_t EnableSubnormals = 0,
-    int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept;
+extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<Wout>
+__spirv_ArbitraryFloatPowNINTEL(sycl::detail::ap_int<WA> A, int32_t MA,
+                                sycl::detail::ap_int<WB> B, bool SignOfB,
+                                int32_t Mout, int32_t EnableSubnormals = 0,
+                                int32_t RoundingMode = 0,
+                                int32_t RoundingAccuracy = 0) noexcept;
 
 template <typename dataT>
-extern SYCL_EXTERNAL int32_t __spirv_ReadPipe(__ocl_RPipeTy<dataT> Pipe,
-                                              dataT *Data, int32_t Size,
-                                              int32_t Alignment) noexcept;
+extern __DPCPP_SYCL_EXTERNAL int32_t
+__spirv_ReadPipe(__ocl_RPipeTy<dataT> Pipe, dataT *Data, int32_t Size,
+                 int32_t Alignment) noexcept;
 template <typename dataT>
-extern SYCL_EXTERNAL int32_t __spirv_WritePipe(__ocl_WPipeTy<dataT> Pipe,
-                                               const dataT *Data, int32_t Size,
-                                               int32_t Alignment) noexcept;
+extern __DPCPP_SYCL_EXTERNAL int32_t
+__spirv_WritePipe(__ocl_WPipeTy<dataT> Pipe, const dataT *Data, int32_t Size,
+                  int32_t Alignment) noexcept;
 template <typename dataT>
-extern SYCL_EXTERNAL void
+extern __DPCPP_SYCL_EXTERNAL void
 __spirv_ReadPipeBlockingINTEL(__ocl_RPipeTy<dataT> Pipe, dataT *Data,
                               int32_t Size, int32_t Alignment) noexcept;
 template <typename dataT>
-extern SYCL_EXTERNAL void
+extern __DPCPP_SYCL_EXTERNAL void
 __spirv_WritePipeBlockingINTEL(__ocl_WPipeTy<dataT> Pipe, const dataT *Data,
                                int32_t Size, int32_t Alignment) noexcept;
 template <typename dataT>
-extern SYCL_EXTERNAL __ocl_RPipeTy<dataT>
+extern __DPCPP_SYCL_EXTERNAL __ocl_RPipeTy<dataT>
 __spirv_CreatePipeFromPipeStorage_read(
     const ConstantPipeStorage *Storage) noexcept;
 template <typename dataT>
-extern SYCL_EXTERNAL __ocl_WPipeTy<dataT>
+extern __DPCPP_SYCL_EXTERNAL __ocl_WPipeTy<dataT>
 __spirv_CreatePipeFromPipeStorage_write(
     const ConstantPipeStorage *Storage) noexcept;
 
-extern SYCL_EXTERNAL void
+extern __DPCPP_SYCL_EXTERNAL void
 __spirv_ocl_prefetch(const __attribute__((opencl_global)) char *Ptr,
                      size_t NumBytes) noexcept;
 
-extern SYCL_EXTERNAL uint16_t __spirv_ConvertFToBF16INTEL(float) noexcept;
-extern SYCL_EXTERNAL float __spirv_ConvertBF16ToFINTEL(uint16_t) noexcept;
+extern __DPCPP_SYCL_EXTERNAL uint16_t
+__spirv_ConvertFToBF16INTEL(float) noexcept;
+extern __DPCPP_SYCL_EXTERNAL float
+    __spirv_ConvertBF16ToFINTEL(uint16_t) noexcept;
 
-__SYCL_CONVERGENT__ extern SYCL_EXTERNAL __SYCL_EXPORT __ocl_vec_t<uint32_t, 4>
-__spirv_GroupNonUniformBallot(uint32_t Execution, bool Predicate) noexcept;
+__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL
+    __SYCL_EXPORT __ocl_vec_t<uint32_t, 4>
+    __spirv_GroupNonUniformBallot(uint32_t Execution, bool Predicate) noexcept;
 
-extern SYCL_EXTERNAL __SYCL_EXPORT void
+extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void
 __clc_BarrierInitialize(int64_t *state, int32_t expected_count) noexcept;
 
-extern SYCL_EXTERNAL __SYCL_EXPORT void
+extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void
 __clc_BarrierInvalidate(int64_t *state) noexcept;
 
-extern SYCL_EXTERNAL __SYCL_EXPORT int64_t
+extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT int64_t
 __clc_BarrierArrive(int64_t *state) noexcept;
 
-extern SYCL_EXTERNAL __SYCL_EXPORT int64_t
+extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT int64_t
 __clc_BarrierArriveAndDrop(int64_t *state) noexcept;
 
-extern SYCL_EXTERNAL __SYCL_EXPORT int64_t
+extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT int64_t
 __clc_BarrierArriveNoComplete(int64_t *state, int32_t count) noexcept;
 
-extern SYCL_EXTERNAL __SYCL_EXPORT int64_t
+extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT int64_t
 __clc_BarrierArriveAndDropNoComplete(int64_t *state, int32_t count) noexcept;
 
-extern SYCL_EXTERNAL __SYCL_EXPORT void
+extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void
 __clc_BarrierCopyAsyncArrive(int64_t *state) noexcept;
 
-extern SYCL_EXTERNAL __SYCL_EXPORT void
+extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void
 __clc_BarrierCopyAsyncArriveNoInc(int64_t *state) noexcept;
 
-__SYCL_CONVERGENT__ extern SYCL_EXTERNAL __SYCL_EXPORT void
+__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void
 __clc_BarrierWait(int64_t *state, int64_t arrival) noexcept;
 
-extern SYCL_EXTERNAL __SYCL_EXPORT bool
+extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT bool
 __clc_BarrierTestWait(int64_t *state, int64_t arrival) noexcept;
 
-__SYCL_CONVERGENT__ extern SYCL_EXTERNAL __SYCL_EXPORT void
+__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void
 __clc_BarrierArriveAndWait(int64_t *state) noexcept;
 
 #ifdef __SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__
 template <typename... Args>
-extern SYCL_EXTERNAL int
+extern __DPCPP_SYCL_EXTERNAL int
 __spirv_ocl_printf(const __attribute__((opencl_constant)) char *Format,
                    Args... args);
 template <typename... Args>
-extern SYCL_EXTERNAL int __spirv_ocl_printf(const char *Format, Args... args);
+extern __DPCPP_SYCL_EXTERNAL int __spirv_ocl_printf(const char *Format,
+                                                    Args... args);
 #else
-extern SYCL_EXTERNAL int
+extern __DPCPP_SYCL_EXTERNAL int
 __spirv_ocl_printf(const __attribute__((opencl_constant)) char *Format, ...);
-extern SYCL_EXTERNAL int __spirv_ocl_printf(const char *Format, ...);
+extern __DPCPP_SYCL_EXTERNAL int __spirv_ocl_printf(const char *Format, ...);
 #endif
 
 // Native builtin extension
 
-extern SYCL_EXTERNAL float __clc_native_tanh(float);
-extern SYCL_EXTERNAL __ocl_vec_t<float, 2>
+extern __DPCPP_SYCL_EXTERNAL float __clc_native_tanh(float);
+extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<float, 2>
     __clc_native_tanh(__ocl_vec_t<float, 2>);
-extern SYCL_EXTERNAL __ocl_vec_t<float, 3>
+extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<float, 3>
     __clc_native_tanh(__ocl_vec_t<float, 3>);
-extern SYCL_EXTERNAL __ocl_vec_t<float, 4>
+extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<float, 4>
     __clc_native_tanh(__ocl_vec_t<float, 4>);
-extern SYCL_EXTERNAL __ocl_vec_t<float, 8>
+extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<float, 8>
     __clc_native_tanh(__ocl_vec_t<float, 8>);
-extern SYCL_EXTERNAL __ocl_vec_t<float, 16>
+extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<float, 16>
     __clc_native_tanh(__ocl_vec_t<float, 16>);
 
-extern SYCL_EXTERNAL _Float16 __clc_native_tanh(_Float16);
-extern SYCL_EXTERNAL __ocl_vec_t<_Float16, 2>
+extern __DPCPP_SYCL_EXTERNAL _Float16 __clc_native_tanh(_Float16);
+extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<_Float16, 2>
     __clc_native_tanh(__ocl_vec_t<_Float16, 2>);
-extern SYCL_EXTERNAL __ocl_vec_t<_Float16, 3>
+extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<_Float16, 3>
     __clc_native_tanh(__ocl_vec_t<_Float16, 3>);
-extern SYCL_EXTERNAL __ocl_vec_t<_Float16, 4>
+extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<_Float16, 4>
     __clc_native_tanh(__ocl_vec_t<_Float16, 4>);
-extern SYCL_EXTERNAL __ocl_vec_t<_Float16, 8>
+extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<_Float16, 8>
     __clc_native_tanh(__ocl_vec_t<_Float16, 8>);
-extern SYCL_EXTERNAL __ocl_vec_t<_Float16, 16>
+extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<_Float16, 16>
     __clc_native_tanh(__ocl_vec_t<_Float16, 16>);
 
-extern SYCL_EXTERNAL _Float16 __clc_native_exp2(_Float16);
-extern SYCL_EXTERNAL __ocl_vec_t<_Float16, 2>
+extern __DPCPP_SYCL_EXTERNAL _Float16 __clc_native_exp2(_Float16);
+extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<_Float16, 2>
     __clc_native_exp2(__ocl_vec_t<_Float16, 2>);
-extern SYCL_EXTERNAL __ocl_vec_t<_Float16, 3>
+extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<_Float16, 3>
     __clc_native_exp2(__ocl_vec_t<_Float16, 3>);
-extern SYCL_EXTERNAL __ocl_vec_t<_Float16, 4>
+extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<_Float16, 4>
     __clc_native_exp2(__ocl_vec_t<_Float16, 4>);
-extern SYCL_EXTERNAL __ocl_vec_t<_Float16, 8>
+extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<_Float16, 8>
     __clc_native_exp2(__ocl_vec_t<_Float16, 8>);
-extern SYCL_EXTERNAL __ocl_vec_t<_Float16, 16>
+extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<_Float16, 16>
     __clc_native_exp2(__ocl_vec_t<_Float16, 16>);
 
 #define __CLC_BF16(...)                                                        \
-  extern SYCL_EXTERNAL __SYCL_EXPORT __VA_ARGS__ __clc_fabs(                   \
+  extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT __VA_ARGS__ __clc_fabs(           \
       __VA_ARGS__) noexcept;                                                   \
-  extern SYCL_EXTERNAL __SYCL_EXPORT __VA_ARGS__ __clc_fmin(                   \
+  extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT __VA_ARGS__ __clc_fmin(           \
       __VA_ARGS__, __VA_ARGS__) noexcept;                                      \
-  extern SYCL_EXTERNAL __SYCL_EXPORT __VA_ARGS__ __clc_fmax(                   \
+  extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT __VA_ARGS__ __clc_fmax(           \
       __VA_ARGS__, __VA_ARGS__) noexcept;                                      \
-  extern SYCL_EXTERNAL __SYCL_EXPORT __VA_ARGS__ __clc_fma(                    \
+  extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT __VA_ARGS__ __clc_fma(            \
       __VA_ARGS__, __VA_ARGS__, __VA_ARGS__) noexcept;
 
 #define __CLC_BF16_SCAL_VEC(TYPE)                                              \
@@ -1052,14 +1080,14 @@ __SYCL_OpGroupAsyncCopyLocalToGlobal(__spv::Scope::Flag, dataT *Dest,
 extern __SYCL_EXPORT void __spirv_ocl_prefetch(const char *Ptr,
                                                size_t NumBytes) noexcept;
 
-__SYCL_CONVERGENT__ extern SYCL_EXTERNAL __SYCL_EXPORT void
+__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void
 __spirv_ControlBarrier(__spv::Scope Execution, __spv::Scope Memory,
                        uint32_t Semantics) noexcept;
 
-__SYCL_CONVERGENT__ extern SYCL_EXTERNAL __SYCL_EXPORT void
+__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void
 __spirv_MemoryBarrier(__spv::Scope Memory, uint32_t Semantics) noexcept;
 
-__SYCL_CONVERGENT__ extern SYCL_EXTERNAL __SYCL_EXPORT void
+__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void
 __spirv_GroupWaitEvents(__spv::Scope Execution, uint32_t NumEvents,
                         __ocl_event_t *WaitEvents) noexcept;
 
diff --git a/sycl/include/CL/__spirv/spirv_vars.hpp b/sycl/include/CL/__spirv/spirv_vars.hpp
index 9aea1192eb37e..7aabb6c056235 100644
--- a/sycl/include/CL/__spirv/spirv_vars.hpp
+++ b/sycl/include/CL/__spirv/spirv_vars.hpp
@@ -17,39 +17,39 @@
 
 #if defined(__NVPTX__) || defined(__AMDGCN__)
 
-SYCL_EXTERNAL size_t __spirv_GlobalInvocationId_x();
-SYCL_EXTERNAL size_t __spirv_GlobalInvocationId_y();
-SYCL_EXTERNAL size_t __spirv_GlobalInvocationId_z();
+__DPCPP_SYCL_EXTERNAL size_t __spirv_GlobalInvocationId_x();
+__DPCPP_SYCL_EXTERNAL size_t __spirv_GlobalInvocationId_y();
+__DPCPP_SYCL_EXTERNAL size_t __spirv_GlobalInvocationId_z();
 
-SYCL_EXTERNAL size_t __spirv_GlobalSize_x();
-SYCL_EXTERNAL size_t __spirv_GlobalSize_y();
-SYCL_EXTERNAL size_t __spirv_GlobalSize_z();
+__DPCPP_SYCL_EXTERNAL size_t __spirv_GlobalSize_x();
+__DPCPP_SYCL_EXTERNAL size_t __spirv_GlobalSize_y();
+__DPCPP_SYCL_EXTERNAL size_t __spirv_GlobalSize_z();
 
-SYCL_EXTERNAL size_t __spirv_GlobalOffset_x();
-SYCL_EXTERNAL size_t __spirv_GlobalOffset_y();
-SYCL_EXTERNAL size_t __spirv_GlobalOffset_z();
+__DPCPP_SYCL_EXTERNAL size_t __spirv_GlobalOffset_x();
+__DPCPP_SYCL_EXTERNAL size_t __spirv_GlobalOffset_y();
+__DPCPP_SYCL_EXTERNAL size_t __spirv_GlobalOffset_z();
 
-SYCL_EXTERNAL size_t __spirv_NumWorkgroups_x();
-SYCL_EXTERNAL size_t __spirv_NumWorkgroups_y();
-SYCL_EXTERNAL size_t __spirv_NumWorkgroups_z();
+__DPCPP_SYCL_EXTERNAL size_t __spirv_NumWorkgroups_x();
+__DPCPP_SYCL_EXTERNAL size_t __spirv_NumWorkgroups_y();
+__DPCPP_SYCL_EXTERNAL size_t __spirv_NumWorkgroups_z();
 
-SYCL_EXTERNAL size_t __spirv_WorkgroupSize_x();
-SYCL_EXTERNAL size_t __spirv_WorkgroupSize_y();
-SYCL_EXTERNAL size_t __spirv_WorkgroupSize_z();
+__DPCPP_SYCL_EXTERNAL size_t __spirv_WorkgroupSize_x();
+__DPCPP_SYCL_EXTERNAL size_t __spirv_WorkgroupSize_y();
+__DPCPP_SYCL_EXTERNAL size_t __spirv_WorkgroupSize_z();
 
-SYCL_EXTERNAL size_t __spirv_WorkgroupId_x();
-SYCL_EXTERNAL size_t __spirv_WorkgroupId_y();
-SYCL_EXTERNAL size_t __spirv_WorkgroupId_z();
+__DPCPP_SYCL_EXTERNAL size_t __spirv_WorkgroupId_x();
+__DPCPP_SYCL_EXTERNAL size_t __spirv_WorkgroupId_y();
+__DPCPP_SYCL_EXTERNAL size_t __spirv_WorkgroupId_z();
 
-SYCL_EXTERNAL size_t __spirv_LocalInvocationId_x();
-SYCL_EXTERNAL size_t __spirv_LocalInvocationId_y();
-SYCL_EXTERNAL size_t __spirv_LocalInvocationId_z();
+__DPCPP_SYCL_EXTERNAL size_t __spirv_LocalInvocationId_x();
+__DPCPP_SYCL_EXTERNAL size_t __spirv_LocalInvocationId_y();
+__DPCPP_SYCL_EXTERNAL size_t __spirv_LocalInvocationId_z();
 
-SYCL_EXTERNAL uint32_t __spirv_SubgroupSize();
-SYCL_EXTERNAL uint32_t __spirv_SubgroupMaxSize();
-SYCL_EXTERNAL uint32_t __spirv_NumSubgroups();
-SYCL_EXTERNAL uint32_t __spirv_SubgroupId();
-SYCL_EXTERNAL uint32_t __spirv_SubgroupLocalInvocationId();
+__DPCPP_SYCL_EXTERNAL uint32_t __spirv_SubgroupSize();
+__DPCPP_SYCL_EXTERNAL uint32_t __spirv_SubgroupMaxSize();
+__DPCPP_SYCL_EXTERNAL uint32_t __spirv_NumSubgroups();
+__DPCPP_SYCL_EXTERNAL uint32_t __spirv_SubgroupId();
+__DPCPP_SYCL_EXTERNAL uint32_t __spirv_SubgroupLocalInvocationId();
 
 #else // defined(__NVPTX__) || defined(__AMDGCN__)
 
@@ -68,89 +68,89 @@ __SPIRV_VAR_QUALIFIERS uint32_t __spirv_BuiltInNumSubgroups;
 __SPIRV_VAR_QUALIFIERS uint32_t __spirv_BuiltInSubgroupId;
 __SPIRV_VAR_QUALIFIERS uint32_t __spirv_BuiltInSubgroupLocalInvocationId;
 
-SYCL_EXTERNAL inline size_t __spirv_GlobalInvocationId_x() {
+__DPCPP_SYCL_EXTERNAL inline size_t __spirv_GlobalInvocationId_x() {
   return __spirv_BuiltInGlobalInvocationId.x;
 }
-SYCL_EXTERNAL inline size_t __spirv_GlobalInvocationId_y() {
+__DPCPP_SYCL_EXTERNAL inline size_t __spirv_GlobalInvocationId_y() {
   return __spirv_BuiltInGlobalInvocationId.y;
 }
-SYCL_EXTERNAL inline size_t __spirv_GlobalInvocationId_z() {
+__DPCPP_SYCL_EXTERNAL inline size_t __spirv_GlobalInvocationId_z() {
   return __spirv_BuiltInGlobalInvocationId.z;
 }
 
-SYCL_EXTERNAL inline size_t __spirv_GlobalSize_x() {
+__DPCPP_SYCL_EXTERNAL inline size_t __spirv_GlobalSize_x() {
   return __spirv_BuiltInGlobalSize.x;
 }
-SYCL_EXTERNAL inline size_t __spirv_GlobalSize_y() {
+__DPCPP_SYCL_EXTERNAL inline size_t __spirv_GlobalSize_y() {
   return __spirv_BuiltInGlobalSize.y;
 }
-SYCL_EXTERNAL inline size_t __spirv_GlobalSize_z() {
+__DPCPP_SYCL_EXTERNAL inline size_t __spirv_GlobalSize_z() {
   return __spirv_BuiltInGlobalSize.z;
 }
 
-SYCL_EXTERNAL inline size_t __spirv_GlobalOffset_x() {
+__DPCPP_SYCL_EXTERNAL inline size_t __spirv_GlobalOffset_x() {
   return __spirv_BuiltInGlobalOffset.x;
 }
-SYCL_EXTERNAL inline size_t __spirv_GlobalOffset_y() {
+__DPCPP_SYCL_EXTERNAL inline size_t __spirv_GlobalOffset_y() {
   return __spirv_BuiltInGlobalOffset.y;
 }
-SYCL_EXTERNAL inline size_t __spirv_GlobalOffset_z() {
+__DPCPP_SYCL_EXTERNAL inline size_t __spirv_GlobalOffset_z() {
   return __spirv_BuiltInGlobalOffset.z;
 }
 
-SYCL_EXTERNAL inline size_t __spirv_NumWorkgroups_x() {
+__DPCPP_SYCL_EXTERNAL inline size_t __spirv_NumWorkgroups_x() {
   return __spirv_BuiltInNumWorkgroups.x;
 }
-SYCL_EXTERNAL inline size_t __spirv_NumWorkgroups_y() {
+__DPCPP_SYCL_EXTERNAL inline size_t __spirv_NumWorkgroups_y() {
   return __spirv_BuiltInNumWorkgroups.y;
 }
-SYCL_EXTERNAL inline size_t __spirv_NumWorkgroups_z() {
+__DPCPP_SYCL_EXTERNAL inline size_t __spirv_NumWorkgroups_z() {
   return __spirv_BuiltInNumWorkgroups.z;
 }
 
-SYCL_EXTERNAL inline size_t __spirv_WorkgroupSize_x() {
+__DPCPP_SYCL_EXTERNAL inline size_t __spirv_WorkgroupSize_x() {
   return __spirv_BuiltInWorkgroupSize.x;
 }
-SYCL_EXTERNAL inline size_t __spirv_WorkgroupSize_y() {
+__DPCPP_SYCL_EXTERNAL inline size_t __spirv_WorkgroupSize_y() {
   return __spirv_BuiltInWorkgroupSize.y;
 }
-SYCL_EXTERNAL inline size_t __spirv_WorkgroupSize_z() {
+__DPCPP_SYCL_EXTERNAL inline size_t __spirv_WorkgroupSize_z() {
   return __spirv_BuiltInWorkgroupSize.z;
 }
 
-SYCL_EXTERNAL inline size_t __spirv_WorkgroupId_x() {
+__DPCPP_SYCL_EXTERNAL inline size_t __spirv_WorkgroupId_x() {
   return __spirv_BuiltInWorkgroupId.x;
 }
-SYCL_EXTERNAL inline size_t __spirv_WorkgroupId_y() {
+__DPCPP_SYCL_EXTERNAL inline size_t __spirv_WorkgroupId_y() {
   return __spirv_BuiltInWorkgroupId.y;
 }
-SYCL_EXTERNAL inline size_t __spirv_WorkgroupId_z() {
+__DPCPP_SYCL_EXTERNAL inline size_t __spirv_WorkgroupId_z() {
   return __spirv_BuiltInWorkgroupId.z;
 }
 
-SYCL_EXTERNAL inline size_t __spirv_LocalInvocationId_x() {
+__DPCPP_SYCL_EXTERNAL inline size_t __spirv_LocalInvocationId_x() {
   return __spirv_BuiltInLocalInvocationId.x;
 }
-SYCL_EXTERNAL inline size_t __spirv_LocalInvocationId_y() {
+__DPCPP_SYCL_EXTERNAL inline size_t __spirv_LocalInvocationId_y() {
   return __spirv_BuiltInLocalInvocationId.y;
 }
-SYCL_EXTERNAL inline size_t __spirv_LocalInvocationId_z() {
+__DPCPP_SYCL_EXTERNAL inline size_t __spirv_LocalInvocationId_z() {
   return __spirv_BuiltInLocalInvocationId.z;
 }
 
-SYCL_EXTERNAL inline uint32_t __spirv_SubgroupSize() {
+__DPCPP_SYCL_EXTERNAL inline uint32_t __spirv_SubgroupSize() {
   return __spirv_BuiltInSubgroupSize;
 }
-SYCL_EXTERNAL inline uint32_t __spirv_SubgroupMaxSize() {
+__DPCPP_SYCL_EXTERNAL inline uint32_t __spirv_SubgroupMaxSize() {
   return __spirv_BuiltInSubgroupMaxSize;
 }
-SYCL_EXTERNAL inline uint32_t __spirv_NumSubgroups() {
+__DPCPP_SYCL_EXTERNAL inline uint32_t __spirv_NumSubgroups() {
   return __spirv_BuiltInNumSubgroups;
 }
-SYCL_EXTERNAL inline uint32_t __spirv_SubgroupId() {
+__DPCPP_SYCL_EXTERNAL inline uint32_t __spirv_SubgroupId() {
   return __spirv_BuiltInSubgroupId;
 }
-SYCL_EXTERNAL inline uint32_t __spirv_SubgroupLocalInvocationId() {
+__DPCPP_SYCL_EXTERNAL inline uint32_t __spirv_SubgroupLocalInvocationId() {
   return __spirv_BuiltInSubgroupLocalInvocationId;
 }
 
diff --git a/sycl/include/sycl/builtins.hpp b/sycl/include/sycl/builtins.hpp
index 7dde0dcdea5b3..fe6eeaaec2e4e 100644
--- a/sycl/include/sycl/builtins.hpp
+++ b/sycl/include/sycl/builtins.hpp
@@ -1953,556 +1953,635 @@ detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> tan(T x) __NOEXC {
 
 #ifdef __SYCL_DEVICE_ONLY__
 extern "C" {
-extern SYCL_EXTERNAL int abs(int x);
-extern SYCL_EXTERNAL long int labs(long int x);
-extern SYCL_EXTERNAL long long int llabs(long long int x);
-
-extern SYCL_EXTERNAL div_t div(int x, int y);
-extern SYCL_EXTERNAL ldiv_t ldiv(long int x, long int y);
-extern SYCL_EXTERNAL lldiv_t lldiv(long long int x, long long int y);
-extern SYCL_EXTERNAL float scalbnf(float x, int n);
-extern SYCL_EXTERNAL double scalbn(double x, int n);
-extern SYCL_EXTERNAL float logf(float x);
-extern SYCL_EXTERNAL double log(double x);
-extern SYCL_EXTERNAL float expf(float x);
-extern SYCL_EXTERNAL double exp(double x);
-extern SYCL_EXTERNAL float log10f(float x);
-extern SYCL_EXTERNAL double log10(double x);
-extern SYCL_EXTERNAL float modff(float x, float *intpart);
-extern SYCL_EXTERNAL double modf(double x, double *intpart);
-extern SYCL_EXTERNAL float exp2f(float x);
-extern SYCL_EXTERNAL double exp2(double x);
-extern SYCL_EXTERNAL float expm1f(float x);
-extern SYCL_EXTERNAL double expm1(double x);
-extern SYCL_EXTERNAL int ilogbf(float x);
-extern SYCL_EXTERNAL int ilogb(double x);
-extern SYCL_EXTERNAL float log1pf(float x);
-extern SYCL_EXTERNAL double log1p(double x);
-extern SYCL_EXTERNAL float log2f(float x);
-extern SYCL_EXTERNAL double log2(double x);
-extern SYCL_EXTERNAL float logbf(float x);
-extern SYCL_EXTERNAL double logb(double x);
-extern SYCL_EXTERNAL float sqrtf(float x);
-extern SYCL_EXTERNAL double sqrt(double x);
-extern SYCL_EXTERNAL float cbrtf(float x);
-extern SYCL_EXTERNAL double cbrt(double x);
-extern SYCL_EXTERNAL float erff(float x);
-extern SYCL_EXTERNAL double erf(double x);
-extern SYCL_EXTERNAL float erfcf(float x);
-extern SYCL_EXTERNAL double erfc(double x);
-extern SYCL_EXTERNAL float tgammaf(float x);
-extern SYCL_EXTERNAL double tgamma(double x);
-extern SYCL_EXTERNAL float lgammaf(float x);
-extern SYCL_EXTERNAL double lgamma(double x);
-extern SYCL_EXTERNAL float fmodf(float x, float y);
-extern SYCL_EXTERNAL double fmod(double x, double y);
-extern SYCL_EXTERNAL float remainderf(float x, float y);
-extern SYCL_EXTERNAL double remainder(double x, double y);
-extern SYCL_EXTERNAL float remquof(float x, float y, int *q);
-extern SYCL_EXTERNAL double remquo(double x, double y, int *q);
-extern SYCL_EXTERNAL float nextafterf(float x, float y);
-extern SYCL_EXTERNAL double nextafter(double x, double y);
-extern SYCL_EXTERNAL float fdimf(float x, float y);
-extern SYCL_EXTERNAL double fdim(double x, double y);
-extern SYCL_EXTERNAL float fmaf(float x, float y, float z);
-extern SYCL_EXTERNAL double fma(double x, double y, double z);
-extern SYCL_EXTERNAL float sinf(float x);
-extern SYCL_EXTERNAL double sin(double x);
-extern SYCL_EXTERNAL float cosf(float x);
-extern SYCL_EXTERNAL double cos(double x);
-extern SYCL_EXTERNAL float tanf(float x);
-extern SYCL_EXTERNAL double tan(double x);
-extern SYCL_EXTERNAL float asinf(float x);
-extern SYCL_EXTERNAL double asin(double x);
-extern SYCL_EXTERNAL float acosf(float x);
-extern SYCL_EXTERNAL double acos(double x);
-extern SYCL_EXTERNAL float atanf(float x);
-extern SYCL_EXTERNAL double atan(double x);
-extern SYCL_EXTERNAL float powf(float x, float y);
-extern SYCL_EXTERNAL double pow(double x, double y);
-extern SYCL_EXTERNAL float atan2f(float x, float y);
-extern SYCL_EXTERNAL double atan2(double x, double y);
-
-extern SYCL_EXTERNAL float sinhf(float x);
-extern SYCL_EXTERNAL double sinh(double x);
-extern SYCL_EXTERNAL float coshf(float x);
-extern SYCL_EXTERNAL double cosh(double x);
-extern SYCL_EXTERNAL float tanhf(float x);
-extern SYCL_EXTERNAL double tanh(double x);
-extern SYCL_EXTERNAL float asinhf(float x);
-extern SYCL_EXTERNAL double asinh(double x);
-extern SYCL_EXTERNAL float acoshf(float x);
-extern SYCL_EXTERNAL double acosh(double x);
-extern SYCL_EXTERNAL float atanhf(float x);
-extern SYCL_EXTERNAL double atanh(double x);
-extern SYCL_EXTERNAL double frexp(double x, int *exp);
-extern SYCL_EXTERNAL double ldexp(double x, int exp);
-extern SYCL_EXTERNAL double hypot(double x, double y);
-
-extern SYCL_EXTERNAL void *memcpy(void *dest, const void *src, size_t n);
-extern SYCL_EXTERNAL void *memset(void *dest, int c, size_t n);
-extern SYCL_EXTERNAL int memcmp(const void *s1, const void *s2, size_t n);
-extern SYCL_EXTERNAL long long int __imf_llmax(long long int x,
-                                               long long int y);
-extern SYCL_EXTERNAL long long int __imf_llmin(long long int x,
-                                               long long int y);
-extern SYCL_EXTERNAL unsigned long long int
+extern __DPCPP_SYCL_EXTERNAL int abs(int x);
+extern __DPCPP_SYCL_EXTERNAL long int labs(long int x);
+extern __DPCPP_SYCL_EXTERNAL long long int llabs(long long int x);
+
+extern __DPCPP_SYCL_EXTERNAL div_t div(int x, int y);
+extern __DPCPP_SYCL_EXTERNAL ldiv_t ldiv(long int x, long int y);
+extern __DPCPP_SYCL_EXTERNAL lldiv_t lldiv(long long int x, long long int y);
+extern __DPCPP_SYCL_EXTERNAL float scalbnf(float x, int n);
+extern __DPCPP_SYCL_EXTERNAL double scalbn(double x, int n);
+extern __DPCPP_SYCL_EXTERNAL float logf(float x);
+extern __DPCPP_SYCL_EXTERNAL double log(double x);
+extern __DPCPP_SYCL_EXTERNAL float expf(float x);
+extern __DPCPP_SYCL_EXTERNAL double exp(double x);
+extern __DPCPP_SYCL_EXTERNAL float log10f(float x);
+extern __DPCPP_SYCL_EXTERNAL double log10(double x);
+extern __DPCPP_SYCL_EXTERNAL float modff(float x, float *intpart);
+extern __DPCPP_SYCL_EXTERNAL double modf(double x, double *intpart);
+extern __DPCPP_SYCL_EXTERNAL float exp2f(float x);
+extern __DPCPP_SYCL_EXTERNAL double exp2(double x);
+extern __DPCPP_SYCL_EXTERNAL float expm1f(float x);
+extern __DPCPP_SYCL_EXTERNAL double expm1(double x);
+extern __DPCPP_SYCL_EXTERNAL int ilogbf(float x);
+extern __DPCPP_SYCL_EXTERNAL int ilogb(double x);
+extern __DPCPP_SYCL_EXTERNAL float log1pf(float x);
+extern __DPCPP_SYCL_EXTERNAL double log1p(double x);
+extern __DPCPP_SYCL_EXTERNAL float log2f(float x);
+extern __DPCPP_SYCL_EXTERNAL double log2(double x);
+extern __DPCPP_SYCL_EXTERNAL float logbf(float x);
+extern __DPCPP_SYCL_EXTERNAL double logb(double x);
+extern __DPCPP_SYCL_EXTERNAL float sqrtf(float x);
+extern __DPCPP_SYCL_EXTERNAL double sqrt(double x);
+extern __DPCPP_SYCL_EXTERNAL float cbrtf(float x);
+extern __DPCPP_SYCL_EXTERNAL double cbrt(double x);
+extern __DPCPP_SYCL_EXTERNAL float erff(float x);
+extern __DPCPP_SYCL_EXTERNAL double erf(double x);
+extern __DPCPP_SYCL_EXTERNAL float erfcf(float x);
+extern __DPCPP_SYCL_EXTERNAL double erfc(double x);
+extern __DPCPP_SYCL_EXTERNAL float tgammaf(float x);
+extern __DPCPP_SYCL_EXTERNAL double tgamma(double x);
+extern __DPCPP_SYCL_EXTERNAL float lgammaf(float x);
+extern __DPCPP_SYCL_EXTERNAL double lgamma(double x);
+extern __DPCPP_SYCL_EXTERNAL float fmodf(float x, float y);
+extern __DPCPP_SYCL_EXTERNAL double fmod(double x, double y);
+extern __DPCPP_SYCL_EXTERNAL float remainderf(float x, float y);
+extern __DPCPP_SYCL_EXTERNAL double remainder(double x, double y);
+extern __DPCPP_SYCL_EXTERNAL float remquof(float x, float y, int *q);
+extern __DPCPP_SYCL_EXTERNAL double remquo(double x, double y, int *q);
+extern __DPCPP_SYCL_EXTERNAL float nextafterf(float x, float y);
+extern __DPCPP_SYCL_EXTERNAL double nextafter(double x, double y);
+extern __DPCPP_SYCL_EXTERNAL float fdimf(float x, float y);
+extern __DPCPP_SYCL_EXTERNAL double fdim(double x, double y);
+extern __DPCPP_SYCL_EXTERNAL float fmaf(float x, float y, float z);
+extern __DPCPP_SYCL_EXTERNAL double fma(double x, double y, double z);
+extern __DPCPP_SYCL_EXTERNAL float sinf(float x);
+extern __DPCPP_SYCL_EXTERNAL double sin(double x);
+extern __DPCPP_SYCL_EXTERNAL float cosf(float x);
+extern __DPCPP_SYCL_EXTERNAL double cos(double x);
+extern __DPCPP_SYCL_EXTERNAL float tanf(float x);
+extern __DPCPP_SYCL_EXTERNAL double tan(double x);
+extern __DPCPP_SYCL_EXTERNAL float asinf(float x);
+extern __DPCPP_SYCL_EXTERNAL double asin(double x);
+extern __DPCPP_SYCL_EXTERNAL float acosf(float x);
+extern __DPCPP_SYCL_EXTERNAL double acos(double x);
+extern __DPCPP_SYCL_EXTERNAL float atanf(float x);
+extern __DPCPP_SYCL_EXTERNAL double atan(double x);
+extern __DPCPP_SYCL_EXTERNAL float powf(float x, float y);
+extern __DPCPP_SYCL_EXTERNAL double pow(double x, double y);
+extern __DPCPP_SYCL_EXTERNAL float atan2f(float x, float y);
+extern __DPCPP_SYCL_EXTERNAL double atan2(double x, double y);
+
+extern __DPCPP_SYCL_EXTERNAL float sinhf(float x);
+extern __DPCPP_SYCL_EXTERNAL double sinh(double x);
+extern __DPCPP_SYCL_EXTERNAL float coshf(float x);
+extern __DPCPP_SYCL_EXTERNAL double cosh(double x);
+extern __DPCPP_SYCL_EXTERNAL float tanhf(float x);
+extern __DPCPP_SYCL_EXTERNAL double tanh(double x);
+extern __DPCPP_SYCL_EXTERNAL float asinhf(float x);
+extern __DPCPP_SYCL_EXTERNAL double asinh(double x);
+extern __DPCPP_SYCL_EXTERNAL float acoshf(float x);
+extern __DPCPP_SYCL_EXTERNAL double acosh(double x);
+extern __DPCPP_SYCL_EXTERNAL float atanhf(float x);
+extern __DPCPP_SYCL_EXTERNAL double atanh(double x);
+extern __DPCPP_SYCL_EXTERNAL double frexp(double x, int *exp);
+extern __DPCPP_SYCL_EXTERNAL double ldexp(double x, int exp);
+extern __DPCPP_SYCL_EXTERNAL double hypot(double x, double y);
+
+extern __DPCPP_SYCL_EXTERNAL void *memcpy(void *dest, const void *src,
+                                          size_t n);
+extern __DPCPP_SYCL_EXTERNAL void *memset(void *dest, int c, size_t n);
+extern __DPCPP_SYCL_EXTERNAL int memcmp(const void *s1, const void *s2,
+                                        size_t n);
+extern __DPCPP_SYCL_EXTERNAL long long int __imf_llmax(long long int x,
+                                                       long long int y);
+extern __DPCPP_SYCL_EXTERNAL long long int __imf_llmin(long long int x,
+                                                       long long int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned long long int
 __imf_ullmax(unsigned long long int x, unsigned long long int y);
-extern SYCL_EXTERNAL unsigned long long int
+extern __DPCPP_SYCL_EXTERNAL unsigned long long int
 __imf_ullmin(unsigned long long int x, unsigned long long int y);
-extern SYCL_EXTERNAL unsigned int __imf_umax(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_umin(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_brev(unsigned int x);
-extern SYCL_EXTERNAL unsigned long long int
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_umax(unsigned int x,
+                                                     unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_umin(unsigned int x,
+                                                     unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_brev(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL unsigned long long int
 __imf_brevll(unsigned long long int x);
-extern SYCL_EXTERNAL unsigned int
+extern __DPCPP_SYCL_EXTERNAL unsigned int
 __imf_byte_perm(unsigned int x, unsigned int y, unsigned int s);
-extern SYCL_EXTERNAL int __imf_ffs(int x);
-extern SYCL_EXTERNAL int __imf_ffsll(long long int x);
-extern SYCL_EXTERNAL int __imf_clz(int x);
-extern SYCL_EXTERNAL int __imf_clzll(long long int x);
-extern SYCL_EXTERNAL int __imf_popc(unsigned int x);
-extern SYCL_EXTERNAL int __imf_popcll(unsigned long long int x);
-extern SYCL_EXTERNAL unsigned int __imf_sad(int x, int y, unsigned int z);
-extern SYCL_EXTERNAL unsigned int __imf_usad(unsigned int x, unsigned int y,
-                                             unsigned int z);
-extern SYCL_EXTERNAL int __imf_rhadd(int x, int y);
-extern SYCL_EXTERNAL unsigned int __imf_urhadd(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_uhadd(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL int __imf_mul24(int x, int y);
-extern SYCL_EXTERNAL unsigned int __imf_umul24(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL int __imf_mulhi(int x, int y);
-extern SYCL_EXTERNAL unsigned int __imf_umulhi(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL long long int __imf_mul64hi(long long int x,
-                                                 long long int y);
-extern SYCL_EXTERNAL unsigned long long int
+extern __DPCPP_SYCL_EXTERNAL int __imf_ffs(int x);
+extern __DPCPP_SYCL_EXTERNAL int __imf_ffsll(long long int x);
+extern __DPCPP_SYCL_EXTERNAL int __imf_clz(int x);
+extern __DPCPP_SYCL_EXTERNAL int __imf_clzll(long long int x);
+extern __DPCPP_SYCL_EXTERNAL int __imf_popc(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL int __imf_popcll(unsigned long long int x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_sad(int x, int y,
+                                                    unsigned int z);
+extern __DPCPP_SYCL_EXTERNAL unsigned int
+__imf_usad(unsigned int x, unsigned int y, unsigned int z);
+extern __DPCPP_SYCL_EXTERNAL int __imf_rhadd(int x, int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_urhadd(unsigned int x,
+                                                       unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_uhadd(unsigned int x,
+                                                      unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL int __imf_mul24(int x, int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_umul24(unsigned int x,
+                                                       unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL int __imf_mulhi(int x, int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_umulhi(unsigned int x,
+                                                       unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL long long int __imf_mul64hi(long long int x,
+                                                         long long int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned long long int
 __imf_umul64hi(unsigned long long int x, unsigned long long int y);
-extern SYCL_EXTERNAL float __imf_saturatef(float x);
-extern SYCL_EXTERNAL float __imf_fmaf(float x, float y, float z);
-extern SYCL_EXTERNAL float __imf_fabsf(float x);
-extern SYCL_EXTERNAL float __imf_floorf(float x);
-extern SYCL_EXTERNAL float __imf_ceilf(float x);
-extern SYCL_EXTERNAL float __imf_truncf(float x);
-extern SYCL_EXTERNAL float __imf_rintf(float x);
-extern SYCL_EXTERNAL float __imf_nearbyintf(float x);
-extern SYCL_EXTERNAL float __imf_sqrtf(float x);
-extern SYCL_EXTERNAL float __imf_rsqrtf(float x);
-extern SYCL_EXTERNAL float __imf_invf(float x);
-extern SYCL_EXTERNAL float __imf_fmaxf(float x, float y);
-extern SYCL_EXTERNAL float __imf_fminf(float x, float y);
-extern SYCL_EXTERNAL float __imf_copysignf(float x, float y);
-extern SYCL_EXTERNAL int __imf_float2int_rd(float x);
-extern SYCL_EXTERNAL int __imf_float2int_rn(float x);
-extern SYCL_EXTERNAL int __imf_float2int_ru(float x);
-extern SYCL_EXTERNAL int __imf_float2int_rz(float x);
-extern SYCL_EXTERNAL unsigned int __imf_float2uint_rd(float x);
-extern SYCL_EXTERNAL unsigned int __imf_float2uint_rn(float x);
-extern SYCL_EXTERNAL unsigned int __imf_float2uint_ru(float x);
-extern SYCL_EXTERNAL unsigned int __imf_float2uint_rz(float x);
-extern SYCL_EXTERNAL long long int __imf_float2ll_rd(float x);
-extern SYCL_EXTERNAL long long int __imf_float2ll_rn(float x);
-extern SYCL_EXTERNAL long long int __imf_float2ll_ru(float x);
-extern SYCL_EXTERNAL long long int __imf_float2ll_rz(float x);
-extern SYCL_EXTERNAL unsigned long long int __imf_float2ull_rd(float x);
-extern SYCL_EXTERNAL unsigned long long int __imf_float2ull_rn(float x);
-extern SYCL_EXTERNAL unsigned long long int __imf_float2ull_ru(float x);
-extern SYCL_EXTERNAL unsigned long long int __imf_float2ull_rz(float x);
-extern SYCL_EXTERNAL int __imf_float_as_int(float x);
-extern SYCL_EXTERNAL unsigned int __imf_float_as_uint(float x);
-extern SYCL_EXTERNAL float __imf_int2float_rd(int x);
-extern SYCL_EXTERNAL float __imf_int2float_rn(int x);
-extern SYCL_EXTERNAL float __imf_int2float_ru(int x);
-extern SYCL_EXTERNAL float __imf_int2float_rz(int x);
-extern SYCL_EXTERNAL float __imf_int_as_float(int x);
-extern SYCL_EXTERNAL float __imf_ll2float_rd(long long int x);
-extern SYCL_EXTERNAL float __imf_ll2float_rn(long long int x);
-extern SYCL_EXTERNAL float __imf_ll2float_ru(long long int x);
-extern SYCL_EXTERNAL float __imf_ll2float_rz(long long int x);
-extern SYCL_EXTERNAL float __imf_uint2float_rd(unsigned int x);
-extern SYCL_EXTERNAL float __imf_uint2float_rn(unsigned int x);
-extern SYCL_EXTERNAL float __imf_uint2float_ru(unsigned int x);
-extern SYCL_EXTERNAL float __imf_uint2float_rz(unsigned int x);
-extern SYCL_EXTERNAL float __imf_uint_as_float(unsigned int x);
-extern SYCL_EXTERNAL float __imf_ull2float_rd(unsigned long long int x);
-extern SYCL_EXTERNAL float __imf_ull2float_rn(unsigned long long int x);
-extern SYCL_EXTERNAL float __imf_ull2float_ru(unsigned long long int x);
-extern SYCL_EXTERNAL float __imf_ull2float_rz(unsigned long long int x);
-extern SYCL_EXTERNAL float __imf_half2float(_Float16 x);
-extern SYCL_EXTERNAL _Float16 __imf_float2half_rd(float x);
-extern SYCL_EXTERNAL _Float16 __imf_float2half_rn(float x);
-extern SYCL_EXTERNAL _Float16 __imf_float2half_ru(float x);
-extern SYCL_EXTERNAL _Float16 __imf_float2half_rz(float x);
-extern SYCL_EXTERNAL int __imf_half2int_rd(_Float16 x);
-extern SYCL_EXTERNAL int __imf_half2int_rn(_Float16 x);
-extern SYCL_EXTERNAL int __imf_half2int_ru(_Float16 x);
-extern SYCL_EXTERNAL int __imf_half2int_rz(_Float16 x);
-extern SYCL_EXTERNAL long long __imf_half2ll_rd(_Float16 x);
-extern SYCL_EXTERNAL long long __imf_half2ll_rn(_Float16 x);
-extern SYCL_EXTERNAL long long __imf_half2ll_ru(_Float16 x);
-extern SYCL_EXTERNAL long long __imf_half2ll_rz(_Float16 x);
-extern SYCL_EXTERNAL short __imf_half2short_rd(_Float16 x);
-extern SYCL_EXTERNAL short __imf_half2short_rn(_Float16 x);
-extern SYCL_EXTERNAL short __imf_half2short_ru(_Float16 x);
-extern SYCL_EXTERNAL short __imf_half2short_rz(_Float16 x);
-extern SYCL_EXTERNAL unsigned int __imf_half2uint_rd(_Float16 x);
-extern SYCL_EXTERNAL unsigned int __imf_half2uint_rn(_Float16 x);
-extern SYCL_EXTERNAL unsigned int __imf_half2uint_ru(_Float16 x);
-extern SYCL_EXTERNAL unsigned int __imf_half2uint_rz(_Float16 x);
-extern SYCL_EXTERNAL unsigned long long __imf_half2ull_rd(_Float16 x);
-extern SYCL_EXTERNAL unsigned long long __imf_half2ull_rn(_Float16 x);
-extern SYCL_EXTERNAL unsigned long long __imf_half2ull_ru(_Float16 x);
-extern SYCL_EXTERNAL unsigned long long __imf_half2ull_rz(_Float16 x);
-extern SYCL_EXTERNAL unsigned short __imf_half2ushort_rd(_Float16 x);
-extern SYCL_EXTERNAL unsigned short __imf_half2ushort_rn(_Float16 x);
-extern SYCL_EXTERNAL unsigned short __imf_half2ushort_ru(_Float16 x);
-extern SYCL_EXTERNAL unsigned short __imf_half2ushort_rz(_Float16 x);
-extern SYCL_EXTERNAL short __imf_half_as_short(_Float16 x);
-extern SYCL_EXTERNAL unsigned short __imf_half_as_ushort(_Float16 x);
-extern SYCL_EXTERNAL _Float16 __imf_int2half_rd(int x);
-extern SYCL_EXTERNAL _Float16 __imf_int2half_rn(int x);
-extern SYCL_EXTERNAL _Float16 __imf_int2half_ru(int x);
-extern SYCL_EXTERNAL _Float16 __imf_int2half_rz(int x);
-extern SYCL_EXTERNAL _Float16 __imf_ll2half_rd(long long x);
-extern SYCL_EXTERNAL _Float16 __imf_ll2half_rn(long long x);
-extern SYCL_EXTERNAL _Float16 __imf_ll2half_ru(long long x);
-extern SYCL_EXTERNAL _Float16 __imf_ll2half_rz(long long x);
-extern SYCL_EXTERNAL _Float16 __imf_short2half_rd(short x);
-extern SYCL_EXTERNAL _Float16 __imf_short2half_rn(short x);
-extern SYCL_EXTERNAL _Float16 __imf_short2half_ru(short x);
-extern SYCL_EXTERNAL _Float16 __imf_short2half_rz(short x);
-extern SYCL_EXTERNAL _Float16 __imf_short_as_half(short x);
-extern SYCL_EXTERNAL _Float16 __imf_uint2half_rd(unsigned int x);
-extern SYCL_EXTERNAL _Float16 __imf_uint2half_rn(unsigned int x);
-extern SYCL_EXTERNAL _Float16 __imf_uint2half_ru(unsigned int x);
-extern SYCL_EXTERNAL _Float16 __imf_uint2half_rz(unsigned int x);
-extern SYCL_EXTERNAL _Float16 __imf_ull2half_rd(unsigned long long x);
-extern SYCL_EXTERNAL _Float16 __imf_ull2half_rn(unsigned long long x);
-extern SYCL_EXTERNAL _Float16 __imf_ull2half_ru(unsigned long long x);
-extern SYCL_EXTERNAL _Float16 __imf_ull2half_rz(unsigned long long x);
-extern SYCL_EXTERNAL _Float16 __imf_ushort2half_rd(unsigned short x);
-extern SYCL_EXTERNAL _Float16 __imf_ushort2half_rn(unsigned short x);
-extern SYCL_EXTERNAL _Float16 __imf_ushort2half_ru(unsigned short x);
-extern SYCL_EXTERNAL _Float16 __imf_ushort2half_rz(unsigned short x);
-extern SYCL_EXTERNAL _Float16 __imf_ushort_as_half(unsigned short x);
-extern SYCL_EXTERNAL _Float16 __imf_double2half(double x);
-
-extern SYCL_EXTERNAL _Float16 __imf_fmaf16(_Float16 x, _Float16 y, _Float16 z);
-extern SYCL_EXTERNAL _Float16 __imf_fabsf16(_Float16 x);
-extern SYCL_EXTERNAL _Float16 __imf_floorf16(_Float16 x);
-extern SYCL_EXTERNAL _Float16 __imf_ceilf16(_Float16 x);
-extern SYCL_EXTERNAL _Float16 __imf_truncf16(_Float16 x);
-extern SYCL_EXTERNAL _Float16 __imf_rintf16(_Float16 x);
-extern SYCL_EXTERNAL _Float16 __imf_nearbyintf16(_Float16 x);
-extern SYCL_EXTERNAL _Float16 __imf_sqrtf16(_Float16 x);
-extern SYCL_EXTERNAL _Float16 __imf_rsqrtf16(_Float16 x);
-extern SYCL_EXTERNAL _Float16 __imf_invf16(_Float16 x);
-extern SYCL_EXTERNAL _Float16 __imf_fmaxf16(_Float16 x, _Float16 y);
-extern SYCL_EXTERNAL _Float16 __imf_fminf16(_Float16 x, _Float16 y);
-extern SYCL_EXTERNAL _Float16 __imf_copysignf16(_Float16 x, _Float16 y);
-extern SYCL_EXTERNAL float __imf_half2float(_Float16 x);
-extern SYCL_EXTERNAL float __imf_bfloat162float(uint16_t x);
-extern SYCL_EXTERNAL unsigned int __imf_bfloat162uint_rd(uint16_t x);
-extern SYCL_EXTERNAL unsigned int __imf_bfloat162uint_rn(uint16_t x);
-extern SYCL_EXTERNAL unsigned int __imf_bfloat162uint_ru(uint16_t x);
-extern SYCL_EXTERNAL unsigned int __imf_bfloat162uint_rz(uint16_t x);
-extern SYCL_EXTERNAL unsigned short __imf_bfloat162ushort_rd(uint16_t x);
-extern SYCL_EXTERNAL unsigned short __imf_bfloat162ushort_rn(uint16_t x);
-extern SYCL_EXTERNAL unsigned short __imf_bfloat162ushort_ru(uint16_t x);
-extern SYCL_EXTERNAL unsigned short __imf_bfloat162ushort_rz(uint16_t x);
-extern SYCL_EXTERNAL unsigned long long __imf_bfloat162ull_rd(uint16_t x);
-extern SYCL_EXTERNAL unsigned long long __imf_bfloat162ull_rn(uint16_t x);
-extern SYCL_EXTERNAL unsigned long long __imf_bfloat162ull_ru(uint16_t x);
-extern SYCL_EXTERNAL unsigned long long __imf_bfloat162ull_rz(uint16_t x);
-extern SYCL_EXTERNAL int __imf_bfloat162int_rd(uint16_t x);
-extern SYCL_EXTERNAL int __imf_bfloat162int_rn(uint16_t x);
-extern SYCL_EXTERNAL int __imf_bfloat162int_ru(uint16_t x);
-extern SYCL_EXTERNAL int __imf_bfloat162int_rz(uint16_t x);
-extern SYCL_EXTERNAL short __imf_bfloat162short_rd(uint16_t x);
-extern SYCL_EXTERNAL short __imf_bfloat162short_rn(uint16_t x);
-extern SYCL_EXTERNAL short __imf_bfloat162short_ru(uint16_t x);
-extern SYCL_EXTERNAL short __imf_bfloat162short_rz(uint16_t x);
-extern SYCL_EXTERNAL long long __imf_bfloat162ll_rd(uint16_t x);
-extern SYCL_EXTERNAL long long __imf_bfloat162ll_rn(uint16_t x);
-extern SYCL_EXTERNAL long long __imf_bfloat162ll_ru(uint16_t x);
-extern SYCL_EXTERNAL long long __imf_bfloat162ll_rz(uint16_t x);
-extern SYCL_EXTERNAL uint16_t __imf_float2bfloat16(float x);
-extern SYCL_EXTERNAL uint16_t __imf_float2bfloat16_rd(float x);
-extern SYCL_EXTERNAL uint16_t __imf_float2bfloat16_rn(float x);
-extern SYCL_EXTERNAL uint16_t __imf_float2bfloat16_ru(float x);
-extern SYCL_EXTERNAL uint16_t __imf_float2bfloat16_rz(float x);
-extern SYCL_EXTERNAL uint16_t __imf_ushort2bfloat16_rd(unsigned short x);
-extern SYCL_EXTERNAL uint16_t __imf_ushort2bfloat16_rn(unsigned short x);
-extern SYCL_EXTERNAL uint16_t __imf_ushort2bfloat16_ru(unsigned short x);
-extern SYCL_EXTERNAL uint16_t __imf_ushort2bfloat16_rz(unsigned short x);
-extern SYCL_EXTERNAL uint16_t __imf_uint2bfloat16_rd(unsigned int x);
-extern SYCL_EXTERNAL uint16_t __imf_uint2bfloat16_rn(unsigned int x);
-extern SYCL_EXTERNAL uint16_t __imf_uint2bfloat16_ru(unsigned int x);
-extern SYCL_EXTERNAL uint16_t __imf_uint2bfloat16_rz(unsigned int x);
-extern SYCL_EXTERNAL uint16_t __imf_ull2bfloat16_rd(unsigned long long x);
-extern SYCL_EXTERNAL uint16_t __imf_ull2bfloat16_rn(unsigned long long x);
-extern SYCL_EXTERNAL uint16_t __imf_ull2bfloat16_ru(unsigned long long x);
-extern SYCL_EXTERNAL uint16_t __imf_ull2bfloat16_rz(unsigned long long x);
-extern SYCL_EXTERNAL uint16_t __imf_short2bfloat16_rd(short x);
-extern SYCL_EXTERNAL uint16_t __imf_short2bfloat16_rn(short x);
-extern SYCL_EXTERNAL uint16_t __imf_short2bfloat16_ru(short x);
-extern SYCL_EXTERNAL uint16_t __imf_short2bfloat16_rz(short x);
-extern SYCL_EXTERNAL uint16_t __imf_int2bfloat16_rd(int x);
-extern SYCL_EXTERNAL uint16_t __imf_int2bfloat16_rn(int x);
-extern SYCL_EXTERNAL uint16_t __imf_int2bfloat16_ru(int x);
-extern SYCL_EXTERNAL uint16_t __imf_int2bfloat16_rz(int x);
-extern SYCL_EXTERNAL uint16_t __imf_ll2bfloat16_rd(long long x);
-extern SYCL_EXTERNAL uint16_t __imf_ll2bfloat16_rn(long long x);
-extern SYCL_EXTERNAL uint16_t __imf_ll2bfloat16_ru(long long x);
-extern SYCL_EXTERNAL uint16_t __imf_ll2bfloat16_rz(long long x);
-extern SYCL_EXTERNAL uint16_t __imf_double2bfloat16(double x);
-extern SYCL_EXTERNAL short __imf_bfloat16_as_short(uint16_t x);
-extern SYCL_EXTERNAL unsigned short __imf_bfloat16_as_ushort(uint16_t x);
-extern SYCL_EXTERNAL uint16_t __imf_short_as_bfloat16(short x);
-extern SYCL_EXTERNAL uint16_t __imf_ushort_as_bfloat16(unsigned short x);
-extern SYCL_EXTERNAL uint16_t __imf_fmabf16(uint16_t x, uint16_t y, uint16_t z);
-extern SYCL_EXTERNAL uint16_t __imf_fmaxbf16(uint16_t x, uint16_t y);
-extern SYCL_EXTERNAL uint16_t __imf_fminbf16(uint16_t x, uint16_t y);
-extern SYCL_EXTERNAL uint16_t __imf_fabsbf16(uint16_t x);
-extern SYCL_EXTERNAL uint16_t __imf_rintbf16(uint16_t x);
-extern SYCL_EXTERNAL uint16_t __imf_floorbf16(uint16_t x);
-extern SYCL_EXTERNAL uint16_t __imf_ceilbf16(uint16_t x);
-extern SYCL_EXTERNAL uint16_t __imf_truncbf16(uint16_t x);
-extern SYCL_EXTERNAL uint16_t __imf_copysignbf16(uint16_t x, uint16_t y);
-extern SYCL_EXTERNAL uint16_t __imf_sqrtbf16(uint16_t x);
-extern SYCL_EXTERNAL uint16_t __imf_rsqrtbf16(uint16_t x);
-extern SYCL_EXTERNAL double __imf_fma(double x, double y, double z);
-extern SYCL_EXTERNAL double __imf_fabs(double x);
-extern SYCL_EXTERNAL double __imf_floor(double x);
-extern SYCL_EXTERNAL double __imf_ceil(double x);
-extern SYCL_EXTERNAL double __imf_trunc(double x);
-extern SYCL_EXTERNAL double __imf_rint(double x);
-extern SYCL_EXTERNAL double __imf_nearbyint(double x);
-extern SYCL_EXTERNAL double __imf_sqrt(double x);
-extern SYCL_EXTERNAL double __imf_rsqrt(double x);
-extern SYCL_EXTERNAL double __imf_inv(double x);
-extern SYCL_EXTERNAL double __imf_fmax(double x, double y);
-extern SYCL_EXTERNAL double __imf_fmin(double x, double y);
-extern SYCL_EXTERNAL double __imf_copysign(double x, double y);
-extern SYCL_EXTERNAL float __imf_double2float_rd(double x);
-extern SYCL_EXTERNAL float __imf_double2float_rn(double x);
-extern SYCL_EXTERNAL float __imf_double2float_ru(double x);
-extern SYCL_EXTERNAL float __imf_double2float_rz(double x);
-extern SYCL_EXTERNAL int __imf_double2hiint(double x);
-extern SYCL_EXTERNAL int __imf_double2loint(double x);
-extern SYCL_EXTERNAL int __imf_double2int_rd(double x);
-extern SYCL_EXTERNAL int __imf_double2int_rn(double x);
-extern SYCL_EXTERNAL int __imf_double2int_ru(double x);
-extern SYCL_EXTERNAL int __imf_double2int_rz(double x);
-extern SYCL_EXTERNAL double __imf_int2double_rn(int x);
-extern SYCL_EXTERNAL unsigned int __imf_double2uint_rd(double x);
-extern SYCL_EXTERNAL unsigned int __imf_double2uint_rn(double x);
-extern SYCL_EXTERNAL unsigned int __imf_double2uint_ru(double x);
-extern SYCL_EXTERNAL unsigned int __imf_double2uint_rz(double x);
-extern SYCL_EXTERNAL long long int __imf_double2ll_rd(double x);
-extern SYCL_EXTERNAL long long int __imf_double2ll_rn(double x);
-extern SYCL_EXTERNAL long long int __imf_double2ll_ru(double x);
-extern SYCL_EXTERNAL long long int __imf_double2ll_rz(double x);
-extern SYCL_EXTERNAL double __imf_ll2double_rd(long long int x);
-extern SYCL_EXTERNAL double __imf_ll2double_rn(long long int x);
-extern SYCL_EXTERNAL double __imf_ll2double_ru(long long int x);
-extern SYCL_EXTERNAL double __imf_ll2double_rz(long long int x);
-extern SYCL_EXTERNAL double __imf_ull2double_rd(unsigned long long int x);
-extern SYCL_EXTERNAL double __imf_ull2double_rn(unsigned long long int x);
-extern SYCL_EXTERNAL double __imf_ull2double_ru(unsigned long long int x);
-extern SYCL_EXTERNAL double __imf_ull2double_rz(unsigned long long int x);
-extern SYCL_EXTERNAL unsigned long long int __imf_double2ull_rd(double x);
-extern SYCL_EXTERNAL unsigned long long int __imf_double2ull_rn(double x);
-extern SYCL_EXTERNAL unsigned long long int __imf_double2ull_ru(double x);
-extern SYCL_EXTERNAL unsigned long long int __imf_double2ull_rz(double x);
-extern SYCL_EXTERNAL long long int __imf_double_as_longlong(double x);
-extern SYCL_EXTERNAL double __imf_longlong_as_double(long long int x);
-extern SYCL_EXTERNAL double __imf_uint2double_rd(unsigned int x);
-extern SYCL_EXTERNAL double __imf_uint2double_rn(unsigned int x);
-extern SYCL_EXTERNAL double __imf_uint2double_ru(unsigned int x);
-extern SYCL_EXTERNAL double __imf_uint2double_rz(unsigned int x);
-extern SYCL_EXTERNAL double __imf_hiloint2double(int hi, int lo);
-
-extern SYCL_EXTERNAL unsigned int __imf_vabs2(unsigned int x);
-extern SYCL_EXTERNAL unsigned int __imf_vabs4(unsigned int x);
-extern SYCL_EXTERNAL unsigned int __imf_vabsss2(unsigned int x);
-extern SYCL_EXTERNAL unsigned int __imf_vabsss4(unsigned int x);
-extern SYCL_EXTERNAL unsigned int __imf_vneg2(unsigned int x);
-extern SYCL_EXTERNAL unsigned int __imf_vneg4(unsigned int x);
-extern SYCL_EXTERNAL unsigned int __imf_vnegss2(unsigned int x);
-extern SYCL_EXTERNAL unsigned int __imf_vnegss4(unsigned int x);
-extern SYCL_EXTERNAL unsigned int __imf_vabsdiffs2(unsigned int x,
-                                                   unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vabsdiffs4(unsigned int x,
-                                                   unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vabsdiffu2(unsigned int x,
-                                                   unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vabsdiffu4(unsigned int x,
-                                                   unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vadd2(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vadd4(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vaddss2(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vaddss4(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vaddus2(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vaddus4(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsub2(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsub4(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsubss2(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsubss4(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsubus2(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsubus4(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vavgs2(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vavgs4(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vavgu2(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vavgu4(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vhaddu2(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vhaddu4(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vcmpeq2(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vcmpeq4(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vcmpne2(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vcmpne4(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vcmpges2(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vcmpges4(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vcmpgeu2(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vcmpgeu4(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vcmpgts2(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vcmpgts4(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vcmpgtu2(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vcmpgtu4(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vcmples2(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vcmples4(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vcmpleu2(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vcmpleu4(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vcmplts2(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vcmplts4(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vcmpltu2(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vcmpltu4(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vmaxs2(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vmaxs4(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vmaxu2(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vmaxu4(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vmins2(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vmins4(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vminu2(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vminu4(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vseteq2(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vseteq4(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsetne2(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsetne4(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsetges2(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsetges4(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsetgeu2(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsetgeu4(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsetgts2(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsetgts4(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsetgtu2(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsetgtu4(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsetles2(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsetles4(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsetleu2(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsetleu4(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsetlts2(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsetlts4(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsetltu2(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsetltu4(unsigned int x,
-                                                 unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsads2(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsads4(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsadu2(unsigned int x, unsigned int y);
-extern SYCL_EXTERNAL unsigned int __imf_vsadu4(unsigned int x, unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL float __imf_saturatef(float x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_fmaf(float x, float y, float z);
+extern __DPCPP_SYCL_EXTERNAL float __imf_fabsf(float x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_floorf(float x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_ceilf(float x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_truncf(float x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_rintf(float x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_nearbyintf(float x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_sqrtf(float x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_rsqrtf(float x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_invf(float x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_fmaxf(float x, float y);
+extern __DPCPP_SYCL_EXTERNAL float __imf_fminf(float x, float y);
+extern __DPCPP_SYCL_EXTERNAL float __imf_copysignf(float x, float y);
+extern __DPCPP_SYCL_EXTERNAL int __imf_float2int_rd(float x);
+extern __DPCPP_SYCL_EXTERNAL int __imf_float2int_rn(float x);
+extern __DPCPP_SYCL_EXTERNAL int __imf_float2int_ru(float x);
+extern __DPCPP_SYCL_EXTERNAL int __imf_float2int_rz(float x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_float2uint_rd(float x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_float2uint_rn(float x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_float2uint_ru(float x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_float2uint_rz(float x);
+extern __DPCPP_SYCL_EXTERNAL long long int __imf_float2ll_rd(float x);
+extern __DPCPP_SYCL_EXTERNAL long long int __imf_float2ll_rn(float x);
+extern __DPCPP_SYCL_EXTERNAL long long int __imf_float2ll_ru(float x);
+extern __DPCPP_SYCL_EXTERNAL long long int __imf_float2ll_rz(float x);
+extern __DPCPP_SYCL_EXTERNAL unsigned long long int __imf_float2ull_rd(float x);
+extern __DPCPP_SYCL_EXTERNAL unsigned long long int __imf_float2ull_rn(float x);
+extern __DPCPP_SYCL_EXTERNAL unsigned long long int __imf_float2ull_ru(float x);
+extern __DPCPP_SYCL_EXTERNAL unsigned long long int __imf_float2ull_rz(float x);
+extern __DPCPP_SYCL_EXTERNAL int __imf_float_as_int(float x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_float_as_uint(float x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_int2float_rd(int x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_int2float_rn(int x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_int2float_ru(int x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_int2float_rz(int x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_int_as_float(int x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_ll2float_rd(long long int x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_ll2float_rn(long long int x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_ll2float_ru(long long int x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_ll2float_rz(long long int x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_uint2float_rd(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_uint2float_rn(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_uint2float_ru(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_uint2float_rz(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_uint_as_float(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_ull2float_rd(unsigned long long int x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_ull2float_rn(unsigned long long int x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_ull2float_ru(unsigned long long int x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_ull2float_rz(unsigned long long int x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_half2float(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_float2half_rd(float x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_float2half_rn(float x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_float2half_ru(float x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_float2half_rz(float x);
+extern __DPCPP_SYCL_EXTERNAL int __imf_half2int_rd(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL int __imf_half2int_rn(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL int __imf_half2int_ru(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL int __imf_half2int_rz(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL long long __imf_half2ll_rd(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL long long __imf_half2ll_rn(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL long long __imf_half2ll_ru(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL long long __imf_half2ll_rz(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL short __imf_half2short_rd(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL short __imf_half2short_rn(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL short __imf_half2short_ru(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL short __imf_half2short_rz(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_half2uint_rd(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_half2uint_rn(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_half2uint_ru(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_half2uint_rz(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL unsigned long long __imf_half2ull_rd(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL unsigned long long __imf_half2ull_rn(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL unsigned long long __imf_half2ull_ru(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL unsigned long long __imf_half2ull_rz(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL unsigned short __imf_half2ushort_rd(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL unsigned short __imf_half2ushort_rn(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL unsigned short __imf_half2ushort_ru(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL unsigned short __imf_half2ushort_rz(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL short __imf_half_as_short(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL unsigned short __imf_half_as_ushort(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_int2half_rd(int x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_int2half_rn(int x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_int2half_ru(int x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_int2half_rz(int x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_ll2half_rd(long long x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_ll2half_rn(long long x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_ll2half_ru(long long x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_ll2half_rz(long long x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_short2half_rd(short x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_short2half_rn(short x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_short2half_ru(short x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_short2half_rz(short x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_short_as_half(short x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_uint2half_rd(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_uint2half_rn(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_uint2half_ru(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_uint2half_rz(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_ull2half_rd(unsigned long long x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_ull2half_rn(unsigned long long x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_ull2half_ru(unsigned long long x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_ull2half_rz(unsigned long long x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_ushort2half_rd(unsigned short x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_ushort2half_rn(unsigned short x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_ushort2half_ru(unsigned short x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_ushort2half_rz(unsigned short x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_ushort_as_half(unsigned short x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_double2half(double x);
+
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_fmaf16(_Float16 x, _Float16 y,
+                                                   _Float16 z);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_fabsf16(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_floorf16(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_ceilf16(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_truncf16(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_rintf16(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_nearbyintf16(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_sqrtf16(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_rsqrtf16(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_invf16(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_fmaxf16(_Float16 x, _Float16 y);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_fminf16(_Float16 x, _Float16 y);
+extern __DPCPP_SYCL_EXTERNAL _Float16 __imf_copysignf16(_Float16 x, _Float16 y);
+extern __DPCPP_SYCL_EXTERNAL float __imf_half2float(_Float16 x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_bfloat162float(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_bfloat162uint_rd(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_bfloat162uint_rn(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_bfloat162uint_ru(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_bfloat162uint_rz(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL unsigned short
+__imf_bfloat162ushort_rd(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL unsigned short
+__imf_bfloat162ushort_rn(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL unsigned short
+__imf_bfloat162ushort_ru(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL unsigned short
+__imf_bfloat162ushort_rz(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL unsigned long long
+__imf_bfloat162ull_rd(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL unsigned long long
+__imf_bfloat162ull_rn(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL unsigned long long
+__imf_bfloat162ull_ru(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL unsigned long long
+__imf_bfloat162ull_rz(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL int __imf_bfloat162int_rd(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL int __imf_bfloat162int_rn(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL int __imf_bfloat162int_ru(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL int __imf_bfloat162int_rz(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL short __imf_bfloat162short_rd(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL short __imf_bfloat162short_rn(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL short __imf_bfloat162short_ru(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL short __imf_bfloat162short_rz(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL long long __imf_bfloat162ll_rd(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL long long __imf_bfloat162ll_rn(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL long long __imf_bfloat162ll_ru(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL long long __imf_bfloat162ll_rz(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_float2bfloat16(float x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_float2bfloat16_rd(float x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_float2bfloat16_rn(float x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_float2bfloat16_ru(float x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_float2bfloat16_rz(float x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t
+__imf_ushort2bfloat16_rd(unsigned short x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t
+__imf_ushort2bfloat16_rn(unsigned short x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t
+__imf_ushort2bfloat16_ru(unsigned short x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t
+__imf_ushort2bfloat16_rz(unsigned short x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_uint2bfloat16_rd(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_uint2bfloat16_rn(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_uint2bfloat16_ru(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_uint2bfloat16_rz(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t
+__imf_ull2bfloat16_rd(unsigned long long x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t
+__imf_ull2bfloat16_rn(unsigned long long x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t
+__imf_ull2bfloat16_ru(unsigned long long x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t
+__imf_ull2bfloat16_rz(unsigned long long x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_short2bfloat16_rd(short x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_short2bfloat16_rn(short x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_short2bfloat16_ru(short x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_short2bfloat16_rz(short x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_int2bfloat16_rd(int x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_int2bfloat16_rn(int x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_int2bfloat16_ru(int x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_int2bfloat16_rz(int x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_ll2bfloat16_rd(long long x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_ll2bfloat16_rn(long long x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_ll2bfloat16_ru(long long x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_ll2bfloat16_rz(long long x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_double2bfloat16(double x);
+extern __DPCPP_SYCL_EXTERNAL short __imf_bfloat16_as_short(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL unsigned short
+__imf_bfloat16_as_ushort(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_short_as_bfloat16(short x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t
+__imf_ushort_as_bfloat16(unsigned short x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_fmabf16(uint16_t x, uint16_t y,
+                                                    uint16_t z);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_fmaxbf16(uint16_t x, uint16_t y);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_fminbf16(uint16_t x, uint16_t y);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_fabsbf16(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_rintbf16(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_floorbf16(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_ceilbf16(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_truncbf16(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_copysignbf16(uint16_t x,
+                                                         uint16_t y);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_sqrtbf16(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL uint16_t __imf_rsqrtbf16(uint16_t x);
+extern __DPCPP_SYCL_EXTERNAL double __imf_fma(double x, double y, double z);
+extern __DPCPP_SYCL_EXTERNAL double __imf_fabs(double x);
+extern __DPCPP_SYCL_EXTERNAL double __imf_floor(double x);
+extern __DPCPP_SYCL_EXTERNAL double __imf_ceil(double x);
+extern __DPCPP_SYCL_EXTERNAL double __imf_trunc(double x);
+extern __DPCPP_SYCL_EXTERNAL double __imf_rint(double x);
+extern __DPCPP_SYCL_EXTERNAL double __imf_nearbyint(double x);
+extern __DPCPP_SYCL_EXTERNAL double __imf_sqrt(double x);
+extern __DPCPP_SYCL_EXTERNAL double __imf_rsqrt(double x);
+extern __DPCPP_SYCL_EXTERNAL double __imf_inv(double x);
+extern __DPCPP_SYCL_EXTERNAL double __imf_fmax(double x, double y);
+extern __DPCPP_SYCL_EXTERNAL double __imf_fmin(double x, double y);
+extern __DPCPP_SYCL_EXTERNAL double __imf_copysign(double x, double y);
+extern __DPCPP_SYCL_EXTERNAL float __imf_double2float_rd(double x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_double2float_rn(double x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_double2float_ru(double x);
+extern __DPCPP_SYCL_EXTERNAL float __imf_double2float_rz(double x);
+extern __DPCPP_SYCL_EXTERNAL int __imf_double2hiint(double x);
+extern __DPCPP_SYCL_EXTERNAL int __imf_double2loint(double x);
+extern __DPCPP_SYCL_EXTERNAL int __imf_double2int_rd(double x);
+extern __DPCPP_SYCL_EXTERNAL int __imf_double2int_rn(double x);
+extern __DPCPP_SYCL_EXTERNAL int __imf_double2int_ru(double x);
+extern __DPCPP_SYCL_EXTERNAL int __imf_double2int_rz(double x);
+extern __DPCPP_SYCL_EXTERNAL double __imf_int2double_rn(int x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_double2uint_rd(double x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_double2uint_rn(double x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_double2uint_ru(double x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_double2uint_rz(double x);
+extern __DPCPP_SYCL_EXTERNAL long long int __imf_double2ll_rd(double x);
+extern __DPCPP_SYCL_EXTERNAL long long int __imf_double2ll_rn(double x);
+extern __DPCPP_SYCL_EXTERNAL long long int __imf_double2ll_ru(double x);
+extern __DPCPP_SYCL_EXTERNAL long long int __imf_double2ll_rz(double x);
+extern __DPCPP_SYCL_EXTERNAL double __imf_ll2double_rd(long long int x);
+extern __DPCPP_SYCL_EXTERNAL double __imf_ll2double_rn(long long int x);
+extern __DPCPP_SYCL_EXTERNAL double __imf_ll2double_ru(long long int x);
+extern __DPCPP_SYCL_EXTERNAL double __imf_ll2double_rz(long long int x);
+extern __DPCPP_SYCL_EXTERNAL double
+__imf_ull2double_rd(unsigned long long int x);
+extern __DPCPP_SYCL_EXTERNAL double
+__imf_ull2double_rn(unsigned long long int x);
+extern __DPCPP_SYCL_EXTERNAL double
+__imf_ull2double_ru(unsigned long long int x);
+extern __DPCPP_SYCL_EXTERNAL double
+__imf_ull2double_rz(unsigned long long int x);
+extern __DPCPP_SYCL_EXTERNAL unsigned long long int
+__imf_double2ull_rd(double x);
+extern __DPCPP_SYCL_EXTERNAL unsigned long long int
+__imf_double2ull_rn(double x);
+extern __DPCPP_SYCL_EXTERNAL unsigned long long int
+__imf_double2ull_ru(double x);
+extern __DPCPP_SYCL_EXTERNAL unsigned long long int
+__imf_double2ull_rz(double x);
+extern __DPCPP_SYCL_EXTERNAL long long int __imf_double_as_longlong(double x);
+extern __DPCPP_SYCL_EXTERNAL double __imf_longlong_as_double(long long int x);
+extern __DPCPP_SYCL_EXTERNAL double __imf_uint2double_rd(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL double __imf_uint2double_rn(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL double __imf_uint2double_ru(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL double __imf_uint2double_rz(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL double __imf_hiloint2double(int hi, int lo);
+
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vabs2(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vabs4(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vabsss2(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vabsss4(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vneg2(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vneg4(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vnegss2(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vnegss4(unsigned int x);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vabsdiffs2(unsigned int x,
+                                                           unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vabsdiffs4(unsigned int x,
+                                                           unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vabsdiffu2(unsigned int x,
+                                                           unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vabsdiffu4(unsigned int x,
+                                                           unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vadd2(unsigned int x,
+                                                      unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vadd4(unsigned int x,
+                                                      unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vaddss2(unsigned int x,
+                                                        unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vaddss4(unsigned int x,
+                                                        unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vaddus2(unsigned int x,
+                                                        unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vaddus4(unsigned int x,
+                                                        unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsub2(unsigned int x,
+                                                      unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsub4(unsigned int x,
+                                                      unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsubss2(unsigned int x,
+                                                        unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsubss4(unsigned int x,
+                                                        unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsubus2(unsigned int x,
+                                                        unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsubus4(unsigned int x,
+                                                        unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vavgs2(unsigned int x,
+                                                       unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vavgs4(unsigned int x,
+                                                       unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vavgu2(unsigned int x,
+                                                       unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vavgu4(unsigned int x,
+                                                       unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vhaddu2(unsigned int x,
+                                                        unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vhaddu4(unsigned int x,
+                                                        unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vcmpeq2(unsigned int x,
+                                                        unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vcmpeq4(unsigned int x,
+                                                        unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vcmpne2(unsigned int x,
+                                                        unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vcmpne4(unsigned int x,
+                                                        unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vcmpges2(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vcmpges4(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vcmpgeu2(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vcmpgeu4(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vcmpgts2(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vcmpgts4(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vcmpgtu2(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vcmpgtu4(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vcmples2(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vcmples4(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vcmpleu2(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vcmpleu4(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vcmplts2(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vcmplts4(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vcmpltu2(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vcmpltu4(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vmaxs2(unsigned int x,
+                                                       unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vmaxs4(unsigned int x,
+                                                       unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vmaxu2(unsigned int x,
+                                                       unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vmaxu4(unsigned int x,
+                                                       unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vmins2(unsigned int x,
+                                                       unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vmins4(unsigned int x,
+                                                       unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vminu2(unsigned int x,
+                                                       unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vminu4(unsigned int x,
+                                                       unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vseteq2(unsigned int x,
+                                                        unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vseteq4(unsigned int x,
+                                                        unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsetne2(unsigned int x,
+                                                        unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsetne4(unsigned int x,
+                                                        unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsetges2(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsetges4(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsetgeu2(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsetgeu4(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsetgts2(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsetgts4(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsetgtu2(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsetgtu4(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsetles2(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsetles4(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsetleu2(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsetleu4(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsetlts2(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsetlts4(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsetltu2(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsetltu4(unsigned int x,
+                                                         unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsads2(unsigned int x,
+                                                       unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsads4(unsigned int x,
+                                                       unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsadu2(unsigned int x,
+                                                       unsigned int y);
+extern __DPCPP_SYCL_EXTERNAL unsigned int __imf_vsadu4(unsigned int x,
+                                                       unsigned int y);
 }
 #ifdef __GLIBC__
 extern "C" {
-extern SYCL_EXTERNAL void __assert_fail(const char *expr, const char *file,
-                                        unsigned int line, const char *func);
-extern SYCL_EXTERNAL float frexpf(float x, int *exp);
-extern SYCL_EXTERNAL float ldexpf(float x, int exp);
-extern SYCL_EXTERNAL float hypotf(float x, float y);
+extern __DPCPP_SYCL_EXTERNAL void __assert_fail(const char *expr,
+                                                const char *file,
+                                                unsigned int line,
+                                                const char *func);
+extern __DPCPP_SYCL_EXTERNAL float frexpf(float x, int *exp);
+extern __DPCPP_SYCL_EXTERNAL float ldexpf(float x, int exp);
+extern __DPCPP_SYCL_EXTERNAL float hypotf(float x, float y);
 
 // MS UCRT supports most of the C standard library but <complex.h> is
 // an exception.
-extern SYCL_EXTERNAL float cimagf(float __complex__ z);
-extern SYCL_EXTERNAL double cimag(double __complex__ z);
-extern SYCL_EXTERNAL float crealf(float __complex__ z);
-extern SYCL_EXTERNAL double creal(double __complex__ z);
-extern SYCL_EXTERNAL float cargf(float __complex__ z);
-extern SYCL_EXTERNAL double carg(double __complex__ z);
-extern SYCL_EXTERNAL float cabsf(float __complex__ z);
-extern SYCL_EXTERNAL double cabs(double __complex__ z);
-extern SYCL_EXTERNAL float __complex__ cprojf(float __complex__ z);
-extern SYCL_EXTERNAL double __complex__ cproj(double __complex__ z);
-extern SYCL_EXTERNAL float __complex__ cexpf(float __complex__ z);
-extern SYCL_EXTERNAL double __complex__ cexp(double __complex__ z);
-extern SYCL_EXTERNAL float __complex__ clogf(float __complex__ z);
-extern SYCL_EXTERNAL double __complex__ clog(double __complex__ z);
-extern SYCL_EXTERNAL float __complex__ cpowf(float __complex__ z);
-extern SYCL_EXTERNAL double __complex__ cpow(double __complex__ z);
-extern SYCL_EXTERNAL float __complex__ csqrtf(float __complex__ z);
-extern SYCL_EXTERNAL double __complex__ csqrt(double __complex__ z);
-extern SYCL_EXTERNAL float __complex__ csinhf(float __complex__ z);
-extern SYCL_EXTERNAL double __complex__ csinh(double __complex__ z);
-extern SYCL_EXTERNAL float __complex__ ccoshf(float __complex__ z);
-extern SYCL_EXTERNAL double __complex__ ccosh(double __complex__ z);
-extern SYCL_EXTERNAL float __complex__ ctanhf(float __complex__ z);
-extern SYCL_EXTERNAL double __complex__ ctanh(double __complex__ z);
-extern SYCL_EXTERNAL float __complex__ csinf(float __complex__ z);
-extern SYCL_EXTERNAL double __complex__ csin(double __complex__ z);
-extern SYCL_EXTERNAL float __complex__ ccosf(float __complex__ z);
-extern SYCL_EXTERNAL double __complex__ ccos(double __complex__ z);
-extern SYCL_EXTERNAL float __complex__ ctanf(float __complex__ z);
-extern SYCL_EXTERNAL double __complex__ ctan(double __complex__ z);
-extern SYCL_EXTERNAL float __complex__ cacosf(float __complex__ z);
-extern SYCL_EXTERNAL double __complex__ cacos(double __complex__ z);
-extern SYCL_EXTERNAL float __complex__ cacoshf(float __complex__ z);
-extern SYCL_EXTERNAL double __complex__ cacosh(double __complex__ z);
-extern SYCL_EXTERNAL float __complex__ casinf(float __complex__ z);
-extern SYCL_EXTERNAL double __complex__ casin(double __complex__ z);
-extern SYCL_EXTERNAL float __complex__ casinhf(float __complex__ z);
-extern SYCL_EXTERNAL double __complex__ casinh(double __complex__ z);
-extern SYCL_EXTERNAL float __complex__ catanf(float __complex__ z);
-extern SYCL_EXTERNAL double __complex__ catan(double __complex__ z);
-extern SYCL_EXTERNAL float __complex__ catanhf(float __complex__ z);
-extern SYCL_EXTERNAL double __complex__ catanh(double __complex__ z);
-extern SYCL_EXTERNAL float __complex__ cpolarf(float rho, float theta);
-extern SYCL_EXTERNAL double __complex__ cpolar(double rho, double theta);
-extern SYCL_EXTERNAL float __complex__ __mulsc3(float a, float b, float c,
-                                                float d);
-extern SYCL_EXTERNAL double __complex__ __muldc3(double a, double b, double c,
-                                                 double d);
-extern SYCL_EXTERNAL float __complex__ __divsc3(float a, float b, float c,
-                                                float d);
-extern SYCL_EXTERNAL double __complex__ __divdc3(float a, float b, float c,
-                                                 float d);
+extern __DPCPP_SYCL_EXTERNAL float cimagf(float __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL double cimag(double __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL float crealf(float __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL double creal(double __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL float cargf(float __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL double carg(double __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL float cabsf(float __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL double cabs(double __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL float __complex__ cprojf(float __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL double __complex__ cproj(double __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL float __complex__ cexpf(float __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL double __complex__ cexp(double __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL float __complex__ clogf(float __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL double __complex__ clog(double __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL float __complex__ cpowf(float __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL double __complex__ cpow(double __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL float __complex__ csqrtf(float __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL double __complex__ csqrt(double __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL float __complex__ csinhf(float __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL double __complex__ csinh(double __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL float __complex__ ccoshf(float __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL double __complex__ ccosh(double __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL float __complex__ ctanhf(float __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL double __complex__ ctanh(double __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL float __complex__ csinf(float __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL double __complex__ csin(double __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL float __complex__ ccosf(float __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL double __complex__ ccos(double __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL float __complex__ ctanf(float __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL double __complex__ ctan(double __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL float __complex__ cacosf(float __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL double __complex__ cacos(double __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL float __complex__ cacoshf(float __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL double __complex__ cacosh(double __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL float __complex__ casinf(float __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL double __complex__ casin(double __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL float __complex__ casinhf(float __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL double __complex__ casinh(double __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL float __complex__ catanf(float __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL double __complex__ catan(double __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL float __complex__ catanhf(float __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL double __complex__ catanh(double __complex__ z);
+extern __DPCPP_SYCL_EXTERNAL float __complex__ cpolarf(float rho, float theta);
+extern __DPCPP_SYCL_EXTERNAL double __complex__ cpolar(double rho,
+                                                       double theta);
+extern __DPCPP_SYCL_EXTERNAL float __complex__ __mulsc3(float a, float b,
+                                                        float c, float d);
+extern __DPCPP_SYCL_EXTERNAL double __complex__ __muldc3(double a, double b,
+                                                         double c, double d);
+extern __DPCPP_SYCL_EXTERNAL float __complex__ __divsc3(float a, float b,
+                                                        float c, float d);
+extern __DPCPP_SYCL_EXTERNAL double __complex__ __divdc3(float a, float b,
+                                                         float c, float d);
 }
 #elif defined(_WIN32)
 extern "C" {
@@ -2512,23 +2591,23 @@ extern "C" {
 //       APIs used by STL, such as _Cosh, are undocumented, even though
 //       they are open-sourced. Recognizing them as builtins is not
 //       straightforward currently.
-extern SYCL_EXTERNAL double _Cosh(double x, double y);
-extern SYCL_EXTERNAL int _dpcomp(double x, double y);
-extern SYCL_EXTERNAL int _dsign(double x);
-extern SYCL_EXTERNAL short _Dtest(double *px);
-extern SYCL_EXTERNAL short _dtest(double *px);
-extern SYCL_EXTERNAL short _Exp(double *px, double y, short eoff);
-extern SYCL_EXTERNAL float _FCosh(float x, float y);
-extern SYCL_EXTERNAL int _fdpcomp(float x, float y);
-extern SYCL_EXTERNAL int _fdsign(float x);
-extern SYCL_EXTERNAL short _FDtest(float *px);
-extern SYCL_EXTERNAL short _fdtest(float *px);
-extern SYCL_EXTERNAL short _FExp(float *px, float y, short eoff);
-extern SYCL_EXTERNAL float _FSinh(float x, float y);
-extern SYCL_EXTERNAL double _Sinh(double x, double y);
-extern SYCL_EXTERNAL float _hypotf(float x, float y);
-extern SYCL_EXTERNAL void _wassert(const wchar_t *wexpr, const wchar_t *wfile,
-                                   unsigned line);
+extern __DPCPP_SYCL_EXTERNAL double _Cosh(double x, double y);
+extern __DPCPP_SYCL_EXTERNAL int _dpcomp(double x, double y);
+extern __DPCPP_SYCL_EXTERNAL int _dsign(double x);
+extern __DPCPP_SYCL_EXTERNAL short _Dtest(double *px);
+extern __DPCPP_SYCL_EXTERNAL short _dtest(double *px);
+extern __DPCPP_SYCL_EXTERNAL short _Exp(double *px, double y, short eoff);
+extern __DPCPP_SYCL_EXTERNAL float _FCosh(float x, float y);
+extern __DPCPP_SYCL_EXTERNAL int _fdpcomp(float x, float y);
+extern __DPCPP_SYCL_EXTERNAL int _fdsign(float x);
+extern __DPCPP_SYCL_EXTERNAL short _FDtest(float *px);
+extern __DPCPP_SYCL_EXTERNAL short _fdtest(float *px);
+extern __DPCPP_SYCL_EXTERNAL short _FExp(float *px, float y, short eoff);
+extern __DPCPP_SYCL_EXTERNAL float _FSinh(float x, float y);
+extern __DPCPP_SYCL_EXTERNAL double _Sinh(double x, double y);
+extern __DPCPP_SYCL_EXTERNAL float _hypotf(float x, float y);
+extern __DPCPP_SYCL_EXTERNAL void _wassert(const wchar_t *wexpr,
+                                           const wchar_t *wfile, unsigned line);
 }
 #endif
 #endif // __SYCL_DEVICE_ONLY__
diff --git a/sycl/include/sycl/detail/assert_happened.hpp b/sycl/include/sycl/detail/assert_happened.hpp
index 43403a5b8a823..ae4e1b707ab48 100644
--- a/sycl/include/sycl/detail/assert_happened.hpp
+++ b/sycl/include/sycl/detail/assert_happened.hpp
@@ -14,7 +14,7 @@
 
 #ifdef __SYCL_DEVICE_ONLY__
 // Reads Flag of AssertHappened on device
-SYCL_EXTERNAL __attribute__((weak)) extern "C" void
+__DPCPP_SYCL_EXTERNAL __attribute__((weak)) extern "C" void
 __devicelib_assert_read(void *);
 #endif
 
diff --git a/sycl/include/sycl/detail/defines_elementary.hpp b/sycl/include/sycl/detail/defines_elementary.hpp
index cffc6076a176a..d894f77917746 100644
--- a/sycl/include/sycl/detail/defines_elementary.hpp
+++ b/sycl/include/sycl/detail/defines_elementary.hpp
@@ -26,9 +26,16 @@
 #endif
 #endif // __SYCL_ALWAYS_INLINE
 
-#ifndef SYCL_EXTERNAL
+#ifdef SYCL_EXTERNAL
+#define __DPCPP_SYCL_EXTERNAL SYCL_EXTERNAL
+#else
+#ifdef __SYCL_DEVICE_ONLY__
+#define __DPCPP_SYCL_EXTERNAL __attribute__((sycl_device))
+#else
+#define __DPCPP_SYCL_EXTERNAL
 #define SYCL_EXTERNAL
 #endif
+#endif
 
 #ifndef __SYCL_ID_QUERIES_FIT_IN_INT__
 #define __SYCL_ID_QUERIES_FIT_IN_INT__ 0
diff --git a/sycl/include/sycl/detail/sycl_fe_intrins.hpp b/sycl/include/sycl/detail/sycl_fe_intrins.hpp
index a79c496c3a225..d0ec6b955452f 100644
--- a/sycl/include/sycl/detail/sycl_fe_intrins.hpp
+++ b/sycl/include/sycl/detail/sycl_fe_intrins.hpp
@@ -19,10 +19,10 @@
 // Post-link tool traces the ID to a string literal it points to and assigns
 // integer ID.
 template <typename T>
-SYCL_EXTERNAL T __sycl_getScalarSpecConstantValue(const char *ID);
+__DPCPP_SYCL_EXTERNAL T __sycl_getScalarSpecConstantValue(const char *ID);
 
 template <typename T>
-SYCL_EXTERNAL T __sycl_getCompositeSpecConstantValue(const char *ID);
+__DPCPP_SYCL_EXTERNAL T __sycl_getCompositeSpecConstantValue(const char *ID);
 
 // The intrinsics below are used to implement support SYCL2020 specialization
 // constants. SYCL2020 version requires more parameters compared to the initial
@@ -36,16 +36,15 @@ SYCL_EXTERNAL T __sycl_getCompositeSpecConstantValue(const char *ID);
 // specialization constant and should be used if native specialization constants
 // are not available.
 template <typename T>
-SYCL_EXTERNAL T __sycl_getScalar2020SpecConstantValue(const char *SymbolicID,
-                                                      const void *DefaultValue,
-                                                      const void *RTBuffer);
+__DPCPP_SYCL_EXTERNAL T __sycl_getScalar2020SpecConstantValue(
+    const char *SymbolicID, const void *DefaultValue, const void *RTBuffer);
 
 template <typename T>
-SYCL_EXTERNAL T __sycl_getComposite2020SpecConstantValue(
+__DPCPP_SYCL_EXTERNAL T __sycl_getComposite2020SpecConstantValue(
     const char *SymbolicID, const void *DefaultValue, const void *RTBuffer);
 
 // Request a fixed-size allocation in local address space at kernel scope.
-extern "C" SYCL_EXTERNAL __attribute__((opencl_local)) std::uint8_t *
+extern "C" __DPCPP_SYCL_EXTERNAL __attribute__((opencl_local)) std::uint8_t *
 __sycl_allocateLocalMemory(std::size_t Size, std::size_t Alignment);
 
 #endif
diff --git a/sycl/include/sycl/ext/intel/esimd/detail/util.hpp b/sycl/include/sycl/ext/intel/esimd/detail/util.hpp
index 36bea945a42ca..bb740842fe22e 100644
--- a/sycl/include/sycl/ext/intel/esimd/detail/util.hpp
+++ b/sycl/include/sycl/ext/intel/esimd/detail/util.hpp
@@ -18,7 +18,7 @@
 #include <type_traits>
 
 #ifdef __SYCL_DEVICE_ONLY__
-#define __ESIMD_INTRIN SYCL_EXTERNAL SYCL_ESIMD_FUNCTION
+#define __ESIMD_INTRIN __DPCPP_SYCL_EXTERNAL SYCL_ESIMD_FUNCTION
 #else
 #define __ESIMD_INTRIN inline
 #endif // __SYCL_DEVICE_ONLY__
diff --git a/sycl/include/sycl/ext/intel/experimental/detail/misc_intrin.hpp b/sycl/include/sycl/ext/intel/experimental/detail/misc_intrin.hpp
index fbae9267158ed..be1862fa7a378 100644
--- a/sycl/include/sycl/ext/intel/experimental/detail/misc_intrin.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/detail/misc_intrin.hpp
@@ -13,7 +13,7 @@
 /// @cond SYCL_DETAIL
 
 #ifdef __SYCL_DEVICE_ONLY__
-#define __SYCL_INTRIN SYCL_EXTERNAL
+#define __SYCL_INTRIN __DPCPP_SYCL_EXTERNAL
 #else
 #define __SYCL_INTRIN inline
 #endif // __SYCL_DEVICE_ONLY__
diff --git a/sycl/include/sycl/ext/oneapi/bfloat16.hpp b/sycl/include/sycl/ext/oneapi/bfloat16.hpp
index 7f2b4c27ab39c..b9b0b8c31fa15 100644
--- a/sycl/include/sycl/ext/oneapi/bfloat16.hpp
+++ b/sycl/include/sycl/ext/oneapi/bfloat16.hpp
@@ -15,9 +15,9 @@
 #include <cmath>
 #endif
 
-extern "C" SYCL_EXTERNAL uint16_t
+extern "C" __DPCPP_SYCL_EXTERNAL uint16_t
 __devicelib_ConvertFToBF16INTEL(const float &) noexcept;
-extern "C" SYCL_EXTERNAL float
+extern "C" __DPCPP_SYCL_EXTERNAL float
 __devicelib_ConvertBF16ToFINTEL(const uint16_t &) noexcept;
 
 namespace sycl {
diff --git a/sycl/include/sycl/ext/oneapi/experimental/invoke_simd.hpp b/sycl/include/sycl/ext/oneapi/experimental/invoke_simd.hpp
index 69ec803fd232e..4bd5ee2946f66 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/invoke_simd.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/invoke_simd.hpp
@@ -41,7 +41,7 @@
 /// target function and the original SPMD arguments passed to invoke_simd.
 template <bool IsFunc, class SpmdRet, class HelperFunc,
           class... UserSimdFuncAndSpmdArgs, class = std::enable_if_t<!IsFunc>>
-SYCL_EXTERNAL __regcall SpmdRet
+__DPCPP_SYCL_EXTERNAL __regcall SpmdRet
 __builtin_invoke_simd(HelperFunc helper, const void *obj,
                       UserSimdFuncAndSpmdArgs... args)
 #ifdef __SYCL_DEVICE_ONLY__
@@ -56,7 +56,7 @@ __builtin_invoke_simd(HelperFunc helper, const void *obj,
 
 template <bool IsFunc, class SpmdRet, class HelperFunc,
           class... UserSimdFuncAndSpmdArgs, class = std::enable_if_t<IsFunc>>
-SYCL_EXTERNAL __regcall SpmdRet
+__DPCPP_SYCL_EXTERNAL __regcall SpmdRet
 __builtin_invoke_simd(HelperFunc helper, UserSimdFuncAndSpmdArgs... args)
 #ifdef __SYCL_DEVICE_ONLY__
     ;
@@ -193,7 +193,7 @@ template <class SimdCallable, class... SpmdArgs> struct sg_size {
   using IsInvocableSgSize = __MP11_NS::mp_bool<std::is_invocable_v<
       SimdCallable, typename spmd2simd<SpmdArgs, N::value>::type...>>;
 
-  SYCL_EXTERNAL constexpr operator int() {
+  __DPCPP_SYCL_EXTERNAL constexpr operator int() {
     using SupportedSgSizes = __MP11_NS::mp_list_c<int, 1, 2, 4, 8, 16, 32>;
     using InvocableSgSizes =
         __MP11_NS::mp_copy_if<SupportedSgSizes, IsInvocableSgSize>;
@@ -234,7 +234,7 @@ static constexpr int get_sg_size() {
 // with captures. Note __regcall - this is needed for efficient argument
 // forwarding.
 template <int N, class Callable, class... T>
-[[intel::device_indirectly_callable]] SYCL_EXTERNAL __regcall detail::
+[[intel::device_indirectly_callable]] __DPCPP_SYCL_EXTERNAL __regcall detail::
     SimdRetType<N, Callable, T...>
     simd_obj_call_helper(const void *obj_ptr,
                          typename detail::spmd2simd<T, N>::type... simd_args) {
@@ -245,7 +245,7 @@ template <int N, class Callable, class... T>
 
 // This function is a wrapper around a call to a function.
 template <int N, class Callable, class... T>
-[[intel::device_indirectly_callable]] SYCL_EXTERNAL __regcall detail::
+[[intel::device_indirectly_callable]] __DPCPP_SYCL_EXTERNAL __regcall detail::
     SimdRetType<N, Callable, T...>
     simd_func_call_helper(Callable f,
                           typename detail::spmd2simd<T, N>::type... simd_args) {
diff --git a/sycl/include/sycl/ext/oneapi/experimental/sycl_complex.hpp b/sycl/include/sycl/ext/oneapi/experimental/sycl_complex.hpp
index d09b9bcead51a..179c24c5a2fed 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/sycl_complex.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/sycl_complex.hpp
@@ -770,13 +770,13 @@ template <class _Tp> struct __libcpp_complex_overload_traits<_Tp, false, true> {
 // real
 
 template <class _Tp>
-SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr _Tp
+__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr _Tp
 real(const complex<_Tp> &__c) {
   return __c.real();
 }
 
 template <class _Tp>
-SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr
+__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr
     typename __libcpp_complex_overload_traits<_Tp>::_ValueType
     real(_Tp __re) {
   return __re;
@@ -785,13 +785,13 @@ SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr
 // imag
 
 template <class _Tp>
-SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr _Tp
+__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr _Tp
 imag(const complex<_Tp> &__c) {
   return __c.imag();
 }
 
 template <class _Tp>
-SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr
+__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr
     typename __libcpp_complex_overload_traits<_Tp>::_ValueType
     imag(_Tp) {
   return 0;
@@ -800,7 +800,7 @@ SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr
 // abs
 
 template <class _Tp>
-SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY _Tp
+__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY _Tp
 abs(const complex<_Tp> &__c) {
   return sycl::hypot(__c.real(), __c.imag());
 }
@@ -808,13 +808,13 @@ abs(const complex<_Tp> &__c) {
 // arg
 
 template <class _Tp>
-SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY _Tp
+__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY _Tp
 arg(const complex<_Tp> &__c) {
   return sycl::atan2(__c.imag(), __c.real());
 }
 
 template <class _Tp>
-SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY
+__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY
     typename enable_if<is_integral<_Tp>::value || is_same<_Tp, double>::value,
                        double>::type
     arg(_Tp __re) {
@@ -822,7 +822,7 @@ SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY
 }
 
 template <class _Tp>
-SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY
+__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY
     typename enable_if<is_same<_Tp, float>::value, float>::type
     arg(_Tp __re) {
   return sycl::atan2(0.F, __re);
@@ -831,7 +831,7 @@ SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY
 // norm
 
 template <class _Tp>
-SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY _Tp
+__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY _Tp
 norm(const complex<_Tp> &__c) {
   if (sycl::isinf(__c.real()))
     return sycl::fabs(__c.real());
@@ -841,7 +841,7 @@ norm(const complex<_Tp> &__c) {
 }
 
 template <class _Tp>
-SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY
+__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY
     typename __libcpp_complex_overload_traits<_Tp>::_ValueType
     norm(_Tp __re) {
   typedef typename __libcpp_complex_overload_traits<_Tp>::_ValueType _ValueType;
@@ -851,13 +851,13 @@ SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY
 // conj
 
 template <class _Tp>
-SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY complex<_Tp>
+__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY complex<_Tp>
 conj(const complex<_Tp> &__c) {
   return complex<_Tp>(__c.real(), -__c.imag());
 }
 
 template <class _Tp>
-SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY
+__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY
     typename __libcpp_complex_overload_traits<_Tp>::_ComplexType
     conj(_Tp __re) {
   typedef
@@ -868,7 +868,7 @@ SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY
 // proj
 
 template <class _Tp>
-SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY complex<_Tp>
+__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY complex<_Tp>
 proj(const complex<_Tp> &__c) {
   complex<_Tp> __r = __c;
   if (sycl::isinf(__c.real()) || sycl::isinf(__c.imag()))
@@ -877,7 +877,7 @@ proj(const complex<_Tp> &__c) {
 }
 
 template <class _Tp>
-SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY typename enable_if<
+__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY typename enable_if<
     is_floating_point<_Tp>::value,
     typename __libcpp_complex_overload_traits<_Tp>::_ComplexType>::type
 proj(_Tp __re) {
@@ -887,7 +887,7 @@ proj(_Tp __re) {
 }
 
 template <class _Tp>
-SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY typename enable_if<
+__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY typename enable_if<
     is_integral<_Tp>::value,
     typename __libcpp_complex_overload_traits<_Tp>::_ComplexType>::type
 proj(_Tp __re) {
@@ -899,7 +899,8 @@ proj(_Tp __re) {
 // polar
 
 template <class _Tp, class = std::enable_if<is_gencomplex<_Tp>::value>>
-SYCL_EXTERNAL complex<_Tp> polar(const _Tp &__rho, const _Tp &__theta = _Tp()) {
+__DPCPP_SYCL_EXTERNAL complex<_Tp> polar(const _Tp &__rho,
+                                         const _Tp &__theta = _Tp()) {
   if (sycl::isnan(__rho) || sycl::signbit(__rho))
     return complex<_Tp>(_Tp(NAN), _Tp(NAN));
   if (sycl::isnan(__theta)) {
@@ -924,7 +925,7 @@ SYCL_EXTERNAL complex<_Tp> polar(const _Tp &__rho, const _Tp &__theta = _Tp()) {
 // log
 
 template <class _Tp, class = std::enable_if<is_gencomplex<_Tp>::value>>
-SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY complex<_Tp>
+__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY complex<_Tp>
 log(const complex<_Tp> &__x) {
   return complex<_Tp>(sycl::log(abs(__x)), arg(__x));
 }
@@ -932,7 +933,7 @@ log(const complex<_Tp> &__x) {
 // log10
 
 template <class _Tp, class = std::enable_if<is_gencomplex<_Tp>::value>>
-SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY complex<_Tp>
+__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY complex<_Tp>
 log10(const complex<_Tp> &__x) {
   return log(__x) / sycl::log(_Tp(10));
 }
@@ -940,7 +941,7 @@ log10(const complex<_Tp> &__x) {
 // sqrt
 
 template <class _Tp, class = std::enable_if<is_gencomplex<_Tp>::value>>
-SYCL_EXTERNAL complex<_Tp> sqrt(const complex<_Tp> &__x) {
+__DPCPP_SYCL_EXTERNAL complex<_Tp> sqrt(const complex<_Tp> &__x) {
   if (sycl::isinf(__x.imag()))
     return complex<_Tp>(_Tp(INFINITY), __x.imag());
   if (sycl::isinf(__x.real())) {
@@ -957,7 +958,7 @@ SYCL_EXTERNAL complex<_Tp> sqrt(const complex<_Tp> &__x) {
 // exp
 
 template <class _Tp, class = std::enable_if<is_gencomplex<_Tp>::value>>
-SYCL_EXTERNAL complex<_Tp> exp(const complex<_Tp> &__x) {
+__DPCPP_SYCL_EXTERNAL complex<_Tp> exp(const complex<_Tp> &__x) {
   _Tp __i = __x.imag();
   if (__i == 0) {
     return complex<_Tp>(sycl::exp(__x.real()),
@@ -980,14 +981,14 @@ SYCL_EXTERNAL complex<_Tp> exp(const complex<_Tp> &__x) {
 // pow
 
 template <class _Tp, class = std::enable_if<is_gencomplex<_Tp>::value>>
-SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY complex<_Tp>
+__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY complex<_Tp>
 pow(const complex<_Tp> &__x, const complex<_Tp> &__y) {
   return exp(__y * log(__x));
 }
 
 template <class _Tp, class _Up,
           class = std::enable_if<is_gencomplex<_Tp>::value>>
-SYCL_EXTERNAL
+__DPCPP_SYCL_EXTERNAL
     _SYCL_EXT_CPLX_INLINE_VISIBILITY complex<typename __promote<_Tp, _Up>::type>
     pow(const complex<_Tp> &__x, const complex<_Up> &__y) {
   typedef complex<typename __promote<_Tp, _Up>::type> result_type;
@@ -997,7 +998,7 @@ SYCL_EXTERNAL
 
 template <class _Tp, class _Up,
           class = std::enable_if<is_gencomplex<_Tp>::value>>
-SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY
+__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY
     typename enable_if<is_genfloat<_Up>::value,
                        complex<typename __promote<_Tp, _Up>::type>>::type
     pow(const complex<_Tp> &__x, const _Up &__y) {
@@ -1008,7 +1009,7 @@ SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY
 
 template <class _Tp, class _Up,
           class = std::enable_if<is_gencomplex<_Tp>::value>>
-SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY
+__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY
     typename enable_if<is_genfloat<_Up>::value,
                        complex<typename __promote<_Tp, _Up>::type>>::type
     pow(const _Tp &__x, const complex<_Up> &__y) {
@@ -1028,7 +1029,7 @@ _SYCL_EXT_CPLX_INLINE_VISIBILITY complex<_Tp> __sqr(const complex<_Tp> &__x) {
 // asinh
 
 template <class _Tp, class = std::enable_if<is_gencomplex<_Tp>::value>>
-SYCL_EXTERNAL complex<_Tp> asinh(const complex<_Tp> &__x) {
+__DPCPP_SYCL_EXTERNAL complex<_Tp> asinh(const complex<_Tp> &__x) {
   const _Tp __pi(sycl::atan2(_Tp(+0.), _Tp(-0.)));
   if (sycl::isinf(__x.real())) {
     if (sycl::isnan(__x.imag()))
@@ -1056,7 +1057,7 @@ SYCL_EXTERNAL complex<_Tp> asinh(const complex<_Tp> &__x) {
 // acosh
 
 template <class _Tp, class = std::enable_if<is_gencomplex<_Tp>::value>>
-SYCL_EXTERNAL complex<_Tp> acosh(const complex<_Tp> &__x) {
+__DPCPP_SYCL_EXTERNAL complex<_Tp> acosh(const complex<_Tp> &__x) {
   const _Tp __pi(sycl::atan2(_Tp(+0.), _Tp(-0.)));
   if (sycl::isinf(__x.real())) {
     if (sycl::isnan(__x.imag()))
@@ -1089,7 +1090,7 @@ SYCL_EXTERNAL complex<_Tp> acosh(const complex<_Tp> &__x) {
 // atanh
 
 template <class _Tp, class = std::enable_if<is_gencomplex<_Tp>::value>>
-SYCL_EXTERNAL complex<_Tp> atanh(const complex<_Tp> &__x) {
+__DPCPP_SYCL_EXTERNAL complex<_Tp> atanh(const complex<_Tp> &__x) {
   const _Tp __pi(sycl::atan2(_Tp(+0.), _Tp(-0.)));
   if (sycl::isinf(__x.imag())) {
     return complex<_Tp>(sycl::copysign(_Tp(0), __x.real()),
@@ -1119,7 +1120,7 @@ SYCL_EXTERNAL complex<_Tp> atanh(const complex<_Tp> &__x) {
 // sinh
 
 template <class _Tp, class = std::enable_if<is_gencomplex<_Tp>::value>>
-SYCL_EXTERNAL complex<_Tp> sinh(const complex<_Tp> &__x) {
+__DPCPP_SYCL_EXTERNAL complex<_Tp> sinh(const complex<_Tp> &__x) {
   if (sycl::isinf(__x.real()) && !sycl::isfinite(__x.imag()))
     return complex<_Tp>(__x.real(), _Tp(NAN));
   if (__x.real() == 0 && !sycl::isfinite(__x.imag()))
@@ -1133,7 +1134,7 @@ SYCL_EXTERNAL complex<_Tp> sinh(const complex<_Tp> &__x) {
 // cosh
 
 template <class _Tp, class = std::enable_if<is_gencomplex<_Tp>::value>>
-SYCL_EXTERNAL complex<_Tp> cosh(const complex<_Tp> &__x) {
+__DPCPP_SYCL_EXTERNAL complex<_Tp> cosh(const complex<_Tp> &__x) {
   if (sycl::isinf(__x.real()) && !sycl::isfinite(__x.imag()))
     return complex<_Tp>(sycl::fabs(__x.real()), _Tp(NAN));
   if (__x.real() == 0 && !sycl::isfinite(__x.imag()))
@@ -1149,7 +1150,7 @@ SYCL_EXTERNAL complex<_Tp> cosh(const complex<_Tp> &__x) {
 // tanh
 
 template <class _Tp, class = std::enable_if<is_gencomplex<_Tp>::value>>
-SYCL_EXTERNAL complex<_Tp> tanh(const complex<_Tp> &__x) {
+__DPCPP_SYCL_EXTERNAL complex<_Tp> tanh(const complex<_Tp> &__x) {
   if (sycl::isinf(__x.real())) {
     if (!sycl::isfinite(__x.imag()))
       return complex<_Tp>(sycl::copysign(_Tp(1), __x.real()), _Tp(0));
@@ -1171,7 +1172,7 @@ SYCL_EXTERNAL complex<_Tp> tanh(const complex<_Tp> &__x) {
 // asin
 
 template <class _Tp, class = std::enable_if<is_gencomplex<_Tp>::value>>
-SYCL_EXTERNAL complex<_Tp> asin(const complex<_Tp> &__x) {
+__DPCPP_SYCL_EXTERNAL complex<_Tp> asin(const complex<_Tp> &__x) {
   complex<_Tp> __z = asinh(complex<_Tp>(-__x.imag(), __x.real()));
   return complex<_Tp>(__z.imag(), -__z.real());
 }
@@ -1179,7 +1180,7 @@ SYCL_EXTERNAL complex<_Tp> asin(const complex<_Tp> &__x) {
 // acos
 
 template <class _Tp, class = std::enable_if<is_gencomplex<_Tp>::value>>
-SYCL_EXTERNAL complex<_Tp> acos(const complex<_Tp> &__x) {
+__DPCPP_SYCL_EXTERNAL complex<_Tp> acos(const complex<_Tp> &__x) {
   const _Tp __pi(sycl::atan2(_Tp(+0.), _Tp(-0.)));
   if (sycl::isinf(__x.real())) {
     if (sycl::isnan(__x.imag()))
@@ -1213,7 +1214,7 @@ SYCL_EXTERNAL complex<_Tp> acos(const complex<_Tp> &__x) {
 // atan
 
 template <class _Tp, class = std::enable_if<is_gencomplex<_Tp>::value>>
-SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY complex<_Tp>
+__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY complex<_Tp>
 atan(const complex<_Tp> &__x) {
   complex<_Tp> __z = atanh(complex<_Tp>(-__x.imag(), __x.real()));
   return complex<_Tp>(__z.imag(), -__z.real());
@@ -1222,7 +1223,7 @@ atan(const complex<_Tp> &__x) {
 // sin
 
 template <class _Tp, class = std::enable_if<is_gencomplex<_Tp>::value>>
-SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY complex<_Tp>
+__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY complex<_Tp>
 sin(const complex<_Tp> &__x) {
   complex<_Tp> __z = sinh(complex<_Tp>(-__x.imag(), __x.real()));
   return complex<_Tp>(__z.imag(), -__z.real());
@@ -1231,7 +1232,7 @@ sin(const complex<_Tp> &__x) {
 // cos
 
 template <class _Tp, class = std::enable_if<is_gencomplex<_Tp>::value>>
-SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY complex<_Tp>
+__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY complex<_Tp>
 cos(const complex<_Tp> &__x) {
   return cosh(complex<_Tp>(-__x.imag(), __x.real()));
 }
@@ -1239,7 +1240,7 @@ cos(const complex<_Tp> &__x) {
 // tan
 
 template <class _Tp, class = std::enable_if<is_gencomplex<_Tp>::value>>
-SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY complex<_Tp>
+__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY complex<_Tp>
 tan(const complex<_Tp> &__x) {
   complex<_Tp> __z = tanh(complex<_Tp>(-__x.imag(), __x.real()));
   return complex<_Tp>(__z.imag(), -__z.real());
@@ -1303,8 +1304,9 @@ basic_ostream<_CharT, _Traits> &operator<<(basic_ostream<_CharT, _Traits> &__os,
 }
 
 template <class _Tp, class = std::enable_if<is_gencomplex<_Tp>::value>>
-SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY inline const sycl::stream &
-operator<<(const sycl::stream &__ss, const complex<_Tp> &_x) {
+__DPCPP_SYCL_EXTERNAL
+    _SYCL_EXT_CPLX_INLINE_VISIBILITY inline const sycl::stream &
+    operator<<(const sycl::stream &__ss, const complex<_Tp> &_x) {
   return __ss << "(" << _x.real() << "," << _x.imag() << ")";
 }
 
diff --git a/sycl/test/basic_tests/macros_no_rdc.cpp b/sycl/test/basic_tests/macros_no_rdc.cpp
new file mode 100644
index 0000000000000..78dc1bb21ce09
--- /dev/null
+++ b/sycl/test/basic_tests/macros_no_rdc.cpp
@@ -0,0 +1,22 @@
+// clang-format off
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsycl-device-only -E -dD -fno-sycl-rdc %s -o %t.device
+// RUN: %clangxx %fsycl-host-only -fno-sycl-rdc -E -dD %s -o %t.host
+//
+// RUN: FileCheck --match-full-lines %s < %t.device --check-prefixes=DEVICE-FULL-LINE --implicit-check-not="#define SYCL_EXTERNAL"
+// RUN: FileCheck --match-full-lines %s < %t.host --check-prefixes=HOST
+//
+// Remove __DPCPP_SYCL_EXTERNAL to simplify regex for DEVICE prefix
+// RUN: sed -i 's|__DPCPP_SYCL_EXTERNAL||g' %t.device
+// RUN: FileCheck %s < %t.device --check-prefixes=DEVICE
+//
+// With -fno-sycl-rdc, device code should not define or use SYCL_EXTERNAL
+// DEVICE-FULL-LINE: #define __DPCPP_SYCL_EXTERNAL __attribute__((sycl_device))
+// DEVICE-NOT:SYCL_EXTERNAL
+//
+// With -fno-sycl-rdc, host code should have SYCL_EXTERNAL defined to empty
+// HOST-DAG: #define SYCL_EXTERNAL
+// HOST-DAG: #define __DPCPP_SYCL_EXTERNAL
+#include <sycl/sycl.hpp>
+#include "ext/oneapi/bfloat16.hpp"
+#include "ext/intel/esimd.hpp"
+#include "ext/oneapi/experimental/sycl_complex.hpp"