[NFC] Drop legacy get/setElement helpers

tests/common/common.h

@@ -27,7 +27,6 @@
 #include <catch2/catch_test_macros.hpp>
 
 #include "../../util/conversion.h"
-#include "../../util/math_vector.h"
 #include "../../util/proxy.h"
 #include "../../util/sycl_enums.h"
 #include "../../util/test_base.h"

tests/common/common_vec.h

@@ -27,7 +27,6 @@
 
 #include "../../util/accuracy.h"
 #include "../../util/math_reference.h"
-#include "../../util/math_vector.h"
 #include "../../util/proxy.h"
 #include "../../util/test_base.h"
 #include "../../util/type_traits.h"
@@ -61,7 +60,7 @@ template <typename vecType, int numOfElems>
 bool check_vector_values(sycl::vec<vecType, numOfElems> vector,
                          vecType* vals) {
   for (int i = 0; i < numOfElems; i++) {
-    if ((vals[i] != getElement(vector, i))) {
+    if ((vals[i] != vector[i])) {
       return false;
     }
   }
@@ -77,7 +76,7 @@ typename std::enable_if<is_sycl_floating_point<vecType>::value, bool>::type
 check_vector_values_div(sycl::vec<vecType, numOfElems> vector,
                         vecType *vals) {
   for (int i = 0; i < numOfElems; i++) {
-    vecType vectorValue = getElement(vector, i);
+    vecType vectorValue = vector[i];
     if (vals[i] == vectorValue)
       continue;
     const vecType ulpsExpected = 2.5; // Min Accuracy for x / y
@@ -115,7 +114,7 @@ bool check_single_vector_op(vectorType vector1, lambdaFunc lambda) {
     return false;
   }
   for (int i = 0; i < vecSize; i++) {
-    if (getElement(vector1, i) != getElement(vector2, i)) {
+    if (vector1[i] != vector2[i]) {
       return false;
     }
   }
@@ -130,8 +129,7 @@ template <typename vecType, int N, typename convertType>
 sycl::vec<convertType, N> convert_vec(sycl::vec<vecType, N> inputVec) {
   sycl::vec<convertType, N> resVec;
   for (size_t i = 0; i < N; ++i) {
-    vecType elem = getElement(inputVec, i);
-    setElement<convertType, N>(resVec, i, convertType(elem));
+    resVec[i] = convertType(inputVec[i]);
   }
   return resVec;
 }
@@ -143,8 +141,7 @@ sycl::vec<convertType, N> rte(sycl::vec<vecType, N> inputVec) {
     sycl::vec<vecType, N> roundedVec = reference::rint(inputVec);
     sycl::vec<convertType, N> resVec;
     for (size_t i = 0; i < N; ++i) {
-      vecType elem = getElement(roundedVec, i);
-      setElement<convertType, N>(resVec, i, static_cast<convertType>(elem));
+      resVec[i] = static_cast<convertType>(roundedVec[i]);
     }
     return resVec;
   }
@@ -158,8 +155,7 @@ sycl::vec<convertType, N> rtz(sycl::vec<vecType, N> inputVec) {
     sycl::vec<vecType, N> roundedVec = reference::trunc(inputVec);
     sycl::vec<convertType, N> resVec;
     for (size_t i = 0; i < N; ++i) {
-      vecType elem = getElement(roundedVec, i);
-      setElement<convertType, N>(resVec, i, static_cast<convertType>(elem));
+      resVec[i] = static_cast<convertType>(roundedVec[i]);
     }
     return resVec;
   }
@@ -173,8 +169,7 @@ sycl::vec<convertType, N> rtp(sycl::vec<vecType, N> inputVec) {
     sycl::vec<vecType, N> roundedVec = reference::ceil(inputVec);
     sycl::vec<convertType, N> resVec;
     for (size_t i = 0; i < N; ++i) {
-      vecType elem = getElement(roundedVec, i);
-      setElement<convertType, N>(resVec, i, static_cast<convertType>(elem));
+      resVec[i] = static_cast<convertType>(roundedVec[i]);
     }
     return resVec;
   }
@@ -188,8 +183,7 @@ sycl::vec<convertType, N> rtn(sycl::vec<vecType, N> inputVec) {
     sycl::vec<vecType, N> roundedVec = reference::floor(inputVec);
     sycl::vec<convertType, N> resVec;
     for (size_t i = 0; i < N; ++i) {
-      vecType elem = getElement(roundedVec, i);
-      setElement<convertType, N>(resVec, i, static_cast<convertType>(elem));
+      resVec[i] = static_cast<convertType>(roundedVec[i]);
     }
     return resVec;
   }
@@ -205,8 +199,8 @@ void handleFPToUnsignedConv(sycl::vec<vecType, N>& inputVec) {
   if constexpr (is_sycl_floating_point<vecType>::value &&
                 std::is_unsigned_v<convertType>) {
     for (size_t i = 0; i < N; ++i) {
-      vecType elem = getElement(inputVec, i);
-      if (elem < 0) setElement<vecType, N>(inputVec, i, -elem);
+      vecType elem = inputVec[i];
+      if (elem < 0) inputVec[i] = -elem;
     }
   }
 }
@@ -383,15 +377,15 @@ bool check_as_result(sycl::vec<vecType, N> inputVec,
                      sycl::vec<asType, asN> asVec) {
   vecType tmp_ptr[N];
   for (size_t i = 0; i < N; ++i) {
-    tmp_ptr[i] = getElement(inputVec, i);
+    tmp_ptr[i] = inputVec[i];
   }
   asType exp_ptr[asN];
   for (size_t i = 0; i < asN; ++i) {
-    exp_ptr[i] = getElement(asVec, i);
+    exp_ptr[i] = asVec[i];
   }
   std::memcpy(exp_ptr, tmp_ptr, std::min(sizeof(exp_ptr), sizeof(tmp_ptr)));
   for (size_t i = 0; i < asN; ++i) {
-    if (exp_ptr[i] != getElement(asVec, i)) {
+    if (exp_ptr[i] != asVec[i]) {
       return false;
     }
   }

tests/hierarchical/hierarchical_non_uniform_local_range.cpp

@@ -28,7 +28,6 @@
 *******************************************************************************/
 
 #include "../common/common.h"
-#include "../../util/math_vector.h"
 
 #define TEST_NAME hierarchical_non_uniform_local_range
 
@@ -41,12 +40,12 @@ using namespace sycl_cts;
 void check_expected(const std::vector<sycl::int3> &data, unsigned local_id,
                     unsigned idx, int dim, bool set, util::logger &log) {
   int expected = set ? local_id : -1;
-  if (getElement(data[idx], dim - 1) != expected) {
+  if (data[idx][dim - 1] != expected) {
     std::string errorMessage =
         std::string("Value for global id ") + std::to_string(idx) +
         std::string(" for dim = ") + std::to_string(dim) +
         std::string(" was not correct (") +
-        std::to_string(getElement(data[idx], dim - 1)) +
+        std::to_string(data[idx][dim - 1]) +
         std::string(" instead of ") + std::to_string(expected) + ")";
     FAIL(log, errorMessage);
   }

util/math_helper.h

@@ -29,7 +29,6 @@
 
 #include "../util/stl.h"
 #include "./../oclmath/mt19937.h"
-#include "./math_vector.h"
 
 namespace sycl_cts {
 /** math utility functions
@@ -92,7 +91,7 @@ T getElement(const T &f, int) {
  * extract an individual element. */
 template <typename T, int dim>
 T getElement(sycl::vec<T, dim> &f, int ix) {
-  return getComponent<T, dim>()(f, ix);
+  return f[ix];
 }
 
 // FIXME: hipSYCL does not support marray
@@ -107,7 +106,7 @@ T getElement(sycl::marray<T, dim> &f, size_t ix) {
 
 template <typename T, int dim>
 void setElement(sycl::vec<T, dim> &f, int ix, T value) {
-  setComponent<T, dim>()(f, ix, value);
+  f[ix] = value;
 }
 
 template <typename R, typename T, int N, typename funT, typename... Args>

util/math_reference.cpp

@@ -579,7 +579,9 @@ sycl::vec<T, N> cross_t(sycl::vec<T, N> a, sycl::vec<T, N> b) {
   temp_res[1] = av[2] * bv[0] - av[0] * bv[2];
   temp_res[2] = av[0] * bv[1] - av[1] * bv[0];
   temp_res[3] = 0.0;
-  for (int i = 0; i < N; i++) setElement<T, N>(res, i, temp_res[i]);
+  for (int i = 0; i < N; i++) res[i] = temp_res[i];
+
+
   return res;
 }
 

util/math_reference.h

@@ -208,7 +208,7 @@ bool any(T x) {
 template <typename T, int N>
 int any(sycl::vec<T, N> a) {
   for (int i = 0; i < N; i++) {
-    if (any(getElement(a, i)) == 1) return true;
+    if (any(a[i]) == 1) return true;
   }
   return false;
 }
@@ -230,7 +230,7 @@ bool all(T x) {
 template <typename T, int N>
 int all(sycl::vec<T, N> a) {
   for (int i = 0; i < N; i++) {
-    if (all(getElement(a, i)) == 0) return false;
+    if (all(a[i]) == 0) return false;
   }
   return true;
 }
@@ -263,10 +263,10 @@ sycl::vec<T, N> select(sycl::vec<T, N> a, sycl::vec<T, N> b,
                        sycl::vec<K, N> c) {
   sycl::vec<T, N> res;
   for (int i = 0; i < N; i++) {
-    if (any(getElement<K, N>(c, i)) == 1)
-      setElement<T, N>(res, i, getElement(b, i));
+    if (any(c[i]) == 1)
+      res[i] = b[i];
     else
-      setElement<T, N>(res, i, getElement(a, i));
+      res[i] = a[i];
   }
   return res;
 }
@@ -365,9 +365,9 @@ sycl_cts::resultRef<sycl::vec<T, N>> clamp(sycl::vec<T, N> a, T b, T c) {
   sycl::vec<T, N> res;
   std::map<int, bool> undefined;
   for (int i = 0; i < N; i++) {
-    sycl_cts::resultRef<T> element = clamp(getElement(a, i), b, c);
+    sycl_cts::resultRef<T> element = clamp(a[i], b, c);
     if (element.undefined.empty())
-      setElement<T, N>(res, i, element.res);
+      res[i] = element.res;
     else
       undefined[i] = true;
   }
@@ -715,9 +715,9 @@ sycl_cts::resultRef<sycl::vec<T, N>> mix(sycl::vec<T, N> a, sycl::vec<T, N> b,
   sycl::vec<T, N> res;
   std::map<int, bool> undefined;
   for (int i = 0; i < N; i++) {
-    sycl_cts::resultRef<T> element = mix(getElement(a, i), getElement(b, i), c);
+    sycl_cts::resultRef<T> element = mix(a[i], b[i], c);
     if (element.undefined.empty())
-      setElement<T, N>(res, i, element.res);
+      res[i] = element.res;
     else
       undefined[i] = true;
   }
@@ -764,7 +764,7 @@ template <typename T, int N>
 sycl::vec<T, N> step(T a, sycl::vec<T, N> b) {
   sycl::vec<T, N> res;
   for (int i = 0; i < N; i++) {
-    setElement<T, N>(res, i, step(a, getElement(b, i)));
+    res[i] = step(a, b[i]);
   }
   return res;
 }
@@ -798,9 +798,9 @@ sycl_cts::resultRef<sycl::vec<T, N>> smoothstep(T a, T b, sycl::vec<T, N> c) {
   sycl::vec<T, N> res;
   std::map<int, bool> undefined;
   for (int i = 0; i < N; i++) {
-    sycl_cts::resultRef<T> element = smoothstep(a, b, getElement(c, i));
+    sycl_cts::resultRef<T> element = smoothstep(a, b, c[i]);
     if (element.undefined.empty())
-      setElement<T, N>(res, i, element.res);
+      res[i] = element.res;
     else
       undefined[i] = true;
   }
@@ -1032,8 +1032,8 @@ sycl::vec<T, N> fract(sycl::vec<T, N> a, sycl::vec<T, N> *b) {
   sycl::vec<T, N> resPtr;
   for (int i = 0; i < N; i++) {
     T value;
-    setElement<T, N>(res, i, fract(getElement(a, i), &value));
-    setElement<T, N>(resPtr, i, value);
+    res[i] = fract(a[i], &value);
+    resPtr[i] = value;
   }
   *b = resPtr;
   return res;
@@ -1061,8 +1061,8 @@ sycl::vec<T, N> frexp(sycl::vec<T, N> a, sycl::vec<int, N> *b) {
   sycl::vec<int, N> resPtr;
   for (int i = 0; i < N; i++) {
     int value;
-    setElement<T, N>(res, i, frexp(getElement(a, i), &value));
-    setElement<int, N>(resPtr, i, value);
+    res[i] = frexp(a[i], &value);
+    resPtr[i] = value;
   }
   *b = resPtr;
   return res;
@@ -1095,7 +1095,7 @@ template <typename T, int N>
 sycl::vec<int, N> ilogb(sycl::vec<T, N> a) {
   sycl::vec<int, N> res;
   for (int i = 0; i < N; i++) {
-    setElement<int, N>(res, i, ilogb(getElement<T, N>(a, i)));
+    res[i] = ilogb(a[i]);
   }
   return res;
 }
@@ -1116,8 +1116,7 @@ template <typename T, int N>
 sycl::vec<T, N> ldexp(sycl::vec<T, N> a, sycl::vec<int, N> b) {
   sycl::vec<T, N> res;
   for (int i = 0; i < N; i++) {
-    setElement<T, N>(res, i,
-                     ldexp(getElement<T, N>(a, i), getElement<int, N>(b, i)));
+    res[i] = ldexp(a[i], b[i]);
   }
   return res;
 }
@@ -1136,7 +1135,7 @@ template <typename T, int N>
 sycl::vec<T, N> ldexp(sycl::vec<T, N> a, int b) {
   sycl::vec<T, N> res;
   for (int i = 0; i < N; i++) {
-    setElement<T, N>(res, i, ldexp(getElement<T, N>(a, i), b));
+    res[i] = ldexp(a[i], b);
   }
   return res;
 }
@@ -1166,8 +1165,8 @@ sycl::vec<T, N> lgamma_r(sycl::vec<T, N> a, sycl::vec<int, N> *b) {
   sycl::vec<int, N> resPtr;
   for (int i = 0; i < N; i++) {
     int value;
-    setElement<T, N>(res, i, lgamma_r(getElement(a, i), &value));
-    setElement<int, N>(resPtr, i, value);
+    res[i] = lgamma_r(a[i], &value);
+    resPtr[i] = value;
   }
   *b = resPtr;
   return res;
@@ -1249,8 +1248,8 @@ sycl::vec<T, N> modf(sycl::vec<T, N> a, sycl::vec<T, N> *b) {
   sycl::vec<T, N> resPtr;
   for (int i = 0; i < N; i++) {
     T value;
-    setElement<T, N>(res, i, modf(getElement(a, i), &value));
-    setElement<T, N>(resPtr, i, value);
+    res[i] = modf(a[i], &value);
+    resPtr[i] = value;
   }
   *b = resPtr;
   return res;
@@ -1324,8 +1323,7 @@ template <typename T, int N>
 sycl::vec<T, N> pown(sycl::vec<T, N> a, sycl::vec<int, N> b) {
   sycl::vec<T, N> res;
   for (int i = 0; i < N; i++) {
-    setElement<T, N>(res, i,
-                     pown(getElement<T, N>(a, i), getElement<int, N>(b, i)));
+    res[i] = pown(a[i], b[i]);
   }
   return res;
 }
@@ -1378,9 +1376,8 @@ sycl::vec<T, N> remquo(sycl::vec<T, N> a, sycl::vec<T, N> b,
   sycl::vec<int, N> resPtr;
   for (int i = 0; i < N; i++) {
     int value;
-    setElement<T, N>(res, i,
-                     remquo(getElement(a, i), getElement(b, i), &value));
-    setElement<int, N>(resPtr, i, value);
+    res[i] = remquo(a[i], b[i], &value);
+    resPtr[i] = value;
   }
   *c = resPtr;
   return res;
@@ -1414,8 +1411,7 @@ template <typename T, int N>
 sycl::vec<T, N> rootn(sycl::vec<T, N> a, sycl::vec<int, N> b) {
   sycl::vec<T, N> res;
   for (int i = 0; i < N; i++) {
-    setElement<T, N>(res, i,
-                     rootn(getElement<T, N>(a, i), getElement<int, N>(b, i)));
+    res[i] = rootn(a[i], b[i]);
   }
   return res;
 }
@@ -1451,8 +1447,8 @@ sycl::vec<T, N> sincos(sycl::vec<T, N> a, sycl::vec<T, N> *b) {
   sycl::vec<T, N> resPtr;
   for (int i = 0; i < N; i++) {
     T value;
-    setElement<T, N>(res, i, sincos(getElement(a, i), &value));
-    setElement<T, N>(resPtr, i, value);
+    res[i] = sincos(a[i], &value);
+    resPtr[i] = value;
   }
   *b = resPtr;
   return res;
@@ -1557,7 +1553,7 @@ template <typename T, int N>
 T dot(sycl::vec<T, N> a, sycl::vec<T, N> b) {
   T res = 0;
   for (int i = 0; i < N; i++)
-    res += getElement<T, N>(a, i) * getElement<T, N>(b, i);
+    res += a[i] * b[i];
   return res;
 }
 // FIXME: hipSYCL does not support marray
@@ -1591,7 +1587,7 @@ sycl::vec<T, N> normalize(sycl::vec<T, N> a) {
   T len_a = reference::length(a);
   if (len_a == 0) return sycl::vec<T, N>(0);
   for (int i = 0; i < N; i++)
-    setElement<T, N>(res, i, getElement<T, N>(a, i) / len_a);
+    res[i] = a[i] / len_a;
   return res;
 }
 // FIXME: hipSYCL does not support marray