From 51cf1f1068466d865e69228d2602a5099c6d6734 Mon Sep 17 00:00:00 2001 From: jpekkila Date: Mon, 19 Aug 2019 18:19:28 +0300 Subject: [PATCH] The C header is now generated from the DSL, stashing the changes just to be sure since I might overwrite something when updating the compilation scripts to work with this new scheme --- acc/compile.sh | 12 +- acc/mhd_solver/stencil_defines.h | 163 ------------------ .../mhd_solver/stencil_definition.sdh | 0 acc/mhd_solver/stencil_process.sps | 2 +- acc/src/acc.l | 2 +- acc/src/acc.y | 3 +- acc/src/code_generator.c | 79 ++++++++- config/astaroth.conf | 2 +- scripts/compile_acc.sh | 22 ++- 9 files changed, 99 insertions(+), 186 deletions(-) delete mode 100644 acc/mhd_solver/stencil_defines.h rename src/core/kernels/stencil_header.hh => acc/mhd_solver/stencil_definition.sdh (100%) 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