coarse_op_kernel.cuh File Reference

#include <color_spinor_field_order.h>
#include <gauge_field_order.h>
#include <clover_field_order.h>
#include <multigrid_helper.cuh>
#include <index_helper.cuh>
#include <gamma.cuh>
#include <linalg.cuh>

Include dependency graph for coarse_op_kernel.cuh:

This graph shows which files directly or indirectly include this file:

Go to the source code of this file.

Classes
struct	quda::CalculateYArg< Float, fineSpin, coarseSpin, fineColor, coarseColor, coarseGauge, coarseGaugeAtomic, fineGauge, fineSpinor, fineSpinorTmp, fineSpinorV, fineClover >

Namespaces
	quda

Macros
#define	max_color_per_block   8

Typedefs
typedef int	quda::storeType

Functions
template<typename Float >
__device__ __host__ void	quda::caxpy (const complex< Float > &a, const complex< Float > &x, complex< Float > &y)
template<bool from_coarse, typename Float , int dim, QudaDirection dir, int fineSpin, int fineColor, int coarseSpin, int coarseColor, typename Wtype , typename Arg >
__device__ __host__ void	quda::computeUV (Arg &arg, const Wtype &W, int parity, int x_cb, int ic_c)
template<bool from_coarse, typename Float , int dim, QudaDirection dir, int fineSpin, int fineColor, int coarseSpin, int coarseColor, typename Arg >
void	quda::ComputeUVCPU (Arg &arg)
template<bool from_coarse, typename Float , int dim, QudaDirection dir, int fineSpin, int fineColor, int coarseSpin, int coarseColor, typename Arg >
__global__ void	quda::ComputeUVGPU (Arg arg)
template<typename Float , int fineSpin, int fineColor, int coarseColor, typename Arg >
__device__ __host__ void	quda::computeAV (Arg &arg, int parity, int x_cb, int ch, int ic_c)
template<typename Float , int fineSpin, int fineColor, int coarseColor, typename Arg >
void	quda::ComputeAVCPU (Arg &arg)
template<typename Float , int fineSpin, int fineColor, int coarseColor, typename Arg >
__global__ void	quda::ComputeAVGPU (Arg arg)
template<typename Float , int fineSpin, int fineColor, int coarseColor, typename Arg >
__device__ __host__ void	quda::computeTMAV (Arg &arg, int parity, int x_cb, int v)
template<typename Float , int fineSpin, int fineColor, int coarseColor, typename Arg >
void	quda::ComputeTMAVCPU (Arg &arg)
template<typename Float , int fineSpin, int fineColor, int coarseColor, typename Arg >
__global__ void	quda::ComputeTMAVGPU (Arg arg)
template<typename Float , int fineSpin, int fineColor, int coarseColor, typename Arg >
__device__ __host__ void	quda::computeTMCAV (Arg &arg, int parity, int x_cb, int ch, int ic_c)
template<typename Float , int fineSpin, int fineColor, int coarseColor, typename Arg >
void	quda::ComputeTMCAVCPU (Arg &arg)
template<typename Float , int fineSpin, int fineColor, int coarseColor, typename Arg >
__global__ void	quda::ComputeTMCAVGPU (Arg arg)
template<bool from_coarse, typename Float , int dim, QudaDirection dir, int fineSpin, int fineColor, int coarseSpin, int coarseColor, typename Arg , typename Gamma >
__device__ __host__ void	quda::multiplyVUV (complex< Float > vuv[], const Arg &arg, const Gamma &gamma, int parity, int x_cb, int ic_c, int jc_c)
	Do a single (AV)^ * UV product, where for preconditioned clover, AV correspond to the clover inverse multiplied by the packed null space vectors, else AV is simply the packed null space vectors. More...
template<typename Arg >
__device__ __host__ int	quda::virtualThreadIdx (const Arg &arg)
template<typename Arg >
__device__ __host__ int	quda::virtualBlockDim (const Arg &arg)
template<typename Arg >
__device__ __host__ int	quda::coarseIndex (const Arg &arg)
template<bool shared_atomic, bool parity_flip, bool from_coarse, typename Float , int dim, QudaDirection dir, int fineSpin, int fineColor, int coarseSpin, int coarseColor, typename Arg , typename Gamma >
__device__ __host__ void	quda::computeVUV (Arg &arg, const Gamma &gamma, int parity, int x_cb, int c_row, int c_col, int parity_coarse_, int coarse_x_cb_)
template<bool from_coarse, typename Float , int dim, QudaDirection dir, int fineSpin, int fineColor, int coarseSpin, int coarseColor, typename Arg >
void	quda::ComputeVUVCPU (Arg arg)
template<bool parity_flip, typename Arg >
__device__ void	quda::getIndicesShared (const Arg &arg, int &parity, int &x_cb, int &parity_coarse, int &x_coarse_cb, int &c_col, int &c_row)
template<bool parity_flip, typename Arg >
__device__ void	quda::getIndicesGlobal (const Arg &arg, int &parity, int &x_cb, int &parity_coarse, int &x_coarse_cb, int &c_col, int &c_row)
template<bool shared_atomic, bool parity_flip, bool from_coarse, typename Float , int dim, QudaDirection dir, int fineSpin, int fineColor, int coarseSpin, int coarseColor, typename Arg >
__global__ void	quda::ComputeVUVGPU (Arg arg)
template<typename Float , int nSpin, int nColor, typename Arg >
__device__ __host__ void	quda::computeYreverse (Arg &arg, int parity, int x_cb, int ic_c, int jc_c)
template<typename Float , int nSpin, int nColor, typename Arg >
void	quda::ComputeYReverseCPU (Arg &arg)
template<typename Float , int nSpin, int nColor, typename Arg >
__global__ void	quda::ComputeYReverseGPU (Arg arg)
template<bool from_coarse, typename Float , int fineSpin, int coarseSpin, int fineColor, int coarseColor, typename Arg >
__device__ __host__ void	quda::computeCoarseClover (Arg &arg, int parity, int x_cb, int ic_c, int jc_c)
template<bool from_coarse, typename Float , int fineSpin, int coarseSpin, int fineColor, int coarseColor, typename Arg >
void	quda::ComputeCoarseCloverCPU (Arg &arg)
template<bool from_coarse, typename Float , int fineSpin, int coarseSpin, int fineColor, int coarseColor, typename Arg >
__global__ void	quda::ComputeCoarseCloverGPU (Arg arg)
template<typename Float , int nSpin, int nColor, typename Arg >
void	quda::AddCoarseDiagonalCPU (Arg &arg)
template<typename Float , int nSpin, int nColor, typename Arg >
__global__ void	quda::AddCoarseDiagonalGPU (Arg arg)
template<typename Float , int nSpin, int nColor, typename Arg >
void	quda::AddCoarseTmDiagonalCPU (Arg &arg)
template<typename Float , int nSpin, int nColor, typename Arg >
__global__ void	quda::AddCoarseTmDiagonalGPU (Arg arg)
template<typename Float , int nSpin, int nColor, typename Arg >
__device__ __host__ void	quda::convert (Arg &arg, int parity, int x_cb, int c_row, int c_col)
template<typename Float , int nSpin, int nColor, typename Arg >
void	quda::ConvertCPU (Arg &arg)
template<typename Float , int nSpin, int nColor, typename Arg >
__global__ void	quda::ConvertGPU (Arg arg)
template<typename Float , int nSpin, int nColor, typename Arg >
__device__ __host__ void	quda::rescaleY (Arg &arg, int parity, int x_cb, int c_row, int c_col)
template<typename Float , int nSpin, int nColor, typename Arg >
void	quda::RescaleYCPU (Arg &arg)
template<typename Float , int nSpin, int nColor, typename Arg >
__global__ void	quda::RescaleYGPU (Arg arg)

Macro Definition Documentation

max_color_per_block

#define max_color_per_block   8

Definition at line 9 of file coarse_op_kernel.cuh.

Referenced by quda::CalculateY< from_coarse, Float, fineSpin, fineColor, coarseSpin, coarseColor, Arg >::advanceAux(), quda::computeVUV(), quda::CalculateY< from_coarse, Float, fineSpin, fineColor, coarseSpin, coarseColor, Arg >::setComputeType(), and quda::CalculateY< from_coarse, Float, fineSpin, fineColor, coarseSpin, coarseColor, Arg >::sharedBytesPerBlock().