Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
Macros | Functions | Variables
ax_helm.c File Reference
#include <CL/cl.h>
#include <stdlib.h>
#include <stdio.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 "ax_helm_kernel.cl.h"
Include dependency graph for ax_helm.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)
 
#define CASE_LARGE(LX)
 
#define STR(X)   #X
 
#define CASE_VECTOR(LX)
 

Functions

void opencl_ax_helm (void *w, void *u, void *dx, void *dy, void *dz, void *dxt, void *dyt, void *dzt, void *h1, void *g11, void *g22, void *g33, void *g12, void *g13, void *g23, int *nelv, int *lx)
 
void opencl_ax_helm_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 *g11, void *g22, void *g33, void *g12, void *g13, void *g23, int *nelv, int *lx)
 
void opencl_ax_helm_vector_part2 (void *au, void *av, void *aw, void *u, void *v, void *w, void *h2, void *B, int *n)
 

Variables

int * autotune_ax_helm = NULL
 

Macro Definition Documentation

◆ CASE

#define CASE (   LX)

◆ CASE_1D

#define CASE_1D (   LX,
  QUEUE,
  EVENT 
)
Value:
{ \
cl_kernel kernel = clCreateKernel(ax_helm_program, \
STR(ax_helm_kernel_lx##LX), &err); \
CL_CHECK(err); \
\
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &w)); \
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &u)); \
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &dx)); \
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &dy)); \
CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &dz)); \
CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &dxt)); \
CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &dyt)); \
CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &dzt)); \
CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &h1)); \
CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_mem), (void *) &g11)); \
CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_mem), (void *) &g22)); \
CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_mem), (void *) &g33)); \
CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_mem), (void *) &g12)); \
CL_CHECK(clSetKernelArg(kernel, 13, sizeof(cl_mem), (void *) &g13)); \
CL_CHECK(clSetKernelArg(kernel, 14, sizeof(cl_mem), (void *) &g23)); \
\
CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) QUEUE, \
kernel, 1, NULL, &global_item_size, \
&local_item_size, 0, NULL, EVENT)); \
CL_CHECK(clReleaseKernel(kernel)); \
}
STR
#define STR(X)
ale_add_kinematics_kernel
__global__ void ale_add_kinematics_kernel(const int n, T *__restrict__ wx, T *__restrict__ wy, T *__restrict__ wz, const T *__restrict__ x_ref, const T *__restrict__ y_ref, const T *__restrict__ z_ref, const T *__restrict__ phi, const T *__restrict__ x, const T *__restrict__ y, const T *__restrict__ z, const kinematics_params_t kin_params)
Definition ale_kinematics_kernel.h:52
w
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ w
Definition ax_helm_full_kernel.h:48
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
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
u
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ u
Definition ax_helm_full_kernel.h:46
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
h1
__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
Definition ax_helm_full_kernel.h:52
g23
__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__ const T *__restrict__ const T *__restrict__ const T *__restrict__ g23
Definition ax_helm_kernel.h:184
g22
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ g22
Definition ax_helm_kernel.h:180
g13
__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__ const T *__restrict__ const T *__restrict__ g13
Definition ax_helm_kernel.h:183
g12
__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__ const T *__restrict__ g12
Definition ax_helm_kernel.h:182
g33
__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__ g33
Definition ax_helm_kernel.h:181
g11
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ g11
Definition ax_helm_kernel.h:179
dyt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dyt
Definition cdtp_kernel.h:115
dzt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dzt
Definition cdtp_kernel.h:116
dxt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dxt
Definition cdtp_kernel.h:114
ax_helm_program
void * ax_helm_program

◆ CASE_KSTEP

#define CASE_KSTEP (   LX,
  QUEUE,
  EVENT 
)
Value:
{ \
cl_kernel kernel = clCreateKernel(ax_helm_program, \
STR(ax_helm_kernel_kstep_lx##LX), &err);\
CL_CHECK(err); \
\
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &w)); \
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &u)); \
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &dx)); \
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &dy)); \
CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &dz)); \
CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &h1)); \
CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &g11)); \
CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &g22)); \
CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &g33)); \
CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_mem), (void *) &g12)); \
CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_mem), (void *) &g13)); \
CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_mem), (void *) &g23)); \
\
CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) QUEUE, \
kernel, 2, NULL, global_kstep, \
local_kstep, 0, NULL, EVENT)); \
CL_CHECK(clReleaseKernel(kernel)); \
}

◆ CASE_LARGE

#define CASE_LARGE (   LX)
Value:
case LX: \
CASE_KSTEP(LX, glb_cmd_queue, NULL); \
break

◆ CASE_VECTOR

#define CASE_VECTOR (   LX)
Value:
case LX: \
{ \
cl_kernel kernel = \
clCreateKernel(ax_helm_program, \
STR(ax_helm_kernel_vector_kstep_lx##LX), &err); \
CL_CHECK(err); \
\
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &au)); \
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &av)); \
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &aw)); \
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &u)); \
CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &v)); \
CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &w)); \
CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &dx)); \
CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &dy)); \
CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &dz)); \
CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_mem), (void *) &h1)); \
CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_mem), (void *) &g11)); \
CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_mem), (void *) &g22)); \
CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_mem), (void *) &g33)); \
CL_CHECK(clSetKernelArg(kernel, 13, sizeof(cl_mem), (void *) &g12)); \
CL_CHECK(clSetKernelArg(kernel, 14, sizeof(cl_mem), (void *) &g13)); \
CL_CHECK(clSetKernelArg(kernel, 15, sizeof(cl_mem), (void *) &g23)); \
\
CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, \
kernel, 2, NULL, global_kstep, \
local_kstep, 0, NULL, NULL)); \
CL_CHECK(clReleaseKernel(kernel)); \
} \
break
aw
__global__ void T *__restrict__ T *__restrict__ aw
Definition ax_helm_full_kernel.h:45
av
__global__ void T *__restrict__ av
Definition ax_helm_full_kernel.h:44
v
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ v
Definition ax_helm_full_kernel.h:47

◆ STR [1/2]

#define STR (   X)    #X

◆ STR [2/2]

#define STR (   X)    #X

Function Documentation

◆ opencl_ax_helm()

void opencl_ax_helm ( void *  w,
void *  u,
void *  dx,
void *  dy,
void *  dz,
void *  dxt,
void *  dyt,
void *  dzt,
void *  h1,
void *  g11,
void *  g22,
void *  g33,
void *  g12,
void *  g13,
void *  g23,
int *  nelv,
int *  lx 
)

Fortran wrapper for device OpenCL Ax

Definition at line 57 of file ax_helm.c.

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

◆ opencl_ax_helm_vector()

void opencl_ax_helm_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 *  g11,
void *  g22,
void *  g33,
void *  g12,
void *  g13,
void *  g23,
int *  nelv,
int *  lx 
)

Fortran wrapper for device OpenCL Ax (vector version)

Definition at line 258 of file ax_helm.c.

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

◆ opencl_ax_helm_vector_part2()

void opencl_ax_helm_vector_part2 ( void *  au,
void *  av,
void *  aw,
void *  u,
void *  v,
void *  w,
void *  h2,
void *  B,
int *  n 
)

Fortran wrapper for device OpenCL Ax (vector version) part2

Definition at line 341 of file ax_helm.c.

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

Variable Documentation

◆ autotune_ax_helm

int* autotune_ax_helm = NULL

Definition at line 52 of file ax_helm.c.