36 #include <hip/hip_runtime.h>
47 void *G22,
void *G23,
void *G33,
54 const dim3 nthrds(1024, 1, 1);
55 const dim3 nblcks((*nel), 1, 1);
57 #define GEO_CASE(LX) \
60 HIP_KERNEL_NAME(coef_generate_geo_kernel<real, LX, 1024>), \
61 nblcks, nthrds, 0, (hipStream_t) glb_cmd_queue, \
62 (real *) G11, (real *) G12, (real *) G13, \
63 (real *) G22, (real *) G23, (real *) G33, \
64 (real *) drdx, (real *) drdy, (real *) drdz, \
65 (real *) dsdx, (real *) dsdy, (real *) dsdz, \
66 (real *) dtdx, (real *) dtdy, (real *) dtdz, \
67 (real *) jacinv, (real *) w3, *gdim); \
68 HIP_CHECK(hipGetLastError()); \
89 fprintf(stderr, __FILE__
": size not supported: %d\n", *lx);
101 void *dxdr,
void *dydr,
void *dzdr,
102 void *dxds,
void *dyds,
void *dzds,
103 void *dxdt,
void *dydt,
void *dzdt,
104 void *
dx,
void *
dy,
void *
dz,
105 void *
x,
void *y,
void *z,
109 const int n = (*nel) * (*lx) * (*lx) * (*lx);
110 const dim3 nthrds(1024, 1, 1);
111 const dim3 nblcks_dxyz((*nel), 1, 1);
112 const dim3 nblcks_drst((n + 1024 - 1)/ 1024, 1, 1);
114 #define DXYZDRST_CASE(LX) \
116 hipLaunchKernelGGL( \
117 HIP_KERNEL_NAME(coef_generate_dxyz_kernel<real, LX, 1024>), \
118 nblcks_dxyz, nthrds, 0, (hipStream_t) glb_cmd_queue, \
119 (real *) dxdr, (real *) dydr, (real *) dzdr, \
120 (real *) dxds, (real *) dyds, (real *) dzds, \
121 (real *) dxdt, (real *) dydt, (real *) dzdt, \
122 (real *) dx, (real *) dy, (real *) dz, \
123 (real *) x, (real *) y, (real *) z); \
124 HIP_CHECK(hipGetLastError()); \
145 fprintf(stderr, __FILE__
": size not supported: %d\n", *lx);
150 hipLaunchKernelGGL(HIP_KERNEL_NAME(coef_generate_drst_kernel<real>),
void hip_coef_generate_geo(void *G11, void *G12, void *G13, void *G22, void *G23, void *G33, void *drdx, void *drdy, void *drdz, void *dsdx, void *dsdy, void *dsdz, void *dtdx, void *dtdy, void *dtdz, void *jacinv, void *w3, int *nel, int *lx, int *gdim)
void hip_coef_generate_dxyzdrst(void *drdx, void *drdy, void *drdz, void *dsdx, void *dsdy, void *dsdz, void *dtdx, void *dtdy, void *dtdz, void *dxdr, void *dydr, void *dzdr, void *dxds, void *dyds, void *dzds, void *dxdt, void *dydt, void *dzdt, void *dx, void *dy, void *dz, void *x, void *y, void *z, void *jacinv, void *jac, int *lx, int *nel)
#define DXYZDRST_CASE(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__ 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__ 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__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dx
__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__ 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__ 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__ x
__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