Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
Macros | Functions
ax_helm_compr.cu File Reference
#include <string.h>
#include <stdlib.h>
#include <stdio.h>
#include "ax_helm_compr_kernel.h"
#include "elem_block_tune.h"
#include <device/device_config.h>
#include <device/cuda/check.h>
#include <common/neko_log.h>

Go to the source code of this file.

Macros

#define CASE_1D(LX, C)
 
#define CASE_1D_SEL(LX, SEL)
 
#define CASE_KSTEP(LX, C)
 
#define CASE_KSTEP_PADDED(LX, C)
 
#define CASE_KSTEP_SEL(LX, SEL)
 
#define CASE_KSTEP_PADDED_SEL(LX, SEL)
 
#define CASE_DMMA(LX, C)
 
#define CASE_DMMA_SEL(LX, SEL)
 
#define CASE_DMMA_TMA(LX, C)
 
#define CASE_DMMA_TMA_SEL(LX, SEL)
 
#define CASE(LX)
 
#define CASE_PADDED(LX)
 
#define CASE_LARGE(LX)
 
#define CASE_LARGE_PADDED(LX)
 
#define CASE_VECTOR_KSTEP(LX, C)
 
#define CASE_VECTOR_KSTEP_PADDED(LX, C)
 
#define CASE_VECTOR_KSTEP_SEL(LX, SEL)
 
#define CASE_VECTOR_KSTEP_PADDED_SEL(LX, SEL)
 
#define CASE_VECTOR_DMMA(LX, C)
 
#define CASE_VECTOR_DMMA_SEL(LX, SEL)
 
#define CASE_VECTOR_DMMA_TMA(LX, C)
 
#define CASE_VECTOR_DMMA_TMA_SEL(LX, SEL)
 
#define CASE_VECTOR_DMMA_TMA_BATCH(LX, C)
 
#define CASE_VECTOR_DMMA_TMA_BATCH_SEL(LX, SEL)
 
#define CASE_VECTOR(LX)
 
#define CASE_VECTOR_PADDED(LX)
 

Functions

template<const int >
int tune_compr (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, void *ind, int *nelv, int *lx, int *eb_sel, int *ch_sel, int *nw_sel, int *tw_sel)
 
template<const int >
int tune_compr_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, void *ind, int *nelv, int *lx, int *eb_sel, int *ch_sel, int *nw_sel, int *tw_sel)
 
template<const int >
int tune_compr_vector (void *au, void *av, void *aw, void *u, void *v, void *w, void *dx, void *dy, void *dz, void *h1, void *g11, void *g22, void *g33, void *g12, void *g13, void *g23, void *ind, int *nelv, int *lx, int *eb_sel, int *nw_sel, int *tw_sel, int *bw_sel)
 
template<const int >
int tune_compr_vector_padded (void *au, void *av, void *aw, void *u, void *v, void *w, void *dx, void *dy, void *dz, void *h1, void *g11, void *g22, void *g33, void *g12, void *g13, void *g23, void *ind, int *nelv, int *lx, int *eb_sel, int *nw_sel, int *tw_sel, int *bw_sel)
 
void cuda_ax_helm_compr (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, void *ind, int *nelv, int *lx)
 
void cuda_ax_helm_compr_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, void *ind, int *nelv, int *lx)
 

Macro Definition Documentation

◆ CASE

#define CASE (   LX)
Value:
case LX: \
if(autotune[LX] == 0 ) { \
autotune[LX]=tune_compr<LX>( w, u, \
dx, dy, dz, \
dxt, dyt, dzt,h1, \
g11, g22, g33, \
g12, g13, g23, ind, nelv, lx, \
&autotune_eb[LX], &autotune_ch[LX], \
&autotune_nw[LX], &autotune_tw[LX]); \
} else if (autotune[LX] == 1 ) { \
CASE_1D_SEL(LX, autotune_ch[LX]); \
} else if (autotune[LX] == 2 ) { \
CASE_KSTEP_SEL(LX, autotune_eb[LX]); \
} else if (autotune[LX] == 3 ) { \
CASE_DMMA_SEL(LX, autotune_nw[LX]); \
} else if (autotune[LX] == 4 ) { \
CASE_DMMA_TMA_SEL(LX, autotune_tw[LX]); \
} \
break
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
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_compr_kernel.h:209
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_compr_kernel.h:205
ind
__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__ const int *__restrict__ ind
Definition ax_helm_compr_kernel.h:210
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_compr_kernel.h:208
w
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ w
Definition ax_helm_compr_kernel.h:978
nelv
__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__ const int *__restrict__ const int nelv
Definition ax_helm_compr_kernel.h:211
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_compr_kernel.h:207
dz
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dz
Definition ax_helm_compr_kernel.h:202
dx
__global__ void const T *__restrict__ const T *__restrict__ dx
Definition ax_helm_compr_kernel.h:200
u
__global__ void const T *__restrict__ u
Definition ax_helm_compr_kernel.h:199
dy
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ dy
Definition ax_helm_compr_kernel.h:201
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_compr_kernel.h:206
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_compr_kernel.h:204
h1
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ h1
Definition ax_helm_compr_kernel.h:203
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

◆ CASE_1D

#define CASE_1D (   LX,
  C 
)
Value:
ax_helm_compr_kernel_1d<real, LX, NEKO_CHUNKS(LX, C)> \
<<<nblcks_1d, NEKO_CHUNKS_NTHRDS(LX, C), 0, stream>>> \
((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, \
(const int *) ind); \
CUDA_CHECK(cudaGetLastError());
ax_helm_compr_kernel_1d
__global__ void ax_helm_compr_kernel_1d(T *__restrict__ w, const T *__restrict__ u, const T *__restrict__ dx, const T *__restrict__ dy, const T *__restrict__ dz, const T *__restrict__ dxt, const T *__restrict__ dyt, const T *__restrict__ dzt, const T *__restrict__ h1, const T *__restrict__ g11, const T *__restrict__ g22, const T *__restrict__ g33, const T *__restrict__ g12, const T *__restrict__ g13, const T *__restrict__ g23, const int *__restrict__ ind)
Definition ax_helm_compr_kernel.h:88
NEKO_CHUNKS
#define NEKO_CHUNKS(LX, C)
Definition elem_block.h:145
NEKO_CHUNKS_NTHRDS
#define NEKO_CHUNKS_NTHRDS(LX, C)
Definition elem_block.h:146
real
double real
Definition device_config.h:12

◆ CASE_1D_SEL

#define CASE_1D_SEL (   LX,
  SEL 
)
Value:
switch (SEL) { \
case 0: CASE_1D(LX, 0); break; \
case 1: CASE_1D(LX, 1); break; \
case 2: CASE_1D(LX, 2); break; \
default: CASE_1D(LX, 3); break; \
}
CASE_1D
#define CASE_1D(LX, C)

◆ CASE_DMMA

#define CASE_DMMA (   LX,
  C 
)
Value:
ax_helm_compr_kernel_dmma<real, LX, NEKO_DMMA_NW(C)> \
<<<NEKO_DMMA_NBLCKS(*nelv, LX), NEKO_DMMA_NTHRDS(C), 0, stream>>> \
((real *) w, (real *) u, \
(real *) dx, (real *) dy, (real *) dz, (real *) h1, \
(real *) g11, (real *) g22, (real *) g33, \
(real *) g12, (real *) g13, (real *) g23, \
(const int *) ind, *nelv); \
CUDA_CHECK(cudaGetLastError());
NEKO_DMMA_NW
#define NEKO_DMMA_NW(C)
Definition dmma_kernel.h:124
NEKO_DMMA_NBLCKS
#define NEKO_DMMA_NBLCKS(NELV, LX)
Definition dmma_kernel.h:150
NEKO_DMMA_NTHRDS
#define NEKO_DMMA_NTHRDS(C)
Definition dmma_kernel.h:125

◆ CASE_DMMA_SEL

#define CASE_DMMA_SEL (   LX,
  SEL 
)
Value:
switch (SEL) { \
case 0: CASE_DMMA(LX, 0); break; \
case 1: CASE_DMMA(LX, 1); break; \
default: CASE_DMMA(LX, 2); break; \
}
CASE_DMMA
#define CASE_DMMA(LX, C)

◆ CASE_DMMA_TMA

#define CASE_DMMA_TMA (   LX,
  C 
)
Value:
ax_helm_compr_kernel_dmma_tma<real, LX, NEKO_DMMA_NW(C)> \
<<<NEKO_DMMA_NBLCKS(*nelv, LX), NEKO_DMMA_NTHRDS(C), 0, stream>>> \
((real *) w, (real *) u, \
(real *) dx, (real *) dy, (real *) dz, (real *) h1, \
(real *) g11, (real *) g22, (real *) g33, \
(real *) g12, (real *) g13, (real *) g23, \
(const int *) ind); \
CUDA_CHECK(cudaGetLastError());

◆ CASE_DMMA_TMA_SEL

#define CASE_DMMA_TMA_SEL (   LX,
  SEL 
)
Value:
switch (SEL) { \
case 0: CASE_DMMA_TMA(LX, 0); break; \
case 1: CASE_DMMA_TMA(LX, 1); break; \
default: CASE_DMMA_TMA(LX, 2); break; \
}
CASE_DMMA_TMA
#define CASE_DMMA_TMA(LX, C)

◆ CASE_KSTEP

#define CASE_KSTEP (   LX,
  C 
)
Value:
ax_helm_compr_kernel_kstep<real, LX, NEKO_EB(LX, C)> \
<<<NEKO_EB_NBLCKS(*nelv, LX, C), NEKO_EB_NTHRDS(LX, C), 0, stream>>> \
((real *) w, (real *) u, \
(real *) dx, (real *) dy, (real *) dz, (real *) h1, \
(real *) g11, (real *) g22, (real *) g33, \
(real *) g12, (real *) g13, (real *) g23, \
(const int *) ind, *nelv); \
CUDA_CHECK(cudaGetLastError());
NEKO_EB
#define NEKO_EB(LX, C)
Definition elem_block.h:104
NEKO_EB_NBLCKS
#define NEKO_EB_NBLCKS(NELV, LX, C)
Definition elem_block.h:106
NEKO_EB_NTHRDS
#define NEKO_EB_NTHRDS(LX, C)
Definition elem_block.h:105

◆ CASE_KSTEP_PADDED

#define CASE_KSTEP_PADDED (   LX,
  C 
)
Value:
ax_helm_compr_kernel_kstep_padded<real, LX, NEKO_EB(LX, C)> \
<<<NEKO_EB_NBLCKS(*nelv, LX, C), NEKO_EB_NTHRDS(LX, C), 0, stream>>> \
((real *) w, (real *) u, \
(real *) dx, (real *) dy, (real *) dz, (real *) h1, \
(real *) g11, (real *) g22, (real *) g33, \
(real *) g12, (real *) g13, (real *) g23, \
(const int *) ind, *nelv); \
CUDA_CHECK(cudaGetLastError());

◆ CASE_KSTEP_PADDED_SEL

#define CASE_KSTEP_PADDED_SEL (   LX,
  SEL 
)
Value:
switch (SEL) { \
case 0: CASE_KSTEP_PADDED(LX, 0); break; \
case 1: CASE_KSTEP_PADDED(LX, 1); break; \
default: CASE_KSTEP_PADDED(LX, 2); break; \
}
CASE_KSTEP_PADDED
#define CASE_KSTEP_PADDED(LX, C)

◆ CASE_KSTEP_SEL

#define CASE_KSTEP_SEL (   LX,
  SEL 
)
Value:
switch (SEL) { \
case 0: CASE_KSTEP(LX, 0); break; \
case 1: CASE_KSTEP(LX, 1); break; \
default: CASE_KSTEP(LX, 2); break; \
}
CASE_KSTEP
#define CASE_KSTEP(LX, C)

◆ CASE_LARGE

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

◆ CASE_LARGE_PADDED

#define CASE_LARGE_PADDED (   LX)
Value:
case LX: \
CASE_KSTEP_PADDED(LX, 0); \
break

◆ CASE_PADDED

#define CASE_PADDED (   LX)
Value:
case LX: \
if(autotune[LX] == 0 ) { \
autotune[LX]=tune_compr_padded<LX>(w, u, \
dx, dy, dz, \
dxt, dyt, dzt,h1, \
g11, g22, g33, \
g12, g13, g23, ind,nelv,lx, \
&autotune_eb[LX], &autotune_ch[LX], \
&autotune_nw[LX], &autotune_tw[LX]); \
} else if (autotune[LX] == 1 ) { \
CASE_1D_SEL(LX, autotune_ch[LX]); \
} else if (autotune[LX] == 2 ) { \
CASE_KSTEP_PADDED_SEL(LX, autotune_eb[LX]); \
} else if (autotune[LX] == 3 ) { \
CASE_DMMA_SEL(LX, autotune_nw[LX]); \
} else if (autotune[LX] == 4 ) { \
CASE_DMMA_TMA_SEL(LX, autotune_tw[LX]); \
} \
break

◆ CASE_VECTOR

#define CASE_VECTOR (   LX)
Value:
case LX: \
if (autotune_vec[LX] == 0 ) { \
autotune_vec[LX] = tune_compr_vector<LX>(au, av, aw, u, v, w, dx, dy, \
dz, h1, g11, g22, g33, g12, \
g13, g23, ind, nelv, lx, \
&autotune_vec_eb[LX], \
&autotune_vec_nw[LX], \
&autotune_vec_tw[LX], \
&autotune_vec_bw[LX]); \
} else if (autotune_vec[LX] == 2 ) { \
CASE_VECTOR_KSTEP_SEL(LX, autotune_vec_eb[LX]); \
} else if (autotune_vec[LX] == 3 ) { \
CASE_VECTOR_DMMA_SEL(LX, autotune_vec_nw[LX]); \
} else if (autotune_vec[LX] == 4 ) { \
CASE_VECTOR_DMMA_TMA_SEL(LX, autotune_vec_tw[LX]); \
} else if (autotune_vec[LX] == 5 ) { \
CASE_VECTOR_DMMA_TMA_BATCH_SEL(LX, autotune_vec_bw[LX]); \
} \
break
aw
__global__ void T *__restrict__ T *__restrict__ aw
Definition ax_helm_compr_kernel.h:975
av
__global__ void T *__restrict__ av
Definition ax_helm_compr_kernel.h:974
v
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ v
Definition ax_helm_compr_kernel.h:977

◆ CASE_VECTOR_DMMA

#define CASE_VECTOR_DMMA (   LX,
  C 
)
Value:
ax_helm_compr_kernel_dmma_vector<real, LX, NEKO_DMMA_NW(C)> \
<<<nblcks_dmma, NEKO_DMMA_NTHRDS(C), 0, stream>>> \
((real *) au, (real *) av, (real *) aw, \
(real *) u, (real *) v, (real *) w, \
(real *) dx, (real *) dy, (real *) dz, \
(real *) h1, (real *) g11, (real *) g22, \
(real *) g33, (real *) g12, (real *) g13, \
(real *) g23, \
(const int *) ind); \
CUDA_CHECK(cudaGetLastError());

◆ CASE_VECTOR_DMMA_SEL

#define CASE_VECTOR_DMMA_SEL (   LX,
  SEL 
)
Value:
switch (SEL) { \
case 0: CASE_VECTOR_DMMA(LX, 0); break; \
case 1: CASE_VECTOR_DMMA(LX, 1); break; \
default: CASE_VECTOR_DMMA(LX, 2); break; \
}
CASE_VECTOR_DMMA
#define CASE_VECTOR_DMMA(LX, C)

◆ CASE_VECTOR_DMMA_TMA

#define CASE_VECTOR_DMMA_TMA (   LX,
  C 
)
Value:
ax_helm_compr_kernel_dmma_tma_vector<real, LX, NEKO_DMMA_NW(C)> \
<<<nblcks_dmma, NEKO_DMMA_NTHRDS(C), 0, stream>>> \
((real *) au, (real *) av, (real *) aw, \
(real *) u, (real *) v, (real *) w, \
(real *) dx, (real *) dy, (real *) dz, \
(real *) h1, (real *) g11, (real *) g22, \
(real *) g33, (real *) g12, (real *) g13, \
(real *) g23, \
(const int *) ind); \
CUDA_CHECK(cudaGetLastError());

◆ CASE_VECTOR_DMMA_TMA_BATCH

#define CASE_VECTOR_DMMA_TMA_BATCH (   LX,
  C 
)
Value:
(void) ax_helm_compr_dmma_tma_batch_optin<real, LX, NEKO_DMMA_NW(C)>(); \
ax_helm_compr_kernel_dmma_tma_batch<real, LX, NEKO_DMMA_NW(C)> \
<<<nblcks_dmma, NEKO_DMMA_NTHRDS(C), \
NEKO_DMMA_TMA_BATCH_SMEM, stream>>> \
((real *) au, (real *) av, (real *) aw, \
(real *) u, (real *) v, (real *) w, \
(real *) dx, (real *) dy, (real *) dz, \
(real *) h1, (real *) g11, (real *) g22, \
(real *) g33, (real *) g12, (real *) g13, \
(real *) g23, \
(const int *) ind); \
CUDA_CHECK(cudaGetLastError());
ax_helm_compr_dmma_tma_batch_optin
static bool ax_helm_compr_dmma_tma_batch_optin()
Definition ax_helm_compr_kernel.h:2289
NEKO_DMMA_TMA_BATCH_SMEM
#define NEKO_DMMA_TMA_BATCH_SMEM
Definition dmma_tma_kernel.h:295

◆ CASE_VECTOR_DMMA_TMA_BATCH_SEL

#define CASE_VECTOR_DMMA_TMA_BATCH_SEL (   LX,
  SEL 
)
Value:
switch (SEL) { \
case 0: CASE_VECTOR_DMMA_TMA_BATCH(LX, 0); break; \
case 1: CASE_VECTOR_DMMA_TMA_BATCH(LX, 1); break; \
default: CASE_VECTOR_DMMA_TMA_BATCH(LX, 2); break; \
}
CASE_VECTOR_DMMA_TMA_BATCH
#define CASE_VECTOR_DMMA_TMA_BATCH(LX, C)

◆ CASE_VECTOR_DMMA_TMA_SEL

#define CASE_VECTOR_DMMA_TMA_SEL (   LX,
  SEL 
)
Value:
switch (SEL) { \
case 0: CASE_VECTOR_DMMA_TMA(LX, 0); break; \
case 1: CASE_VECTOR_DMMA_TMA(LX, 1); break; \
default: CASE_VECTOR_DMMA_TMA(LX, 2); break; \
}
CASE_VECTOR_DMMA_TMA
#define CASE_VECTOR_DMMA_TMA(LX, C)

◆ CASE_VECTOR_KSTEP

#define CASE_VECTOR_KSTEP (   LX,
  C 
)
Value:
ax_helm_compr_kernel_vector_kstep<real, LX, NEKO_EB(LX, C)> \
<<<NEKO_EB_NBLCKS(*nelv, LX, C), NEKO_EB_NTHRDS(LX, C), 0, stream>>> \
((real *) au, (real *) av, (real *) aw, \
(real *) u, (real *) v, (real *) w, \
(real *) dx, (real *) dy, (real *) dz, \
(real *) h1, (real *) g11, (real *) g22, \
(real *) g33, (real *) g12, (real *) g13, \
(real *) g23, \
(const int *) ind, *nelv); \
CUDA_CHECK(cudaGetLastError());

◆ CASE_VECTOR_KSTEP_PADDED

#define CASE_VECTOR_KSTEP_PADDED (   LX,
  C 
)
Value:
ax_helm_compr_kernel_vector_kstep_padded<real, LX, NEKO_EB(LX, C)> \
<<<NEKO_EB_NBLCKS(*nelv, LX, C), NEKO_EB_NTHRDS(LX, C), 0, stream>>> \
((real *) au, (real *) av, (real *) aw, \
(real *) u, (real *) v, (real *) w, \
(real *) dx, (real *) dy, (real *) dz, \
(real *) h1, (real *) g11, (real *) g22, \
(real *) g33, (real *) g12, (real *) g13, \
(real *) g23, \
(const int *) ind, *nelv); \
CUDA_CHECK(cudaGetLastError());

◆ CASE_VECTOR_KSTEP_PADDED_SEL

#define CASE_VECTOR_KSTEP_PADDED_SEL (   LX,
  SEL 
)
Value:
switch (SEL) { \
case 0: CASE_VECTOR_KSTEP_PADDED(LX, 0); break; \
case 1: CASE_VECTOR_KSTEP_PADDED(LX, 1); break; \
default: CASE_VECTOR_KSTEP_PADDED(LX, 2); break; \
}
CASE_VECTOR_KSTEP_PADDED
#define CASE_VECTOR_KSTEP_PADDED(LX, C)

◆ CASE_VECTOR_KSTEP_SEL

#define CASE_VECTOR_KSTEP_SEL (   LX,
  SEL 
)
Value:
switch (SEL) { \
case 0: CASE_VECTOR_KSTEP(LX, 0); break; \
case 1: CASE_VECTOR_KSTEP(LX, 1); break; \
default: CASE_VECTOR_KSTEP(LX, 2); break; \
}
CASE_VECTOR_KSTEP
#define CASE_VECTOR_KSTEP(LX, C)

◆ CASE_VECTOR_PADDED

#define CASE_VECTOR_PADDED (   LX)
Value:
case LX: \
if (autotune_vec[LX] == 0 ) { \
autotune_vec[LX] = tune_compr_vector_padded<LX>(au, av, aw, u, v, w, \
dx, dy, dz, h1, \
g11, g22, g33, \
g12, g13, g23, ind, \
nelv, lx, \
&autotune_vec_eb[LX], \
&autotune_vec_nw[LX], \
&autotune_vec_tw[LX], \
&autotune_vec_bw[LX]); \
} else if (autotune_vec[LX] == 2 ) { \
CASE_VECTOR_KSTEP_PADDED_SEL(LX, autotune_vec_eb[LX]); \
} else if (autotune_vec[LX] == 3 ) { \
CASE_VECTOR_DMMA_SEL(LX, autotune_vec_nw[LX]); \
} else if (autotune_vec[LX] == 4 ) { \
CASE_VECTOR_DMMA_TMA_SEL(LX, autotune_vec_tw[LX]); \
} else if (autotune_vec[LX] == 5 ) { \
CASE_VECTOR_DMMA_TMA_BATCH_SEL(LX, autotune_vec_bw[LX]); \
} \
break

Function Documentation

◆ cuda_ax_helm_compr()

void cuda_ax_helm_compr ( 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,
void *  ind,
int *  nelv,
int *  lx 
)

Fortran wrapper for device CUDA Ax, compressed geometric factors

Definition at line 97 of file ax_helm_compr.cu.

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

◆ cuda_ax_helm_compr_vector()

void cuda_ax_helm_compr_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,
void *  ind,
int *  nelv,
int *  lx 
)

Fortran wrapper for device CUDA Ax vector version, compressed geometric factors

Definition at line 309 of file ax_helm_compr.cu.

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

◆ tune_compr()

template<const int >
int tune_compr ( 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,
void *  ind,
int *  nelv,
int *  lx,
int *  eb_sel,
int *  ch_sel,
int *  nw_sel,
int *  tw_sel 
)

Definition at line 522 of file ax_helm_compr.cu.

Here is the call graph for this function:

◆ tune_compr_padded()

template<const int >
int tune_compr_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,
void *  ind,
int *  nelv,
int *  lx,
int *  eb_sel,
int *  ch_sel,
int *  nw_sel,
int *  tw_sel 
)

Definition at line 768 of file ax_helm_compr.cu.

Here is the call graph for this function:

◆ tune_compr_vector()

template<const int >
int tune_compr_vector ( void *  au,
void *  av,
void *  aw,
void *  u,
void *  v,
void *  w,
void *  dx,
void *  dy,
void *  dz,
void *  h1,
void *  g11,
void *  g22,
void *  g33,
void *  g12,
void *  g13,
void *  g23,
void *  ind,
int *  nelv,
int *  lx,
int *  eb_sel,
int *  nw_sel,
int *  tw_sel,
int *  bw_sel 
)

Definition at line 1020 of file ax_helm_compr.cu.

Here is the call graph for this function:

◆ tune_compr_vector_padded()

template<const int >
int tune_compr_vector_padded ( void *  au,
void *  av,
void *  aw,
void *  u,
void *  v,
void *  w,
void *  dx,
void *  dy,
void *  dz,
void *  h1,
void *  g11,
void *  g22,
void *  g33,
void *  g12,
void *  g13,
void *  g23,
void *  ind,
int *  nelv,
int *  lx,
int *  eb_sel,
int *  nw_sel,
int *  tw_sel,
int *  bw_sel 
)

Definition at line 1258 of file ax_helm_compr.cu.

Here is the call graph for this function: