math_kernel.h File Reference

#include "wave.h"

Include dependency graph for math_kernel.h:

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

Go to the source code of this file.

Functions
template<typename T >
__global__ void	cmult_kernel (T *__restrict__ a, const T c, const int n)
template<typename T >
__global__ void	masked_gather_copy_kernel (T *__restrict__ a, T *__restrict__ b, int *__restrict__ mask, const int n, const int n_mask)
template<typename T >
__global__ void	masked_gather_copy_aligned_kernel (T *__restrict__ a, T *__restrict__ b, int *__restrict__ mask, const int n, const int n_mask)
__device__ __forceinline__ void	face_gather_nonlinear_index (int *index, const int idx, const int lx, const int ly, const int lz)
__device__ __forceinline__ int	face_gather_idx (const int i, const int j, const int k, const int l, const int n1, const int n2, const int nf)
template<typename T >
__global__ void	face_masked_gather_copy_kernel (T *__restrict__ a, const T *__restrict__ b, const int *__restrict__ mask, const int *__restrict__ facet, const int n1, const int n2, const int lx, const int ly, const int lz, const int n_mask)
template<typename T >
__global__ void	masked_scatter_copy_kernel (T *__restrict__ a, T *__restrict__ b, int *__restrict__ mask, const int n, const int n_mask)
template<typename T >
__global__ void	masked_scatter_copy_aligned_kernel (T *__restrict__ a, T *__restrict__ b, int *__restrict__ mask, const int n, const int n_mask)
template<typename T >
__global__ void	masked_atomic_reduction_kernel (T *__restrict__ a, T *__restrict__ b, int *__restrict__ mask, const int n, const int n_mask)
template<typename T >
__global__ void	masked_copy_kernel_0 (T *__restrict__ a, T *__restrict__ b, int *__restrict__ mask, const int n, const int n_mask)
template<typename T >
__global__ void	masked_copy_kernel_aligned (T *__restrict__ a, T *__restrict__ b, int *__restrict__ mask, const int n, const int n_mask)
template<typename T >
__global__ void	cfill_mask_kernel (T *__restrict__ a, const T c, const int n, int *__restrict__ mask, const int n_mask)
template<typename T >
__global__ void	cmult2_kernel (T *__restrict__ a, T *__restrict__ b, const T c, const int n)
template<typename T >
__global__ void	cdiv_kernel (T *__restrict__ a, const T c, const int n)
template<typename T >
__global__ void	cdiv2_kernel (T *__restrict__ a, T *__restrict__ b, const T c, const int n)
template<typename T >
__global__ void	cadd_kernel (T *__restrict__ a, const T c, const int n)
template<typename T >
__global__ void	cadd2_kernel (T *__restrict__ a, T *__restrict__ b, const T c, const int n)
template<typename T >
__global__ void	cwrap_kernel (T *__restrict__ a, const T min_val, const T max_val, const int n)
template<typename T >
__global__ void	sqrt_inplace_kernel (T *__restrict__ a, const int n)
template<typename T >
__global__ void	power_kernel (T *__restrict__ ap, const T *__restrict__ a, const T p, const int n)
template<typename T >
__global__ void	cfill_kernel (T *__restrict__ a, const T c, const int n)
template<typename T >
__global__ void	copy_kernel (T *__restrict__ a, const T *__restrict__ b, const int n)
template<typename T >
__global__ void	add2_kernel (T *__restrict__ a, const T *__restrict__ b, const int n)
template<typename T >
__global__ void	add3_kernel (T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ c, const int n)
template<typename T >
__global__ void	add4_kernel (T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ c, const T *__restrict__ d, const int n)
template<typename T >
__global__ void	add2s1_kernel (T *__restrict__ a, const T *__restrict__ b, const T c1, const int n)
template<typename T >
__global__ void	add2s2_many_kernel (T *__restrict__ x, const T **p, const T *alpha, const int p_cur, const int n)
template<typename T >
__global__ void	add2s2_kernel (T *__restrict__ a, const T *__restrict__ b, const T c1, const int n)
template<typename T >
__global__ void	addsqr2s2_kernel (T *__restrict__ a, const T *__restrict__ b, const T c1, const int n)
template<typename T >
__global__ void	add3s2_kernel (T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ c, const T c1, const T c2, const int n)
template<typename T >
__global__ void	add4s3_kernel (T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ c, const T *__restrict__ d, const T c1, const T c2, const T c3, const int n)
template<typename T >
__global__ void	add5s4_kernel (T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ c, const T *__restrict__ d, const T *__restrict__ e, const T c1, const T c2, const T c3, const T c4, const int n)
template<typename T >
__global__ void	invcol1_kernel (T *__restrict__ a, const int n)
template<typename T >
__global__ void	invcol2_kernel (T *__restrict__ a, const T *__restrict__ b, const int n)
template<typename T >
__global__ void	invcol3_kernel (T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ c, const int n)
template<typename T >
__global__ void	col2_kernel (T *__restrict__ a, const T *__restrict__ b, const int n)
template<typename T >
__global__ void	col3_kernel (T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ c, const int n)
template<typename T >
__global__ void	subcol3_kernel (T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ c, const int n)
template<typename T >
__global__ void	sub2_kernel (T *__restrict__ a, const T *__restrict__ b, const int n)
template<typename T >
__global__ void	sub3_kernel (T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ c, const int n)
template<typename T >
__global__ void	addcol3_kernel (T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ c, const int n)
template<typename T >
__global__ void	addcol4_kernel (T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ c, const T *__restrict__ d, const int n)
template<typename T >
__global__ void	addcol3s2_kernel (T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ c, const T s, const int n)
template<typename T >
__global__ void	vdot3_kernel (T *__restrict__ dot, const T *__restrict__ u1, const T *__restrict__ u2, const T *__restrict__ u3, const T *__restrict__ v1, const T *__restrict__ v2, const T *__restrict__ v3, const int n)
template<typename T >
__global__ void	vcross_kernel (T *__restrict__ u1, T *__restrict__ u2, T *__restrict__ u3, const T *__restrict__ v1, const T *__restrict__ v2, const T *__restrict__ v3, const T *__restrict__ w1, const T *__restrict__ w2, const T *__restrict__ w3, const int n)
template<typename T >
__inline__ __device__ T	reduce_warp (T val)
template<typename T >
__inline__ __device__ T	reduce_max_warp (T val)
template<typename T >
__inline__ __device__ T	reduce_min_warp (T val)
template<typename T >
__global__ void	reduce_kernel (T *bufred, const int n)
template<typename T >
__global__ void	reduce_max_kernel (T *bufred, const T ninf, const int n)
template<typename T >
__global__ void	reduce_min_kernel (T *bufred, const T pinf, const int n)
template<typename T_acc >
__global__ void	glsc3_reduce_kernel (T_acc *bufred, const int n, const int j)
template<typename T >
__global__ void	vlsc3_kernel (const T *a, const T *b, const T *c, T *buf_h, const int n)
template<typename T , typename T_acc >
__global__ void	glsc3_kernel (const T *a, const T *b, const T *c, T_acc *buf_h, const int n)
template<typename T , typename T_acc >
__global__ void	glsc3_many_kernel (const T *a, const T **b, const T *c, T_acc *buf_h, const int j, const int n)
template<typename T , typename T_acc >
__global__ void	glsc2_kernel (const T *a, const T *b, T_acc *buf_h, const int n)
template<typename T , typename T_acc >
__global__ void	glsubnorm2_kernel (const T *a, const T *b, T_acc *buf_h, const int n)
template<typename T , typename T_acc >
__global__ void	glsum_kernel (const T *a, T_acc *buf_h, const int n)
template<typename T >
__global__ void	glmax_kernel (const T *a, const T ninf, T *buf_h, const int n)
template<typename T >
__global__ void	glmin_kernel (const T *a, const T pinf, T *buf_h, const int n)
template<typename T >
__global__ void	absval_kernel (T *__restrict__ a, const int n)
template<typename T >
__global__ void	pwmax_vec2_kernel (T *__restrict__ a, const T *__restrict__ b, const int n)
template<typename T >
__global__ void	pwmax_vec3_kernel (T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ c, const int n)
template<typename T >
__global__ void	pwmax_sca2_kernel (T *__restrict__ a, const T c, const int n)
template<typename T >
__global__ void	pwmax_sca3_kernel (T *__restrict__ a, const T *__restrict b, const T c, const int n)
template<typename T >
__global__ void	pwmin_vec2_kernel (T *__restrict__ a, const T *__restrict__ b, const int n)
template<typename T >
__global__ void	pwmin_vec3_kernel (T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ c, const int n)
template<typename T >
__global__ void	pwmin_sca2_kernel (T *__restrict__ a, const T c, const int n)
template<typename T >
__global__ void	pwmin_sca3_kernel (T *__restrict__ a, const T *__restrict b, const T c, const int n)

Function Documentation

absval_kernel()

template<typename T >

__global__ void absval_kernel	(	T *__restrict__	a,
		const int	n
	)

Device kernel for abs_value

Definition at line 1321 of file math_kernel.h.

Here is the call graph for this function:

add2_kernel()

template<typename T >

__global__ void add2_kernel	(	T *__restrict__	a,
		const T *__restrict__	b,
		const int	n
	)

Device kernel for add2

Definition at line 428 of file math_kernel.h.

Here is the call graph for this function:

add2s1_kernel()

template<typename T >

__global__ void add2s1_kernel	(	T *__restrict__	a,
		const T *__restrict__	b,
		const T	c1,
		const int	n
	)

Device kernel for add2s1

Definition at line 479 of file math_kernel.h.

Here is the call graph for this function:

add2s2_kernel()

template<typename T >

__global__ void add2s2_kernel	(	T *__restrict__	a,
		const T *__restrict__	b,
		const T	c1,
		const int	n
	)

Device kernel for add2s2

Definition at line 519 of file math_kernel.h.

Here is the call graph for this function:

add2s2_many_kernel()

template<typename T >

__global__ void add2s2_many_kernel	(	T *__restrict__	x,
		const T **	p,
		const T *	alpha,
		const int	p_cur,
		const int	n
	)

Device kernel for add2s2

Definition at line 496 of file math_kernel.h.

Here is the call graph for this function:

add3_kernel()

template<typename T >

__global__ void add3_kernel	(	T *__restrict__	a,
		const T *__restrict__	b,
		const T *__restrict__	c,
		const int	n
	)

Device kernel for add3

Definition at line 444 of file math_kernel.h.

Here is the call graph for this function:

add3s2_kernel()

template<typename T >

__global__ void add3s2_kernel	(	T *__restrict__	a,
		const T *__restrict__	b,
		const T *__restrict__	c,
		const T	c1,
		const T	c2,
		const int	n
	)

Device kernel for add3s2

Definition at line 553 of file math_kernel.h.

Here is the call graph for this function:

add4_kernel()

template<typename T >

__global__ void add4_kernel	(	T *__restrict__	a,
		const T *__restrict__	b,
		const T *__restrict__	c,
		const T *__restrict__	d,
		const int	n
	)

Device kernel for add4

Definition at line 461 of file math_kernel.h.

Here is the call graph for this function:

add4s3_kernel()

template<typename T >

__global__ void add4s3_kernel	(	T *__restrict__	a,
		const T *__restrict__	b,
		const T *__restrict__	c,
		const T *__restrict__	d,
		const T	c1,
		const T	c2,
		const T	c3,
		const int	n
	)

Device kernel for add4s3

Definition at line 572 of file math_kernel.h.

Here is the call graph for this function:

add5s4_kernel()

template<typename T >

__global__ void add5s4_kernel	(	T *__restrict__	a,
		const T *__restrict__	b,
		const T *__restrict__	c,
		const T *__restrict__	d,
		const T *__restrict__	e,
		const T	c1,
		const T	c2,
		const T	c3,
		const T	c4,
		const int	n
	)

Device kernel for add5s4

Definition at line 593 of file math_kernel.h.

Here is the call graph for this function:

addcol3_kernel()

template<typename T >

__global__ void addcol3_kernel	(	T *__restrict__	a,
		const T *__restrict__	b,
		const T *__restrict__	c,
		const int	n
	)

Device kernel for addcol3

Definition at line 748 of file math_kernel.h.

Here is the call graph for this function:

addcol3s2_kernel()

template<typename T >

__global__ void addcol3s2_kernel	(	T *__restrict__	a,
		const T *__restrict__	b,
		const T *__restrict__	c,
		const T	s,
		const int	n
	)

Device kernel for addcol3s2

Definition at line 785 of file math_kernel.h.

Here is the call graph for this function:

addcol4_kernel()

template<typename T >

__global__ void addcol4_kernel	(	T *__restrict__	a,
		const T *__restrict__	b,
		const T *__restrict__	c,
		const T *__restrict__	d,
		const int	n
	)

Device kernel for addcol4

Definition at line 766 of file math_kernel.h.

Here is the call graph for this function:

addsqr2s2_kernel()

template<typename T >

__global__ void addsqr2s2_kernel	(	T *__restrict__	a,
		const T *__restrict__	b,
		const T	c1,
		const int	n
	)

Device kernel for addsqr2s2

Definition at line 536 of file math_kernel.h.

Here is the call graph for this function:

cadd2_kernel()

template<typename T >

__global__ void cadd2_kernel	(	T *__restrict__	a,
		T *__restrict__	b,
		const T	c,
		const int	n
	)

Device kernel for cadd2

Definition at line 329 of file math_kernel.h.

Here is the call graph for this function:

cadd_kernel()

template<typename T >

__global__ void cadd_kernel	(	T *__restrict__	a,
		const T	c,
		const int	n
	)

Device kernel for cadd

Definition at line 313 of file math_kernel.h.

Here is the call graph for this function:

cdiv2_kernel()

template<typename T >

__global__ void cdiv2_kernel	(	T *__restrict__	a,
		T *__restrict__	b,
		const T	c,
		const int	n
	)

Device kernel for cdiv2

Definition at line 296 of file math_kernel.h.

Here is the call graph for this function:

cdiv_kernel()

template<typename T >

__global__ void cdiv_kernel	(	T *__restrict__	a,
		const T	c,
		const int	n
	)

Device kernel for cdiv

Definition at line 280 of file math_kernel.h.

Here is the call graph for this function:

cfill_kernel()

template<typename T >

__global__ void cfill_kernel	(	T *__restrict__	a,
		const T	c,
		const int	n
	)

Device kernel for cmult

Definition at line 396 of file math_kernel.h.

Here is the call graph for this function:

cfill_mask_kernel()

template<typename T >

__global__ void cfill_mask_kernel	(	T *__restrict__	a,
		const T	c,
		const int	n,
		int *__restrict__	mask,
		const int	n_mask
	)

Device kernel for cfill_mask

Definition at line 247 of file math_kernel.h.

Here is the call graph for this function:

cmult2_kernel()

template<typename T >

__global__ void cmult2_kernel	(	T *__restrict__	a,
		T *__restrict__	b,
		const T	c,
		const int	n
	)

Device kernel for cmult2

Definition at line 263 of file math_kernel.h.

Here is the call graph for this function:

cmult_kernel()

template<typename T >

__global__ void cmult_kernel	(	T *__restrict__	a,
		const T	c,
		const int	n
	)

Device kernel for cmult

Definition at line 43 of file math_kernel.h.

Here is the call graph for this function:

col2_kernel()

template<typename T >

__global__ void col2_kernel	(	T *__restrict__	a,
		const T *__restrict__	b,
		const int	n
	)

Device kernel for col2

Definition at line 665 of file math_kernel.h.

Here is the call graph for this function:

col3_kernel()

template<typename T >

__global__ void col3_kernel	(	T *__restrict__	a,
		const T *__restrict__	b,
		const T *__restrict__	c,
		const int	n
	)

Device kernel for col3

Definition at line 681 of file math_kernel.h.

Here is the call graph for this function:

copy_kernel()

template<typename T >

__global__ void copy_kernel	(	T *__restrict__	a,
		const T *__restrict__	b,
		const int	n
	)

Device kernel for copy

Definition at line 412 of file math_kernel.h.

Here is the call graph for this function:

cwrap_kernel()

template<typename T >

__global__ void cwrap_kernel	(	T *__restrict__	a,
		const T	min_val,
		const T	max_val,
		const int	n
	)

Device kernel for cwrap

Definition at line 346 of file math_kernel.h.

Here is the call graph for this function:

face_gather_idx()

__device__ __forceinline__ int face_gather_idx	(	const int	i,
		const int	j,
		const int	k,
		const int	l,
		const int	n1,
		const int	n2,
		const int	nf
	)

Definition at line 107 of file math_kernel.h.

Here is the call graph for this function:

face_gather_nonlinear_index()

__device__ __forceinline__ void face_gather_nonlinear_index	(	int *	index,
		const int	idx,
		const int	lx,
		const int	ly,
		const int	lz
	)

Definition at line 92 of file math_kernel.h.

Here is the call graph for this function:

face_masked_gather_copy_kernel()

template<typename T >

__global__ void face_masked_gather_copy_kernel	(	T *__restrict__	a,
		const T *__restrict__	b,
		const int *__restrict__	mask,
		const int *__restrict__	facet,
		const int	n1,
		const int	n2,
		const int	lx,
		const int	ly,
		const int	lz,
		const int	n_mask
	)

Device kernel for masked gather copy from a face-local field

Definition at line 116 of file math_kernel.h.

Here is the call graph for this function:

glmax_kernel()

template<typename T >

__global__ void glmax_kernel	(	const T *	a,
		const T	ninf,
		T *	buf_h,
		const int	n
	)

Device kernel for glmax

Definition at line 1249 of file math_kernel.h.

Here is the call graph for this function:

glmin_kernel()

template<typename T >

__global__ void glmin_kernel	(	const T *	a,
		const T	pinf,
		T *	buf_h,
		const int	n
	)

Device kernel for glmin

Definition at line 1285 of file math_kernel.h.

Here is the call graph for this function:

glsc2_kernel()

template<typename T , typename T_acc >

__global__ void glsc2_kernel	(	const T *	a,
		const T *	b,
		T_acc *	buf_h,
		const int	n
	)

Device kernel for glsc2

Definition at line 1144 of file math_kernel.h.

Here is the call graph for this function:

glsc3_kernel()

template<typename T , typename T_acc >

__global__ void glsc3_kernel	(	const T *	a,
		const T *	b,
		const T *	c,
		T_acc *	buf_h,
		const int	n
	)

Device kernel for glsc3

Definition at line 1068 of file math_kernel.h.

Here is the call graph for this function:

glsc3_many_kernel()

template<typename T , typename T_acc >

__global__ void glsc3_many_kernel	(	const T *	a,
		const T **	b,
		const T *	c,
		T_acc *	buf_h,
		const int	j,
		const int	n
	)

Device kernel for glsc3 many

Definition at line 1103 of file math_kernel.h.

Here is the call graph for this function:

glsc3_reduce_kernel()

template<typename T_acc >

__global__ void glsc3_reduce_kernel	(	T_acc *	bufred,
		const int	n,
		const int	j
	)

Reduction kernel for glsc3

Definition at line 999 of file math_kernel.h.

Here is the call graph for this function:

glsubnorm2_kernel()

template<typename T , typename T_acc >

__global__ void glsubnorm2_kernel	(	const T *	a,
		const T *	b,
		T_acc *	buf_h,
		const int	n
	)

Device kernel for glsubnorm

Definition at line 1179 of file math_kernel.h.

Here is the call graph for this function:

glsum_kernel()

template<typename T , typename T_acc >

__global__ void glsum_kernel	(	const T *	a,
		T_acc *	buf_h,
		const int	n
	)

Device kernel for glsum

Definition at line 1214 of file math_kernel.h.

Here is the call graph for this function:

invcol1_kernel()

template<typename T >

__global__ void invcol1_kernel	(	T *__restrict__	a,
		const int	n
	)

Device kernel for invcol1

Definition at line 616 of file math_kernel.h.

Here is the call graph for this function:

invcol2_kernel()

template<typename T >

__global__ void invcol2_kernel	(	T *__restrict__	a,
		const T *__restrict__	b,
		const int	n
	)

Device kernel for invcol2

Definition at line 632 of file math_kernel.h.

Here is the call graph for this function:

invcol3_kernel()

template<typename T >

__global__ void invcol3_kernel	(	T *__restrict__	a,
		const T *__restrict__	b,
		const T *__restrict__	c,
		const int	n
	)

Device kernel for invcol3

Definition at line 648 of file math_kernel.h.

Here is the call graph for this function:

masked_atomic_reduction_kernel()

template<typename T >

__global__ void masked_atomic_reduction_kernel	(	T *__restrict__	a,
		T *__restrict__	b,
		int *__restrict__	mask,
		const int	n,
		const int	n_mask
	)

Device kernel for masked atomic update

Definition at line 192 of file math_kernel.h.

Here is the call graph for this function:

masked_copy_kernel_0()

template<typename T >

__global__ void masked_copy_kernel_0	(	T *__restrict__	a,
		T *__restrict__	b,
		int *__restrict__	mask,
		const int	n,
		const int	n_mask
	)

Device kernel for masked copy with BC style mask

Definition at line 211 of file math_kernel.h.

Here is the call graph for this function:

masked_copy_kernel_aligned()

template<typename T >

__global__ void masked_copy_kernel_aligned	(	T *__restrict__	a,
		T *__restrict__	b,
		int *__restrict__	mask,
		const int	n,
		const int	n_mask
	)

Device kernel for masked copy with point zone style mask

Definition at line 229 of file math_kernel.h.

Here is the call graph for this function:

masked_gather_copy_aligned_kernel()

template<typename T >

__global__ void masked_gather_copy_aligned_kernel	(	T *__restrict__	a,
		T *__restrict__	b,
		int *__restrict__	mask,
		const int	n,
		const int	n_mask
	)

Device kernel for masked gather copy aligned

Definition at line 77 of file math_kernel.h.

Here is the call graph for this function:

masked_gather_copy_kernel()

template<typename T >

__global__ void masked_gather_copy_kernel	(	T *__restrict__	a,
		T *__restrict__	b,
		int *__restrict__	mask,
		const int	n,
		const int	n_mask
	)

Device kernel for masked gather copy

Definition at line 59 of file math_kernel.h.

Here is the call graph for this function:

masked_scatter_copy_aligned_kernel()

template<typename T >

__global__ void masked_scatter_copy_aligned_kernel	(	T *__restrict__	a,
		T *__restrict__	b,
		int *__restrict__	mask,
		const int	n,
		const int	n_mask
	)

Device kernel for masked scatter copy with aligned mask

Definition at line 174 of file math_kernel.h.

Here is the call graph for this function:

masked_scatter_copy_kernel()

template<typename T >

__global__ void masked_scatter_copy_kernel	(	T *__restrict__	a,
		T *__restrict__	b,
		int *__restrict__	mask,
		const int	n,
		const int	n_mask
	)

Device kernel for masked scatter copy

Definition at line 156 of file math_kernel.h.

Here is the call graph for this function:

power_kernel()

template<typename T >

__global__ void power_kernel	(	T *__restrict__	ap,
		const T *__restrict__	a,
		const T	p,
		const int	n
	)

Device kernel for power

Definition at line 379 of file math_kernel.h.

Here is the call graph for this function:

pwmax_sca2_kernel()

template<typename T >

__global__ void pwmax_sca2_kernel	(	T *__restrict__	a,
		const T	c,
		const int	n
	)

Device kernel for point-wise max of vector and scalar a = max(a, c)

Definition at line 1369 of file math_kernel.h.

Here is the call graph for this function:

pwmax_sca3_kernel()

template<typename T >

__global__ void pwmax_sca3_kernel	(	T *__restrict__	a,
		const T *__restrict	b,
		const T	c,
		const int	n
	)

Device kernel for point-wise max of vector and scalar a = max(b, c)

Definition at line 1382 of file math_kernel.h.

Here is the call graph for this function:

pwmax_vec2_kernel()

template<typename T >

__global__ void pwmax_vec2_kernel	(	T *__restrict__	a,
		const T *__restrict__	b,
		const int	n
	)

Device kernel for point-wise max of two vectors a = max(a, b)

Definition at line 1341 of file math_kernel.h.

Here is the call graph for this function:

pwmax_vec3_kernel()

template<typename T >

__global__ void pwmax_vec3_kernel	(	T *__restrict__	a,
		const T *__restrict__	b,
		const T *__restrict__	c,
		const int	n
	)

Device kernel for point-wise max of two vectors a = max(b, c)

Definition at line 1354 of file math_kernel.h.

Here is the call graph for this function:

pwmin_sca2_kernel()

template<typename T >

__global__ void pwmin_sca2_kernel	(	T *__restrict__	a,
		const T	c,
		const int	n
	)

Device kernel for point-wise min of vector and scalar a = min(a, c)

Definition at line 1425 of file math_kernel.h.

Here is the call graph for this function:

pwmin_sca3_kernel()

template<typename T >

__global__ void pwmin_sca3_kernel	(	T *__restrict__	a,
		const T *__restrict	b,
		const T	c,
		const int	n
	)

Device kernel for point-wise min of vector and scalar a = min(b, c)

Definition at line 1438 of file math_kernel.h.

Here is the call graph for this function:

pwmin_vec2_kernel()

template<typename T >

__global__ void pwmin_vec2_kernel	(	T *__restrict__	a,
		const T *__restrict__	b,
		const int	n
	)

Device kernel for point-wise min of two vectors a = min(a, b)

Definition at line 1397 of file math_kernel.h.

Here is the call graph for this function:

pwmin_vec3_kernel()

template<typename T >

__global__ void pwmin_vec3_kernel	(	T *__restrict__	a,
		const T *__restrict__	b,
		const T *__restrict__	c,
		const int	n
	)

Device kernel for point-wise min of two vectors a = min(b, c)

Definition at line 1410 of file math_kernel.h.

Here is the call graph for this function:

reduce_kernel()

template<typename T >

__global__ void reduce_kernel	(	T *	bufred,
		const int	n
	)

Vector reduction kernel

Definition at line 905 of file math_kernel.h.

Here is the call graph for this function:

reduce_max_kernel()

template<typename T >

__global__ void reduce_max_kernel	(	T *	bufred,
		const T	ninf,
		const int	n
	)

Vector reduction maximisation kernel

Definition at line 936 of file math_kernel.h.

Here is the call graph for this function:

reduce_max_warp()

template<typename T >

__inline__ __device__ T reduce_max_warp

(

T

val

)

Warp shuffle reduction of maximisation

Definition at line 873 of file math_kernel.h.

Here is the call graph for this function:

reduce_min_kernel()

template<typename T >

__global__ void reduce_min_kernel	(	T *	bufred,
		const T	pinf,
		const int	n
	)

Vector reduction minimisation kernel

Definition at line 967 of file math_kernel.h.

Here is the call graph for this function:

reduce_min_warp()

template<typename T >

__inline__ __device__ T reduce_min_warp

(

T

val

)

Warp shuffle reduction of minimisation

Definition at line 889 of file math_kernel.h.

Here is the call graph for this function:

reduce_warp()

template<typename T >

__inline__ __device__ T reduce_warp

(

T

val

)

Warp shuffle reduction

The ladder is stepped down from NEKO_WAVE_SIZE/2, so the leading half wave shuffle is only emitted on a 64 lane wavefront. At wave64 this preprocesses to exactly the sequence it always has; at wave32 the 32 step would have shuffled past the end of the wave, where __shfl_down returns the lane's own value and the reduction silently doubles it.

Definition at line 857 of file math_kernel.h.

Here is the call graph for this function:

sqrt_inplace_kernel()

template<typename T >

__global__ void sqrt_inplace_kernel	(	T *__restrict__	a,
		const int	n
	)

Device kernel for sqrt_inplace

Definition at line 364 of file math_kernel.h.

Here is the call graph for this function:

sub2_kernel()

template<typename T >

__global__ void sub2_kernel	(	T *__restrict__	a,
		const T *__restrict__	b,
		const int	n
	)

Device kernel for sub2

Definition at line 715 of file math_kernel.h.

Here is the call graph for this function:

sub3_kernel()

template<typename T >

__global__ void sub3_kernel	(	T *__restrict__	a,
		const T *__restrict__	b,
		const T *__restrict__	c,
		const int	n
	)

Device kernel for sub3

Definition at line 731 of file math_kernel.h.

Here is the call graph for this function:

subcol3_kernel()

template<typename T >

__global__ void subcol3_kernel	(	T *__restrict__	a,
		const T *__restrict__	b,
		const T *__restrict__	c,
		const int	n
	)

Device kernel for subcol3

Definition at line 698 of file math_kernel.h.

Here is the call graph for this function:

vcross_kernel()

template<typename T >

__global__ void vcross_kernel	(	T *__restrict__	u1,
		T *__restrict__	u2,
		T *__restrict__	u3,
		const T *__restrict__	v1,
		const T *__restrict__	v2,
		const T *__restrict__	v3,
		const T *__restrict__	w1,
		const T *__restrict__	w2,
		const T *__restrict__	w3,
		const int	n
	)

Device kernel for vcross

Definition at line 826 of file math_kernel.h.

Here is the call graph for this function:

vdot3_kernel()

template<typename T >

__global__ void vdot3_kernel	(	T *__restrict__	dot,
		const T *__restrict__	u1,
		const T *__restrict__	u2,
		const T *__restrict__	u3,
		const T *__restrict__	v1,
		const T *__restrict__	v2,
		const T *__restrict__	v3,
		const int	n
	)

Device kernel for vdot3

Definition at line 804 of file math_kernel.h.

Here is the call graph for this function:

vlsc3_kernel()

template<typename T >

__global__ void vlsc3_kernel	(	const T *	a,
		const T *	b,
		const T *	c,
		T *	buf_h,
		const int	n
	)

Device kernel for vlsc3

Definition at line 1033 of file math_kernel.h.

Here is the call graph for this function: