enable HCC printf when using hip-clang (#1947)

This allows printf to work with hip-clang and HCC runtime. See comments under #1919 for a reported bug and feature request.

[ROCm/hip commit: 899c878703]

projects/hip/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

projects/hip/include/hip/hcc_detail/hip_runtime.h

@@ -311,20 +311,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__