Revert "added full data type support for __ldg"

This reverts commit 552acd5ed5.

Conflicts:
	include/hcc_detail/hip_ldg.h


[ROCm/hip commit: 75532471b2]

projects/hip/src/hip_ldg.cpp

@@ -19,300 +19,43 @@ THE SOFTWARE.
 
 #include"hcc_detail/hip_ldg.h"
 
-__device__ char                 __ldg(const char* ptr)
+__device__ char __ldg(const char* ptr)
+{
+    return *ptr;
+}
+
+__device__ signed char __ldg(const signed char* ptr)
 {
     return ptr[0];
 }
 
-__device__ char1                __ldg(const char1* ptr)
+__device__ short __ldg(const short* ptr)
 {
     return ptr[0];
 }
 
-__device__ char2                __ldg(const char2* ptr)
+__device__ int __ldg(const int* 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)
+__device__ long long __ldg(const long long* ptr)
 {
     return ptr[0];
 }
 
 
-__device__ short                __ldg(const short* ptr)
+__device__ int2 __ldg(const int2* ptr)
 {
     return ptr[0];
 }
 
-__device__ short1               __ldg(const short1* ptr)
+__device__ int4 __ldg(const int4* ptr)
 {
     return ptr[0];
 }
 
-__device__ short2               __ldg(const short2* ptr)
+__device__ float __ldg(const float* 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)
-{
-    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__ longlong1            __ldg(const longlong1* ptr)
-{
-    return ptr[0];
-}
-
-__device__ longlong2            __ldg(const longlong2* ptr)
-{
-    return ptr[0];
-}
-
-__device__ longlong3            __ldg(const longlong3* ptr)
-{
-    return ptr[0];
-}
-
-__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)
-{
-    return ptr[0];
-}
-
-__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];
-}
-
-
-

projects/hip/tests/src/hip_test_ldg.cpp

@@ -149,151 +149,14 @@ int main() {
   cout << " System major " << devProp.major << endl;
   cout << " agent prop name " << devProp.name << endl;
 
-    int errors;
-
-    errors =    dataTypesRun<char>() &
-                dataTypesRun<char1>() &
-                dataTypesRun<char2>() &
-                dataTypesRun<char3>() &
-                dataTypesRun<char4>() &
-                dataTypesRun<signed char>() &
-                dataTypesRun<unsigned char>();
-
-    if(errors == 1){
-        errors = 0;
-    }else{
-        std::cout<<"Failed Char"<<std::endl;
-        return -1;
-    }
-
-    errors =    dataTypesRun<short>() &
-                dataTypesRun<short1>() &
-                dataTypesRun<short2>() &
-                dataTypesRun<short3>() &
-                dataTypesRun<short4>() &
-                dataTypesRun<unsigned short>();
-
-    if(errors == 1){
-        errors = 0;
-    }else{
-        std::cout<<"Failed Short"<<std::endl;
-        return -1;
-    }
-
-    errors =    dataTypesRun<int>() &
-                dataTypesRun<int1>() &
-                dataTypesRun<int2>() &
-                dataTypesRun<int3>() &
-                dataTypesRun<int4>() &
-                dataTypesRun<unsigned int>();
-
-    if(errors == 1){
-        errors = 0;
-    }else{
-        std::cout<<"Failed Int"<<std::endl;
-        return -1;
-    }
-
-    errors =    dataTypesRun<long>() &
-                dataTypesRun<long1>() &
-                dataTypesRun<long2>() &
-                dataTypesRun<long3>() &
-                dataTypesRun<long4>() &
-                dataTypesRun<unsigned long>();
-
-    if(errors == 1){
-        errors = 0;
-    }else{
-        std::cout<<"Failed Long"<<std::endl;
-        return -1;
-    }
-
-    errors =    dataTypesRun<long long>() &
-                dataTypesRun<longlong1>() &
-                dataTypesRun<longlong2>() &
-                dataTypesRun<longlong3>() &
-                dataTypesRun<longlong4>() &
-                dataTypesRun<unsigned long long>();
-
-    if(errors == 1){
-        errors = 0;
-    }else{
-        std::cout<<"Failed Long Long"<<std::endl;
-        return -1;
-    }
-
-    errors =    dataTypesRun<uchar1>() &
-                dataTypesRun<uchar2>() &
-                dataTypesRun<uchar3>() &
-                dataTypesRun<uchar4>();
-
-    if(errors == 1){
-        errors = 0;
-    }else{
-        std::cout<<"Failed Unsigned Char"<<std::endl;
-        return -1;
-    }
-
-    errors =    dataTypesRun<ushort1>() &
-                dataTypesRun<ushort2>() &
-                dataTypesRun<ushort3>() &
-                dataTypesRun<ushort4>();
-
-    if(errors == 1){
-        errors = 0;
-    }else{
-        std::cout<<"Failed Unsigned Short"<<std::endl;
-        return -1;
-    }
-
-    errors =    dataTypesRun<uint1>() &
-                dataTypesRun<uint2>() &
-                dataTypesRun<uint3>() &
-                dataTypesRun<uint4>();
-
-    if(errors == 1){
-        errors = 0;
-    }else{
-        std::cout<<"Failed Unsigned Int"<<std::endl;
-        return -1;
-    }
-
-    errors =    dataTypesRun<ulonglong1>() &
-                dataTypesRun<ulonglong2>() &
-                dataTypesRun<ulonglong3>() &
-                dataTypesRun<ulonglong4>();
-
-    if(errors == 1){
-        errors = 0;
-    }else{
-        std::cout<<"Failed Unsigned Long Long"<<std::endl;
-        return -1;
-    }
-
-    errors =    dataTypesRun<float>() &
-                dataTypesRun<float1>() &
-                dataTypesRun<float2>() &
-                dataTypesRun<float3>() &
-                dataTypesRun<float4>();
-
-    if(errors == 1){
-        errors = 0;
-    }else{
-        std::cout<<"Failed Float"<<std::endl;
-        return -1;
-    }
-
-    errors =    dataTypesRun<double>() &
-                dataTypesRun<double1>() &
-                dataTypesRun<double2>() &
-                dataTypesRun<double3>() &
-                dataTypesRun<double4>();
-
+   int errors = dataTypesRun<char>() &
+           dataTypesRun<signed char>() &
+           dataTypesRun<short>() &
+           dataTypesRun<int>();
   //hipResetDefaultAccelerator();
     if(errors == 1){
         passed();
-    }else{
-        std::cout<<"Failed Float"<<std::endl;
-        return -1;
+        return 0;
     }
+
 }