Neko 1.99.5
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
Macros | Functions | Variables
opr_cdtp.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 "cdtp_kernel.cl.h"
Include dependency graph for opr_cdtp.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)
 

Functions

void opencl_cdtp (void *dtx, void *x, void *dr, void *ds, void *dt, void *dxt, void *dyt, void *dzt, void *w3, int *nel, int *lx)
 

Variables

int * autotune_cdtp = NULL
 

Macro Definition Documentation

◆ CASE

#define CASE (   LX)

◆ CASE_1D

#define CASE_1D (   LX,
  QUEUE,
  EVENT 
)
Value:
{ \
cl_kernel kernel = clCreateKernel(cdtp_program, \
STR(cdtp_kernel_lx##LX), &err); \
CL_CHECK(err); \
\
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &dtx)); \
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &x)); \
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &dr)); \
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &ds)); \
CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &dt)); \
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 *) &w3)); \
\
CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) QUEUE, \
kernel, 1, NULL, &global_item_size, \
&local_item_size, 0, NULL, EVENT)); \
CL_CHECK(clReleaseKernel(kernel)); \
}
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
dr
__global__ void const T *__restrict__ const T *__restrict__ dr
Definition cdtp_kernel.h:107
ds
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ ds
Definition cdtp_kernel.h:108
x
__global__ void const T *__restrict__ x
Definition cdtp_kernel.h:106
dt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dt
Definition cdtp_kernel.h:109
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
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:112
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:111
dxt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dxt
Definition cdtp_kernel.h:110
STR
#define STR(X)
cdtp_program
void * cdtp_program

◆ CASE_KSTEP

#define CASE_KSTEP (   LX,
  QUEUE,
  EVENT 
)
Value:
{ \
cl_kernel kernel = clCreateKernel(cdtp_program, \
STR(cdtp_kernel_kstep_lx##LX), &err); \
CL_CHECK(err); \
\
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &dtx)); \
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &x)); \
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &dr)); \
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &ds)); \
CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &dt)); \
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 *) &w3)); \
\
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

◆ STR

#define STR (   X)    #X

Function Documentation

◆ opencl_cdtp()

void opencl_cdtp ( void *  dtx,
void *  x,
void *  dr,
void *  ds,
void *  dt,
void *  dxt,
void *  dyt,
void *  dzt,
void *  w3,
int *  nel,
int *  lx 
)

Fortran wrapper for device OpenCL \( D^T X \)

Definition at line 57 of file opr_cdtp.c.

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

Variable Documentation

◆ autotune_cdtp

int* autotune_cdtp = NULL

Definition at line 52 of file opr_cdtp.c.