Deprecate doubled vector as much as possible.

Arm complex instructions on M3/M4 NEON v8.3 and simplify A64FX code paths/broaden.
Grid_vector_types and Simd.h mainly reorg and prep for sComplexD and sComplexF alternat Nsimd=1 types

Grid/simd/Grid_a64fx-2.h

@@ -933,6 +933,8 @@ typedef Optimization::Div            DivSIMD;
 typedef Optimization::Mult           MultSIMD;
 typedef Optimization::MultComplex    MultComplexSIMD;
 typedef Optimization::MultAddComplex MultAddComplexSIMD;
+// Asserts MultAddComplexSIMD and a three argument MultSIMD operator()
+#define GRID_ARCH_HAS_COMPLEX_MULT_ADD 1
 typedef Optimization::MultRealPart   MultRealPartSIMD;
 typedef Optimization::MaddRealPart   MaddRealPartSIMD;
 typedef Optimization::Conj           ConjSIMD;

Grid/simd/Grid_a64fx-fixedsize.h

@@ -760,6 +760,8 @@ typedef Optimization::Div            DivSIMD;
 typedef Optimization::Mult           MultSIMD;
 typedef Optimization::MultComplex    MultComplexSIMD;
 typedef Optimization::MultAddComplex MultAddComplexSIMD;
+// Asserts MultAddComplexSIMD and a three argument MultSIMD operator()
+#define GRID_ARCH_HAS_COMPLEX_MULT_ADD 1
 typedef Optimization::MultRealPart   MultRealPartSIMD;
 typedef Optimization::MaddRealPart   MaddRealPartSIMD;
 typedef Optimization::Conj           ConjSIMD;

Grid/simd/Grid_doubled_vector.h

@@ -339,54 +339,6 @@ typedef Grid_simd2<complex<double>  , vComplexD>  vComplexD2;
 typedef Grid_simd2<double           , vRealD>     vRealD2;
 
 
-
-/////////////////////////////////////////
-// Some traits to recognise the types
-/////////////////////////////////////////
-template <typename T>
-struct is_simd : public std::false_type {};
-template <> struct is_simd<vRealF>     : public std::true_type {};
-template <> struct is_simd<vRealD>     : public std::true_type {};
-template <> struct is_simd<vRealH>     : public std::true_type {};
-template <> struct is_simd<vComplexF>  : public std::true_type {};
-template <> struct is_simd<vComplexD>  : public std::true_type {};
-template <> struct is_simd<vComplexH>  : public std::true_type {};
-template <> struct is_simd<vInteger>   : public std::true_type {};
-template <> struct is_simd<vRealD2>    : public std::true_type {};
-template <> struct is_simd<vComplexD2> : public std::true_type {};
-
-template <typename T> using IfSimd    = Invoke<std::enable_if<is_simd<T>::value, int> >;
-template <typename T> using IfNotSimd = Invoke<std::enable_if<!is_simd<T>::value, unsigned> >;
-
-///////////////////////////////////////////////
-// insert / extract with complex support
-///////////////////////////////////////////////
-template <class S, class V>
-accelerator_inline S getlane(const Grid_simd<S, V> &in,int lane) {
-  return in.getlane(lane);
-}
-template <class S, class V>
-accelerator_inline void putlane(Grid_simd<S, V> &vec,const S &_S, int lane){
-  vec.putlane(_S,lane);
-}
-template <class S,IfNotSimd<S> = 0 >
-accelerator_inline S getlane(const S &in,int lane) {
-  return in;
-}
-template <class S,IfNotSimd<S> = 0 >
-accelerator_inline void putlane(S &vec,const S &_S, int lane){
-  vec = _S;
-}
-template <class S, class V>
-accelerator_inline S getlane(const Grid_simd2<S, V> &in,int lane) {
-  return in.getlane(lane);
-}
-template <class S, class V>
-accelerator_inline void putlane(Grid_simd2<S, V> &vec,const S &_S, int lane){
-  vec.putlane(_S,lane);
-}
-
-
 ////////////////////////////////////////////////////////////////////
 // General rotate
 ////////////////////////////////////////////////////////////////////

Grid/simd/Grid_neon.h

@@ -42,6 +42,16 @@
 #define GEN_SIMD_WIDTH 16u
 #endif
 
+//////////////////////////////////////////////////////////////////////////////
+// ARMv8.3 complex assist (FCMLA). Kept separate from the ACLE
+// __ARM_FEATURE_COMPLEX: the instructions are also present on A64FX, where
+// their throughput is 1:4 and this path would be a regression. Revisit the
+// rule once measured per target.
+//////////////////////////////////////////////////////////////////////////////
+#ifdef __ARM_FEATURE_COMPLEX
+#define GRID_ARM_FEATURE_COMPLEX 1
+#endif
+
 #include "Grid_generic_types.h"
 #include <arm_neon.h>
 
@@ -226,6 +236,22 @@ struct Div{
 };
 
 struct MultComplex{
+#ifdef GRID_ARM_FEATURE_COMPLEX
+  // fcmla #0  : r += (ar,ar)*(br,bi) = ( ar*br, ar*bi)
+  // fcmla #90 : r += (-ai,ai)*(bi,br) = (-ai*bi, ai*br)
+  // Complex float: two products, two instructions
+  inline float32x4_t operator()(float32x4_t a, float32x4_t b){
+    float32x4_t r = vdupq_n_f32(0.0f);
+    r = vcmlaq_f32(r, a, b);
+    return vcmlaq_rot90_f32(r, a, b);
+  }
+  // Complex double
+  inline float64x2_t operator()(float64x2_t a, float64x2_t b){
+    float64x2_t r = vdupq_n_f64(0.0);
+    r = vcmlaq_f64(r, a, b);
+    return vcmlaq_rot90_f64(r, a, b);
+  }
+#else
   // Complex float
   inline float32x4_t operator()(float32x4_t a, float32x4_t b){
 
@@ -273,8 +299,23 @@ struct MultComplex{
     // r5 = vmulq_f64(r0, a);
     // return vaddq_f64(r4, r5);
   }
+#endif
 };
 
+#ifdef GRID_ARM_FEATURE_COMPLEX
+struct MultAddComplex{
+  // Complex a*b+c
+  inline float32x4_t operator()(float32x4_t a, float32x4_t b, float32x4_t c){
+    c = vcmlaq_f32(c, a, b);
+    return vcmlaq_rot90_f32(c, a, b);
+  }
+  inline float64x2_t operator()(float64x2_t a, float64x2_t b, float64x2_t c){
+    c = vcmlaq_f64(c, a, b);
+    return vcmlaq_rot90_f64(c, a, b);
+  }
+};
+#endif
+
 struct Mult{
   // Real float
   inline float32x4_t mac(float32x4_t a, float32x4_t b, float32x4_t c){
@@ -285,6 +326,13 @@ struct Mult{
     //return vaddq_f64(vmulq_f64(b,c),a);
     return vfmaq_f64(a, b, c);
   }
+  // Real a*b+c
+  inline float32x4_t operator()(float32x4_t a, float32x4_t b, float32x4_t c){
+    return vfmaq_f32(c, a, b);
+  }
+  inline float64x2_t operator()(float64x2_t a, float64x2_t b, float64x2_t c){
+    return vfmaq_f64(c, a, b);
+  }
   inline float32x4_t operator()(float32x4_t a, float32x4_t b){
     return vmulq_f32(a,b);
   }
@@ -585,6 +633,11 @@ typedef Optimization::Sub         SubSIMD;
 typedef Optimization::Div         DivSIMD;
 typedef Optimization::Mult        MultSIMD;
 typedef Optimization::MultComplex MultComplexSIMD;
+#ifdef GRID_ARM_FEATURE_COMPLEX
+typedef Optimization::MultAddComplex MultAddComplexSIMD;
+// Asserts MultAddComplexSIMD and a three argument MultSIMD operator()
+#define GRID_ARCH_HAS_COMPLEX_MULT_ADD 1
+#endif
 typedef Optimization::MultRealPart MultRealPartSIMD;
 typedef Optimization::MaddRealPart MaddRealPartSIMD;
 typedef Optimization::Conj        ConjSIMD;

Grid/simd/Grid_vector_types.h

@@ -36,8 +36,7 @@ Author: Michael Marshall <michael.marshall@ed.ac.au>
 */
 // Time-stamp: <2015-07-10 17:45:33 neo>
 //---------------------------------------------------------------------------
-#ifndef GRID_VECTOR_TYPES
+#pragma once
-#define GRID_VECTOR_TYPES
 
 // PAB - Lifted and adapted from Eigen, which is GPL V2
 struct Grid_half {
@@ -169,78 +168,6 @@ accelerator_inline Grid_half sfw_float_to_half(float ff) {
 
 NAMESPACE_BEGIN(Grid);
 
-
-//////////////////////////////////////
-// To take the floating point type of real/complex type
-//////////////////////////////////////
-template <typename T>
-struct RealPart {
-  typedef T type;
-};
-template <typename T>
-struct RealPart<complex<T> > {
-  typedef T type;
-};
-
-#include <type_traits>
-
-//////////////////////////////////////
-// demote a vector to real type
-//////////////////////////////////////
-// type alias used to simplify the syntax of std::enable_if
-template <typename T> using Invoke = typename T::type;
-template <typename Condition, typename ReturnType = void> using EnableIf    = Invoke<std::enable_if<Condition::value, ReturnType> >;
-template <typename Condition, typename ReturnType = void> using NotEnableIf = Invoke<std::enable_if<!Condition::value, ReturnType> >;
-
-////////////////////////////////////////////////////////
-// Check for complexity with type traits
-template <typename T> struct is_complex : public std::false_type {};
-template <> struct is_complex<ComplexD> : public std::true_type {};
-template <> struct is_complex<ComplexF> : public std::true_type {};
-
-template <typename T> struct is_ComplexD : public std::false_type {};
-template <> struct is_ComplexD<ComplexD> : public std::true_type {};
-
-template <typename T> struct is_ComplexF : public std::false_type {};
-template <> struct is_ComplexF<ComplexF> : public std::true_type {};
-
-template<typename T, typename V=void> struct is_real : public std::false_type {};
-template<typename T> struct is_real<T, typename std::enable_if<std::is_floating_point<T>::value,
-  void>::type> : public std::true_type {};
-
-template<typename T, typename V=void> struct is_integer : public std::false_type {};
-template<typename T> struct is_integer<T, typename std::enable_if<std::is_integral<T>::value,
-  void>::type> : public std::true_type {};
-
-template <typename T>              using IfReal    = Invoke<std::enable_if<is_real<T>::value, int> >;
-template <typename T>              using IfComplex = Invoke<std::enable_if<is_complex<T>::value, int> >;
-template <typename T>              using IfInteger = Invoke<std::enable_if<is_integer<T>::value, int> >;
-template <typename T1,typename T2> using IfSame    = Invoke<std::enable_if<std::is_same<T1,T2>::value, int> >;
-
-template <typename T>              using IfNotReal    = Invoke<std::enable_if<!is_real<T>::value, int> >;
-template <typename T>              using IfNotComplex = Invoke<std::enable_if<!is_complex<T>::value, int> >;
-template <typename T>              using IfNotInteger = Invoke<std::enable_if<!is_integer<T>::value, int> >;
-template <typename T1,typename T2> using IfNotSame    = Invoke<std::enable_if<!std::is_same<T1,T2>::value, int> >;
-
-////////////////////////////////////////////////////////
-// Define the operation templates functors
-// general forms to allow for vsplat syntax
-// need explicit declaration of types when used since
-// clang cannot automatically determine the output type sometimes
-template <class Out, class Input1, class Input2, class Input3, class Operation>
-Out accelerator_inline  trinary(Input1 src_1, Input2 src_2, Input3 src_3, Operation op) {
-  return op(src_1, src_2, src_3);
-}
-template <class Out, class Input1, class Input2, class Operation>
-Out accelerator_inline binary(Input1 src_1, Input2 src_2, Operation op) {
-  return op(src_1, src_2);
-}
-template <class Out, class Input, class Operation>
-Out accelerator_inline  unary(Input src, Operation op) {
-  return op(src);
-}
-///////////////////////////////////////////////
-
 /*
   @brief Grid_simd class for the SIMD vector type operations
 */
@@ -378,11 +305,18 @@ public:
 
   // FIXME -- alias this to an accelerator_inline MAC struct.
 
-  #if defined(A64FX) || defined(A64FXFIXEDSIZE)
+  #ifdef GRID_ARCH_HAS_COMPLEX_MULT_ADD
+  // Architecture has a fused complex multiply-add (FCMLA and friends).
+  // Defining GRID_ARCH_HAS_COMPLEX_MULT_ADD asserts BOTH MultAddComplexSIMD
+  // and a three argument MultSIMD operator() computing a*b+c.
   friend accelerator_inline void mac(Grid_simd *__restrict__ y,
 				     const Grid_simd *__restrict__ a,
 				     const Grid_simd *__restrict__ x) {
-    *y = fxmac((*a), (*x), (*y));
+    if constexpr ( is_complex<Scalar_type>::value ) {
+      y->v = trinary<Vector_type>(a->v, x->v, y->v, MultAddComplexSIMD());
+    } else {
+      y->v = trinary<Vector_type>(a->v, x->v, y->v, MultSIMD());
+    }
   };
   #else
   friend accelerator_inline void mac(Grid_simd *__restrict__ y,
@@ -670,7 +604,6 @@ typedef Grid_simd<complex<uint16_t>, SIMD_Htype> vComplexH;
 typedef Grid_simd<complex<float>   , SIMD_Ftype> vComplexF;
 typedef Grid_simd<complex<double>  , SIMD_Dtype> vComplexD;
 #endif
-
 /////////////////////////////////////////
 // Pointer type to use on extractLane
 /////////////////////////////////////////
@@ -885,25 +818,7 @@ accelerator_inline Grid_simd<S, V> operator*(Grid_simd<S, V> a, Grid_simd<S, V>
   return ret;
 };
 
-// ---------------- A64FX MAC -------------------
+// mac() selects MultAddComplexSIMD directly where the architecture has it.
-// Distinguish between complex types and others
-#if defined(A64FX) || defined(A64FXFIXEDSIZE)
-template <class S, class V, IfComplex<S> = 0>
-accelerator_inline Grid_simd<S, V> fxmac(Grid_simd<S, V> a, Grid_simd<S, V> b, Grid_simd<S, V> c) {
-  Grid_simd<S, V> ret;
-  ret.v = trinary<V>(a.v, b.v, c.v, MultAddComplexSIMD());
-  return ret;
-};
-
-// Real/Integer types
-template <class S, class V, IfNotComplex<S> = 0>
-accelerator_inline Grid_simd<S, V> fxmac(Grid_simd<S, V> a, Grid_simd<S, V> b, Grid_simd<S, V> c) {
-  Grid_simd<S, V> ret;
-  ret.v = trinary<V>(a.v, b.v, c.v, MultSIMD());
-  return ret;
-};
-#endif
-// ----------------------------------------------
 
 
 ///////////////////////
@@ -1141,13 +1056,3 @@ template<class vobj> void gpermute(vobj & inout,int perm){
 
 NAMESPACE_END(Grid);
 
-#ifdef GRID_SYCL
-template<> struct sycl::is_device_copyable<Grid::vComplexF> : public std::true_type {};
-template<> struct sycl::is_device_copyable<Grid::vComplexD> : public std::true_type {};
-template<> struct sycl::is_device_copyable<Grid::vRealF   > : public std::true_type {};
-template<> struct sycl::is_device_copyable<Grid::vRealD   > : public std::true_type {};
-template<> struct sycl::is_device_copyable<Grid::vInteger > : public std::true_type {};
-#endif
-
-
-#endif

Grid/simd/Simd.h

@@ -28,216 +28,163 @@ See the full license in the file "LICENSE" in the top level distribution
 directory
 *************************************************************************************/
 /*  END LEGAL */
-#ifndef GRID_SIMD_H
+#pragma once
-#define GRID_SIMD_H
 
-#if defined(GRID_CUDA) || defined(GRID_HIP)
+#include <type_traits>
-#include <thrust/complex.h>
-#endif
 
-////////////////////////////////////////////////////////////////////////
+#include <Grid/simd/Grid_scalar_support.h>
-// Define scalar and vector floating point types
-//
-// Scalar:   RealF, RealD, ComplexF, ComplexD
-//
-// Vector:  vRealF, vRealD, vComplexF, vComplexD
-//
-// Vector types are arch dependent
-////////////////////////////////////////////////////////////////////////
-
-#define _MM_SELECT_FOUR_FOUR(A,B,C,D) ((A<<6)|(B<<4)|(C<<2)|(D))
-#define _MM_SELECT_FOUR_FOUR_STRING(A,B,C,D) "((" #A "<<6)|(" #B "<<4)|(" #C "<<2)|(" #D "))"
-#define _MM_SELECT_EIGHT_TWO(A,B,C,D,E,F,G,H) ((A<<7)|(B<<6)|(C<<5)|(D<<4)|(E<<3)|(F<<2)|(G<<4)|(H))
-#define _MM_SELECT_FOUR_TWO (A,B,C,D) _MM_SELECT_EIGHT_TWO(0,0,0,0,A,B,C,D)
-#define _MM_SELECT_TWO_TWO  (A,B)     _MM_SELECT_FOUR_TWO(0,0,A,B)
-
-#define RotateBit (0x100)
 
 NAMESPACE_BEGIN(Grid);
-
+//////////////////////////////////////
-typedef uint32_t Integer;
+// To take the floating point type of real/complex type
-
+//////////////////////////////////////
-typedef  float  RealF;
+template <typename T>
-typedef  double RealD;
+struct RealPart {
-#ifdef GRID_DEFAULT_PRECISION_DOUBLE
+  typedef T type;
-typedef RealD   Real;
+};
-#else
+template <typename T>
-typedef RealF  Real;
+struct RealPart<std::complex<T> > {
-#endif
+  typedef T type;
-
+};
-#if defined(GRID_CUDA) || defined(GRID_HIP)
-typedef thrust::complex<RealF> ComplexF;
-typedef thrust::complex<RealD> ComplexD;
-typedef thrust::complex<Real>  Complex;
-typedef thrust::complex<uint16_t>  ComplexH;
-template<class T> using complex = thrust::complex<T>;
-
-accelerator_inline ComplexD pow(const ComplexD& r,RealD y){ return(thrust::pow(r,(double)y)); }
-accelerator_inline ComplexF pow(const ComplexF& r,RealF y){ return(thrust::pow(r,(float)y)); }
-#else 
-typedef std::complex<RealF> ComplexF;
-typedef std::complex<RealD> ComplexD;
-typedef std::complex<Real>  Complex;
-typedef std::complex<uint16_t>  ComplexH; // Hack
-template<class T> using complex = std::complex<T>;
-
-accelerator_inline ComplexD pow(const ComplexD& r,RealD y){ return(std::pow(r,y)); }
-accelerator_inline ComplexF pow(const ComplexF& r,RealF y){ return(std::pow(r,y)); }
-#endif
-
-//accelerator_inline RealD pow(const RealD& r,RealD y){ return(std::pow(r,y)); }
-//accelerator_inline RealD sqrt(const RealD  & r){ return std::sqrt(r); }
-
-// This comes from ::pow already from math.h and CUDA
-// Calls either Grid::pow for complex, or std::pow for real
-// Problem is CUDA math_functions is exposing ::pow, and I can't define
-
-using std::abs;
-using std::pow;
-using std::sqrt;
-using std::log;
-using std::exp;
-using std::sin;
-using std::cos;
-using std::asin;
-using std::acos;
 
 
-accelerator_inline RealF    conjugate(const RealF  & r){ return r; }
+// type alias used to simplify the syntax of std::enable_if
-accelerator_inline RealD    conjugate(const RealD  & r){ return r; }
+template <typename T> using Invoke = typename T::type;
-accelerator_inline ComplexD conjugate(const ComplexD& r){ return(conj(r)); }
+template <typename Condition, typename ReturnType = void> using EnableIf    = Invoke<std::enable_if<Condition::value, ReturnType> >;
-accelerator_inline ComplexF conjugate(const ComplexF& r ){ return(conj(r)); }
+template <typename Condition, typename ReturnType = void> using NotEnableIf = Invoke<std::enable_if<!Condition::value, ReturnType> >;
 
-accelerator_inline RealF    adj(const RealF  & r){ return r; }
+////////////////////////////////////////////////////////
-accelerator_inline RealD    adj(const RealD  & r){ return r; }
+// Check for complexity with type traits
-accelerator_inline ComplexD adj(const ComplexD& r){ return(conjugate(r)); }
+template <typename T> struct is_complex : public std::false_type {};
-accelerator_inline ComplexF adj(const ComplexF& r ){ return(conjugate(r)); }
+template <> struct is_complex<ComplexD> : public std::true_type {};
+template <> struct is_complex<ComplexF> : public std::true_type {};
 
-#if defined(GRID_CUDA) || defined(GRID_HIP)
+template <typename T> struct is_ComplexD : public std::false_type {};
-//Provide for convenience
+template <> struct is_ComplexD<ComplexD> : public std::true_type {};
-inline std::complex<double> conjugate(const std::complex<double>& r){ return(conj(r)); }
-inline std::complex<float>  conjugate(const std::complex<float>& r) { return(conj(r)); }
-inline std::complex<double> adj(const std::complex<double>& r)      { return(conj(r)); }
-inline std::complex<float>  adj(const std::complex<float>& r)       { return(conj(r)); }
-#endif
 
-accelerator_inline RealF real(const RealF  & r){ return r; }
+template <typename T> struct is_ComplexF : public std::false_type {};
-accelerator_inline RealD real(const RealD  & r){ return r; }
+template <> struct is_ComplexF<ComplexF> : public std::true_type {};
-accelerator_inline RealF real(const ComplexF  & r){ return r.real(); }
-accelerator_inline RealD real(const ComplexD  & r){ return r.real(); }
 
-accelerator_inline RealF imag(const ComplexF  & r){ return r.imag(); }
+template<typename T, typename V=void> struct is_real : public std::false_type {};
-accelerator_inline RealD imag(const ComplexD  & r){ return r.imag(); }
+template<typename T> struct is_real<T, typename std::enable_if<std::is_floating_point<T>::value,
+  void>::type> : public std::true_type {};
 
-accelerator_inline ComplexD innerProduct(const ComplexD & l, const ComplexD & r) { return conjugate(l)*r; }
+template<typename T, typename V=void> struct is_integer : public std::false_type {};
-accelerator_inline ComplexF innerProduct(const ComplexF & l, const ComplexF & r) { return conjugate(l)*r; }
+template<typename T> struct is_integer<T, typename std::enable_if<std::is_integral<T>::value,
-accelerator_inline RealD innerProduct(const RealD & l, const RealD & r) { return l*r; }
+  void>::type> : public std::true_type {};
-accelerator_inline RealF innerProduct(const RealF & l, const RealF & r) { return l*r; }
 
-accelerator_inline ComplexD Reduce(const ComplexD& r){ return r; }
+template <typename T>              using IfReal    = Invoke<std::enable_if<is_real<T>::value, int> >;
-accelerator_inline ComplexF Reduce(const ComplexF& r){ return r; }
+template <typename T>              using IfComplex = Invoke<std::enable_if<is_complex<T>::value, int> >;
-accelerator_inline RealD Reduce(const RealD& r){ return r; }
+template <typename T>              using IfInteger = Invoke<std::enable_if<is_integer<T>::value, int> >;
-accelerator_inline RealF Reduce(const RealF& r){ return r; }
+template <typename T1,typename T2> using IfSame    = Invoke<std::enable_if<std::is_same<T1,T2>::value, int> >;
 
-accelerator_inline RealD toReal(const ComplexD& r){ return r.real(); }
+template <typename T>              using IfNotReal    = Invoke<std::enable_if<!is_real<T>::value, int> >;
-accelerator_inline RealF toReal(const ComplexF& r){ return r.real(); }
+template <typename T>              using IfNotComplex = Invoke<std::enable_if<!is_complex<T>::value, int> >;
-accelerator_inline RealD toReal(const RealD& r){ return r; }
+template <typename T>              using IfNotInteger = Invoke<std::enable_if<!is_integer<T>::value, int> >;
-accelerator_inline RealF toReal(const RealF& r){ return r; }
+template <typename T1,typename T2> using IfNotSame    = Invoke<std::enable_if<!std::is_same<T1,T2>::value, int> >;
-  
-////////////////////////////////////////////////////////////////////////////////
-//Provide support functions for basic real and complex data types required by Grid
-//Single and double precision versions. Should be able to template this once only.
-////////////////////////////////////////////////////////////////////////////////
-accelerator_inline void mac (ComplexD * __restrict__ y,const ComplexD * __restrict__ a,const ComplexD *__restrict__ x){ *y = (*a) * (*x)+(*y); };
-accelerator_inline void mult(ComplexD * __restrict__ y,const ComplexD * __restrict__ l,const ComplexD *__restrict__ r){ *y = (*l) * (*r);}
-accelerator_inline void sub (ComplexD * __restrict__ y,const ComplexD * __restrict__ l,const ComplexD *__restrict__ r){ *y = (*l) - (*r);}
-accelerator_inline void add (ComplexD * __restrict__ y,const ComplexD * __restrict__ l,const ComplexD *__restrict__ r){ *y = (*l) + (*r);}
-// conjugate already supported for complex
-  
-accelerator_inline void mac (ComplexF * __restrict__ y,const ComplexF * __restrict__ a,const ComplexF *__restrict__ x){ *y = (*a) * (*x)+(*y); }
-accelerator_inline void mult(ComplexF * __restrict__ y,const ComplexF * __restrict__ l,const ComplexF *__restrict__ r){ *y = (*l) * (*r); }
-accelerator_inline void sub (ComplexF * __restrict__ y,const ComplexF * __restrict__ l,const ComplexF *__restrict__ r){ *y = (*l) - (*r); }
-accelerator_inline void add (ComplexF * __restrict__ y,const ComplexF * __restrict__ l,const ComplexF *__restrict__ r){ *y = (*l) + (*r); }
-  
-//conjugate already supported for complex
-accelerator_inline ComplexF timesI(const ComplexF &r)     { return(ComplexF(-r.imag(),r.real()));}
-accelerator_inline ComplexD timesI(const ComplexD &r)     { return(ComplexD(-r.imag(),r.real()));}
-accelerator_inline ComplexF timesMinusI(const ComplexF &r){ return(ComplexF(r.imag(),-r.real()));}
-accelerator_inline ComplexD timesMinusI(const ComplexD &r){ return(ComplexD(r.imag(),-r.real()));}
-//accelerator_inline ComplexF timesI(const ComplexF &r)     { return(r*ComplexF(0.0,1.0));}
-//accelerator_inline ComplexD timesI(const ComplexD &r)     { return(r*ComplexD(0.0,1.0));}
-//accelerator_inline ComplexF timesMinusI(const ComplexF &r){ return(r*ComplexF(0.0,-1.0));}
-//accelerator_inline ComplexD timesMinusI(const ComplexD &r){ return(r*ComplexD(0.0,-1.0));}
-
-// define projections to real and imaginay parts
-accelerator_inline ComplexF projReal(const ComplexF &r){return( ComplexF(r.real(), 0.0));}
-accelerator_inline ComplexD projReal(const ComplexD &r){return( ComplexD(r.real(), 0.0));}
-accelerator_inline ComplexF projImag(const ComplexF &r){return (ComplexF(r.imag(), 0.0 ));}
-accelerator_inline ComplexD projImag(const ComplexD &r){return (ComplexD(r.imag(), 0.0));}
-
-// define auxiliary functions for complex computations
-accelerator_inline void timesI(ComplexF &ret,const ComplexF &r)     { ret = timesI(r);}
-accelerator_inline void timesI(ComplexD &ret,const ComplexD &r)     { ret = timesI(r);}
-accelerator_inline void timesMinusI(ComplexF &ret,const ComplexF &r){ ret = timesMinusI(r);}
-accelerator_inline void timesMinusI(ComplexD &ret,const ComplexD &r){ ret = timesMinusI(r);}
-  
-accelerator_inline void mac (RealD * __restrict__ y,const RealD * __restrict__ a,const RealD *__restrict__ x){ *y = (*a) * (*x)+(*y);}
-accelerator_inline void mult(RealD * __restrict__ y,const RealD * __restrict__ l,const RealD *__restrict__ r){ *y = (*l) * (*r);}
-accelerator_inline void sub (RealD * __restrict__ y,const RealD * __restrict__ l,const RealD *__restrict__ r){ *y = (*l) - (*r);}
-accelerator_inline void add (RealD * __restrict__ y,const RealD * __restrict__ l,const RealD *__restrict__ r){ *y = (*l) + (*r);}
-  
-accelerator_inline void mac (RealF * __restrict__ y,const RealF * __restrict__ a,const RealF *__restrict__ x){  *y = (*a) * (*x)+(*y); }
-accelerator_inline void mult(RealF * __restrict__ y,const RealF * __restrict__ l,const RealF *__restrict__ r){ *y = (*l) * (*r); }
-accelerator_inline void sub (RealF * __restrict__ y,const RealF * __restrict__ l,const RealF *__restrict__ r){ *y = (*l) - (*r); }
-accelerator_inline void add (RealF * __restrict__ y,const RealF * __restrict__ l,const RealF *__restrict__ r){ *y = (*l) + (*r); }
-  
-accelerator_inline void vstream(ComplexF &l, const ComplexF &r){ l=r;}
-accelerator_inline void vstream(ComplexD &l, const ComplexD &r){ l=r;}
-accelerator_inline void vstream(RealF &l, const RealF &r){ l=r;}
-accelerator_inline void vstream(RealD &l, const RealD &r){ l=r;}
-  
-accelerator_inline ComplexD toComplex(const RealD &in) { return ComplexD(in);}
-accelerator_inline ComplexF toComplex(const RealF &in) { return ComplexF(in);}
-  
-class Zero{};
-//static Zero Zero();
-template<class itype> accelerator_inline void zeroit(itype &arg)   { arg=Zero();};
-template<>            accelerator_inline void zeroit(ComplexF &arg){ arg=0; };
-template<>            accelerator_inline void zeroit(ComplexD &arg){ arg=0; };
-template<>            accelerator_inline void zeroit(RealF &arg)   { arg=0; };
-template<>            accelerator_inline void zeroit(RealD &arg)   { arg=0; };
-
-// More limited Integer support  
-accelerator_inline Integer Reduce(const Integer& r){ return r; }
-accelerator_inline void mac (Integer * __restrict__ y,const Integer * __restrict__ a,const Integer *__restrict__ x){  *y = (*a) * (*x)+(*y); }
-accelerator_inline void mult(Integer * __restrict__ y,const Integer * __restrict__ l,const Integer *__restrict__ r){ *y = (*l) * (*r); }
-accelerator_inline void sub (Integer * __restrict__ y,const Integer * __restrict__ l,const Integer *__restrict__ r){ *y = (*l) - (*r); }
-accelerator_inline void add (Integer * __restrict__ y,const Integer * __restrict__ l,const Integer *__restrict__ r){ *y = (*l) + (*r); }
-accelerator_inline void vstream(Integer &l, const RealD &r){ l=r;}
-template<>            accelerator_inline void zeroit(Integer &arg)   { arg=0; };
-
-accelerator_inline Integer mod (Integer a,Integer y) { return a%y;}
-accelerator_inline Integer div (Integer a,Integer y) { return a/y;}
-//accelerator_inline Integer abs (Integer &a) { return a%y;}
-
-//////////////////////////////////////////////////////////
-// Permute
-// Permute 0 every ABCDEFGH -> BA DC FE HG
-// Permute 1 every ABCDEFGH -> CD AB GH EF
-// Permute 2 every ABCDEFGH -> EFGH ABCD
-// Permute 3 possible on longer iVector lengths (512bit = 8 double = 16 single)
-// Permute 4 possible on half precision @512bit vectors.
-//
-// Defined inside SIMD specialization files
-//////////////////////////////////////////////////////////
-template<class VectorSIMD>
-accelerator_inline void Gpermute(VectorSIMD &y,const VectorSIMD &b,int perm);
 
+////////////////////////////////////////////////////////
+// Define the operation templates functors
+// general forms to allow for vsplat syntax
+// need explicit declaration of types when used since
+// clang cannot automatically determine the output type sometimes
+template <class Out, class Input1, class Input2, class Input3, class Operation>
+Out accelerator_inline  trinary(Input1 src_1, Input2 src_2, Input3 src_3, Operation op) {
+  return op(src_1, src_2, src_3);
+}
+template <class Out, class Input1, class Input2, class Operation>
+Out accelerator_inline binary(Input1 src_1, Input2 src_2, Operation op) {
+  return op(src_1, src_2);
+}
+template <class Out, class Input, class Operation>
+Out accelerator_inline  unary(Input src, Operation op) {
+  return op(src);
+}
 NAMESPACE_END(Grid);
+///////////////////////////////////////////////
 
 #include <Grid/simd/Grid_vector_types.h>
 #include <Grid/simd/Grid_doubled_vector.h>
+#include <Grid/simd/Grid_scalar_types.h>
+
+#ifdef GRID_SYCL
+template<> struct sycl::is_device_copyable<Grid::vComplexF> : public std::true_type {};
+template<> struct sycl::is_device_copyable<Grid::vComplexD> : public std::true_type {};
+template<> struct sycl::is_device_copyable<Grid::vRealF   > : public std::true_type {};
+template<> struct sycl::is_device_copyable<Grid::vRealD   > : public std::true_type {};
+template<> struct sycl::is_device_copyable<Grid::vInteger > : public std::true_type {};
+template<> struct sycl::is_device_copyable<Grid::sComplexF> : public std::true_type {};
+template<> struct sycl::is_device_copyable<Grid::sComplexD> : public std::true_type {};
+template<> struct sycl::is_device_copyable<Grid::sRealF   > : public std::true_type {};
+template<> struct sycl::is_device_copyable<Grid::sRealD   > : public std::true_type {};
+template<> struct sycl::is_device_copyable<Grid::sInteger > : public std::true_type {};
+#endif
+
+/////////////////////////////////////////
+// Detect vector types
+/////////////////////////////////////////
+NAMESPACE_BEGIN(Grid);
+template <typename T>
+struct is_simd : public std::false_type {};
+template <> struct is_simd<vRealF>     : public std::true_type {};
+template <> struct is_simd<vRealD>     : public std::true_type {};
+template <> struct is_simd<vRealH>     : public std::true_type {};
+template <> struct is_simd<vComplexF>  : public std::true_type {};
+template <> struct is_simd<vComplexD>  : public std::true_type {};
+template <> struct is_simd<vComplexH>  : public std::true_type {};
+template <> struct is_simd<vInteger>   : public std::true_type {};
+template <> struct is_simd<sRealF>     : public std::true_type {};
+template <> struct is_simd<sRealD>     : public std::true_type {};
+template <> struct is_simd<sComplexF>  : public std::true_type {};
+template <> struct is_simd<sComplexD>  : public std::true_type {};
+template <> struct is_simd<sInteger>   : public std::true_type {};
+
+template <typename T> using IfSimd    = Invoke<std::enable_if<is_simd<T>::value, int> >;
+template <typename T> using IfNotSimd = Invoke<std::enable_if<!is_simd<T>::value, unsigned> >;
+
+
+///////////////////////////////////////////////
+// insert / extract with complex support
+///////////////////////////////////////////////
+template <class S, class V>
+accelerator_inline S getlane(const Grid_simd<S, V> &in,int lane) {
+  return in.getlane(lane);
+}
+template <class S, class V>
+accelerator_inline void putlane(Grid_simd<S, V> &vec,const S &_S, int lane){
+  vec.putlane(_S,lane);
+}
+template <class S,IfNotSimd<S> = 0 >
+accelerator_inline S getlane(const S &in,int lane) {
+  return in;
+}
+template <class S,IfNotSimd<S> = 0 >
+accelerator_inline void putlane(S &vec,const S &_S, int lane){
+  vec = _S;
+}
+template <class S, class V>
+accelerator_inline S getlane(const Grid_simd2<S, V> &in,int lane) {
+  return in.getlane(lane);
+}
+template <class S, class V>
+accelerator_inline void putlane(Grid_simd2<S, V> &vec,const S &_S, int lane){
+  vec.putlane(_S,lane);
+}
+template <class S>
+accelerator_inline S getlane(const Grid_simd1<S> &in,int lane) {
+  return in.getlane(lane);
+}
+template <class S>
+accelerator_inline void putlane(Grid_simd1<S> &vec,const S &_S, int lane){
+  vec.putlane(_S,lane);
+}
+
+NAMESPACE_END(Grid);
+
+
+
 #include <Grid/simd/Grid_vector_unops.h>
+#include <Grid/simd/Grid_scalar_unops.h>
 
 NAMESPACE_BEGIN(Grid);
 
@@ -317,4 +264,4 @@ inline std::ostream& operator<< (std::ostream& stream, const vInteger &o){
 }
 
 NAMESPACE_END(Grid)
-#endif
+