36 #include <OpenCL/cl.h>
49 #include "cfl_kernel.cl.h"
58 void *dr_inv,
void *ds_inv,
void *dt_inv,
59 void *
jacinv,
int *nel,
int *lx) {
66 const size_t global_item_size = 256 * (*nel);
67 const size_t local_item_size = 256;
70 cl_mem
cfl_d = clCreateBuffer(
glb_ctx, CL_MEM_READ_WRITE,
71 (*nel) *
sizeof(
real), NULL, &err);
78 cl_kernel kernel = clCreateKernel(cfl_program, \
79 STR(cfl_kernel_lx##LX), &err); \
82 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(real), dt)); \
83 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &u)); \
84 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &v)); \
85 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &w)); \
86 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &drdx)); \
87 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &dsdx)); \
88 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &dtdx)); \
89 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &drdy)); \
90 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &dsdy)); \
91 CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_mem), (void *) &dtdy)); \
92 CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_mem), (void *) &drdz)); \
93 CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_mem), (void *) &dsdz)); \
94 CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_mem), (void *) &dtdz)); \
95 CL_CHECK(clSetKernelArg(kernel, 13, sizeof(cl_mem), (void *) &dr_inv)); \
96 CL_CHECK(clSetKernelArg(kernel, 14, sizeof(cl_mem), (void *) &ds_inv)); \
97 CL_CHECK(clSetKernelArg(kernel, 15, sizeof(cl_mem), (void *) &dt_inv)); \
98 CL_CHECK(clSetKernelArg(kernel, 16, sizeof(cl_mem), (void *) &jacinv)); \
99 CL_CHECK(clSetKernelArg(kernel, 17, sizeof(cl_mem), (void *) &cfl_d)); \
101 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, \
102 kernel, 1, NULL, &global_item_size, \
103 &local_item_size, 0, NULL, &kern_wait)); \
122 cfl_d, CL_TRUE, 0, (*nel) *
sizeof(
real),
123 cfl, 1, &kern_wait, NULL));
126 for (
i = 0;
i < (*nel);
i++) {
127 cfl_max = fmax(cfl_max,
cfl[
i]);
__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__ 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__ 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__ 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__ dt
__global__ void cfl_kernel(const T dt, const T *__restrict__ u, const T *__restrict__ v, const T *__restrict__ w, const T *__restrict__ drdx, const T *__restrict__ dsdx, const T *__restrict__ dtdx, const T *__restrict__ drdy, const T *__restrict__ dsdy, const T *__restrict__ dtdy, const T *__restrict__ drdz, const T *__restrict__ dsdz, const T *__restrict__ dtdz, const T *__restrict__ dr_inv, const T *__restrict__ ds_inv, const T *__restrict__ dt_inv, const T *__restrict__ jacinv, T *__restrict__ cfl_h)
void opencl_kernel_jit(const char *kernel, cl_program *program)
real(kind=rp) function, public cfl(dt, u, v, w, Xh, coef, nelv, gdim)
real opencl_cfl(real *dt, void *u, void *v, void *w, void *drdx, void *dsdx, void *dtdx, void *drdy, void *dsdy, void *dtdy, void *drdz, void *dsdz, void *dtdz, void *dr_inv, void *ds_inv, void *dt_inv, void *jacinv, int *nel, int *lx)