bicgstab_kernel.h File Reference

#include <math/bcknd/device/hip/math_kernel.h>

Include dependency graph for bicgstab_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	bicgstab_update_p_kernel (T *__restrict__ p, const T *__restrict__ r, const T *__restrict__ v, const T beta, const T omega, const int n)
template<typename T , typename T_acc >
__global__ void	bicgstab_product_and_norm_kernel (const T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ mult, T_acc *__restrict__ buf_h, const int n)
template<typename T , typename T_acc >
__global__ void	bicgstab_part1_kernel (T *__restrict__ s, const T *__restrict__ r, const T *__restrict__ v, const T *__restrict__ mult, T_acc *__restrict__ buf_h, const T alpha, const int n)
template<typename T , typename T_acc >
__global__ void	bicgstab_part2_kernel (T *__restrict__ x, T *__restrict__ r, const T *__restrict__ p_hat, const T *__restrict__ s_hat, const T *__restrict__ s, const T *__restrict__ t, const T *__restrict__ f, const T *__restrict__ mult, T_acc *__restrict__ buf_h, const T alpha, const T omega, const int n)

Function Documentation

bicgstab_part1_kernel()

template<typename T , typename T_acc >

__global__ void bicgstab_part1_kernel	(	T *__restrict__	s,
		const T *__restrict__	r,
		const T *__restrict__	v,
		const T *__restrict__	mult,
		T_acc *__restrict__	buf_h,
		const T	alpha,
		const int	n
	)

Kernel for BiCGStab part 1, \( s = r - \alpha v \), reducing \( s^T M s \)

Definition at line 116 of file bicgstab_kernel.h.

Here is the call graph for this function:

bicgstab_part2_kernel()

template<typename T , typename T_acc >

__global__ void bicgstab_part2_kernel	(	T *__restrict__	x,
		T *__restrict__	r,
		const T *__restrict__	p_hat,
		const T *__restrict__	s_hat,
		const T *__restrict__	s,
		const T *__restrict__	t,
		const T *__restrict__	f,
		const T *__restrict__	mult,
		T_acc *__restrict__	buf_h,
		const T	alpha,
		const T	omega,
		const int	n
	)

Kernel for BiCGStab part 2, \( x = x + \alpha \hat{p} + \omega \hat{s} \) and \( r = s - \omega t \), reducing \( (r^T M r, f^T M r) \)

The second value is the rho inner product of the next iteration, which the recurrence would otherwise reduce separately at the top of the loop.

Definition at line 160 of file bicgstab_kernel.h.

Here is the call graph for this function:

bicgstab_product_and_norm_kernel()

template<typename T , typename T_acc >

__global__ void bicgstab_product_and_norm_kernel	(	const T *__restrict__	a,
		const T *__restrict__	b,
		const T *__restrict__	mult,
		T_acc *__restrict__	buf_h,
		const int	n
	)

Kernel for a weighted inner product and squared norm in one pass, \( (a^T M b, b^T M b) \)

The two partial sums of each block are stored interleaved, so that glsc3_reduce_kernel with j = 2 leaves the results in buf_h[0..1].

Definition at line 68 of file bicgstab_kernel.h.

Here is the call graph for this function:

bicgstab_update_p_kernel()

template<typename T >

__global__ void bicgstab_update_p_kernel	(	T *__restrict__	p,
		const T *__restrict__	r,
		const T *__restrict__	v,
		const T	beta,
		const T	omega,
		const int	n
	)

Kernel for the BiCGStab search direction update \( p = r + \beta (p - \omega v) \)

Definition at line 45 of file bicgstab_kernel.h.

Here is the call graph for this function: