QUDA  1.0.0
All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros
Public Types | Public Member Functions | Static Public Member Functions | Public Attributes | List of all members
quda::blas::HeavyQuarkResidualNorm_< ReduceType, Float2, FloatN > Struct Template Reference
Inheritance diagram for quda::blas::HeavyQuarkResidualNorm_< ReduceType, Float2, FloatN >:
Inheritance graph
[legend]
Collaboration diagram for quda::blas::HeavyQuarkResidualNorm_< ReduceType, Float2, FloatN >:
Collaboration graph
[legend]

Public Types

typedef scalar< ReduceType >::type real
 

Public Member Functions

 HeavyQuarkResidualNorm_ (const Float2 &a, const Float2 &b)
 
__device__ __host__ void pre ()
 pre-computation routine called before the "M-loop" More...
 
__device__ __host__ void operator() (ReduceType &sum, FloatN &x, FloatN &y, FloatN &z, FloatN &w, FloatN &v)
 where the reduction is usually computed and any auxiliary operations More...
 
__device__ __host__ void post (ReduceType &sum)
 sum the solution and residual norms, and compute the heavy-quark norm More...
 

Static Public Member Functions

static int streams ()
 
static int flops ()
 total number of input and output streams More...
 

Public Attributes

Float2 a
 
Float2 b
 
ReduceType aux
 

Detailed Description

template<typename ReduceType, typename Float2, typename FloatN>
struct quda::blas::HeavyQuarkResidualNorm_< ReduceType, Float2, FloatN >

This kernel returns (x, x) and (r,r) and also returns the so-called heavy quark norm as used by MILC: 1 / N * (r, r)_i / (x, x)_i, where i is site index and N is the number of sites. When this kernel is launched, we must enforce that the parameter M in the launcher corresponds to the number of FloatN fields used to represent the spinor, e.g., M=6 for Wilson and M=3 for staggered. This is only the case for half-precision kernels by default. To enable this, the siteUnroll template parameter must be set true when reduceCuda is instantiated.

Definition at line 471 of file reduce_core.cuh.

Member Typedef Documentation

◆ real

template<typename ReduceType , typename Float2 , typename FloatN >
typedef scalar<ReduceType>::type quda::blas::HeavyQuarkResidualNorm_< ReduceType, Float2, FloatN >::real

Definition at line 472 of file reduce_core.cuh.

Constructor & Destructor Documentation

◆ HeavyQuarkResidualNorm_()

template<typename ReduceType , typename Float2 , typename FloatN >
quda::blas::HeavyQuarkResidualNorm_< ReduceType, Float2, FloatN >::HeavyQuarkResidualNorm_ ( const Float2 &  a,
const Float2 &  b 
)
inline

Definition at line 476 of file reduce_core.cuh.

Member Function Documentation

◆ flops()

template<typename ReduceType , typename Float2 , typename FloatN >
static int quda::blas::HeavyQuarkResidualNorm_< ReduceType, Float2, FloatN >::flops ( )
inlinestatic

total number of input and output streams

Definition at line 499 of file reduce_core.cuh.

◆ operator()()

template<typename ReduceType , typename Float2 , typename FloatN >
__device__ __host__ void quda::blas::HeavyQuarkResidualNorm_< ReduceType, Float2, FloatN >::operator() ( ReduceType &  sum,
FloatN &  x,
FloatN &  y,
FloatN &  z,
FloatN &  w,
FloatN &  v 
)
inlinevirtual

where the reduction is usually computed and any auxiliary operations

Implements quda::blas::ReduceFunctor< ReduceType, Float2, FloatN >.

Definition at line 484 of file reduce_core.cuh.

◆ post()

template<typename ReduceType , typename Float2 , typename FloatN >
__device__ __host__ void quda::blas::HeavyQuarkResidualNorm_< ReduceType, Float2, FloatN >::post ( ReduceType &  sum)
inlinevirtual

sum the solution and residual norms, and compute the heavy-quark norm

Reimplemented from quda::blas::ReduceFunctor< ReduceType, Float2, FloatN >.

Definition at line 491 of file reduce_core.cuh.

◆ pre()

template<typename ReduceType , typename Float2 , typename FloatN >
__device__ __host__ void quda::blas::HeavyQuarkResidualNorm_< ReduceType, Float2, FloatN >::pre ( )
inlinevirtual

pre-computation routine called before the "M-loop"

Reimplemented from quda::blas::ReduceFunctor< ReduceType, Float2, FloatN >.

Definition at line 478 of file reduce_core.cuh.

◆ streams()

template<typename ReduceType , typename Float2 , typename FloatN >
static int quda::blas::HeavyQuarkResidualNorm_< ReduceType, Float2, FloatN >::streams ( )
inlinestatic

Definition at line 498 of file reduce_core.cuh.

Member Data Documentation

◆ a

template<typename ReduceType , typename Float2 , typename FloatN >
Float2 quda::blas::HeavyQuarkResidualNorm_< ReduceType, Float2, FloatN >::a

Definition at line 473 of file reduce_core.cuh.

◆ aux

template<typename ReduceType , typename Float2 , typename FloatN >
ReduceType quda::blas::HeavyQuarkResidualNorm_< ReduceType, Float2, FloatN >::aux

Definition at line 475 of file reduce_core.cuh.

◆ b

template<typename ReduceType , typename Float2 , typename FloatN >
Float2 quda::blas::HeavyQuarkResidualNorm_< ReduceType, Float2, FloatN >::b

Definition at line 474 of file reduce_core.cuh.


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