Neko 1.99.1
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
Macros | Functions | Variables
opr_opgrad.c File Reference
#include <CL/cl.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <device/device_config.h>
#include <device/opencl/jit.h>
#include <device/opencl/prgm_lib.h>
#include <device/opencl/check.h>
#include <common/neko_log.h>
#include "opgrad_kernel.cl.h"
Include dependency graph for opr_opgrad.c:

Go to the source code of this file.

Macros

#define STR(X)   #X
 
#define CASE_1D(LX, QUEUE, EVENT)
 
#define CASE_KSTEP(LX, QUEUE, EVENT)
 
#define CASE(LX)
 

Functions

void opencl_opgrad (void *ux, void *uy, void *uz, void *u, void *dx, void *dy, void *dz, void *drdx, void *dsdx, void *dtdx, void *drdy, void *dsdy, void *dtdy, void *drdz, void *dsdz, void *dtdz, void *w3, int *nel, int *lx)
 

Variables

int * autotune_opgrad = NULL
 

Macro Definition Documentation

◆ CASE

#define CASE (   LX)

◆ CASE_1D

#define CASE_1D (   LX,
  QUEUE,
  EVENT 
)
Value:
{ \
cl_kernel kernel = clCreateKernel(opgrad_program, \
STR(opgrad_kernel_lx##LX), &err); \
CL_CHECK(err); \
\
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &ux)); \
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &uy)); \
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &uz)); \
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &u)); \
CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &dx)); \
CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &dy)); \
CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &dz)); \
CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &drdx)); \
CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &dsdx)); \
CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_mem), (void *) &dtdx)); \
CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_mem), (void *) &drdy)); \
CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_mem), (void *) &dsdy)); \
CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_mem), (void *) &dtdy)); \
CL_CHECK(clSetKernelArg(kernel, 13, sizeof(cl_mem), (void *) &drdz)); \
CL_CHECK(clSetKernelArg(kernel, 14, sizeof(cl_mem), (void *) &dsdz)); \
CL_CHECK(clSetKernelArg(kernel, 15, sizeof(cl_mem), (void *) &dtdz)); \
CL_CHECK(clSetKernelArg(kernel, 16, sizeof(cl_mem), (void *) &w3)); \
\
CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) QUEUE, \
kernel, 1, NULL, &global_item_size, \
&local_item_size,0, NULL, EVENT)); \
}
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__ drdy
Definition ax_helm_full_kernel.h:54
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__ drdz
Definition ax_helm_full_kernel.h:55
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__ const T *__restrict__ dsdz
Definition ax_helm_full_kernel.h:58
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__ dsdy
Definition ax_helm_full_kernel.h:57
dtdy
__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
Definition ax_helm_full_kernel.h:60
u
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ u
Definition ax_helm_full_kernel.h:46
dx
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dx
Definition ax_helm_full_kernel.h:49
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__ drdx
Definition ax_helm_full_kernel.h:53
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__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dtdz
Definition ax_helm_full_kernel.h:61
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__ dsdx
Definition ax_helm_full_kernel.h:56
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__ dtdx
Definition ax_helm_full_kernel.h:59
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__ dz
Definition ax_helm_full_kernel.h:51
dy
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dy
Definition ax_helm_full_kernel.h:50
w3
__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
Definition cdtp_kernel.h:113
dirichlet_apply_scalar_kernel
__global__ void dirichlet_apply_scalar_kernel(const int *__restrict__ msk, T *__restrict__ x, const T g, const int m)
Definition dirichlet_kernel.h:42
uy
__global__ void T *__restrict__ uy
Definition opgrad_kernel.h:126
uz
__global__ void T *__restrict__ T *__restrict__ uz
Definition opgrad_kernel.h:127
STR
#define STR(X)
opgrad_program
void * opgrad_program

◆ CASE_KSTEP

#define CASE_KSTEP (   LX,
  QUEUE,
  EVENT 
)
Value:
{ \
cl_kernel kernel = clCreateKernel(opgrad_program, \
STR(opgrad_kernel_kstep_lx##LX), &err); \
CL_CHECK(err); \
\
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &ux)); \
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &uy)); \
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &uz)); \
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &u)); \
CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &dx)); \
CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &dy)); \
CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &dz)); \
CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &drdx)); \
CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &dsdx)); \
CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_mem), (void *) &dtdx)); \
CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_mem), (void *) &drdy)); \
CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_mem), (void *) &dsdy)); \
CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_mem), (void *) &dtdy)); \
CL_CHECK(clSetKernelArg(kernel, 13, sizeof(cl_mem), (void *) &drdz)); \
CL_CHECK(clSetKernelArg(kernel, 14, sizeof(cl_mem), (void *) &dsdz)); \
CL_CHECK(clSetKernelArg(kernel, 15, sizeof(cl_mem), (void *) &dtdz)); \
CL_CHECK(clSetKernelArg(kernel, 16, sizeof(cl_mem), (void *) &w3)); \
\
CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) QUEUE, \
kernel, 2, NULL, global_kstep, \
local_kstep, 0, NULL, EVENT)); \
}

◆ STR

#define STR (   X)    #X

Function Documentation

◆ opencl_opgrad()

void opencl_opgrad ( void *  ux,
void *  uy,
void *  uz,
void *  u,
void *  dx,
void *  dy,
void *  dz,
void *  drdx,
void *  dsdx,
void *  dtdx,
void *  drdy,
void *  dsdy,
void *  dtdy,
void *  drdz,
void *  dsdz,
void *  dtdz,
void *  w3,
int *  nel,
int *  lx 
)

Fortran wrapper for device OpenCL gradients

Definition at line 57 of file opr_opgrad.c.

Here is the call graph for this function:
Here is the caller graph for this function:

Variable Documentation

◆ autotune_opgrad

int* autotune_opgrad = NULL

Definition at line 52 of file opr_opgrad.c.