diff --git a/hipamd/include/hcc_detail/hip_ldg.h b/hipamd/include/hcc_detail/hip_ldg.h index 69ed0de823..4d128dbb02 100644 --- a/hipamd/include/hcc_detail/hip_ldg.h +++ b/hipamd/include/hcc_detail/hip_ldg.h @@ -21,10 +21,16 @@ THE SOFTWARE. #define HIP_LDG_H #if __HCC__ +<<<<<<< HEAD +#include"hip_vector_types.h" +#include"host_defines.h" +#if __hcc_workweek__ >= 16164 +======= #include"hip/hip_vector_types.h" #include"hip/hcc_detail/host_defines.h" +>>>>>>> 75532471b2d124c4af823e838bf983a20debe6c4 __device__ char __ldg(const char* ); __device__ char1 __ldg(const char1* ); __device__ char2 __ldg(const char2* ); @@ -94,6 +100,14 @@ __device__ double2 __ldg(const double2* ); __device__ double3 __ldg(const double3* ); __device__ double4 __ldg(const double4* ); +<<<<<<< HEAD +#endif // __hcc_workweek__ + +#endif // __HCC__ + +#endif // HIP_LDG_H +======= #endif #endif +>>>>>>> 75532471b2d124c4af823e838bf983a20debe6c4 diff --git a/hipamd/src/hip_ldg.cpp b/hipamd/src/hip_ldg.cpp index 1e35a5cd49..c376e0170a 100644 --- a/hipamd/src/hip_ldg.cpp +++ b/hipamd/src/hip_ldg.cpp @@ -18,6 +18,128 @@ THE SOFTWARE. */ #include"hcc_detail/hip_ldg.h" +<<<<<<< HEAD +#if __hcc_workweek__ >= 16164 +__device__ char __ldg(const char* ptr) +{ + return ptr[0]; +} + +__device__ char1 __ldg(const char1* ptr) +{ + return ptr[0]; +} + +__device__ char2 __ldg(const char2* ptr) +{ + return ptr[0]; +} + +__device__ char3 __ldg(const char3* ptr) +{ + return ptr[0]; +} + +__device__ char4 __ldg(const char4* ptr) +{ + return ptr[0]; +} + +__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__ short1 __ldg(const short1* ptr) +{ + return ptr[0]; +} + +__device__ short2 __ldg(const short2* ptr) +{ + return ptr[0]; +} + +__device__ short3 __ldg(const short3* 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__ int1 __ldg(const int1* ptr) +{ + return ptr[0]; +} + +__device__ int2 __ldg(const int2* ptr) +{ + return ptr[0]; +} + +__device__ int3 __ldg(const int3* 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__ long1 __ldg(const long1* ptr) +{ + return ptr[0]; +} + +__device__ long2 __ldg(const long2* ptr) +{ + return ptr[0]; +} + +__device__ long3 __ldg(const long3* ptr) +{ + return ptr[0]; +} + +__device__ long4 __ldg(const long4* ptr) +======= __device__ char __ldg(const char* ptr) { @@ -25,37 +147,210 @@ __device__ char __ldg(const char* ptr) } __device__ signed char __ldg(const signed char* ptr) +>>>>>>> 75532471b2d124c4af823e838bf983a20debe6c4 { return ptr[0]; } +<<<<<<< HEAD +__device__ unsigned long __ldg(const unsigned long* ptr) +======= __device__ short __ldg(const short* ptr) +>>>>>>> 75532471b2d124c4af823e838bf983a20debe6c4 { return ptr[0]; } +<<<<<<< HEAD + +__device__ long long __ldg(const long long* ptr) +======= __device__ int __ldg(const int* ptr) +>>>>>>> 75532471b2d124c4af823e838bf983a20debe6c4 { return ptr[0]; } +<<<<<<< HEAD +__device__ longlong1 __ldg(const longlong1* ptr) +======= __device__ long long __ldg(const long long* ptr) +>>>>>>> 75532471b2d124c4af823e838bf983a20debe6c4 { return ptr[0]; } +<<<<<<< HEAD +__device__ longlong2 __ldg(const longlong2* ptr) +{ + return ptr[0]; +} + +__device__ longlong3 __ldg(const longlong3* ptr) +======= __device__ int2 __ldg(const int2* ptr) +>>>>>>> 75532471b2d124c4af823e838bf983a20debe6c4 { return ptr[0]; } +<<<<<<< HEAD +__device__ longlong4 __ldg(const longlong4* ptr) +{ + return ptr[0]; +} + +__device__ unsigned long long __ldg(const unsigned long long* ptr) +{ + return ptr[0]; +} + + +__device__ uchar1 __ldg(const uchar1* ptr) +======= __device__ int4 __ldg(const int4* ptr) +>>>>>>> 75532471b2d124c4af823e838bf983a20debe6c4 { return ptr[0]; } +<<<<<<< HEAD +__device__ uchar2 __ldg(const uchar2* ptr) +{ + return ptr[0]; +} + +__device__ uchar3 __ldg(const uchar3* ptr) +{ + return ptr[0]; +} + +__device__ uchar4 __ldg(const uchar4* ptr) +{ + return ptr[0]; +} + + +__device__ ushort1 __ldg(const ushort1* ptr) +{ + return ptr[0]; +} + +__device__ ushort2 __ldg(const ushort2* ptr) +{ + return ptr[0]; +} + +__device__ ushort3 __ldg(const ushort3* ptr) +{ + return ptr[0]; +} + +__device__ ushort4 __ldg(const ushort4* ptr) +{ + return ptr[0]; +} + + +__device__ uint1 __ldg(const uint1* ptr) +{ + return ptr[0]; +} + +__device__ uint2 __ldg(const uint2* ptr) +{ + return ptr[0]; +} + +__device__ uint3 __ldg(const uint3* ptr) +{ + return ptr[0]; +} + +__device__ uint4 __ldg(const uint4* ptr) +{ + return ptr[0]; +} + + +__device__ ulonglong1 __ldg(const ulonglong1* ptr) +{ + return ptr[0]; +} + +__device__ ulonglong2 __ldg(const ulonglong2* ptr) +{ + return ptr[0]; +} + +__device__ ulonglong3 __ldg(const ulonglong3* ptr) +{ + return ptr[0]; +} + +__device__ ulonglong4 __ldg(const ulonglong4* ptr) +{ + return ptr[0]; +} + + +__device__ float __ldg(const float* ptr) +{ + return ptr[0]; +} + +__device__ float1 __ldg(const float1* ptr) +{ + return ptr[0]; +} + +__device__ float2 __ldg(const float2* ptr) +{ + return ptr[0]; +} + +__device__ float3 __ldg(const float3* ptr) +{ + return ptr[0]; +} + +__device__ float4 __ldg(const float4* ptr) +{ + return ptr[0]; +} + + +__device__ double __ldg(const double* ptr) +{ + return ptr[0]; +} + +__device__ double1 __ldg(const double1* ptr) +{ + return ptr[0]; +} + +__device__ double2 __ldg(const double2* ptr) +{ + return ptr[0]; +} + +__device__ double3 __ldg(const double3* ptr) +{ + return ptr[0]; +} + +__device__ double4 __ldg(const double4* ptr) +{ + return ptr[0]; +} + +#endif + +======= __device__ float __ldg(const float* ptr) { return ptr[0]; } +>>>>>>> 75532471b2d124c4af823e838bf983a20debe6c4