Neko  0.8.1
A portable framework for high-order spectral element flow simulations
Macros | Functions
ax_helm.hip File Reference
#include <string.h>
#include <stdlib.h>
#include <stdio.h>
#include <hip/hip_runtime.h>
#include <device/device_config.h>
#include <device/hip/check.h>
#include "ax_helm_kernel.h"
#include <common/neko_log.h>
Include dependency graph for ax_helm.hip:

Go to the source code of this file.

Macros

#define CASE_1D(LX)
 
#define CASE_KSTEP(LX)
 
#define CASE_KSTEP_PADDED(LX)
 
#define CASE(LX)
 
#define CASE_PADDED(LX)
 

Functions

template<const int >
int tune (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)
 
template<const int >
int tune_padded (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 hip_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)
 

Macro Definition Documentation

◆ CASE

#define CASE (   LX)
Value:
case LX: \
if(autotune[LX] == 0 ) { \
autotune[LX]=tune<LX>( w, u, \
dx, dy, dz, \
dxt, dyt, dzt,h1, \
g11, g22, g33, \
g12, g13, g23, nelv, lx); \
} else if (autotune[LX] == 1 ) { \
CASE_1D(LX); \
} else if (autotune[LX] == 2 ) { \
CASE_KSTEP(LX); \
} \
break
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:115
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:114
dxt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dxt
Definition: cdtp_kernel.h:113
dz
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dz
Definition: conv1_kernel.h:138
dx
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dx
Definition: conv1_kernel.h:136
u
__global__ void const T *__restrict__ u
Definition: conv1_kernel.h:132
dy
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dy
Definition: conv1_kernel.h:137
w
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ w
Definition: lambda2_kernel.h:244
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:160
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:159
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:157
h1
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ h1
Definition: ax_helm_kernel.h:155
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:158
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:161
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:156

◆ CASE_1D

#define CASE_1D (   LX)
Value:
hipLaunchKernelGGL(HIP_KERNEL_NAME( ax_helm_kernel_1d<real, LX, 1024> ), \
nblcks_1d, nthrds_1d, 0, \
(hipStream_t) glb_cmd_queue, \
(real *) w, (real *) u, \
(real *) dx, (real *) dy, (real *) dz, \
(real *) dxt, (real *) dyt, (real *) dzt, (real *) h1, \
(real *) g11, (real *) g22, (real *) g33, \
(real *) g12, (real *) g13, (real *) g23); \
HIP_CHECK(hipGetLastError());
real
double real
Definition: device_config.h:12
glb_cmd_queue
void * glb_cmd_queue

◆ CASE_KSTEP

#define CASE_KSTEP (   LX)
Value:
hipLaunchKernelGGL( HIP_KERNEL_NAME( ax_helm_kernel_kstep<real, LX> ), \
nblcks_kstep, nthrds_kstep, 0, \
(hipStream_t) glb_cmd_queue, \
(real *) w, (real *) u, \
(real *) dx, (real *) dy, (real *) dz, (real *) h1, \
(real *) g11, (real *) g22, (real *) g33, \
(real *) g12, (real *) g13, (real *) g23); \
HIP_CHECK(hipGetLastError());

◆ CASE_KSTEP_PADDED

#define CASE_KSTEP_PADDED (   LX)
Value:
hipLaunchKernelGGL( HIP_KERNEL_NAME(ax_helm_kernel_kstep_padded<real, LX> ),\
nblcks_kstep, nthrds_kstep, 0, \
(hipStream_t) glb_cmd_queue, \
(real *) w, (real *) u, \
(real *) dx, (real *) dy, (real *) dz, (real *) h1, \
(real *) g11, (real *) g22, (real *) g33, \
(real *) g12, (real *) g13, (real *) g23); \
HIP_CHECK(hipGetLastError());

◆ CASE_PADDED

#define CASE_PADDED (   LX)
Value:
case LX: \
if(autotune[LX] == 0 ) { \
autotune[LX]=tune_padded<LX>(w, u, \
dx, dy, dz, \
dxt, dyt, dzt,h1, \
g11, g22, g33, \
g12, g13, g23,nelv,lx); \
} else if (autotune[LX] == 1 ) { \
CASE_1D(LX); \
} else if (autotune[LX] == 2 ) { \
CASE_KSTEP_PADDED(LX); \
} \
break

Function Documentation

◆ hip_ax_helm()

void hip_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 HIP Ax

Definition at line 64 of file ax_helm.hip.

◆ tune()

template<const int >
int tune ( 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 
)

Definition at line 164 of file ax_helm.hip.

Here is the call graph for this function:

◆ tune_padded()

template<const int >
int tune_padded ( 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 
)

Definition at line 241 of file ax_helm.hip.

Here is the call graph for this function: