From 23a91cd8aaf4b342bbf57bfa05581bc240ae2e53 Mon Sep 17 00:00:00 2001
From: Aditya Atluri <adityaavinash1@gmail.com>
Date: Thu, 21 Apr 2016 11:17:26 -0500
Subject: [PATCH] added full data type support for __ldg

[ROCm/clr commit: 3aac6d45adb968545589adda72265e12ce889065]
---
 .../clr/hipamd/include/hcc_detail/hip_ldg.h   |  98 ++++--
 projects/clr/hipamd/src/hip_ldg.cpp           | 281 +++++++++++++++++-
 .../clr/hipamd/tests/src/hip_test_ldg.cpp     | 149 +++++++++-
 3 files changed, 481 insertions(+), 47 deletions(-)

diff --git a/projects/clr/hipamd/include/hcc_detail/hip_ldg.h b/projects/clr/hipamd/include/hcc_detail/hip_ldg.h
index d996a978f2..50d5bcbf7a 100644
--- a/projects/clr/hipamd/include/hcc_detail/hip_ldg.h
+++ b/projects/clr/hipamd/include/hcc_detail/hip_ldg.h
@@ -23,35 +23,75 @@ THE SOFTWARE.
 #if __HCC__
 #include"hip_vector_types.h"
 #include"host_defines.h"
-__device__ char            __ldg(const char* );
-__device__ signed char     __ldg(const signed char* );
-__device__ short           __ldg(const short* );
-__device__ int             __ldg(const int* );
-__device__ long            __ldg(const long* );
-__device__ long long       __ldg(const long long* );
-__device__ char2           __ldg(const char2* );
-__device__ char4           __ldg(const char4* );
-__device__ short2          __ldg(const short2* );
-__device__ short4          __ldg(const short4* );
-__device__ int2            __ldg(const int2* );
-__device__ int4            __ldg(const int4* );
-__device__ longlong2       __ldg(const longlong2* );
-__device__ unsigned char   __ldg(const unsigned char* );
-__device__ unsigned short  __ldg(const unsigned short* );
-__device__ unsigned int    __ldg(const unsigned int* );
-__device__ unsigned long   __ldg(const unsigned long* );
-__device__ unsigned long long __ldg(const unsigned long long* );
-__device__ uchar2          __ldg(const uchar2* );
-__device__ uchar4          __ldg(const uchar4* );
-__device__ ushort2         __ldg(const ushort2* );
-__device__ uint2           __ldg(const uint2* );
-__device__ uint4           __ldg(const uint4* );
-__device__ ulonglong2      __ldg(const ulonglong2* );
-__device__ float           __ldg(const float* );
-__device__ double          __ldg(const double* );
-__device__ float2          __ldg(const float2* );
-__device__ float4          __ldg(const float4* );
-__device__ double2         __ldg(const double2* );
+
+__device__ char                 __ldg(const char* );
+__device__ char1                __ldg(const char1* );
+__device__ char2                __ldg(const char2* );
+__device__ char3                __ldg(const char3* );
+__device__ char4                __ldg(const char4* );
+__device__ signed char          __ldg(const signed char* );
+__device__ unsigned char        __ldg(const unsigned char* );
+
+__device__ short                __ldg(const short* );
+__device__ short1               __ldg(const short1* );
+__device__ short2               __ldg(const short2* );
+__device__ short3               __ldg(const short3* );
+__device__ short4               __ldg(const short4* );
+__device__ unsigned short       __ldg(const unsigned short* );
+
+__device__ int                  __ldg(const int* );
+__device__ int1                 __ldg(const int1* );
+__device__ int2                 __ldg(const int2* );
+__device__ int3                 __ldg(const int3* );
+__device__ int4                 __ldg(const int4* );
+__device__ unsigned int         __ldg(const unsigned int* );
+
+
+__device__ long                 __ldg(const long* );
+__device__ long1                __ldg(const long1* );
+__device__ long2                __ldg(const long2* );
+__device__ long3                __ldg(const long3* );
+__device__ long4                __ldg(const long4* );
+__device__ unsigned long        __ldg(const unsigned long* );
+
+__device__ long long            __ldg(const long long* );
+__device__ longlong1            __ldg(const longlong1* );
+__device__ longlong2            __ldg(const longlong2* );
+__device__ longlong3            __ldg(const longlong3* );
+__device__ longlong4            __ldg(const longlong4* );
+__device__ unsigned long long   __ldg(const unsigned long long* );
+
+__device__ uchar1               __ldg(const uchar1* );
+__device__ uchar2               __ldg(const uchar2* );
+__device__ uchar3               __ldg(const uchar3* );
+__device__ uchar4               __ldg(const uchar4* );
+
+__device__ ushort1              __ldg(const ushort1* );
+__device__ ushort2              __ldg(const ushort2* );
+__device__ ushort3              __ldg(const ushort3* );
+__device__ ushort4              __ldg(const ushort4* );
+
+__device__ uint1                __ldg(const uint1* );
+__device__ uint2                __ldg(const uint2* );
+__device__ uint3                __ldg(const uint3* );
+__device__ uint4                __ldg(const uint4* );
+
+__device__ ulonglong1           __ldg(const ulonglong1* );
+__device__ ulonglong2           __ldg(const ulonglong2* );
+__device__ ulonglong3           __ldg(const ulonglong3* );
+__device__ ulonglong4           __ldg(const ulonglong4* );
+
+__device__ float                __ldg(const float* );
+__device__ float1               __ldg(const float1* );
+__device__ float2               __ldg(const float2* );
+__device__ float3               __ldg(const float3* );
+__device__ float4               __ldg(const float4* );
+
+__device__ double               __ldg(const double* );
+__device__ double1              __ldg(const double1* );
+__device__ double2              __ldg(const double2* );
+__device__ double3              __ldg(const double3* );
+__device__ double4              __ldg(const double4* );
 
 #endif
 
diff --git a/projects/clr/hipamd/src/hip_ldg.cpp b/projects/clr/hipamd/src/hip_ldg.cpp
index 1e35a5cd49..f59c3e962d 100644
--- a/projects/clr/hipamd/src/hip_ldg.cpp
+++ b/projects/clr/hipamd/src/hip_ldg.cpp
@@ -19,43 +19,300 @@ THE SOFTWARE.
 
 #include"hcc_detail/hip_ldg.h"
 
-__device__ char __ldg(const char* ptr)
-{
-    return *ptr;
-}
-
-__device__ signed char __ldg(const signed char* ptr)
+__device__ char                 __ldg(const char* ptr)
 {
     return ptr[0];
 }
 
-__device__ short __ldg(const short* ptr)
+__device__ char1                __ldg(const char1* ptr)
 {
     return ptr[0];
 }
 
-__device__ int __ldg(const int* ptr)
+__device__ char2                __ldg(const char2* ptr)
 {
     return ptr[0];
 }
 
-__device__ long long __ldg(const long long* ptr)
+__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__ int2 __ldg(const int2* ptr)
+__device__ short                __ldg(const short* ptr)
 {
     return ptr[0];
 }
 
-__device__ int4 __ldg(const int4* ptr)
+__device__ short1               __ldg(const short1* ptr)
 {
     return ptr[0];
 }
 
-__device__ float __ldg(const float* ptr)
+__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)
+{
+    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];
+}
+
+
+
diff --git a/projects/clr/hipamd/tests/src/hip_test_ldg.cpp b/projects/clr/hipamd/tests/src/hip_test_ldg.cpp
index a58652d240..9d605fea5c 100644
--- a/projects/clr/hipamd/tests/src/hip_test_ldg.cpp
+++ b/projects/clr/hipamd/tests/src/hip_test_ldg.cpp
@@ -149,14 +149,151 @@ int main() {
   cout << " System major " << devProp.major << endl;
   cout << " agent prop name " << devProp.name << endl;
 
-   int errors = dataTypesRun<char>() &
-           dataTypesRun<signed char>() &
-           dataTypesRun<short>() &
-           dataTypesRun<int>();
+    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>();
+
   //hipResetDefaultAccelerator();
     if(errors == 1){
         passed();
-        return 0;
+    }else{
+        std::cout<<"Failed Float"<<std::endl;
+        return -1;
     }
-
 }