diff --git a/hipamd/bin/hipcc b/hipamd/bin/hipcc index 83774a996b..013179b49e 100755 --- a/hipamd/bin/hipcc +++ b/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) { diff --git a/hipamd/include/hip/hcc_detail/device_functions.h b/hipamd/include/hip/hcc_detail/device_functions.h index b96bbba07f..1be853bc99 100644 --- a/hipamd/include/hip/hcc_detail/device_functions.h +++ b/hipamd/include/hip/hcc_detail/device_functions.h @@ -33,6 +33,15 @@ THE SOFTWARE. #include #include #include + +#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 */ diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h index e86611b213..c1ad5b2fe5 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/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 static inline __device__ void printf(const char* format, All... all) { hc::printf(format, all...); } -#elif defined(__HCC_ACCELERATOR__) || __HIP__ +#else template static inline __device__ void printf(const char* format, All... all) {} -#endif +#endif // HC_FEATURE_PRINTF +#endif // __HCC_ACCELERATOR__ #endif //__HCC_OR_HIP_CLANG__