9.2
general documentation
Macros | Functions | Variables
cs_sles_it_hip.cpp File Reference
#include "base/cs_defs.h"
#include <assert.h>
#include <algorithm>
#include <list>
#include <hip/hip_runtime.h>
#include <hip/hip_runtime_api.h>
#include <cfloat>
#include "bft/bft_error.h"
#include "base/cs_base.h"
#include "base/cs_base_accel.h"
#include "alge/cs_blas_hip.h"
#include "base/cs_dispatch.h"
#include "base/cs_log.h"
#include "alge/cs_matrix.h"
#include "alge/cs_matrix_priv.h"
#include "alge/cs_matrix_spmv_hip.h"
#include "base/cs_mem.h"
#include "base/cs_profiling.h"
#include "alge/cs_sles_it.h"
#include "alge/cs_sles_pc.h"
#include "alge/cs_sles_it_priv.h"
#include "alge/cs_sles_it_hip.h"
+ Include dependency graph for cs_sles_it_hip.cpp:

Macros

#define HAVE_GRAPH_CAPTURE   0
 
#define CS_SIMD_SIZE(s)   (((s-1)/CS_HIP_WARP_SIZE+1)*CS_HIP_WARP_SIZE)
 
#define CS_BLOCKSIZE   256
 

Functions

template<size_t block_size, typename T >
static __device__ void __forceinline__ _warp_reduce_sum (volatile T *stmp, size_t tid)
 Kernel for sum reduction within a warp (for warp size 32). More...
 
template<size_t block_size>
static __global__ void _jacobi_compute_vx (cs_lnum_t n_rows, const cs_real_t *__restrict__ ad_inv, const cs_real_t *__restrict__ rhs, cs_real_t *__restrict__ vx, cs_real_t *__restrict__ rk)
 
template<size_t block_size>
static __global__ void _jacobi_compute_vx_ini0 (cs_lnum_t n_rows, const cs_real_t *__restrict__ ad_inv, const cs_real_t *__restrict__ rhs, cs_real_t *__restrict__ vx, cs_real_t *__restrict__ rk)
 
template<size_t block_size>
static __global__ void _jacobi_compute_vx_and_residual (cs_lnum_t n_rows, const cs_real_t *__restrict__ ad_inv, const cs_real_t *__restrict__ ad, const cs_real_t *__restrict__ rhs, cs_real_t *__restrict__ vx, cs_real_t *__restrict__ rk, double *__restrict__ sum_block)
 
template<size_t block_size>
static __global__ void _jacobi_compute_vx_and_residual_ini0 (cs_lnum_t n_rows, const cs_real_t *__restrict__ ad_inv, const cs_real_t *__restrict__ rhs, cs_real_t *__restrict__ vx, cs_real_t *__restrict__ rk, double *__restrict__ sum_block)
 
static __global__ void _block_3_jacobi_compute_vx (cs_lnum_t n_rows, const cs_real_t *__restrict__ ad_inv, const cs_real_t *__restrict__ rhs, const cs_real_t *__restrict__ vxx, cs_real_t *__restrict__ vx)
 
static __global__ void _block_3_jacobi_compute_vx0 (cs_lnum_t n_rows, const cs_real_t *__restrict__ ad_inv, const cs_real_t *__restrict__ rhs, cs_real_t *__restrict__ rk, cs_real_t *__restrict__ vx)
 
template<size_t block_size>
static __global__ void _block_3_jacobi_compute_residual (cs_lnum_t n_rows, const cs_real_t *__restrict__ ad, const cs_real_t *__restrict__ vx, const cs_real_t *__restrict__ rk, double *__restrict__ sum_block)
 
static __global__ void _block_jacobi_compute_vx (cs_lnum_t n_rows, cs_lnum_t db_size, const cs_real_t *__restrict__ ad_inv, const cs_real_t *__restrict__ rhs, const cs_real_t *__restrict__ vxx, cs_real_t *__restrict__ vx)
 
static __global__ void _block_jacobi_compute_vx0 (cs_lnum_t n_rows, cs_lnum_t db_size, const cs_real_t *__restrict__ ad_inv, const cs_real_t *__restrict__ rhs, cs_real_t *__restrict__ rk, cs_real_t *__restrict__ vx)
 
template<size_t block_size>
static __global__ void _block_jacobi_compute_residual (cs_lnum_t n_rows, cs_lnum_t db_size, const cs_real_t *__restrict__ ad, cs_real_t *__restrict__ vx, cs_real_t *__restrict__ rk, double *__restrict__ sum_block)
 
static __global__ void _fcg_init (cs_lnum_t n, const cs_real_t *__restrict__ rhs, cs_real_t *__restrict__ rk, cs_real_t *__restrict__ qk)
 
static __global__ void _fcg_init_vx0 (cs_lnum_t n, const cs_real_t *__restrict__ rhs, cs_real_t *__restrict__ vx, cs_real_t *__restrict__ rk, cs_real_t *__restrict__ qk)
 
static __global__ void _fcg_update_n (cs_lnum_t n, cs_real_t gk_rk1, cs_real_t ak_rk, const cs_real_t *__restrict__ wk, cs_real_t *__restrict__ dk, cs_real_t *__restrict__ qk, cs_real_t *__restrict__ rk, cs_real_t *__restrict__ vk, cs_real_t *__restrict__ vx)
 
static __global__ void _fcg_update_0 (cs_lnum_t n, cs_real_t ak_rk, const cs_real_t *__restrict__ wk, cs_real_t *__restrict__ dk, cs_real_t *__restrict__ qk, cs_real_t *__restrict__ rk, cs_real_t *__restrict__ vk, cs_real_t *__restrict__ vx)
 
template<size_t block_size>
static __global__ void _ymx_dot_yy (cs_lnum_t n, const cs_real_t *__restrict__ x, cs_real_t *__restrict__ y, double *__restrict__ sum_block)
 
template<size_t block_size>
static __global__ void _x0_ymr_dot_yy (cs_lnum_t n, const cs_real_t *__restrict__ r, cs_real_t *__restrict__ x, cs_real_t *__restrict__ y, double *__restrict__ sum_block)
 
template<size_t block_size>
static __global__ void _smaxpy_dot_yy (cs_lnum_t n, cs_real_t alpha, const cs_real_t *__restrict__ x, cs_real_t *__restrict__ y, double *__restrict__ sum_block)
 
template<size_t block_size>
static __global__ void _y_scale_dot_xy (cs_lnum_t n, cs_real_t alpha, const cs_real_t *__restrict__ x, cs_real_t *__restrict__ y, double *__restrict__ sum_block)
 
static __global__ void _gcr_gkj_inv (int n_c_iter, int n_gkj, const cs_real_t *__restrict__ gkj, cs_real_t *__restrict__ gkj_inv)
 
static __global__ void _gcr_gkj_inv_1 (const cs_real_t *__restrict__ gkj, cs_real_t *__restrict__ gkj_inv)
 
static __global__ void _gcr_update_vx (cs_lnum_t n_rows, int n_c_iter, size_t wa_size, const cs_real_t *__restrict__ alpha, const cs_real_t *__restrict__ gkj_inv, const cs_real_t *__restrict__ zk, cs_real_t *__restrict__ vx)
 
static void _prefetch_h2d (const void *dst, size_t size, int device_id, hipStream_t stream)
 Copy data from host to device. More...
 
static void _sync_reduction_sum (const cs_sles_it_t *c, hipStream_t stream, cs_lnum_t tuple_size, double res[])
 
static void _sync_reduction_sum (const cs_sles_it_t *c, hipStream_t stream, cs_lnum_t tuple_size, double res_d[], double res_h[])
 
static void _sync_reduction_sum_h (const cs_sles_it_t *c, hipStream_t stream, cs_lnum_t tuple_size, double res[])
 
static double _dot_product (const cs_sles_it_t *c, hipStream_t stream, const cs_real_t *x, const cs_real_t *y)
 
template<size_t blockSize, typename T >
static __global__ void _dot_product_xx_stage_1_of_2 (cs_lnum_t n, const T *x, double *b_res)
 Compute dot product x.x, summing result over all threads of a block. More...
 
template<size_t blockSize, typename T >
static __global__ void _dot_product_xy_stage_1_of_2 (cs_lnum_t n, const T *x, const T *y, double *b_res)
 Compute dot product x.y, summing result over all threads of a block. More...
 
template<size_t blockSize, typename T >
static __global__ void _dot_products_xx_xy_stage_1_of_2 (cs_lnum_t n, const T *x, const T *y, double *b_res)
 Compute dot products x.x, x.y, summing result over all threads of a block. More...
 
template<size_t blockSize, typename T >
static __global__ void _dot_products_xy_yz_stage_1_of_2 (cs_lnum_t n, const T *x, const T *y, const T *z, double *b_res)
 Compute dot products x.y, y.z, summing result over all threads of a block. More...
 
template<size_t blockSize, typename T >
static __global__ void _dot_products_xx_xy_yz_stage_1_of_2 (cs_lnum_t n, const T *x, const T *y, const T *z, double *b_res)
 Compute dot products x.y, x.y, y.z, summing result over all threads of a block. More...
 
template<size_t blockSize, typename T >
static __global__ void _dot_products_xx_yy_xy_xz_yz_stage_1_of_2 (cs_lnum_t n, const T *x, const T *y, const T *z, double *b_res)
 Compute dot products x.y, y.y, x.y, x.z, y.z, summing result over all threads of a block. More...
 
template<size_t blockSize, typename T >
static __global__ void _dot_products_vr_vw_vq_rr_stage_1_of_2 (cs_lnum_t n, const T *v, const T *w, const T *q, const T *r, double *b_res)
 Compute dot product x.y, summing result over all threads of a block. More...
 
static void _dot_products_vr_vw_vq_rr (const cs_sles_it_t *c, hipStream_t stream, const cs_real_t *v, const cs_real_t *r, const cs_real_t *w, const cs_real_t *q, double *s1, double *s2, double *s3, double *s4)
 
double cs_sles_it_dot_product (const cs_sles_it_t *c, hipStream_t stream, const cs_real_t *x, const cs_real_t *y)
 
double cs_sles_it_dot_product_xx (const cs_sles_it_t *c, hipStream_t stream, const cs_real_t *x)
 
void cs_sles_it_dot_products_xx_xy (const cs_sles_it_t *c, hipStream_t stream, const cs_real_t *x, const cs_real_t *y, double *xx, double *xy)
 
void cs_sles_it_dot_products_xy_yz (const cs_sles_it_t *c, hipStream_t stream, const cs_real_t *x, const cs_real_t *y, const cs_real_t *z, double *xy, double *yz)
 
void cs_sles_it_dot_products_xx_xy_yz (const cs_sles_it_t *c, hipStream_t stream, const cs_real_t *x, const cs_real_t *y, const cs_real_t *z, double *xx, double *xy, double *yz)
 
void cs_sles_it_dot_products_xx_yy_xy_xz_yz (const cs_sles_it_t *c, hipStream_t stream, const cs_real_t *x, const cs_real_t *y, const cs_real_t *z, double *xx, double *yy, double *xy, double *xz, double *yz)
 
cs_sles_convergence_state_t cs_sles_it_hip_jacobi (cs_sles_it_t *c, const cs_matrix_t *a, cs_lnum_t diag_block_size, cs_sles_it_convergence_t *convergence, const cs_real_t *rhs, cs_real_t *vx_ini, cs_real_t *vx, size_t aux_size, void *aux_vectors)
 
cs_sles_convergence_state_t cs_sles_it_hip_block_jacobi (cs_sles_it_t *c, const cs_matrix_t *a, cs_lnum_t diag_block_size, cs_sles_it_convergence_t *convergence, const cs_real_t *rhs, cs_real_t *__restrict__ vx_ini, cs_real_t *__restrict__ vx, size_t aux_size, void *aux_vectors)
 
cs_sles_convergence_state_t cs_sles_it_hip_fcg (cs_sles_it_t *c, const cs_matrix_t *a, cs_lnum_t diag_block_size, cs_sles_it_convergence_t *convergence, const cs_real_t *rhs, cs_real_t *restrict vx_ini, cs_real_t *restrict vx, size_t aux_size, void *aux_vectors)
 
cs_sles_convergence_state_t cs_sles_it_hip_gcr (cs_sles_it_t *c, const cs_matrix_t *a, cs_lnum_t diag_block_size, cs_sles_it_convergence_t *convergence, const cs_real_t *rhs, cs_real_t *vx_ini, cs_real_t *vx, size_t aux_size, void *aux_vectors)
 

Variables

static bool _use_hipblas = false
 

Macro Definition Documentation

◆ CS_BLOCKSIZE

#define CS_BLOCKSIZE   256

◆ CS_SIMD_SIZE

#define CS_SIMD_SIZE (   s)    (((s-1)/CS_HIP_WARP_SIZE+1)*CS_HIP_WARP_SIZE)

◆ HAVE_GRAPH_CAPTURE

#define HAVE_GRAPH_CAPTURE   0

Function Documentation

◆ _block_3_jacobi_compute_residual()

static __global__ void _block_3_jacobi_compute_residual ( cs_lnum_t  n_rows,
const cs_real_t *__restrict__  ad,
const cs_real_t *__restrict__  vx,
const cs_real_t *__restrict__  rk,
double *__restrict__  sum_block 
)
static

◆ _block_3_jacobi_compute_vx()

static __global__ void _block_3_jacobi_compute_vx ( cs_lnum_t  n_rows,
const cs_real_t *__restrict__  ad_inv,
const cs_real_t *__restrict__  rhs,
const cs_real_t *__restrict__  vxx,
cs_real_t *__restrict__  vx 
)
static

◆ _block_3_jacobi_compute_vx0()

static __global__ void _block_3_jacobi_compute_vx0 ( cs_lnum_t  n_rows,
const cs_real_t *__restrict__  ad_inv,
const cs_real_t *__restrict__  rhs,
cs_real_t *__restrict__  rk,
cs_real_t *__restrict__  vx 
)
static

◆ _block_jacobi_compute_residual()

static __global__ void _block_jacobi_compute_residual ( cs_lnum_t  n_rows,
cs_lnum_t  db_size,
const cs_real_t *__restrict__  ad,
cs_real_t *__restrict__  vx,
cs_real_t *__restrict__  rk,
double *__restrict__  sum_block 
)
static

◆ _block_jacobi_compute_vx()

static __global__ void _block_jacobi_compute_vx ( cs_lnum_t  n_rows,
cs_lnum_t  db_size,
const cs_real_t *__restrict__  ad_inv,
const cs_real_t *__restrict__  rhs,
const cs_real_t *__restrict__  vxx,
cs_real_t *__restrict__  vx 
)
static

◆ _block_jacobi_compute_vx0()

static __global__ void _block_jacobi_compute_vx0 ( cs_lnum_t  n_rows,
cs_lnum_t  db_size,
const cs_real_t *__restrict__  ad_inv,
const cs_real_t *__restrict__  rhs,
cs_real_t *__restrict__  rk,
cs_real_t *__restrict__  vx 
)
static

◆ _dot_product()

static double _dot_product ( const cs_sles_it_t *  c,
hipStream_t  stream,
const cs_real_t *  x,
const cs_real_t *  y 
)
inlinestatic

◆ _dot_product_xx_stage_1_of_2()

static __global__ void _dot_product_xx_stage_1_of_2 ( cs_lnum_t  n,
const T *  x,
double *  b_res 
)
static

Compute dot product x.x, summing result over all threads of a block.

blockSize must be a power of 2.

Parameters
[in]narray size
[in]xx vector
[out]b_resresult of s = x.x

◆ _dot_product_xy_stage_1_of_2()

static __global__ void _dot_product_xy_stage_1_of_2 ( cs_lnum_t  n,
const T *  x,
const T *  y,
double *  b_res 
)
static

Compute dot product x.y, summing result over all threads of a block.

blockSize must be a power of 2.

Parameters
[in]narray size
[in]xx vector
[in]yy vector
[out]b_resresult of s = x.x

◆ _dot_products_vr_vw_vq_rr()

static void _dot_products_vr_vw_vq_rr ( const cs_sles_it_t *  c,
hipStream_t  stream,
const cs_real_t *  v,
const cs_real_t *  r,
const cs_real_t *  w,
const cs_real_t *  q,
double *  s1,
double *  s2,
double *  s3,
double *  s4 
)
static

◆ _dot_products_vr_vw_vq_rr_stage_1_of_2()

static __global__ void _dot_products_vr_vw_vq_rr_stage_1_of_2 ( cs_lnum_t  n,
const T *  v,
const T *  w,
const T *  q,
const T *  r,
double *  b_res 
)
static

Compute dot product x.y, summing result over all threads of a block.

blockSize must be a power of 2.

Parameters
[in]narray size
[in]vv vector
[in]ww vector
[in]qq vector
[in]rr vector
[out]b_resresult of s = x.x

◆ _dot_products_xx_xy_stage_1_of_2()

static __global__ void _dot_products_xx_xy_stage_1_of_2 ( cs_lnum_t  n,
const T *  x,
const T *  y,
double *  b_res 
)
static

Compute dot products x.x, x.y, summing result over all threads of a block.

blockSize must be a power of 2.

Parameters
[in]narray size
[in]xx vector
[in]yy vector
[out]b_resresult of s = x.x

◆ _dot_products_xx_xy_yz_stage_1_of_2()

static __global__ void _dot_products_xx_xy_yz_stage_1_of_2 ( cs_lnum_t  n,
const T *  x,
const T *  y,
const T *  z,
double *  b_res 
)
static

Compute dot products x.y, x.y, y.z, summing result over all threads of a block.

blockSize must be a power of 2.

Parameters
[in]narray size
[in]xx vector
[in]yy vector
[in]zz vector
[out]b_resresult of s = x.x

◆ _dot_products_xx_yy_xy_xz_yz_stage_1_of_2()

static __global__ void _dot_products_xx_yy_xy_xz_yz_stage_1_of_2 ( cs_lnum_t  n,
const T *  x,
const T *  y,
const T *  z,
double *  b_res 
)
static

Compute dot products x.y, y.y, x.y, x.z, y.z, summing result over all threads of a block.

blockSize must be a power of 2.

Parameters
[in]narray size
[in]xx vector
[in]yy vector
[in]zz vector
[out]b_resresult of s = x.x

◆ _dot_products_xy_yz_stage_1_of_2()

static __global__ void _dot_products_xy_yz_stage_1_of_2 ( cs_lnum_t  n,
const T *  x,
const T *  y,
const T *  z,
double *  b_res 
)
static

Compute dot products x.y, y.z, summing result over all threads of a block.

blockSize must be a power of 2.

Parameters
[in]narray size
[in]xx vector
[in]yy vector
[in]zz vector
[out]b_resresult of s = x.x

◆ _fcg_init()

static __global__ void _fcg_init ( cs_lnum_t  n,
const cs_real_t *__restrict__  rhs,
cs_real_t *__restrict__  rk,
cs_real_t *__restrict__  qk 
)
static

◆ _fcg_init_vx0()

static __global__ void _fcg_init_vx0 ( cs_lnum_t  n,
const cs_real_t *__restrict__  rhs,
cs_real_t *__restrict__  vx,
cs_real_t *__restrict__  rk,
cs_real_t *__restrict__  qk 
)
static

◆ _fcg_update_0()

static __global__ void _fcg_update_0 ( cs_lnum_t  n,
cs_real_t  ak_rk,
const cs_real_t *__restrict__  wk,
cs_real_t *__restrict__  dk,
cs_real_t *__restrict__  qk,
cs_real_t *__restrict__  rk,
cs_real_t *__restrict__  vk,
cs_real_t *__restrict__  vx 
)
static

◆ _fcg_update_n()

static __global__ void _fcg_update_n ( cs_lnum_t  n,
cs_real_t  gk_rk1,
cs_real_t  ak_rk,
const cs_real_t *__restrict__  wk,
cs_real_t *__restrict__  dk,
cs_real_t *__restrict__  qk,
cs_real_t *__restrict__  rk,
cs_real_t *__restrict__  vk,
cs_real_t *__restrict__  vx 
)
static

◆ _gcr_gkj_inv()

static __global__ void _gcr_gkj_inv ( int  n_c_iter,
int  n_gkj,
const cs_real_t *__restrict__  gkj,
cs_real_t *__restrict__  gkj_inv 
)
static

◆ _gcr_gkj_inv_1()

static __global__ void _gcr_gkj_inv_1 ( const cs_real_t *__restrict__  gkj,
cs_real_t *__restrict__  gkj_inv 
)
static

◆ _gcr_update_vx()

static __global__ void _gcr_update_vx ( cs_lnum_t  n_rows,
int  n_c_iter,
size_t  wa_size,
const cs_real_t *__restrict__  alpha,
const cs_real_t *__restrict__  gkj_inv,
const cs_real_t *__restrict__  zk,
cs_real_t *__restrict__  vx 
)
static

◆ _jacobi_compute_vx()

static __global__ void _jacobi_compute_vx ( cs_lnum_t  n_rows,
const cs_real_t *__restrict__  ad_inv,
const cs_real_t *__restrict__  rhs,
cs_real_t *__restrict__  vx,
cs_real_t *__restrict__  rk 
)
static

◆ _jacobi_compute_vx_and_residual()

static __global__ void _jacobi_compute_vx_and_residual ( cs_lnum_t  n_rows,
const cs_real_t *__restrict__  ad_inv,
const cs_real_t *__restrict__  ad,
const cs_real_t *__restrict__  rhs,
cs_real_t *__restrict__  vx,
cs_real_t *__restrict__  rk,
double *__restrict__  sum_block 
)
static

◆ _jacobi_compute_vx_and_residual_ini0()

static __global__ void _jacobi_compute_vx_and_residual_ini0 ( cs_lnum_t  n_rows,
const cs_real_t *__restrict__  ad_inv,
const cs_real_t *__restrict__  rhs,
cs_real_t *__restrict__  vx,
cs_real_t *__restrict__  rk,
double *__restrict__  sum_block 
)
static

◆ _jacobi_compute_vx_ini0()

static __global__ void _jacobi_compute_vx_ini0 ( cs_lnum_t  n_rows,
const cs_real_t *__restrict__  ad_inv,
const cs_real_t *__restrict__  rhs,
cs_real_t *__restrict__  vx,
cs_real_t *__restrict__  rk 
)
static

◆ _prefetch_h2d()

static void _prefetch_h2d ( const void *  dst,
size_t  size,
int  device_id,
hipStream_t  stream 
)
static

Copy data from host to device.

This is simply a wrapper over hipMemcpy.

A safety check is added.

Parameters
[out]dstpointer to data
[in]sizesize of data to copy
[in]device_idHIP device id
[in]streamsize of data to copy
Returns
pointer to allocated memory.

◆ _smaxpy_dot_yy()

static __global__ void _smaxpy_dot_yy ( cs_lnum_t  n,
cs_real_t  alpha,
const cs_real_t *__restrict__  x,
cs_real_t *__restrict__  y,
double *__restrict__  sum_block 
)
static

◆ _sync_reduction_sum() [1/2]

static void _sync_reduction_sum ( const cs_sles_it_t *  c,
hipStream_t  stream,
cs_lnum_t  tuple_size,
double  res[] 
)
static

◆ _sync_reduction_sum() [2/2]

static void _sync_reduction_sum ( const cs_sles_it_t *  c,
hipStream_t  stream,
cs_lnum_t  tuple_size,
double  res_d[],
double  res_h[] 
)
static

◆ _sync_reduction_sum_h()

static void _sync_reduction_sum_h ( const cs_sles_it_t *  c,
hipStream_t  stream,
cs_lnum_t  tuple_size,
double  res[] 
)
static

◆ _warp_reduce_sum()

static __device__ void __forceinline__ _warp_reduce_sum ( volatile T *  stmp,
size_t  tid 
)
static

Kernel for sum reduction within a warp (for warp size 32).

Parameters
[in,out]stmpshared value to reduce
[in,out]tidthread id

◆ _x0_ymr_dot_yy()

static __global__ void _x0_ymr_dot_yy ( cs_lnum_t  n,
const cs_real_t *__restrict__  r,
cs_real_t *__restrict__  x,
cs_real_t *__restrict__  y,
double *__restrict__  sum_block 
)
static

◆ _y_scale_dot_xy()

static __global__ void _y_scale_dot_xy ( cs_lnum_t  n,
cs_real_t  alpha,
const cs_real_t *__restrict__  x,
cs_real_t *__restrict__  y,
double *__restrict__  sum_block 
)
static

◆ _ymx_dot_yy()

static __global__ void _ymx_dot_yy ( cs_lnum_t  n,
const cs_real_t *__restrict__  x,
cs_real_t *__restrict__  y,
double *__restrict__  sum_block 
)
static

◆ cs_sles_it_dot_product()

double cs_sles_it_dot_product ( const cs_sles_it_t *  c,
hipStream_t  stream,
const cs_real_t *  x,
const cs_real_t *  y 
)

◆ cs_sles_it_dot_product_xx()

double cs_sles_it_dot_product_xx ( const cs_sles_it_t *  c,
hipStream_t  stream,
const cs_real_t *  x 
)

◆ cs_sles_it_dot_products_xx_xy()

void cs_sles_it_dot_products_xx_xy ( const cs_sles_it_t *  c,
hipStream_t  stream,
const cs_real_t *  x,
const cs_real_t *  y,
double *  xx,
double *  xy 
)

◆ cs_sles_it_dot_products_xx_xy_yz()

void cs_sles_it_dot_products_xx_xy_yz ( const cs_sles_it_t *  c,
hipStream_t  stream,
const cs_real_t *  x,
const cs_real_t *  y,
const cs_real_t *  z,
double *  xx,
double *  xy,
double *  yz 
)

◆ cs_sles_it_dot_products_xx_yy_xy_xz_yz()

void cs_sles_it_dot_products_xx_yy_xy_xz_yz ( const cs_sles_it_t *  c,
hipStream_t  stream,
const cs_real_t *  x,
const cs_real_t *  y,
const cs_real_t *  z,
double *  xx,
double *  yy,
double *  xy,
double *  xz,
double *  yz 
)

◆ cs_sles_it_dot_products_xy_yz()

void cs_sles_it_dot_products_xy_yz ( const cs_sles_it_t *  c,
hipStream_t  stream,
const cs_real_t *  x,
const cs_real_t *  y,
const cs_real_t *  z,
double *  xy,
double *  yz 
)

◆ cs_sles_it_hip_block_jacobi()

cs_sles_convergence_state_t cs_sles_it_hip_block_jacobi ( cs_sles_it_t *  c,
const cs_matrix_t *  a,
cs_lnum_t  diag_block_size,
cs_sles_it_convergence_t *  convergence,
const cs_real_t *  rhs,
cs_real_t *__restrict__  vx_ini,
cs_real_t *__restrict__  vx,
size_t  aux_size,
void *  aux_vectors 
)

◆ cs_sles_it_hip_fcg()

cs_sles_convergence_state_t cs_sles_it_hip_fcg ( cs_sles_it_t *  c,
const cs_matrix_t *  a,
cs_lnum_t  diag_block_size,
cs_sles_it_convergence_t *  convergence,
const cs_real_t *  rhs,
cs_real_t *restrict  vx_ini,
cs_real_t *restrict  vx,
size_t  aux_size,
void *  aux_vectors 
)

◆ cs_sles_it_hip_gcr()

cs_sles_convergence_state_t cs_sles_it_hip_gcr ( cs_sles_it_t *  c,
const cs_matrix_t *  a,
cs_lnum_t  diag_block_size,
cs_sles_it_convergence_t *  convergence,
const cs_real_t *  rhs,
cs_real_t *  vx_ini,
cs_real_t *  vx,
size_t  aux_size,
void *  aux_vectors 
)

◆ cs_sles_it_hip_jacobi()

cs_sles_convergence_state_t cs_sles_it_hip_jacobi ( cs_sles_it_t *  c,
const cs_matrix_t *  a,
cs_lnum_t  diag_block_size,
cs_sles_it_convergence_t *  convergence,
const cs_real_t *  rhs,
cs_real_t *  vx_ini,
cs_real_t *  vx,
size_t  aux_size,
void *  aux_vectors 
)

Variable Documentation

◆ _use_hipblas

bool _use_hipblas = false
static