separate printf declaration for vdi/clang

There are now two implementations of printf in HIP:

1. The implemenation for HCC is controlled by the HC_FEATURE_PRINTF
   macro, and it works only with the HCC compiler used in combination
   with the HCC runtime.

2. The implementation for hip-clang requires the VDI runtime, and is
   always enabled with that combination.

hipamd/bin/hipcc

@@ -133,6 +133,7 @@ if (defined $HIP_RUNTIME and $HIP_RUNTIME eq "VDI" and !defined $HIP_VDI_HOME) {
     } else {
         $HIP_VDI_HOME = $HIP_PATH; # use HIP_PATH
     }
+    $HIPCXXFLAGS .= "-D__HIP_VDI__";
 }
 
 if (defined $HIP_VDI_HOME) {

hipamd/include/hip/hcc_detail/device_functions.h

@@ -33,6 +33,15 @@ THE SOFTWARE.
 #include <hip/hip_vector_types.h>
 #include <hip/hcc_detail/device_library_decls.h>
 #include <hip/hcc_detail/llvm_intrinsics.h>
+
+#if __HIP_CLANG_ONLY__
+#if __HIP_VDI__
+extern "C" __device__ int printf(const char *fmt, ...);
+#else
+static inline __device__ void printf(const char* format, All... all) {}
+#endif
+#endif
+
 /*
 Integer Intrinsics
 */

hipamd/include/hip/hcc_detail/hip_runtime.h

@@ -311,15 +311,19 @@ 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); }
 
-#if defined(__HCC_ACCELERATOR__) && defined(HC_FEATURE_PRINTF)
+// 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...);
 }
-#elif defined(__HCC_ACCELERATOR__) || __HIP__
+#else
 template <typename... All>
 static inline __device__ void printf(const char* format, All... all) {}
-#endif
+#endif // HC_FEATURE_PRINTF
+#endif // __HCC_ACCELERATOR__
 
 #endif //__HCC_OR_HIP_CLANG__