38template<
typename T, const
int LX >
55 const int k =
jk /
LX;
63 for (
int l = 0; l <
LX; l++) {
69 for (
int l = 0; l <
LX; l++) {
75 for (
int l = 0; l <
LX; l++) {
82 if (
i == 0 ||
i ==
LX-1) {
87 if (
j == 0 ||
j ==
LX-1) {
92 if (
k == 0 ||
k ==
LX-1) {
105 int *nel,
int *
lxp) {
114 jacobi_kernel<real, N><<<blocks, threads, 0, stream>>>( \
116 (real*)dxt, (real*)dyt, (real*)dzt,\
117 (real*)G11, (real*)G22, (real*)G33,\
118 (real*)G12, (real*)G13, (real*)G23,\
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dzt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dyt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dxt
__global__ void dirichlet_apply_scalar_kernel(const int *__restrict__ msk, T *__restrict__ x, const T g, const int m)
__global__ void jacobi_kernel(T *__restrict__ du, const T *__restrict__ dxt, const T *__restrict__ dyt, const T *__restrict__ dzt, const T *__restrict__ G11, const T *__restrict__ G22, const T *__restrict__ G33, const T *__restrict__ G12, const T *__restrict__ G13, const T *__restrict__ G23, const int nel)
void cuda_jacobi_update(void *d, void *dxt, void *dyt, void *dzt, void *G11, void *G22, void *G33, void *G12, void *G13, void *G23, int *nel, int *lxp)