From 971b0f3546e0038b2f9c0e19e4285e6b91296d0d Mon Sep 17 00:00:00 2001 From: Peter Boyle Date: Wed, 19 Aug 2026 20:00:26 -0400 Subject: [PATCH] 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 | 2 + Grid/simd/Grid_a64fx-fixedsize.h | 2 + Grid/simd/Grid_doubled_vector.h | 48 ----- Grid/simd/Grid_neon.h | 53 +++++ Grid/simd/Grid_vector_types.h | 117 +---------- Grid/simd/Simd.h | 331 +++++++++++++------------------ 6 files changed, 207 insertions(+), 346 deletions(-) diff --git a/Grid/simd/Grid_a64fx-2.h b/Grid/simd/Grid_a64fx-2.h index a4ef70ae8..b538f2da1 100644 --- a/Grid/simd/Grid_a64fx-2.h +++ b/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; diff --git a/Grid/simd/Grid_a64fx-fixedsize.h b/Grid/simd/Grid_a64fx-fixedsize.h index 5bf1b0a39..e4768823c 100644 --- a/Grid/simd/Grid_a64fx-fixedsize.h +++ b/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; diff --git a/Grid/simd/Grid_doubled_vector.h b/Grid/simd/Grid_doubled_vector.h index ee6047507..43ccecde2 100644 --- a/Grid/simd/Grid_doubled_vector.h +++ b/Grid/simd/Grid_doubled_vector.h @@ -339,54 +339,6 @@ typedef Grid_simd2 , vComplexD> vComplexD2; typedef Grid_simd2 vRealD2; - -///////////////////////////////////////// -// Some traits to recognise the types -///////////////////////////////////////// -template -struct is_simd : public std::false_type {}; -template <> struct is_simd : public std::true_type {}; -template <> struct is_simd : public std::true_type {}; -template <> struct is_simd : public std::true_type {}; -template <> struct is_simd : public std::true_type {}; -template <> struct is_simd : public std::true_type {}; -template <> struct is_simd : public std::true_type {}; -template <> struct is_simd : public std::true_type {}; -template <> struct is_simd : public std::true_type {}; -template <> struct is_simd : public std::true_type {}; - -template using IfSimd = Invoke::value, int> >; -template using IfNotSimd = Invoke::value, unsigned> >; - -/////////////////////////////////////////////// -// insert / extract with complex support -/////////////////////////////////////////////// -template -accelerator_inline S getlane(const Grid_simd &in,int lane) { - return in.getlane(lane); -} -template -accelerator_inline void putlane(Grid_simd &vec,const S &_S, int lane){ - vec.putlane(_S,lane); -} -template = 0 > -accelerator_inline S getlane(const S &in,int lane) { - return in; -} -template = 0 > -accelerator_inline void putlane(S &vec,const S &_S, int lane){ - vec = _S; -} -template -accelerator_inline S getlane(const Grid_simd2 &in,int lane) { - return in.getlane(lane); -} -template -accelerator_inline void putlane(Grid_simd2 &vec,const S &_S, int lane){ - vec.putlane(_S,lane); -} - - //////////////////////////////////////////////////////////////////// // General rotate //////////////////////////////////////////////////////////////////// diff --git a/Grid/simd/Grid_neon.h b/Grid/simd/Grid_neon.h index 16ee4f805..59c90b3c7 100644 --- a/Grid/simd/Grid_neon.h +++ b/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 @@ -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; diff --git a/Grid/simd/Grid_vector_types.h b/Grid/simd/Grid_vector_types.h index 976d55689..86b71e16c 100644 --- a/Grid/simd/Grid_vector_types.h +++ b/Grid/simd/Grid_vector_types.h @@ -36,8 +36,7 @@ Author: Michael Marshall */ // Time-stamp: <2015-07-10 17:45:33 neo> //--------------------------------------------------------------------------- -#ifndef GRID_VECTOR_TYPES -#define GRID_VECTOR_TYPES +#pragma once // 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 -struct RealPart { - typedef T type; -}; -template -struct RealPart > { - typedef T type; -}; - -#include - -////////////////////////////////////// -// demote a vector to real type -////////////////////////////////////// -// type alias used to simplify the syntax of std::enable_if -template using Invoke = typename T::type; -template using EnableIf = Invoke >; -template using NotEnableIf = Invoke >; - -//////////////////////////////////////////////////////// -// Check for complexity with type traits -template struct is_complex : public std::false_type {}; -template <> struct is_complex : public std::true_type {}; -template <> struct is_complex : public std::true_type {}; - -template struct is_ComplexD : public std::false_type {}; -template <> struct is_ComplexD : public std::true_type {}; - -template struct is_ComplexF : public std::false_type {}; -template <> struct is_ComplexF : public std::true_type {}; - -template struct is_real : public std::false_type {}; -template struct is_real::value, - void>::type> : public std::true_type {}; - -template struct is_integer : public std::false_type {}; -template struct is_integer::value, - void>::type> : public std::true_type {}; - -template using IfReal = Invoke::value, int> >; -template using IfComplex = Invoke::value, int> >; -template using IfInteger = Invoke::value, int> >; -template using IfSame = Invoke::value, int> >; - -template using IfNotReal = Invoke::value, int> >; -template using IfNotComplex = Invoke::value, int> >; -template using IfNotInteger = Invoke::value, int> >; -template using IfNotSame = Invoke::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 -Out accelerator_inline trinary(Input1 src_1, Input2 src_2, Input3 src_3, Operation op) { - return op(src_1, src_2, src_3); -} -template -Out accelerator_inline binary(Input1 src_1, Input2 src_2, Operation op) { - return op(src_1, src_2); -} -template -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::value ) { + y->v = trinary(a->v, x->v, y->v, MultAddComplexSIMD()); + } else { + y->v = trinary(a->v, x->v, y->v, MultSIMD()); + } }; #else friend accelerator_inline void mac(Grid_simd *__restrict__ y, @@ -670,7 +604,6 @@ typedef Grid_simd, SIMD_Htype> vComplexH; typedef Grid_simd , SIMD_Ftype> vComplexF; typedef Grid_simd , SIMD_Dtype> vComplexD; #endif - ///////////////////////////////////////// // Pointer type to use on extractLane ///////////////////////////////////////// @@ -885,25 +818,7 @@ accelerator_inline Grid_simd operator*(Grid_simd a, Grid_simd return ret; }; -// ---------------- A64FX MAC ------------------- -// Distinguish between complex types and others -#if defined(A64FX) || defined(A64FXFIXEDSIZE) -template = 0> -accelerator_inline Grid_simd fxmac(Grid_simd a, Grid_simd b, Grid_simd c) { - Grid_simd ret; - ret.v = trinary(a.v, b.v, c.v, MultAddComplexSIMD()); - return ret; -}; - -// Real/Integer types -template = 0> -accelerator_inline Grid_simd fxmac(Grid_simd a, Grid_simd b, Grid_simd c) { - Grid_simd ret; - ret.v = trinary(a.v, b.v, c.v, MultSIMD()); - return ret; -}; -#endif -// ---------------------------------------------- +// mac() selects MultAddComplexSIMD directly where the architecture has it. /////////////////////// @@ -1141,13 +1056,3 @@ template void gpermute(vobj & inout,int perm){ NAMESPACE_END(Grid); -#ifdef GRID_SYCL -template<> struct sycl::is_device_copyable : public std::true_type {}; -template<> struct sycl::is_device_copyable : public std::true_type {}; -template<> struct sycl::is_device_copyable : public std::true_type {}; -template<> struct sycl::is_device_copyable : public std::true_type {}; -template<> struct sycl::is_device_copyable : public std::true_type {}; -#endif - - -#endif diff --git a/Grid/simd/Simd.h b/Grid/simd/Simd.h index b9e706b6b..6b35649f8 100644 --- a/Grid/simd/Simd.h +++ b/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 -#define GRID_SIMD_H +#pragma once -#if defined(GRID_CUDA) || defined(GRID_HIP) -#include -#endif +#include -//////////////////////////////////////////////////////////////////////// -// 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) +#include NAMESPACE_BEGIN(Grid); - -typedef uint32_t Integer; - -typedef float RealF; -typedef double RealD; -#ifdef GRID_DEFAULT_PRECISION_DOUBLE -typedef RealD Real; -#else -typedef RealF Real; -#endif - -#if defined(GRID_CUDA) || defined(GRID_HIP) -typedef thrust::complex ComplexF; -typedef thrust::complex ComplexD; -typedef thrust::complex Complex; -typedef thrust::complex ComplexH; -template using complex = thrust::complex; - -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 ComplexF; -typedef std::complex ComplexD; -typedef std::complex Complex; -typedef std::complex ComplexH; // Hack -template using complex = std::complex; - -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; +////////////////////////////////////// +// To take the floating point type of real/complex type +////////////////////////////////////// +template +struct RealPart { + typedef T type; +}; +template +struct RealPart > { + typedef T type; +}; -accelerator_inline RealF conjugate(const RealF & r){ return r; } -accelerator_inline RealD conjugate(const RealD & r){ return r; } -accelerator_inline ComplexD conjugate(const ComplexD& r){ return(conj(r)); } -accelerator_inline ComplexF conjugate(const ComplexF& r ){ return(conj(r)); } +// type alias used to simplify the syntax of std::enable_if +template using Invoke = typename T::type; +template using EnableIf = Invoke >; +template using NotEnableIf = Invoke >; -accelerator_inline RealF adj(const RealF & r){ return r; } -accelerator_inline RealD adj(const RealD & r){ return r; } -accelerator_inline ComplexD adj(const ComplexD& r){ return(conjugate(r)); } -accelerator_inline ComplexF adj(const ComplexF& r ){ return(conjugate(r)); } +//////////////////////////////////////////////////////// +// Check for complexity with type traits +template struct is_complex : public std::false_type {}; +template <> struct is_complex : public std::true_type {}; +template <> struct is_complex : public std::true_type {}; -#if defined(GRID_CUDA) || defined(GRID_HIP) -//Provide for convenience -inline std::complex conjugate(const std::complex& r){ return(conj(r)); } -inline std::complex conjugate(const std::complex& r) { return(conj(r)); } -inline std::complex adj(const std::complex& r) { return(conj(r)); } -inline std::complex adj(const std::complex& r) { return(conj(r)); } -#endif +template struct is_ComplexD : public std::false_type {}; +template <> struct is_ComplexD : public std::true_type {}; -accelerator_inline RealF real(const RealF & r){ return r; } -accelerator_inline RealD real(const RealD & r){ return r; } -accelerator_inline RealF real(const ComplexF & r){ return r.real(); } -accelerator_inline RealD real(const ComplexD & r){ return r.real(); } +template struct is_ComplexF : public std::false_type {}; +template <> struct is_ComplexF : public std::true_type {}; -accelerator_inline RealF imag(const ComplexF & r){ return r.imag(); } -accelerator_inline RealD imag(const ComplexD & r){ return r.imag(); } +template struct is_real : public std::false_type {}; +template struct is_real::value, + void>::type> : public std::true_type {}; -accelerator_inline ComplexD innerProduct(const ComplexD & l, const ComplexD & r) { return conjugate(l)*r; } -accelerator_inline ComplexF innerProduct(const ComplexF & l, const ComplexF & r) { return conjugate(l)*r; } -accelerator_inline RealD innerProduct(const RealD & l, const RealD & r) { return l*r; } -accelerator_inline RealF innerProduct(const RealF & l, const RealF & r) { return l*r; } +template struct is_integer : public std::false_type {}; +template struct is_integer::value, + void>::type> : public std::true_type {}; -accelerator_inline ComplexD Reduce(const ComplexD& r){ return r; } -accelerator_inline ComplexF Reduce(const ComplexF& r){ return r; } -accelerator_inline RealD Reduce(const RealD& r){ return r; } -accelerator_inline RealF Reduce(const RealF& r){ return r; } +template using IfReal = Invoke::value, int> >; +template using IfComplex = Invoke::value, int> >; +template using IfInteger = Invoke::value, int> >; +template using IfSame = Invoke::value, int> >; -accelerator_inline RealD toReal(const ComplexD& r){ return r.real(); } -accelerator_inline RealF toReal(const ComplexF& r){ return r.real(); } -accelerator_inline RealD toReal(const RealD& r){ return r; } -accelerator_inline RealF toReal(const RealF& r){ return r; } - -//////////////////////////////////////////////////////////////////////////////// -//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 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 -accelerator_inline void Gpermute(VectorSIMD &y,const VectorSIMD &b,int perm); +template using IfNotReal = Invoke::value, int> >; +template using IfNotComplex = Invoke::value, int> >; +template using IfNotInteger = Invoke::value, int> >; +template using IfNotSame = Invoke::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 +Out accelerator_inline trinary(Input1 src_1, Input2 src_2, Input3 src_3, Operation op) { + return op(src_1, src_2, src_3); +} +template +Out accelerator_inline binary(Input1 src_1, Input2 src_2, Operation op) { + return op(src_1, src_2); +} +template +Out accelerator_inline unary(Input src, Operation op) { + return op(src); +} NAMESPACE_END(Grid); +/////////////////////////////////////////////// #include #include +#include + +#ifdef GRID_SYCL +template<> struct sycl::is_device_copyable : public std::true_type {}; +template<> struct sycl::is_device_copyable : public std::true_type {}; +template<> struct sycl::is_device_copyable : public std::true_type {}; +template<> struct sycl::is_device_copyable : public std::true_type {}; +template<> struct sycl::is_device_copyable : public std::true_type {}; +template<> struct sycl::is_device_copyable : public std::true_type {}; +template<> struct sycl::is_device_copyable : public std::true_type {}; +template<> struct sycl::is_device_copyable : public std::true_type {}; +template<> struct sycl::is_device_copyable : public std::true_type {}; +template<> struct sycl::is_device_copyable : public std::true_type {}; +#endif + +///////////////////////////////////////// +// Detect vector types +///////////////////////////////////////// +NAMESPACE_BEGIN(Grid); +template +struct is_simd : public std::false_type {}; +template <> struct is_simd : public std::true_type {}; +template <> struct is_simd : public std::true_type {}; +template <> struct is_simd : public std::true_type {}; +template <> struct is_simd : public std::true_type {}; +template <> struct is_simd : public std::true_type {}; +template <> struct is_simd : public std::true_type {}; +template <> struct is_simd : public std::true_type {}; +template <> struct is_simd : public std::true_type {}; +template <> struct is_simd : public std::true_type {}; +template <> struct is_simd : public std::true_type {}; +template <> struct is_simd : public std::true_type {}; +template <> struct is_simd : public std::true_type {}; + +template using IfSimd = Invoke::value, int> >; +template using IfNotSimd = Invoke::value, unsigned> >; + + +/////////////////////////////////////////////// +// insert / extract with complex support +/////////////////////////////////////////////// +template +accelerator_inline S getlane(const Grid_simd &in,int lane) { + return in.getlane(lane); +} +template +accelerator_inline void putlane(Grid_simd &vec,const S &_S, int lane){ + vec.putlane(_S,lane); +} +template = 0 > +accelerator_inline S getlane(const S &in,int lane) { + return in; +} +template = 0 > +accelerator_inline void putlane(S &vec,const S &_S, int lane){ + vec = _S; +} +template +accelerator_inline S getlane(const Grid_simd2 &in,int lane) { + return in.getlane(lane); +} +template +accelerator_inline void putlane(Grid_simd2 &vec,const S &_S, int lane){ + vec.putlane(_S,lane); +} +template +accelerator_inline S getlane(const Grid_simd1 &in,int lane) { + return in.getlane(lane); +} +template +accelerator_inline void putlane(Grid_simd1 &vec,const S &_S, int lane){ + vec.putlane(_S,lane); +} + +NAMESPACE_END(Grid); + + + #include +#include NAMESPACE_BEGIN(Grid); @@ -317,4 +264,4 @@ inline std::ostream& operator<< (std::ostream& stream, const vInteger &o){ } NAMESPACE_END(Grid) -#endif +