Neko 1.99.1
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
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)
 

Variables

intautotune_ax_helm = NULL
 

Macro Definition Documentation

◆ CASE

#define CASE (   LX)

◆ CASE_1D

#define CASE_1D (   LX,
  QUEUE,
  EVENT 
)
Value:
{ \
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)); \
\
}
#define STR(X)
__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__ u
__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__ 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 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
__global__ void dirichlet_apply_scalar_kernel(const int *__restrict__ msk, T *__restrict__ x, const T g, const int m)
__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
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ g22
__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
__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
__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
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ g11
void * ax_helm_program

◆ CASE_KSTEP

#define CASE_KSTEP (   LX,
  QUEUE,
  EVENT 
)
Value:
{ \
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)); \
\
}

◆ 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_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)); \
local_kstep, 0, NULL, NULL)); \
\
} \
break
__global__ void T *__restrict__ T *__restrict__ aw
__global__ void T *__restrict__ av
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ v

◆ 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:

Variable Documentation

◆ autotune_ax_helm

int* autotune_ax_helm = NULL

Definition at line 52 of file ax_helm.c.