SWDEV-393139 - Add mipmap tests

Add 1D/2D/3D mipmap tests covering necessary data types and
config parameters.

Change-Id: I6079d9f3a2ddf4e42b8a6f7f3902322cfca02cff

catch/include/hip_array_common.hh

@@ -128,7 +128,12 @@ template<
 const std::string getString(const T& t)
 {
   std::ostringstream os;
-  os<< "(" << t.x << ")";
+  if constexpr (std::is_same<decltype(T::x), char>::value ||
+                std::is_same<decltype(T::x), unsigned char>::value) {
+    os << "(" << static_cast<int>(t.x) << ")";
+  } else {
+    os << "(" << t.x << ")";
+  }
   return os.str();
 }
 
@@ -138,7 +143,12 @@ template<
 const std::string getString(const T& t)
 {
   std::ostringstream os;
-  os<< "(" << t.x << ", " << t.y << ")";
+  if constexpr (std::is_same<decltype(T::x), char>::value ||
+                std::is_same<decltype(T::x), unsigned char>::value) {
+    os << "(" << static_cast<int>(t.x) << ", " << static_cast<int>(t.y) << ")";
+  } else {
+    os << "(" << t.x << ", " << t.y << ")";
+  }
   return os.str();
 }
 
@@ -148,7 +158,13 @@ template<
 const std::string getString(const T& t)
 {
   std::ostringstream os;
-  os<< "(" << t.x << ", " << t.y << ", " << t.z << ")";
+  if constexpr (std::is_same<decltype(T::x), char>::value ||
+                std::is_same<decltype(T::x), unsigned char>::value) {
+    os << "(" << static_cast<int>(t.x) << ", " << static_cast<int>(t.y) << ", " <<
+        static_cast<int>(t.z) << ")";
+  } else {
+    os << "(" << t.x << ", " << t.y << ", " << t.z << ")";
+  }
   return os.str();
 }
 
@@ -158,7 +174,13 @@ template<
 const std::string getString(const T& t)
 {
   std::ostringstream os;
-  os<< "(" << t.x << ", " << t.y << ", " << t.z << ", " << t.w << ")";
+  if constexpr (std::is_same<decltype(T::x), char>::value ||
+                std::is_same<decltype(T::x), unsigned char>::value) {
+    os << "(" << static_cast<int>(t.x) << ", " << static_cast<int>(t.y) << ", " <<
+        static_cast<int>(t.z) << ", " << static_cast<int>(t.w) << ")";
+  } else {
+    os << "(" << t.x << ", " << t.y << ", " << t.z << ", " << t.w << ")";
+  }
   return os.str();
 }
 
@@ -168,7 +190,12 @@ template<
 std::string getString(const T& t)
 {
   std::ostringstream os;
-  os << t;
+  if constexpr (std::is_same<T, char>::value ||
+                std::is_same<T, unsigned char>::value) {
+    os << static_cast<int>(t);
+  } else {
+    os << t;
+  }
   return os.str();
 }
 
@@ -180,7 +207,13 @@ static inline T getRandom() {
   } else {
     r = std::rand() / (RAND_MAX + 1.);
   }
-  return static_cast<T>(std::numeric_limits < T > ::max() * r);
+  if constexpr (std::is_floating_point<T>::value) {
+    // Restrict any float within (-1000, 1000)
+    // to prevent too big float value that would make caculation sick
+    return static_cast<T>(r * 1000.);
+  } else {
+    return static_cast<T>(std::numeric_limits<T>::max() * r);
+  }
 }
 
 template<
@@ -214,3 +247,62 @@ template<
 static inline void initVal(T &val) {
   val = getRandom<T>();
 }
+
+/*Convert normalized floatx to typex*/
+template <typename T, typename F> inline __device__ T getTypeFromNormalizedFloat(const F &f) {
+  T t;
+  if constexpr (std::is_scalar<T>::value)
+    t = static_cast<T>(f.x * std::numeric_limits<T>::max());
+  else {
+    if constexpr (rank<T>() > 0)
+      t.x = static_cast<decltype(T::x)>(f.x * std::numeric_limits<decltype(T::x)>::max());
+    if constexpr (rank<T>() > 1)
+      t.y = static_cast<decltype(T::y)>(f.y * std::numeric_limits<decltype(T::y)>::max());
+    if constexpr (rank<T>() > 2)
+      t.z = static_cast<decltype(T::z)>(f.z * std::numeric_limits<decltype(T::z)>::max());
+    if constexpr (rank<T>() > 3)
+      t.w = static_cast<decltype(T::w)>(f.w * std::numeric_limits<decltype(T::w)>::max());
+  }
+  return t;
+}
+
+/*Convert typex to normalized floatx*/
+template <class T>
+inline auto getNormalizedFloatType(const T &t) {
+  if constexpr (std::is_scalar<T>::value)
+    return static_cast<float>(t) / std::numeric_limits<T>::max();
+  else {
+    if constexpr (rank<T>() == 1) {
+      float1 f{static_cast<float>(t.x) / std::numeric_limits<decltype(T::x)>::max()};
+      return f;
+    }
+    if constexpr (rank<T>() == 2) {
+      float2 f{static_cast<float>(t.x) / std::numeric_limits<decltype(T::x)>::max(), 
+               static_cast<float>(t.y) / std::numeric_limits<decltype(T::y)>::max()};
+      return f;
+    }
+    if constexpr (rank<T>() == 3) {
+      float3 f{static_cast<float>(t.x) / std::numeric_limits<decltype(T::x)>::max(),
+               static_cast<float>(t.y) / std::numeric_limits<decltype(T::y)>::max(),
+               static_cast<float>(t.z) / std::numeric_limits<decltype(T::z)>::max()};
+      return f;
+    }
+    if constexpr (rank<T>() == 4) {
+      float4 f{static_cast<float>(t.x) / std::numeric_limits<decltype(T::x)>::max(),
+               static_cast<float>(t.y) / std::numeric_limits<decltype(T::y)>::max(),
+               static_cast<float>(t.z) / std::numeric_limits<decltype(T::z)>::max(),
+               static_cast<float>(t.w) / std::numeric_limits<decltype(T::w)>::max()};
+      return f;
+    }
+  }
+}
+
+/*Check if T is floatx*/
+template <typename T> inline bool constexpr isFloat() {
+  if constexpr (std::is_scalar<T>::value)
+    return std::is_floating_point<T>::value;
+  else {
+    return std::is_floating_point<decltype(T::x)>::value;
+  }
+  return false;
+}