diff --git a/acc/compile.sh b/acc/compile.sh
index 7f7c143..0fe6b4c 100755
--- a/acc/compile.sh
+++ b/acc/compile.sh
@@ -8,17 +8,21 @@ FILENAME="${FULL_NAME%.*}"
EXTENSION="${FULL_NAME##*.}"
if [ "${EXTENSION}" = "sas" ]; then
- echo "Generating stencil assembly stage ${FILENAME}.sas -> stencil_assembly.cuh"
COMPILE_FLAGS="-sas" # Generate stencil assembly stage
CUH_FILENAME="stencil_assembly.cuh"
+ echo "Generating stencil assembly stage ${FILENAME}.sas -> ${CUH_FILENAME}"
elif [ "${EXTENSION}" = "sps" ]; then
- echo "Generating stencil processing stage: ${FILENAME}.sps -> stencil_process.cuh"
COMPILE_FLAGS="-sps" # Generate stencil processing stage
CUH_FILENAME="stencil_process.cuh"
+ echo "Generating stencil processing stage: ${FILENAME}.sps -> ${CUH_FILENAME}"
+elif [ "${EXTENSION}" = "sdh" ]; then
+ COMPILE_FLAGS="-sdh" # Generate stencil definition header
+ CUH_FILENAME="stencil_defines.h"
+ echo "Generating stencil definition header: ${FILENAME}.sdh -> ${CUH_FILENAME}"
else
echo "Error: unknown extension" ${EXTENSION} "of file" ${FULL_NAME}
- echo "Extension should be either .sas or .sps"
+ echo "Extension should be either .sas, .sps or .sdh"
exit
fi
-${ACC_DIR}/preprocess.sh $2 $1 | ${ACC_DIR}/build/acc ${COMPILE_FLAGS} > ${CUH_FILENAME}
+${ACC_DIR}/preprocess.sh $1 | ${ACC_DIR}/build/acc ${COMPILE_FLAGS} > ${CUH_FILENAME}
diff --git a/acc/mhd_solver/stencil_defines.h b/acc/mhd_solver/stencil_defines.h
deleted file mode 100644
index b4bc622..0000000
--- a/acc/mhd_solver/stencil_defines.h
+++ /dev/null
@@ -1,163 +0,0 @@
-/*
- Copyright (C) 2014-2019, Johannes Pekkilae, Miikka Vaeisalae.
-
- This file is part of Astaroth.
-
- Astaroth is free software: you can redistribute it and/or modify
- it under the terms of the GNU General Public License as published by
- the Free Software Foundation, either version 3 of the License, or
- (at your option) any later version.
-
- Astaroth is distributed in the hope that it will be useful,
- but WITHOUT ANY WARRANTY; without even the implied warranty of
- MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
- GNU General Public License for more details.
-
- You should have received a copy of the GNU General Public License
- along with Astaroth. If not, see .
-*/
-#pragma once
-
-/*
- * =============================================================================
- * Logical switches
- * =============================================================================
- */
-#define STENCIL_ORDER (6)
-#define NGHOST (STENCIL_ORDER / 2)
-#define LDENSITY (1)
-#define LHYDRO (1)
-#define LMAGNETIC (1)
-#define LENTROPY (1)
-#define LTEMPERATURE (0)
-#define LFORCING (1)
-#define LUPWD (1)
-
-#define AC_THERMAL_CONDUCTIVITY (AcReal(0.001)) // TODO: make an actual config parameter
-
-/*
- * =============================================================================
- * User-defined parameters
- * =============================================================================
- */
-// clang-format off
-#define AC_FOR_USER_INT_PARAM_TYPES(FUNC)\
- /* Other */\
- FUNC(AC_max_steps), \
- FUNC(AC_save_steps), \
- FUNC(AC_bin_steps), \
- FUNC(AC_bc_type),
-
-#define AC_FOR_USER_INT3_PARAM_TYPES(FUNC)
-
-#define AC_FOR_USER_REAL_PARAM_TYPES(FUNC)\
- /* cparams */\
- FUNC(AC_dsx), \
- FUNC(AC_dsy), \
- FUNC(AC_dsz), \
- FUNC(AC_dsmin), \
- /* physical grid*/\
- FUNC(AC_xlen), \
- FUNC(AC_ylen), \
- FUNC(AC_zlen), \
- FUNC(AC_xorig), \
- FUNC(AC_yorig), \
- FUNC(AC_zorig), \
- /*Physical units*/\
- FUNC(AC_unit_density),\
- FUNC(AC_unit_velocity),\
- FUNC(AC_unit_length),\
- /* properties of gravitating star*/\
- FUNC(AC_star_pos_x),\
- FUNC(AC_star_pos_y),\
- FUNC(AC_star_pos_z),\
- FUNC(AC_M_star),\
- /* Run params */\
- FUNC(AC_cdt), \
- FUNC(AC_cdtv), \
- FUNC(AC_cdts), \
- FUNC(AC_nu_visc), \
- FUNC(AC_cs_sound), \
- FUNC(AC_eta), \
- FUNC(AC_mu0), \
- FUNC(AC_cp_sound), \
- FUNC(AC_gamma), \
- FUNC(AC_cv_sound), \
- FUNC(AC_lnT0), \
- FUNC(AC_lnrho0), \
- FUNC(AC_zeta), \
- FUNC(AC_trans),\
- /* Other */\
- FUNC(AC_bin_save_t), \
- /* Initial condition params */\
- FUNC(AC_ampl_lnrho), \
- FUNC(AC_ampl_uu), \
- FUNC(AC_angl_uu), \
- FUNC(AC_lnrho_edge),\
- FUNC(AC_lnrho_out),\
- /* Forcing parameters. User configured. */\
- FUNC(AC_forcing_magnitude),\
- FUNC(AC_relhel), \
- FUNC(AC_kmin), \
- FUNC(AC_kmax), \
- /* Forcing parameters. Set by the generator. */\
- FUNC(AC_forcing_phase),\
- FUNC(AC_k_forcex),\
- FUNC(AC_k_forcey),\
- FUNC(AC_k_forcez),\
- FUNC(AC_kaver),\
- FUNC(AC_ff_hel_rex),\
- FUNC(AC_ff_hel_rey),\
- FUNC(AC_ff_hel_rez),\
- FUNC(AC_ff_hel_imx),\
- FUNC(AC_ff_hel_imy),\
- FUNC(AC_ff_hel_imz),\
- /* Additional helper params */\
- /* (deduced from other params do not set these directly!) */\
- FUNC(AC_G_CONST),\
- FUNC(AC_GM_star),\
- FUNC(AC_sq2GM_star),\
- FUNC(AC_cs2_sound), \
- FUNC(AC_inv_dsx), \
- FUNC(AC_inv_dsy), \
- FUNC(AC_inv_dsz),
-
-#define AC_FOR_USER_REAL3_PARAM_TYPES(FUNC)
-// clang-format on
-
-/*
- * =============================================================================
- * User-defined vertex buffers
- * =============================================================================
- */
-// clang-format off
-#if LENTROPY
-#define AC_FOR_VTXBUF_HANDLES(FUNC) \
- FUNC(VTXBUF_LNRHO), \
- FUNC(VTXBUF_UUX), \
- FUNC(VTXBUF_UUY), \
- FUNC(VTXBUF_UUZ), \
- FUNC(VTXBUF_AX), \
- FUNC(VTXBUF_AY), \
- FUNC(VTXBUF_AZ), \
- FUNC(VTXBUF_ENTROPY),
-#elif LMAGNETIC
-#define AC_FOR_VTXBUF_HANDLES(FUNC) \
- FUNC(VTXBUF_LNRHO), \
- FUNC(VTXBUF_UUX), \
- FUNC(VTXBUF_UUY), \
- FUNC(VTXBUF_UUZ), \
- FUNC(VTXBUF_AX), \
- FUNC(VTXBUF_AY), \
- FUNC(VTXBUF_AZ),
-#elif LHYDRO
-#define AC_FOR_VTXBUF_HANDLES(FUNC) \
- FUNC(VTXBUF_LNRHO), \
- FUNC(VTXBUF_UUX), \
- FUNC(VTXBUF_UUY), \
- FUNC(VTXBUF_UUZ),
-#else
-#define AC_FOR_VTXBUF_HANDLES(FUNC) \
- FUNC(VTXBUF_LNRHO),
-#endif
-// clang-format on
diff --git a/src/core/kernels/stencil_header.hh b/acc/mhd_solver/stencil_definition.sdh
similarity index 100%
rename from src/core/kernels/stencil_header.hh
rename to acc/mhd_solver/stencil_definition.sdh
diff --git a/acc/mhd_solver/stencil_process.sps b/acc/mhd_solver/stencil_process.sps
index a2219b6..420193e 100644
--- a/acc/mhd_solver/stencil_process.sps
+++ b/acc/mhd_solver/stencil_process.sps
@@ -1,4 +1,4 @@
-#include "stencil_header.hh"
+#include "stencil_definition.sdh"
Vector
diff --git a/acc/src/acc.l b/acc/src/acc.l
index b180ae8..76104d8 100644
--- a/acc/src/acc.l
+++ b/acc/src/acc.l
@@ -15,7 +15,7 @@ L [a-zA-Z_]
"void" { return VOID; } /* Rest of the types inherited from C */
"int" { return INT; }
"int3" { return INT3; }
-"ScalarField" { return SCALAR; }
+"ScalarField" { return SCALARFIELD; }
"VectorField" { return VECTOR; }
"Kernel" { return KERNEL; } /* Function specifiers */
diff --git a/acc/src/acc.y b/acc/src/acc.y
index 7138b1b..0bd1d19 100644
--- a/acc/src/acc.y
+++ b/acc/src/acc.y
@@ -16,7 +16,7 @@ int yyget_lineno();
%token CONSTANT IN OUT UNIFORM
%token IDENTIFIER NUMBER
%token RETURN
-%token SCALAR VECTOR MATRIX
+%token SCALAR VECTOR MATRIX SCALARFIELD
%token VOID INT INT3
%token IF ELSE FOR WHILE ELIF
%token LEQU LAND LOR LLEQU
@@ -209,6 +209,7 @@ type_specifier: VOID
| SCALAR { $$ = astnode_create(NODE_TYPE_SPECIFIER, NULL, NULL); $$->token = SCALAR; }
| VECTOR { $$ = astnode_create(NODE_TYPE_SPECIFIER, NULL, NULL); $$->token = VECTOR; }
| MATRIX { $$ = astnode_create(NODE_TYPE_SPECIFIER, NULL, NULL); $$->token = MATRIX; }
+ | SCALARFIELD { $$ = astnode_create(NODE_TYPE_SPECIFIER, NULL, NULL); $$->token = SCALARFIELD; }
;
identifier: IDENTIFIER { $$ = astnode_create(NODE_IDENTIFIER, NULL, NULL); astnode_set_buffer(yytext, $$); }
diff --git a/acc/src/code_generator.c b/acc/src/code_generator.c
index 01cf1d1..c31ad86 100644
--- a/acc/src/code_generator.c
+++ b/acc/src/code_generator.c
@@ -54,12 +54,13 @@ static const char* translation_table[TRANSLATION_TABLE_SIZE] = {
[WHILE] = "while",
[FOR] = "for",
// Type specifiers
- [VOID] = "void",
- [INT] = "int",
- [INT3] = "int3",
- [SCALAR] = "AcReal",
- [VECTOR] = "AcReal3",
- [MATRIX] = "AcMatrix",
+ [VOID] = "void",
+ [INT] = "int",
+ [INT3] = "int3",
+ [SCALAR] = "AcReal",
+ [VECTOR] = "AcReal3",
+ [MATRIX] = "AcMatrix",
+ [SCALARFIELD] = "AcReal",
// Type qualifiers
[KERNEL] = "template static "
"__global__", //__launch_bounds__(RK_THREADBLOCK_SIZE,
@@ -555,6 +556,68 @@ generate_preprocessed_structures(void)
");
}
+static void
+generate_header(void)
+{
+ printf("\n#pragma once\n");
+
+ // Int params
+ printf("#define AC_FOR_USER_INT_PARAM_TYPES(FUNC)");
+ for (int i = 0; i < num_symbols; ++i) {
+ if (symbol_table[i].type_specifier == INT) {
+ printf("\\\nFUNC(%s),", symbol_table[i].identifier);
+ }
+ }
+ printf("\n\n");
+
+ // Int3 params
+ printf("#define AC_FOR_USER_INT3_PARAM_TYPES(FUNC)");
+ for (int i = 0; i < num_symbols; ++i) {
+ if (symbol_table[i].type_specifier == INT3) {
+ printf("\\\nFUNC(%s),", symbol_table[i].identifier);
+ }
+ }
+ printf("\n\n");
+
+ // Scalar params
+ printf("#define AC_FOR_USER_REAL_PARAM_TYPES(FUNC)");
+ for (int i = 0; i < num_symbols; ++i) {
+ if (symbol_table[i].type_specifier == SCALAR) {
+ printf("\\\nFUNC(%s),", symbol_table[i].identifier);
+ }
+ }
+ printf("\n\n");
+
+ // Vector params
+ printf("#define AC_FOR_USER_REAL3_PARAM_TYPES(FUNC)");
+ for (int i = 0; i < num_symbols; ++i) {
+ if (symbol_table[i].type_specifier == VECTOR) {
+ printf("\\\nFUNC(%s),", symbol_table[i].identifier);
+ }
+ }
+ printf("\n\n");
+
+ // Scalar fields
+ printf("#define AC_FOR_VTXBUF_HANDLES(FUNC)");
+ for (int i = 0; i < num_symbols; ++i) {
+ if (symbol_table[i].type_specifier == SCALARFIELD) {
+ printf("\\\nFUNC(%s),", symbol_table[i].identifier);
+ }
+ }
+ printf("\n\n");
+
+ /*
+ printf("\n");
+ printf("typedef struct {\n");
+ for (int i = 0; i < num_symbols; ++i) {
+ if (symbol_table[i].type_qualifier == PREPROCESSED)
+ printf("%s %s;\n", translate(symbol_table[i].type_specifier),
+ symbol_table[i].identifier);
+ }
+ printf("} %sData;\n", translate(SCALAR));
+ */
+}
+
int
main(int argc, char** argv)
{
@@ -563,7 +626,7 @@ main(int argc, char** argv)
compilation_type = STENCIL_ASSEMBLY;
else if (!strcmp(argv[1], "-sps"))
compilation_type = STENCIL_PROCESS;
- else if (!strcmp(argv[1], "-hh"))
+ else if (!strcmp(argv[1], "-sdh"))
compilation_type = STENCIL_HEADER;
else
printf("Unknown flag %s. Generating stencil assembly.\n", argv[1]);
@@ -590,6 +653,8 @@ main(int argc, char** argv)
traverse(root);
if (compilation_type == STENCIL_ASSEMBLY)
generate_preprocessed_structures();
+ else if (compilation_type == STENCIL_HEADER)
+ generate_header();
// print_symbol_table();
diff --git a/config/astaroth.conf b/config/astaroth.conf
index 32f50a3..944364b 100644
--- a/config/astaroth.conf
+++ b/config/astaroth.conf
@@ -7,7 +7,7 @@
*/
AC_nx = 128
AC_ny = 128
-AC_nz = 128
+AC_nz = 6
AC_dsx = 0.04908738521
AC_dsy = 0.04908738521
diff --git a/scripts/compile_acc.sh b/scripts/compile_acc.sh
index 8648079..4649bf5 100755
--- a/scripts/compile_acc.sh
+++ b/scripts/compile_acc.sh
@@ -8,17 +8,18 @@ fi
KERNEL_DIR=${AC_HOME}"/src/core/kernels"
ACC_DIR=${AC_HOME}"/acc"
-ACC_DEFAULT_HEADER="mhd_solver/stencil_defines.h"
ACC_DEFAULT_SAS="mhd_solver/stencil_assembly.sas"
ACC_DEFAULT_SPS="mhd_solver/stencil_process.sps"
+ACC_DEFAULT_HEADER="mhd_solver/stencil_definition.sdh"
+ACC_DEFAULT_INCLUDE_DIR="mhd_solver"
${ACC_DIR}/clean.sh
${ACC_DIR}/build_acc.sh
-
-ACC_HEADER=${ACC_DEFAULT_HEADER}
ACC_SAS=${ACC_DEFAULT_SAS}
ACC_SPS=${ACC_DEFAULT_SPS}
+ACC_HEADER=${ACC_DEFAULT_HEADER}
+ACC_INCLUDE=${ACC_DEFAULT_INCLUDE_DIR}
while [ "$#" -gt 0 ]
do
@@ -56,9 +57,14 @@ echo "Header file:" ${ACC_DIR}/${ACC_HEADER}
echo "Assembly file: ${ACC_DIR}/${ACC_SAS}"
echo "Process file: ${ACC_DIR}/${ACC_SPS}"
-cd ${KERNEL_DIR}
-${ACC_DIR}/compile.sh ${ACC_DIR}/${ACC_SAS} ${ACC_DIR}/${ACC_HEADER}
-${ACC_DIR}/compile.sh ${ACC_DIR}/${ACC_SPS} ${ACC_DIR}/${ACC_HEADER}
+cd ${ACC_DIR}/${ACC_INCLUDE_DIR}
+${ACC_DIR}/compile.sh ${ACC_DIR}/${ACC_SAS}
+${ACC_DIR}/compile.sh ${ACC_DIR}/${ACC_SPS}
+${ACC_DIR}/compile.sh ${ACC_DIR}/${ACC_HEADER}
-echo "Linking: " ${ACC_DIR}/${ACC_HEADER} " -> " ${AC_HOME}/include/stencil_defines.h
-ln -sf ${ACC_DIR}/${ACC_HEADER} ${AC_HOME}/include/stencil_defines.h
+#mv ${ACC_SAS} ${AC_HOME}/src/core/kernels
+#mv ${ACC_SPS} ${AC_HOME}/src/core/kernels
+#mv ${ACC_HEADER} ${AC_HOME}/include
+
+#echo "Linking: " ${ACC_DIR}/${ACC_HEADER} " -> " ${AC_HOME}/include/stencil_defines.h
+#ln -sf ${ACC_DIR}/${ACC_HEADER} ${AC_HOME}/include/stencil_defines.h