diff --git a/clang/lib/Sema/SPIRVBuiltins.td b/clang/lib/Sema/SPIRVBuiltins.td
index 2c3410c984e86..30872777e3102 100644
--- a/clang/lib/Sema/SPIRVBuiltins.td
+++ b/clang/lib/Sema/SPIRVBuiltins.td
@@ -722,12 +722,6 @@ foreach VSize1 = [Vec2, Vec4, Vec8, Vec16] in {
 
 // 2.8. Misc instructions
 
-let IsVariadic = 1 in {
-  foreach name = ["printf"] in {
-    def : OCLSPVBuiltin<name, [Int, PointerType<ConstType<TrueChar>, ConstantAS>]>;
-  }
-}
-
 foreach name = ["prefetch"] in {
   def : OCLSPVBuiltin<name, [Void, PointerType<ConstType<AGenTypeN>, GlobalAS>, Size]>;
 }
diff --git a/sycl/include/CL/__spirv/spirv_ops.hpp b/sycl/include/CL/__spirv/spirv_ops.hpp
index 52a2ee5a7af75..d57b28d53791b 100644
--- a/sycl/include/CL/__spirv/spirv_ops.hpp
+++ b/sycl/include/CL/__spirv/spirv_ops.hpp
@@ -636,6 +636,16 @@ extern 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;
 
+#ifdef __SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__
+template <typename... Args>
+extern SYCL_EXTERNAL int
+__spirv_ocl_printf(const __attribute__((opencl_constant)) char *Format,
+                   Args... args);
+#else
+extern SYCL_EXTERNAL int
+__spirv_ocl_printf(const __attribute__((opencl_constant)) char *Format, ...);
+#endif
+
 #else // if !__SYCL_DEVICE_ONLY__
 
 template <typename dataT>
diff --git a/sycl/test/extensions/experimental-printf.cpp b/sycl/test/extensions/experimental-printf.cpp
new file mode 100644
index 0000000000000..652b94c2a06e9
--- /dev/null
+++ b/sycl/test/extensions/experimental-printf.cpp
@@ -0,0 +1,40 @@
+// This test is intended to check that internal
+// __SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__ works as expected, i.e. we can
+// see printf ExtInst regardless of the macro presence and that argument
+// promotion is disabled if the macro is present.
+//
+// RUN: %clangxx -fsycl -fsycl-device-only -fno-sycl-use-bitcode %s -o %t.spv
+// RUN: llvm-spirv -to-text %t.spv -o %t.spt
+// RUN: FileCheck %s --check-prefixes CHECK,CHECK-DOUBLE < %t.spt
+//
+// RUN: %clangxx -fsycl -fsycl-device-only -fno-sycl-use-bitcode -D__SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__ %s -o %t.spv
+// RUN: llvm-spirv -to-text %t.spv -o %t.spt
+// RUN: FileCheck %s --check-prefixes CHECK,CHECK-FLOAT < %t.spt
+
+// CHECK-FLOAT: TypeFloat [[#TYPE:]] 32
+// CHECK-DOUBLE: TypeFloat [[#TYPE:]] 64
+// CHECK: Constant [[#TYPE]] [[#CONST:]]
+// CHECK: ExtInst [[#]] [[#]] [[#]] printf [[#]] [[#CONST]]
+
+#include <CL/sycl.hpp>
+
+#ifdef __SYCL_DEVICE_ONLY__
+#define __SYCL_CONSTANT_AS __attribute__((opencl_constant))
+#else
+#define __SYCL_CONSTANT_AS
+#endif
+
+const __SYCL_CONSTANT_AS char fmt[] = "Hello, World! %f\n";
+
+int main() {
+  sycl::queue q;
+
+  q.submit([&](sycl::handler &cgh) {
+    cgh.single_task([=]() {
+      float f = 3.14;
+      sycl::ext::oneapi::experimental::printf(fmt, f);
+    });
+  });
+
+  return 0;
+}