[SYCL][COMPAT] Adds sycl::length wrapper to SYCLcompat (#12968)

Adds a templated wrapper to sycl::length in syclcompat, that wraps over
`sycl::vec<ValueT,N>` up to N == 4, and calculates the length with N >
4.

---------

Signed-off-by: Alberto Cabrera <[email protected]>

Author: Alcpz

sycl/doc/syclcompat/README.md

@@ -1199,7 +1199,8 @@ static kernel_function_info get_kernel_function_info(const void *function);
 
 `syclcompat::fast_length` provides a wrapper to SYCL's
 `fast_length(sycl::vec<float,N>)` that accepts arguments for a C++ array and a
-length.
+length. `syclcompat::length` provides a templated version that wraps over
+`sycl::length`.
 
 `vectorized_max` and `vectorized_min` are binary operations returning the
 max/min of two arguments, where each argument is treated as a `sycl::vec` type.
@@ -1213,6 +1214,9 @@ which accept `sycl::vec<T,2>` arguments representing complex values.
 ```cpp
 inline float fast_length(const float *a, int len);
 
+template <typename ValueT>
+inline ValueT length(const ValueT *a, const int len);
+
 template <typename S, typename T> inline T vectorized_max(T a, T b);
 
 template <typename S, typename T> inline T vectorized_min(T a, T b);

sycl/include/syclcompat/math.hpp

@@ -73,6 +73,29 @@ inline float fast_length(const float *a, int len) {
   }
 }
 
+/// Calculate the square root of the input array.
+/// \param [in] a The array pointer
+/// \param [in] len Length of the array
+/// \returns The square root
+template <typename ValueT>
+inline ValueT length(const ValueT *a, const int len) {
+  switch (len) {
+  case 1:
+    return a[0];
+  case 2:
+    return sycl::length(sycl::vec<ValueT, 2>(a[0], a[1]));
+  case 3:
+    return sycl::length(sycl::vec<ValueT, 3>(a[0], a[1], a[2]));
+  case 4:
+    return sycl::length(sycl::vec<ValueT, 4>(a[0], a[1], a[2], a[3]));
+  default:
+    ValueT ret = 0;
+    for (int i = 0; i < len; ++i)
+      ret += a[i] * a[i];
+    return sycl::sqrt(ret);
+  }
+}
+
 /// Compute vectorized max for two values, with each value treated as a vector
 /// type \p S
 /// \param [in] S The type of the vector

sycl/test-e2e/syclcompat/math/math_fast_length_test.cpp renamed to sycl/test-e2e/syclcompat/math/math_length_test.cpp

@@ -14,10 +14,10 @@
  *
  *  SYCLcompat API
  *
- *  math_fast_length_test.cpp
+ *  math_length_test.cpp
  *
  *  Description:
- *    Fast length tests
+ *    vector length tests
  **************************************************************************/
 
 // The original source was under the license below:
@@ -41,22 +41,26 @@
 
 #define MAX_LEN 5
 
-void compute_length(float *d_A, size_t n, float *ans) {
+void compute_fast_length(float *d_A, size_t n, float *ans) {
   *ans = syclcompat::fast_length(d_A, n);
 }
 
-class FastLengthLauncher {
+void compute_length(float *d_A, size_t n, float *ans) {
+  *ans = syclcompat::length(d_A, n);
+}
+
+class LengthLauncher {
 protected:
   float *data_;
   float *result_;
   float host_result_{0.0};
 
 public:
-  FastLengthLauncher() {
+  LengthLauncher() {
     data_ = (float *)syclcompat::malloc(MAX_LEN * sizeof(float));
     result_ = (float *)syclcompat::malloc(sizeof(float));
   };
-  ~FastLengthLauncher() {
+  ~LengthLauncher() {
     syclcompat::free(data_);
     syclcompat::free(result_);
   }
@@ -68,13 +72,13 @@ class FastLengthLauncher {
     assert(diff <= 1.e-5);
   }
 
-  void launch(std::vector<float> vec) {
+  template <auto F> void launch(std::vector<float> vec) {
     size_t n = vec.size();
     syclcompat::memcpy(data_, vec.data(), sizeof(float) * n);
     auto data = data_;
     auto result = result_;
     syclcompat::get_default_queue().single_task(
-        [data, result, n]() { compute_length(data, n, result); });
+        [data, result, n]() { F(data, n, result); });
     syclcompat::memcpy(&host_result_, result_, sizeof(float));
     check_result(vec);
   }
@@ -83,18 +87,36 @@ class FastLengthLauncher {
 void test_fast_length() {
   std::cout << __PRETTY_FUNCTION__ << std::endl;
 
-  auto launcher = FastLengthLauncher();
-  launcher.launch(std::vector<float>{0.8970062715});
-  launcher.launch(std::vector<float>{0.8335529744, 0.7346600673});
-  launcher.launch(std::vector<float>{0.1658983906, 0.590226484, 0.4891553616});
-  launcher.launch(std::vector<float>{0.6041178723, 0.7760620605, 0.2944284976,
-                                     0.6851913766});
-  launcher.launch(std::vector<float>{0.6041178723, 0.7760620605, 0.2944284976,
-                                     0.6851913766, 0.6851913766});
+  auto launcher = LengthLauncher();
+  launcher.launch<compute_fast_length>(std::vector<float>{0.8970062715});
+  launcher.launch<compute_fast_length>(
+      std::vector<float>{0.8335529744, 0.7346600673});
+  launcher.launch<compute_fast_length>(
+      std::vector<float>{0.1658983906, 0.590226484, 0.4891553616});
+  launcher.launch<compute_fast_length>(std::vector<float>{
+      0.6041178723, 0.7760620605, 0.2944284976, 0.6851913766});
+  launcher.launch<compute_fast_length>(std::vector<float>{
+      0.6041178723, 0.7760620605, 0.2944284976, 0.6851913766, 0.6851913766});
+}
+
+void test_length() {
+  std::cout << __PRETTY_FUNCTION__ << std::endl;
+
+  auto launcher = LengthLauncher();
+  launcher.launch<compute_length>(std::vector<float>{0.8970062715});
+  launcher.launch<compute_length>(
+      std::vector<float>{0.8335529744, 0.7346600673});
+  launcher.launch<compute_length>(
+      std::vector<float>{0.1658983906, 0.590226484, 0.4891553616});
+  launcher.launch<compute_length>(std::vector<float>{
+      0.6041178723, 0.7760620605, 0.2944284976, 0.6851913766});
+  launcher.launch<compute_length>(std::vector<float>{
+      0.6041178723, 0.7760620605, 0.2944284976, 0.6851913766, 0.6851913766});
 }
 
 int main() {
   test_fast_length();
+  test_length();
 
   return 0;
 }