From 9a0c5d0653defaa629987f10c0afbaf771f75624 Mon Sep 17 00:00:00 2001
From: Sameer Sahasrabuddhe <41661541+ssahasra@users.noreply.github.com>
Date: Tue, 17 Mar 2020 14:03:27 +0530
Subject: [PATCH] enable HCC printf when using hip-clang

This is cherry-picked from PR#1947 that was committed to the
github repo. It allows printf to work with hip-clang and HCC
runtime.

Change-Id: I754753250ea1e694cf3441722e2d4c9d25fa75bc
---
 include/hip/hcc_detail/device_functions.h | 14 ++++++++++----
 include/hip/hcc_detail/hip_runtime.h      | 14 --------------
 2 files changed, 10 insertions(+), 18 deletions(-)

diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h
index 1be853bc99..e6549dde0d 100644
--- a/include/hip/hcc_detail/device_functions.h
+++ b/include/hip/hcc_detail/device_functions.h
@@ -34,13 +34,19 @@ THE SOFTWARE.
 #include <hip/hcc_detail/device_library_decls.h>
 #include <hip/hcc_detail/llvm_intrinsics.h>
 
-#if __HIP_CLANG_ONLY__
-#if __HIP_VDI__
+#if __HIP_CLANG_ONLY__ && __HIP_VDI__
 extern "C" __device__ int printf(const char *fmt, ...);
 #else
+#if HC_FEATURE_PRINTF
+template <typename... All>
+static inline __device__ void printf(const char* format, All... all) {
+    hc::printf(format, all...);
+}
+#else
+template <typename... All>
 static inline __device__ void printf(const char* format, All... all) {}
-#endif
-#endif
+#endif // HC_FEATURE_PRINTF
+#endif // __HIP_CLANG_ONLY__ && __HIP_VDI__
 
 /*
 Integer Intrinsics
diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h
index cbb3a0e99a..ef591b8c51 100644
--- a/include/hip/hcc_detail/hip_runtime.h
+++ b/include/hip/hcc_detail/hip_runtime.h
@@ -314,20 +314,6 @@ extern "C" __device__ void* __hip_free(void* ptr);
 static inline __device__ void* malloc(size_t size) { return __hip_malloc(size); }
 static inline __device__ void* free(void* ptr) { return __hip_free(ptr); }
 
-// Declare printf only for the HCC compiler. hip-clang is handled in
-// device_functions.h
-#if __HCC_ACCELERATOR__
-#if HC_FEATURE_PRINTF
-template <typename... All>
-static inline __device__ void printf(const char* format, All... all) {
-    hc::printf(format, all...);
-}
-#else
-template <typename... All>
-static inline __device__ void printf(const char* format, All... all) {}
-#endif // HC_FEATURE_PRINTF
-#endif // __HCC_ACCELERATOR__
-
 #endif //__HCC_OR_HIP_CLANG__
 
 #ifdef __HCC__