Implement hip_ldg Functions into HIP header

Move all the function definitions for hip_ldg.cpp into hip_ldg.h header and enable for HIP clang path.

include/hip/hcc_detail/hip_ldg.h

@@ -23,54 +23,81 @@ THE SOFTWARE.
 #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_LDG_H
 #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_LDG_H
 
-#if defined __HCC__
-#if __hcc_workweek__ >= 16164
+#if defined(__HCC_OR_HIP_CLANG__)
+#if __hcc_workweek__ >= 16164 || defined(__HIP_CLANG_ONLY__)
 #include "hip_vector_types.h"
 #include "host_defines.h"
 
-__device__ char __ldg(const char*);
-__device__ char2 __ldg(const char2*);
-__device__ char4 __ldg(const char4*);
-__device__ signed char __ldg(const signed char*);
-__device__ unsigned char __ldg(const unsigned char*);
+__device__ inline static char __ldg(const char* ptr) { return *ptr; }
 
-__device__ short __ldg(const short*);
-__device__ short2 __ldg(const short2*);
-__device__ short4 __ldg(const short4*);
-__device__ unsigned short __ldg(const unsigned short*);
+__device__ inline static char2 __ldg(const char2* ptr) { return *ptr; }
 
-__device__ int __ldg(const int*);
-__device__ int2 __ldg(const int2*);
-__device__ int4 __ldg(const int4*);
-__device__ unsigned int __ldg(const unsigned int*);
+__device__ inline static char4 __ldg(const char4* ptr) { return *ptr; }
+
+__device__ inline static signed char __ldg(const signed char* ptr) { return ptr[0]; }
+
+__device__ inline static unsigned char __ldg(const unsigned char* ptr) { return ptr[0]; }
 
 
-__device__ long __ldg(const long*);
-__device__ unsigned long __ldg(const unsigned long*);
+__device__ inline static short __ldg(const short* ptr) { return ptr[0]; }
 
-__device__ long long __ldg(const long long*);
-__device__ longlong2 __ldg(const longlong2*);
-__device__ unsigned long long __ldg(const unsigned long long*);
+__device__ inline static short2 __ldg(const short2* ptr) { return ptr[0]; }
 
-__device__ uchar2 __ldg(const uchar2*);
-__device__ uchar4 __ldg(const uchar4*);
+__device__ inline static short4 __ldg(const short4* ptr) { return ptr[0]; }
 
-__device__ ushort2 __ldg(const ushort2*);
+__device__ inline static unsigned short __ldg(const unsigned short* ptr) { return ptr[0]; }
 
-__device__ uint2 __ldg(const uint2*);
-__device__ uint4 __ldg(const uint4*);
 
-__device__ ulonglong2 __ldg(const ulonglong2*);
+__device__ inline static int __ldg(const int* ptr) { return ptr[0]; }
 
-__device__ float __ldg(const float*);
-__device__ float2 __ldg(const float2*);
-__device__ float4 __ldg(const float4*);
+__device__ inline static int2 __ldg(const int2* ptr) { return ptr[0]; }
 
-__device__ double __ldg(const double*);
-__device__ double2 __ldg(const double2*);
+__device__ inline static int4 __ldg(const int4* ptr) { return ptr[0]; }
 
-#endif  // __hcc_workweek__
+__device__ inline static unsigned int __ldg(const unsigned int* ptr) { return ptr[0]; }
 
-#endif  // __HCC__
+
+__device__ inline static long __ldg(const long* ptr) { return ptr[0]; }
+
+__device__ inline static unsigned long __ldg(const unsigned long* ptr) { return ptr[0]; }
+
+
+__device__ inline static long long __ldg(const long long* ptr) { return ptr[0]; }
+
+__device__ inline static longlong2 __ldg(const longlong2* ptr) { return ptr[0]; }
+
+__device__ inline static unsigned long long __ldg(const unsigned long long* ptr) { return ptr[0]; }
+
+
+__device__ inline static uchar2 __ldg(const uchar2* ptr) { return ptr[0]; }
+
+__device__ inline static uchar4 __ldg(const uchar4* ptr) { return ptr[0]; }
+
+
+__device__ inline static ushort2 __ldg(const ushort2* ptr) { return ptr[0]; }
+
+
+__device__ inline static uint2 __ldg(const uint2* ptr) { return ptr[0]; }
+
+__device__ inline static uint4 __ldg(const uint4* ptr) { return ptr[0]; }
+
+
+__device__ inline static ulonglong2 __ldg(const ulonglong2* ptr) { return ptr[0]; }
+
+
+__device__ inline static float __ldg(const float* ptr) { return ptr[0]; }
+
+__device__ inline static float2 __ldg(const float2* ptr) { return ptr[0]; }
+
+__device__ inline static float4 __ldg(const float4* ptr) { return ptr[0]; }
+
+
+__device__ inline static double __ldg(const double* ptr) { return ptr[0]; }
+
+__device__ inline static double2 __ldg(const double2* ptr) { return ptr[0]; }
+
+#endif  // __hcc_workweek__ || defined(__HIP_CLANG_ONLY__)
+
+#endif  // defined(__HCC_OR_HIP_CLANG__)
 
 #endif  // HIP_LDG_H