38 #include <hip/hip_runtime.h>
53 void *
u,
void *
v,
void *
w,
54 void *
dx,
void *
dy,
void *
dz,
60 void *
jacinv,
void *
w3,
int *nelv,
int *lx) {
62 const dim3 nthrds((*lx), (*lx), 1);
63 const dim3 nblcks((*nelv), 1, 1);
65 #define CASE_VECTOR_KSTEP(LX) \
66 hipLaunchKernelGGL( HIP_KERNEL_NAME( ax_helm_stress_kernel_vector_kstep<real, LX> ), \
68 (hipStream_t) glb_cmd_queue, \
69 (real *) au, (real *) av, (real *) aw, \
70 (real *) u, (real *) v, (real *) w, \
71 (real *) dx, (real *) dy, (real *) dz, (real *) h1, \
72 (real *) drdx, (real *) drdy, (real *) drdz, \
73 (real *) dsdx, (real *) dsdy, (real *) dsdz, \
74 (real *) dtdx, (real *) dtdy, (real *) dtdz, \
75 (real *) jacinv, (real *) w3); \
76 HIP_CHECK(hipGetLastError());
78 #define CASE_VECTOR_KSTEP_PADDED(LX) \
79 hipLaunchKernelGGL( HIP_KERNEL_NAME( ax_helm_stress_kernel_vector_kstep_padded<real, LX> ), \
81 (hipStream_t) glb_cmd_queue, \
82 (real *) au, (real *) av, (real *) aw, \
83 (real *) u, (real *) v, (real *) w, \
84 (real *) dx, (real *) dy, (real *) dz, (real *) h1, \
85 (real *) drdx, (real *) drdy, (real *) drdz, \
86 (real *) dsdx, (real *) dsdy, (real *) dsdz, \
87 (real *) dtdx, (real *) dtdy, (real *) dtdz, \
88 (real *) jacinv, (real *) w3); \
89 HIP_CHECK(hipGetLastError());
91 #define CASE_VECTOR(LX) \
93 CASE_VECTOR_KSTEP(LX); \
96 #define CASE_VECTOR_PADDED(LX) \
98 CASE_VECTOR_KSTEP_PADDED(LX); \
119 fprintf(stderr, __FILE__
": size not supported: %d\n", *lx);
129 void *
u,
void *
v,
void *
w,
130 void *h2,
void *B,
int *n) {
132 const dim3 nthrds(1024, 1, 1);
133 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
136 hipLaunchKernelGGL( HIP_KERNEL_NAME( ax_helm_stress_kernel_vector_part2<real> ),
137 nblcks, nthrds, 0, stream,
void hip_ax_helm_stress_vector(void *au, void *av, void *aw, void *u, void *v, void *w, void *dx, void *dy, void *dz, void *dxt, void *dyt, void *dzt, void *h1, void *drdx, void *drdy, void *drdz, void *dsdx, void *dsdy, void *dsdz, void *dtdx, void *dtdy, void *dtdz, void *jacinv, void *w3, int *nelv, int *lx)
void hip_ax_helm_stress_vector_part2(void *au, void *av, void *aw, void *u, void *v, void *w, void *h2, void *B, int *n)
#define CASE_VECTOR_PADDED(LX)
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ drdy
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ drdz
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dsdz
__global__ void T *__restrict__ T *__restrict__ aw
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dsdy
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ w
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dtdy
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ u
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dx
__global__ void T *__restrict__ av
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ drdx
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dtdz
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dsdx
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ v
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dtdx
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dz
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ h1
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dy
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ jacinv
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ w3
__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