From 0e1d1b9fb4834b798879b8159a743c45816ed9eb Mon Sep 17 00:00:00 2001 From: jpekkila Date: Mon, 7 Oct 2019 16:33:24 +0300 Subject: [PATCH] Some optimizations for DSL compilation. Also a new feature: Inplace addition and subtraction += and -= are now allowed --- acc/src/acc.l | 2 ++ acc/src/acc.y | 4 +++- acc/src/code_generator.c | 6 +++++- acc/stdlib/stdderiv.h | 10 +++++----- src/core/kernels/integration.cuh | 8 ++++---- 5 files changed, 19 insertions(+), 11 deletions(-) diff --git a/acc/src/acc.l b/acc/src/acc.l index e04c045..92d51ad 100644 --- a/acc/src/acc.l +++ b/acc/src/acc.l @@ -51,6 +51,8 @@ L [a-zA-Z_] "++" { return INPLACE_INC; } "--" { return INPLACE_DEC; } +"+=" { return INPLACE_ADD; } +"-=" { return INPLACE_SUB; } [-+*/;=\[\]{}(),\.<>] { return yytext[0]; } /* Characters */ diff --git a/acc/src/acc.y b/acc/src/acc.y index 94abf7c..7c41f02 100644 --- a/acc/src/acc.y +++ b/acc/src/acc.y @@ -21,7 +21,7 @@ int yyget_lineno(); %token IF ELSE FOR WHILE ELIF %token LEQU LAND LOR LLEQU %token KERNEL DEVICE PREPROCESSED -%token INPLACE_INC INPLACE_DEC +%token INPLACE_INC INPLACE_DEC INPLACE_ADD INPLACE_SUB %% @@ -188,6 +188,8 @@ binary_operator: '+' | LAND { $$ = astnode_create(NODE_UNKNOWN, NULL, NULL); astnode_set_buffer(yytext, $$); } | LOR { $$ = astnode_create(NODE_UNKNOWN, NULL, NULL); astnode_set_buffer(yytext, $$); } | LLEQU { $$ = astnode_create(NODE_UNKNOWN, NULL, NULL); astnode_set_buffer(yytext, $$); } + | INPLACE_ADD { $$ = astnode_create(NODE_UNKNOWN, NULL, NULL); astnode_set_buffer(yytext, $$); } + | INPLACE_SUB { $$ = astnode_create(NODE_UNKNOWN, NULL, NULL); astnode_set_buffer(yytext, $$); } ; unary_operator: '-' /* C-style casts are disallowed, would otherwise be defined here */ { $$ = astnode_create(NODE_UNKNOWN, NULL, NULL); $$->infix = yytext[0]; } diff --git a/acc/src/code_generator.c b/acc/src/code_generator.c index 2b8d12a..5ba7ed5 100644 --- a/acc/src/code_generator.c +++ b/acc/src/code_generator.c @@ -66,7 +66,7 @@ static const char* translation_table[TRANSLATION_TABLE_SIZE] = { [COMPLEX] = "acComplex", // Type qualifiers [KERNEL] = "template static __global__", - [DEVICE] = "static __device__", + [DEVICE] = "static __device__ __forceinline__", [PREPROCESSED] = "static __device__ __forceinline__", [CONSTANT] = "const", [IN] = "in", @@ -138,6 +138,9 @@ static size_t current_nest = 0; static Symbol* symboltable_lookup(const char* identifier) { + // TODO assert tha symbol not function! cannot be since we allow overloads->conflicts if not + // explicit + if (!identifier) return NULL; @@ -379,6 +382,7 @@ traverse(const ASTNode* node) tmp = tmp->parent; assert(tmp->type = NODE_FUNCTION_DECLARATION); + // TODO FIX not to use symboltable_lookup const Symbol* parent_function = symboltable_lookup(tmp->lhs->rhs->buffer); assert(parent_function); diff --git a/acc/stdlib/stdderiv.h b/acc/stdlib/stdderiv.h index 1fbefd9..0de0cd3 100644 --- a/acc/stdlib/stdderiv.h +++ b/acc/stdlib/stdderiv.h @@ -27,7 +27,7 @@ first_derivative(Scalar pencil[], Scalar inv_ds) Scalar res = 0; for (int i = 1; i <= MID; ++i) { - res = res + coefficients[i] * (pencil[MID + i] - pencil[MID - i]); + res += coefficients[i] * (pencil[MID + i] - pencil[MID - i]); } return res * inv_ds; @@ -50,7 +50,7 @@ second_derivative(Scalar pencil[], Scalar inv_ds) Scalar res = coefficients[0] * pencil[MID]; for (int i = 1; i <= MID; ++i) { - res = res + coefficients[i] * (pencil[MID + i] + pencil[MID - i]); + res += coefficients[i] * (pencil[MID + i] + pencil[MID - i]); } return res * inv_ds * inv_ds; @@ -76,8 +76,8 @@ cross_derivative(Scalar pencil_a[], Scalar pencil_b[], Scalar inv_ds_a, Scalar i Scalar res = 0.0; for (int i = 1; i <= MID; ++i) { - res = res + coefficients[i] * (pencil_a[MID + i] + pencil_a[MID - i] - pencil_b[MID + i] - - pencil_b[MID - i]); + res += coefficients[i] * + (pencil_a[MID + i] + pencil_a[MID - i] - pencil_b[MID + i] - pencil_b[MID - i]); } return res * inv_ds_a * inv_ds_b; } @@ -311,7 +311,7 @@ contract(const Matrix mat) Scalar res = 0; for (int i = 0; i < 3; ++i) { - res = res + dot(mat.row[i], mat.row[i]); + res += dot(mat.row[i], mat.row[i]); } return res; diff --git a/src/core/kernels/integration.cuh b/src/core/kernels/integration.cuh index 35de310..5aad56e 100644 --- a/src/core/kernels/integration.cuh +++ b/src/core/kernels/integration.cuh @@ -29,7 +29,7 @@ #include -static __device__ __forceinline__ int +static __device__ constexpr int IDX(const int i) { return i; @@ -95,7 +95,7 @@ write(AcReal* __restrict__ out[], const int handle, const int idx, const AcReal out[handle][idx] = value; } -static __device__ void +static __device__ __forceinline__ void write(AcReal* __restrict__ out[], const int3 vec, const int idx, const AcReal3 value) { write(out, vec.x, idx, value.x); @@ -103,13 +103,13 @@ write(AcReal* __restrict__ out[], const int3 vec, const int idx, const AcReal3 v write(out, vec.z, idx, value.z); } -static __device__ AcReal +static __device__ __forceinline__ AcReal read_out(const int idx, AcReal* __restrict__ field[], const int handle) { return field[handle][idx]; } -static __device__ AcReal3 +static __device__ __forceinline__ AcReal3 read_out(const int idx, AcReal* __restrict__ field[], const int3 handle) { return (AcReal3){read_out(idx, field, handle.x), read_out(idx, field, handle.y),