35#include <hip/hip_runtime.h>
53 void *
u,
int *n,
void *
gd,
int *
nb,
54 void *b,
void *
bo,
int *op,
57 if ((*m) == 0)
return;
60 const dim3 nblcks(((*m)+ 1024 - 1)/ 1024, 1, 1);
68 *
nb, (
int *) b, (
int *)
bo);
76 *
nb, (
int *) b, (
int *)
bo);
84 *
nb, (
int *) b, (
int *)
bo);
92 *
nb, (
int *) b, (
int *)
bo);
102 void *
u,
int *n,
void *
gd,
103 int *
nb,
void *b,
void *
bo,
106 if ((*m) == 0)
return;
109 const dim3 nblcks(((*m)+1024 - 1)/ 1024, 1, 1);
115 *
nb, (
int *) b, (
int *)
bo);
128 if (stream ==
NULL) {
131 (
real *) u_d, (
real *) buf_d + offset,
132 (
int *) dof_d + offset, n);
137 (
real *) u_d, (
real *) buf_d + offset,
138 (
int *) dof_d + offset, n);
154 if (stream ==
NULL) {
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ u
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ v
__global__ void dirichlet_apply_scalar_kernel(const int *__restrict__ msk, T *__restrict__ x, const T g, const int m)
void hip_gs_unpack(real *u_d, int op, real *buf_d, int *dof_d, int offset, int n, hipStream_t stream)
void hip_scatter_kernel(void *v, int *m, void *dg, void *u, int *n, void *gd, int *nb, void *b, void *bo, hipStream_t stream)
void hip_gs_pack(void *u_d, void *buf_d, void *dof_d, int offset, int n, hipStream_t stream)
void hip_gather_kernel(void *v, int *m, int *o, void *dg, void *u, int *n, void *gd, int *nb, void *b, void *bo, int *op, hipStream_t stream)