QUDA  v1.1.0
A library for QCD on GPUs
Public Types | Public Member Functions | Public Attributes | Static Public Attributes | List of all members
quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc > Struct Template Reference

Accessor routine for ColorSpinorFields in native field order. More...

#include <color_spinor_field_order.h>

Public Types

using Accessor = FloatNOrder< Float, Ns, Nc, N, spin_project, huge_alloc >
 
using real = typename mapper< Float >::type
 
using complex = complex< real >
 
using Vector = typename VectorType< Float, N >::type
 
using GhostVector = typename VectorType< Float, N_ghost >::type
 
using AllocInt = typename AllocType< huge_alloc >::type
 
using norm_type = float
 

Public Member Functions

 FloatNOrder (const ColorSpinorField &a, int nFace=1, Float *field_=0, norm_type *norm_=0, Float **ghost_=0, bool override=false)
 
void resetGhost (const ColorSpinorField &a, void *const *ghost_) const
 
__device__ __host__ void load (complex out[length/2], int x, int parity=0) const
 
__device__ __host__ void save (const complex in[length/2], int x, int parity=0)
 
__device__ __host__ colorspinor_wrapper< real, Accessoroperator() (int x_cb, int parity)
 This accessor routine returns a colorspinor_wrapper to this object, allowing us to overload various operators for manipulating at the site level interms of matrix operations. More...
 
__device__ __host__ const colorspinor_wrapper< real, Accessoroperator() (int x_cb, int parity) const
 This accessor routine returns a const colorspinor_wrapper to this object, allowing us to overload various operators for manipulating at the site level interms of matrix operations. More...
 
__device__ __host__ void loadGhost (complex out[length_ghost/2], int x, int dim, int dir, int parity=0) const
 
__device__ __host__ void saveGhost (const complex in[length_ghost/2], int x, int dim, int dir, int parity=0) const
 
__device__ __host__ colorspinor_ghost_wrapper< real, AccessorGhost (int dim, int dir, int ghost_idx, int parity)
 This accessor routine returns a colorspinor_ghost_wrapper to this object, allowing us to overload various operators for manipulating at the site level interms of matrix operations. More...
 
__device__ __host__ const colorspinor_ghost_wrapper< real, AccessorGhost (int dim, int dir, int ghost_idx, int parity) const
 This accessor routine returns a const colorspinor_ghost_wrapper to this object, allowing us to overload various operators for manipulating at the site level interms of matrix operations. More...
 
void save ()
 Backup the field to the host when tuning. More...
 
void load ()
 Restore the field from the host after tuning. More...
 
size_t Bytes () const
 

Public Attributes

Float * field
 
norm_typenorm
 
const AllocInt offset
 
const AllocInt norm_offset
 
int volumeCB
 
int faceVolumeCB [4]
 
int stride
 
Float * ghost [8]
 
norm_typeghost_norm [8]
 
int nParity
 
void * backup_h
 
size_t bytes
 host memory for backing up the field when tuning More...
 

Static Public Attributes

static constexpr int length = 2 * Ns * Nc
 
static constexpr int length_ghost = spin_project ? length / 2 : length
 
static constexpr int N = N_
 
static constexpr int M = length / N
 
static constexpr int N_ghost = !spin_project ? N : (Ns * Nc) % N == 0 ? N : N / 2
 
static constexpr int M_ghost = length_ghost / N_ghost
 

Detailed Description

template<typename Float, int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
struct quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >

Accessor routine for ColorSpinorFields in native field order.

Template Parameters
FloatUnderlying storage data type of the field
NsNumber of spin components
NcNumber of colors
NNumber of real numbers per short vector
spin_projectWhether the ghosts are spin projected or not
huge_allocTemplate parameter that enables 64-bit pointer arithmetic for huge allocations (e.g., packed set of vectors). Default is to use 32-bit pointer arithmetic.

Definition at line 988 of file color_spinor_field_order.h.

Member Typedef Documentation

◆ Accessor

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
using quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::Accessor = FloatNOrder<Float, Ns, Nc, N, spin_project, huge_alloc>

Definition at line 997 of file color_spinor_field_order.h.

◆ AllocInt

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
using quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::AllocInt = typename AllocType<huge_alloc>::type

Definition at line 1002 of file color_spinor_field_order.h.

◆ complex

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
using quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::complex = complex<real>

Definition at line 999 of file color_spinor_field_order.h.

◆ GhostVector

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
using quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::GhostVector = typename VectorType<Float, N_ghost>::type

Definition at line 1001 of file color_spinor_field_order.h.

◆ norm_type

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
using quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::norm_type = float

Definition at line 1003 of file color_spinor_field_order.h.

◆ real

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
using quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::real = typename mapper<Float>::type

Definition at line 998 of file color_spinor_field_order.h.

◆ Vector

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
using quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::Vector = typename VectorType<Float, N>::type

Definition at line 1000 of file color_spinor_field_order.h.

Constructor & Destructor Documentation

◆ FloatNOrder()

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::FloatNOrder ( const ColorSpinorField a,
int  nFace = 1,
Float *  field_ = 0,
norm_type norm_ = 0,
Float **  ghost_ = 0,
bool  override = false 
)
inline

Definition at line 1017 of file color_spinor_field_order.h.

Member Function Documentation

◆ Bytes()

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
size_t quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::Bytes ( ) const
inline

Definition at line 1243 of file color_spinor_field_order.h.

◆ Ghost() [1/2]

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
__device__ __host__ colorspinor_ghost_wrapper<real, Accessor> quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::Ghost ( int  dim,
int  dir,
int  ghost_idx,
int  parity 
)
inline

This accessor routine returns a colorspinor_ghost_wrapper to this object, allowing us to overload various operators for manipulating at the site level interms of matrix operations.

Parameters
[in]dimDimensions of the ghost we are requesting
[in]ghost_idxCheckerboarded space-time ghost index we are requesting
[in]parityParity we are requesting
Returns
Instance of a colorspinor_ghost_wrapper that curries in access to this field at the above coordinates.

Definition at line 1203 of file color_spinor_field_order.h.

◆ Ghost() [2/2]

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
__device__ __host__ const colorspinor_ghost_wrapper<real, Accessor> quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::Ghost ( int  dim,
int  dir,
int  ghost_idx,
int  parity 
) const
inline

This accessor routine returns a const colorspinor_ghost_wrapper to this object, allowing us to overload various operators for manipulating at the site level interms of matrix operations.

Parameters
[in]dimDimensions of the ghost we are requesting
[in]ghost_idxCheckerboarded space-time ghost index we are requesting
[in]parityParity we are requesting
Returns
Instance of a colorspinor_ghost+wrapper that curries in access to this field at the above coordinates.

Definition at line 1219 of file color_spinor_field_order.h.

◆ load() [1/2]

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
void quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::load ( )
inline

Restore the field from the host after tuning.

Definition at line 1237 of file color_spinor_field_order.h.

◆ load() [2/2]

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
__device__ __host__ void quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::load ( complex  out[length/2],
int  x,
int  parity = 0 
) const
inline

Definition at line 1046 of file color_spinor_field_order.h.

◆ loadGhost()

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
__device__ __host__ void quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::loadGhost ( complex  out[length_ghost/2],
int  x,
int  dim,
int  dir,
int  parity = 0 
) const
inline

Definition at line 1133 of file color_spinor_field_order.h.

◆ operator()() [1/2]

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
__device__ __host__ colorspinor_wrapper<real, Accessor> quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::operator() ( int  x_cb,
int  parity 
)
inline

This accessor routine returns a colorspinor_wrapper to this object, allowing us to overload various operators for manipulating at the site level interms of matrix operations.

Parameters
[in]x_cbCheckerboarded space-time index we are requesting
[in]parityParity we are requesting
Returns
Instance of a colorspinor_wrapper that curries in access to this field at the above coordinates.

Definition at line 1114 of file color_spinor_field_order.h.

◆ operator()() [2/2]

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
__device__ __host__ const colorspinor_wrapper<real, Accessor> quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::operator() ( int  x_cb,
int  parity 
) const
inline

This accessor routine returns a const colorspinor_wrapper to this object, allowing us to overload various operators for manipulating at the site level interms of matrix operations.

Parameters
[in]x_cbCheckerboarded space-time index we are requesting
[in]parityParity we are requesting
Returns
Instance of a colorspinor_wrapper that curries in access to this field at the above coordinates.

Definition at line 1128 of file color_spinor_field_order.h.

◆ resetGhost()

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
void quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::resetGhost ( const ColorSpinorField a,
void *const *  ghost_ 
) const
inline

Definition at line 1033 of file color_spinor_field_order.h.

◆ save() [1/2]

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
void quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::save ( )
inline

Backup the field to the host when tuning.

Definition at line 1228 of file color_spinor_field_order.h.

◆ save() [2/2]

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
__device__ __host__ void quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::save ( const complex  in[length/2],
int  x,
int  parity = 0 
)
inline

Definition at line 1065 of file color_spinor_field_order.h.

◆ saveGhost()

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
__device__ __host__ void quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::saveGhost ( const complex  in[length_ghost/2],
int  x,
int  dim,
int  dir,
int  parity = 0 
) const
inline

Definition at line 1151 of file color_spinor_field_order.h.

Member Data Documentation

◆ backup_h

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
void* quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::backup_h

Definition at line 1014 of file color_spinor_field_order.h.

◆ bytes

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
size_t quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::bytes

host memory for backing up the field when tuning

Definition at line 1015 of file color_spinor_field_order.h.

◆ faceVolumeCB

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
int quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::faceVolumeCB[4]

Definition at line 1009 of file color_spinor_field_order.h.

◆ field

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
Float* quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::field

Definition at line 1004 of file color_spinor_field_order.h.

◆ ghost

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
Float* quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::ghost[8]
mutable

Definition at line 1011 of file color_spinor_field_order.h.

◆ ghost_norm

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
norm_type* quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::ghost_norm[8]
mutable

Definition at line 1012 of file color_spinor_field_order.h.

◆ length

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
constexpr int quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::length = 2 * Ns * Nc
staticconstexpr

Definition at line 990 of file color_spinor_field_order.h.

◆ length_ghost

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
constexpr int quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::length_ghost = spin_project ? length / 2 : length
staticconstexpr

Definition at line 991 of file color_spinor_field_order.h.

◆ M

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
constexpr int quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::M = length / N
staticconstexpr

Definition at line 993 of file color_spinor_field_order.h.

◆ M_ghost

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
constexpr int quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::M_ghost = length_ghost / N_ghost
staticconstexpr

Definition at line 996 of file color_spinor_field_order.h.

◆ N

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
constexpr int quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::N = N_
staticconstexpr

Definition at line 992 of file color_spinor_field_order.h.

◆ N_ghost

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
constexpr int quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::N_ghost = !spin_project ? N : (Ns * Nc) % N == 0 ? N : N / 2
staticconstexpr

Definition at line 995 of file color_spinor_field_order.h.

◆ norm

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
norm_type* quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::norm

Definition at line 1005 of file color_spinor_field_order.h.

◆ norm_offset

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
const AllocInt quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::norm_offset

Definition at line 1007 of file color_spinor_field_order.h.

◆ nParity

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
int quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::nParity

Definition at line 1013 of file color_spinor_field_order.h.

◆ offset

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
const AllocInt quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::offset

Definition at line 1006 of file color_spinor_field_order.h.

◆ stride

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
int quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::stride

Definition at line 1010 of file color_spinor_field_order.h.

◆ volumeCB

template<typename Float , int Ns, int Nc, int N_, bool spin_project = false, bool huge_alloc = false>
int quda::colorspinor::FloatNOrder< Float, Ns, Nc, N_, spin_project, huge_alloc >::volumeCB

Definition at line 1008 of file color_spinor_field_order.h.


The documentation for this struct was generated from the following file: