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.


[ROCm/clr commit: d39508b4fd]

projects/clr/hipamd/CMakeLists.txt

@@ -193,8 +193,7 @@ if(HIP_PLATFORM STREQUAL "hcc")
         src/program_state.cpp)
 
     set(SOURCE_FILES_DEVICE
-        src/device_util.cpp
-        src/hip_ldg.cpp)
+        src/device_util.cpp)
 
     execute_process(COMMAND ${HCC_HOME}/bin/hcc-config --ldflags OUTPUT_VARIABLE HCC_LD_FLAGS)
     set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} ${HCC_LD_FLAGS} -Wl,-Bsymbolic")

projects/clr/hipamd/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

projects/clr/hipamd/src/hip_ldg.cpp

@@ -1,83 +0,0 @@
-/*
-Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
-
-Permission is hereby granted, free of charge, to any person obtaining a copy
-of this software and associated documentation files (the "Software"), to deal
-in the Software without restriction, including without limitation the rights
-to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
-copies of the Software, and to permit persons to whom the Software is
-furnished to do so, subject to the following conditions:
-
-The above copyright notice and this permission notice shall be included in
-all copies or substantial portions of the Software.
-
-THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
-IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
-FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
-AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
-LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
-OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
-THE SOFTWARE.
-*/
-
-#include "hip/hcc_detail/hip_ldg.h"
-#include "hip/hcc_detail/hip_vector_types.h"
-
-__device__ char __ldg(const char* ptr) { return *ptr; }
-
-__device__ char2 __ldg(const char2* ptr) { return *ptr; }
-
-__device__ char4 __ldg(const char4* ptr) { return *ptr; }
-
-__device__ signed char __ldg(const signed char* ptr) { return ptr[0]; }
-
-__device__ unsigned char __ldg(const unsigned char* ptr) { return ptr[0]; }
-
-__device__ short __ldg(const short* ptr) { return ptr[0]; }
-
-__device__ short2 __ldg(const short2* ptr) { return ptr[0]; }
-
-__device__ short4 __ldg(const short4* ptr) { return ptr[0]; }
-
-__device__ unsigned short __ldg(const unsigned short* ptr) { return ptr[0]; }
-
-__device__ int __ldg(const int* ptr) { return ptr[0]; }
-
-__device__ int2 __ldg(const int2* ptr) { return ptr[0]; }
-
-__device__ int4 __ldg(const int4* ptr) { return ptr[0]; }
-
-__device__ unsigned int __ldg(const unsigned int* ptr) { return ptr[0]; }
-
-
-__device__ long __ldg(const long* ptr) { return ptr[0]; }
-
-__device__ unsigned long __ldg(const unsigned long* ptr) { return ptr[0]; }
-
-__device__ long long __ldg(const long long* ptr) { return ptr[0]; }
-
-__device__ longlong2 __ldg(const longlong2* ptr) { return ptr[0]; }
-
-__device__ unsigned long long __ldg(const unsigned long long* ptr) { return ptr[0]; }
-
-__device__ uchar2 __ldg(const uchar2* ptr) { return ptr[0]; }
-
-__device__ uchar4 __ldg(const uchar4* ptr) { return ptr[0]; }
-
-__device__ ushort2 __ldg(const ushort2* ptr) { return ptr[0]; }
-
-__device__ uint2 __ldg(const uint2* ptr) { return ptr[0]; }
-
-__device__ uint4 __ldg(const uint4* ptr) { return ptr[0]; }
-
-__device__ ulonglong2 __ldg(const ulonglong2* ptr) { return ptr[0]; }
-
-__device__ float __ldg(const float* ptr) { return ptr[0]; }
-
-__device__ float2 __ldg(const float2* ptr) { return ptr[0]; }
-
-__device__ float4 __ldg(const float4* ptr) { return ptr[0]; }
-
-__device__ double __ldg(const double* ptr) { return ptr[0]; }
-
-__device__ double2 __ldg(const double2* ptr) { return ptr[0]; }

projects/clr/hipamd/tests/src/deviceLib/hip_test_ldg.cpp

@@ -35,7 +35,7 @@ THE SOFTWARE.
 #include "hip/hip_vector_types.h"
 #include "test_common.h"
 
-#if (__hcc_workweek__ >= 16164) || defined(__HIP_PLATFORM_NVCC__)
+#if (__hcc_workweek__ >= 16164) || defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_CLANG_ONLY__)
 
 #define HIP_ASSERT(x) (assert((x) == hipSuccess))