Neko  0.8.99
A portable framework for high-order spectral element flow simulations
Macros | Functions
opr_cdtp.cu File Reference
#include <string.h>
#include <stdlib.h>
#include <stdio.h>
#include "cdtp_kernel.h"
#include <device/device_config.h>
#include <device/cuda/check.h>
#include <common/neko_log.h>
Include dependency graph for opr_cdtp.cu:

Go to the source code of this file.

Macros

#define CASE_1D(LX)
 
#define CASE_KSTEP(LX)
 
#define CASE(LX)
 
#define CASE_LARGE(LX)
 

Functions

template<const int >
int tune_cdtp (void *dtx, void *x, void *dr, void *ds, void *dt, void *dxt, void *dyt, void *dzt, void *w3, int *nel, int *lx)
 
void cuda_cdtp (void *dtx, void *x, void *dr, void *ds, void *dt, void *dxt, void *dyt, void *dzt, void *w3, int *nel, int *lx)
 

Macro Definition Documentation

◆ CASE

#define CASE (   LX)
Value:
case LX: \
if(autotune[LX] == 0 ) { \
autotune[LX]=tune_cdtp<LX>(dtx, x, \
dr, ds, dt, \
dxt, dyt, dzt, \
w3, nel, lx); \
} else if (autotune[LX] == 1 ) { \
CASE_1D(LX); \
} else if (autotune[LX] == 2 ) { \
CASE_KSTEP(LX); \
} \
break
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

◆ CASE_1D

#define CASE_1D (   LX)
Value:
cdtp_kernel_1d<real, LX, 1024> \
<<<nblcks, nthrds_1d, 0, stream>>>((real *) dtx, (real *) x, \
(real *) dr, (real *) ds, (real *) dt, \
(real *) dxt, (real *) dyt, (real *) dzt, \
(real *) w3); \
CUDA_CHECK(cudaGetLastError());
real
double real
Definition: device_config.h:12

◆ CASE_KSTEP

#define CASE_KSTEP (   LX)
Value:
cdtp_kernel_kstep<real, LX> \
<<<nblcks, nthrds_kstep, 0, stream>>>((real *) dtx, (real *) x, \
(real *) dr, (real *) ds, (real *) dt, \
(real *) dxt, (real *) dyt, (real *) dzt, \
(real *) w3); \
CUDA_CHECK(cudaGetLastError());

◆ CASE_LARGE

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

Function Documentation

◆ cuda_cdtp()

void cuda_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 cuda \( D^T X \)

Definition at line 57 of file opr_cdtp.cu.

Here is the caller graph for this function:

◆ tune_cdtp()

template<const int >
int tune_cdtp ( void *  dtx,
void *  x,
void *  dr,
void *  ds,
void *  dt,
void *  dxt,
void *  dyt,
void *  dzt,
void *  w3,
int *  nel,
int *  lx 
)

Definition at line 142 of file opr_cdtp.cu.

Here is the call graph for this function: