[SYCL] Change the way of accessing single element in sycl::vec

Returning scalars from sycl::vec class as a swizzled vec causes problems
with template resolution and ambiguity in overload resolution. This
commit implements workaround for this problem. Scalars are returned by
value for const vec and by reference for vec. Assumed that it is
possible to get a reference the element of the vector, test is added to
check that behaviour is not changed.

Signed-off-by: Artur Gainullin <[email protected]>

Author: againull

sycl/include/CL/sycl/swizzles.def

@@ -64,35 +64,45 @@
     return __SYCL_ACCESS_RETURN;                                    \
   }
 
+#define __SYCL_SCALAR_ACCESS(_COND, _NAME, _INDEX)                             \
+  template <int N = getNumElements()>                                          \
+  typename std::enable_if<(_COND), DataT &>::type _NAME() {                    \
+    return (*__SYCL_ACCESS_RETURN)[_INDEX];                                    \
+  }                                                                            \
+  template <int N = getNumElements()>                                          \
+  typename std::enable_if<(_COND), const DataT &>::type _NAME() const {        \
+    return (*__SYCL_ACCESS_RETURN)[_INDEX];                                    \
+  }
+
 //__swizzled_vec__ XYZW_ACCESS() const;
-__SYCL_ACCESS(N <= 4, x, 0)
-__SYCL_ACCESS(N == 2 || N == 3 || N == 4, y, 1)
-__SYCL_ACCESS(N == 3 || N == 4, z, 2)
-__SYCL_ACCESS(N == 4, w, 3)
+__SYCL_SCALAR_ACCESS(N <= 4, x, 0)
+__SYCL_SCALAR_ACCESS(N == 2 || N == 3 || N == 4, y, 1)
+__SYCL_SCALAR_ACCESS(N == 3 || N == 4, z, 2)
+__SYCL_SCALAR_ACCESS(N == 4, w, 3)
 
 //__swizzled_vec__ RGBA_ACCESS() const;
-__SYCL_ACCESS(N == 4, r, 0)
-__SYCL_ACCESS(N == 4, g, 1)
-__SYCL_ACCESS(N == 4, b, 2)
-__SYCL_ACCESS(N == 4, a, 3)
+__SYCL_SCALAR_ACCESS(N == 4, r, 0)
+__SYCL_SCALAR_ACCESS(N == 4, g, 1)
+__SYCL_SCALAR_ACCESS(N == 4, b, 2)
+__SYCL_SCALAR_ACCESS(N == 4, a, 3)
 
 //__swizzled_vec__ INDEX_ACCESS() const;
-__SYCL_ACCESS(N > 0, s0, 0)
-__SYCL_ACCESS(N > 1, s1, 1)
-__SYCL_ACCESS(N > 2, s2, 2)
-__SYCL_ACCESS(N > 2, s3, 3)
-__SYCL_ACCESS(N > 4, s4, 4)
-__SYCL_ACCESS(N > 4, s5, 5)
-__SYCL_ACCESS(N > 4, s6, 6)
-__SYCL_ACCESS(N > 4, s7, 7)
-__SYCL_ACCESS(N == 16, s8, 8)
-__SYCL_ACCESS(N == 16, s9, 9)
-__SYCL_ACCESS(N == 16, sA, 10)
-__SYCL_ACCESS(N == 16, sB, 11)
-__SYCL_ACCESS(N == 16, sC, 12)
-__SYCL_ACCESS(N == 16, sD, 13)
-__SYCL_ACCESS(N == 16, sE, 14)
-__SYCL_ACCESS(N == 16, sF, 15)
+__SYCL_SCALAR_ACCESS(N > 0, s0, 0)
+__SYCL_SCALAR_ACCESS(N > 1, s1, 1)
+__SYCL_SCALAR_ACCESS(N > 2, s2, 2)
+__SYCL_SCALAR_ACCESS(N > 2, s3, 3)
+__SYCL_SCALAR_ACCESS(N > 4, s4, 4)
+__SYCL_SCALAR_ACCESS(N > 4, s5, 5)
+__SYCL_SCALAR_ACCESS(N > 4, s6, 6)
+__SYCL_SCALAR_ACCESS(N > 4, s7, 7)
+__SYCL_SCALAR_ACCESS(N == 16, s8, 8)
+__SYCL_SCALAR_ACCESS(N == 16, s9, 9)
+__SYCL_SCALAR_ACCESS(N == 16, sA, 10)
+__SYCL_SCALAR_ACCESS(N == 16, sB, 11)
+__SYCL_SCALAR_ACCESS(N == 16, sC, 12)
+__SYCL_SCALAR_ACCESS(N == 16, sD, 13)
+__SYCL_SCALAR_ACCESS(N == 16, sE, 14)
+__SYCL_SCALAR_ACCESS(N == 16, sF, 15)
 
 #ifdef SYCL_SIMPLE_SWIZZLES
 //__swizzled_vec__ XYZW_SWIZZLE() const;

sycl/include/CL/sycl/types.hpp

@@ -605,6 +605,24 @@ template <typename Type, int NumElements> class vec {
     return this;
   }
 
+  // ext_vector_type is used as an underlying type for sycl::vec on device.
+  // The problem is that for clang vector types the return of operator[] is a
+  // temporary and not a reference to the element in the vector. In practice
+  // reinterpret_cast<DataT *>(&m_Data)[i]; is working. According to
+  // http://llvm.org/docs/GetElementPtr.html#can-gep-index-into-vector-elements
+  // this is not disallowed now. But could probably be disallowed in the future.
+  // That is why tests are added to check that behavior of the compiler has
+  // not changed.
+  //
+  // Implement operator [] in the same way for host and device.
+  // TODO: change host side implementation when underlying type for host side
+  // will be changed to std::array.
+  const DataT &operator[](int i) const {
+    return reinterpret_cast<const DataT *>(&m_Data)[i];
+  }
+
+  DataT &operator[](int i) { return reinterpret_cast<DataT *>(&m_Data)[i]; }
+
   // Begin hi/lo, even/odd, xyzw, and rgba swizzles.
 private:
   // Indexer used in the swizzles.def

sycl/test/basic_tests/scalar_vec_access.cpp

@@ -0,0 +1,52 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: env SYCL_DEVICE_TYPE=HOST %t.out | FileCheck %s
+// RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
+// RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER
+// RUN: %ACC_RUN_PLACEHOLDER %t.out %ACC_CHECK_PLACEHOLDER
+//==------- scalar_vec_access.cpp - SYCL scalar access to vec test ---------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+// CHECK-NOT: Error: unexpected behavior because of accessing element of the vector by reference
+
+#include <CL/sycl.hpp>
+#include <iostream>
+
+typedef float float4_t __attribute__((ext_vector_type(4)));
+
+int main() {
+
+  cl::sycl::queue Q;
+
+  Q.submit([=](cl::sycl::handler &cgh) {
+    cl::sycl::stream out(1024, 100, cgh);
+    cgh.single_task<class K>([=]() {
+      // Test that it is possible to get a reference to single element of the
+      // vector type. This behavior could possibly change in the future, this
+      // test is necessary to track that.
+      float4_t my_float4 = {0.0, 1.0, 2.0, 3.0};
+      float f[4];
+      for (int i = 0; i < 4; ++i) {
+        f[i] = reinterpret_cast<float *>(&my_float4)[i];
+        if (f[i] != i)
+          out << "Error: unexpected behavior because of accessing element of "
+                 "the vector by reference"
+              << cl::sycl::endl;
+      }
+
+      // Test that there is no template resolution error
+      cl::sycl::float4 a = {1.0, 2.0, 3.0, 4.0};
+      out << cl::sycl::native::recip(a.x()) << cl::sycl::endl;
+    });
+  });
+
+  // Test that there is no ambiguity in overload resolution.
+  cl::sycl::float4 a = {1.0, 2.0, 3.0, 4.0};
+  std::cout << a.x() << std::endl;
+
+  return 0;
+}