QUDA  v0.7.0
A library for QCD on GPUs
All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Pages
Classes | Namespaces | Functions | Variables
reduce_mixed_core.h File Reference
#include <launch_kernel.cuh>

Go to the source code of this file.

Classes

struct  mixed::ReduceArg< ReduceType, SpinorX, SpinorY, SpinorZ, SpinorW, SpinorV, Reducer >
 
class  mixed::ReduceCuda< doubleN, ReduceType, ReduceSimpleType, FloatN, M, SpinorX, SpinorY, SpinorZ, SpinorW, SpinorV, Reducer >
 

Namespaces

 mixed
 

Functions

__host__ __device__ void mixed::zero (double &x)
 
__host__ __device__ void mixed::zero (double2 &x)
 
__host__ __device__ void mixed::zero (double3 &x)
 
__device__ void mixed::copytoshared (double *s, const int i, const double x, const int block)
 
__device__ void mixed::copytoshared (double *s, const int i, const double2 x, const int block)
 
__device__ void mixed::copytoshared (double *s, const int i, const double3 x, const int block)
 
__device__ void mixed::copytoshared (volatile double *s, const int i, const double x, const int block)
 
__device__ void mixed::copytoshared (volatile double *s, const int i, const double2 x, const int block)
 
__device__ void mixed::copytoshared (volatile double *s, const int i, const double3 x, const int block)
 
__device__ void mixed::copyfromshared (double &x, const double *s, const int i, const int block)
 
__device__ void mixed::copyfromshared (double2 &x, const double *s, const int i, const int block)
 
__device__ void mixed::copyfromshared (double3 &x, const double *s, const int i, const int block)
 
template<typename ReduceType , typename ReduceSimpleType >
__device__ void mixed::add (ReduceType &sum, ReduceSimpleType *s, const int i, const int block)
 
template<>
__device__ void mixed::add< double, double > (double &sum, double *s, const int i, const int block)
 
template<>
__device__ void mixed::add< double2, double > (double2 &sum, double *s, const int i, const int block)
 
template<>
__device__ void mixed::add< double3, double > (double3 &sum, double *s, const int i, const int block)
 
template<typename ReduceType , typename ReduceSimpleType >
__device__ void mixed::add (ReduceSimpleType *s, const int i, const int j, const int block)
 
template<typename ReduceType , typename ReduceSimpleType >
__device__ void mixed::add (volatile ReduceSimpleType *s, const int i, const int j, const int block)
 
template<>
__device__ void mixed::add< double, double > (double *s, const int i, const int j, const int block)
 
template<>
__device__ void mixed::add< double, double > (volatile double *s, const int i, const int j, const int block)
 
template<>
__device__ void mixed::add< double2, double > (double *s, const int i, const int j, const int block)
 
template<>
__device__ void mixed::add< double2, double > (volatile double *s, const int i, const int j, const int block)
 
template<>
__device__ void mixed::add< double3, double > (double *s, const int i, const int j, const int block)
 
template<>
__device__ void mixed::add< double3, double > (volatile double *s, const int i, const int j, const int block)
 
__host__ __device__ void mixed::zero (doublesingle &x)
 
__host__ __device__ void mixed::zero (doublesingle2 &x)
 
__host__ __device__ void mixed::zero (doublesingle3 &x)
 
__device__ void mixed::copytoshared (doublesingle *s, const int i, const doublesingle x, const int block)
 
__device__ void mixed::copytoshared (doublesingle *s, const int i, const doublesingle2 x, const int block)
 
__device__ void mixed::copytoshared (doublesingle *s, const int i, const doublesingle3 x, const int block)
 
__device__ void mixed::copytoshared (volatile doublesingle *s, const int i, const doublesingle x, const int block)
 
__device__ void mixed::copytoshared (volatile doublesingle *s, const int i, const doublesingle2 x, const int block)
 
__device__ void mixed::copytoshared (volatile doublesingle *s, const int i, const doublesingle3 x, const int block)
 
__device__ void mixed::copyfromshared (doublesingle &x, const doublesingle *s, const int i, const int block)
 
__device__ void mixed::copyfromshared (doublesingle2 &x, const doublesingle *s, const int i, const int block)
 
__device__ void mixed::copyfromshared (doublesingle3 &x, const doublesingle *s, const int i, const int block)
 
template<>
__device__ void mixed::add< doublesingle, doublesingle > (doublesingle &sum, doublesingle *s, const int i, const int block)
 
template<>
__device__ void mixed::add< doublesingle2, doublesingle > (doublesingle2 &sum, doublesingle *s, const int i, const int block)
 
template<>
__device__ void mixed::add< doublesingle3, doublesingle > (doublesingle3 &sum, doublesingle *s, const int i, const int block)
 
template<>
__device__ void mixed::add< doublesingle, doublesingle > (doublesingle *s, const int i, const int j, const int block)
 
template<>
__device__ void mixed::add< doublesingle, doublesingle > (volatile doublesingle *s, const int i, const int j, const int block)
 
template<>
__device__ void mixed::add< doublesingle2, doublesingle > (doublesingle *s, const int i, const int j, const int block)
 
template<>
__device__ void mixed::add< doublesingle2, doublesingle > (volatile doublesingle *s, const int i, const int j, const int block)
 
template<>
__device__ void mixed::add< doublesingle3, doublesingle > (doublesingle *s, const int i, const int j, const int block)
 
template<>
__device__ void mixed::add< doublesingle3, doublesingle > (volatile doublesingle *s, const int i, const int j, const int block)
 
template<int block_size, typename ReduceType , typename ReduceSimpleType , typename FloatN , int M, typename SpinorX , typename SpinorY , typename SpinorZ , typename SpinorW , typename SpinorV , typename Reducer >
__global__ void mixed::reduceKernel (ReduceArg< ReduceType, SpinorX, SpinorY, SpinorZ, SpinorW, SpinorV, Reducer > arg)
 
template<typename doubleN , typename ReduceType , typename ReduceSimpleType , typename FloatN , int M, typename SpinorX , typename SpinorY , typename SpinorZ , typename SpinorW , typename SpinorV , typename Reducer >
doubleN mixed::reduceLaunch (ReduceArg< ReduceType, SpinorX, SpinorY, SpinorZ, SpinorW, SpinorV, Reducer > &arg, const TuneParam &tp, const cudaStream_t &stream)
 
template<typename doubleN , typename ReduceType , typename ReduceSimpleType , template< typename ReducerType, typename Float, typename FloatN > class Reducer, int writeX, int writeY, int writeZ, int writeW, int writeV, bool siteUnroll>
doubleN mixed::reduceCuda (const double2 &a, const double2 &b, cudaColorSpinorField &x, cudaColorSpinorField &y, cudaColorSpinorField &z, cudaColorSpinorField &w, cudaColorSpinorField &v)
 

Variables

__device__ unsigned int mixed::count = 0
 
__shared__ bool mixed::isLastBlockDone