cwpearson/astaroth
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge branch 'master' into cmakelist_rewrite_and_C_API_conformity_07-26

Browse Source
This commit is contained in:
jpekkila
2019-08-06 15:14:33 +03:00
parent e4b981fc62 f3de2fa03c
commit 280804a438
2 changed files with 9 additions and 6 deletions
Download Patch File Download Diff File Expand all files Collapse all files

13
acc/src/code_generator.c
View File

@@ -335,7 +335,8 @@ traverse(const ASTNode* node)
// Preprocessed parameter boilerplate
if (node->type == NODE_TYPE_QUALIFIER && node->token == PREPROCESSED)
inside_preprocessed = true;
static const char preprocessed_parameter_boilerplate[] = "const int3 vertexIdx, ";
static const char
preprocessed_parameter_boilerplate[] = "const int3 vertexIdx, const int3 globalVertexIdx, ";
if (inside_preprocessed && node->type == NODE_FUNCTION_PARAMETER_DECLARATION)
printf("%s ", preprocessed_parameter_boilerplate);
// BOILERPLATE END////////////////////////////////////////////////////////
@@ -491,6 +492,7 @@ generate_preprocessed_structures(void)
// FILLING THE DATA STRUCT
printf("static __device__ __forceinline__ AcRealData\
read_data(const int3 vertexIdx,\
const int3 globalVertexIdx,\
AcReal* __restrict__ buf[], const int handle)\
{\n\
%sData data;\n",
@@ -498,7 +500,7 @@ generate_preprocessed_structures(void)
for (int i = 0; i < num_symbols; ++i) {
if (symbol_table[i].type_qualifier == PREPROCESSED)
printf("data.%s = preprocessed_%s(vertexIdx, buf[handle]);\n",
printf("data.%s = preprocessed_%s(vertexIdx, globalVertexIdx, buf[handle]);\n",
symbol_table[i].identifier, symbol_table[i].identifier);
}
printf("return data;\n");
@@ -526,13 +528,14 @@ generate_preprocessed_structures(void)
\
static __device__ __forceinline__ AcReal3Data\
read_data(const int3 vertexIdx,\
const int3 globalVertexIdx,\
AcReal* __restrict__ buf[], const int3& handle)\
{\
AcReal3Data data;\
\
data.x = read_data(vertexIdx, buf, handle.x);\
data.y = read_data(vertexIdx, buf, handle.y);\
data.z = read_data(vertexIdx, buf, handle.z);\
data.x = read_data(vertexIdx, globalVertexIdx, buf, handle.x);\
data.y = read_data(vertexIdx, globalVertexIdx, buf, handle.y);\
data.z = read_data(vertexIdx, globalVertexIdx, buf, handle.z);\
\
return data;\
}\

2
src/core/kernels/kernels.cuh
View File

@@ -671,7 +671,7 @@ read_out(const int idx, AcReal* __restrict__ field[], const int3 handle)
}
#define WRITE_OUT(handle, value) (write(buffer.out, handle, idx, value))
#define READ(handle) (read_data(vertexIdx, buffer.in, handle))
#define READ(handle) (read_data(vertexIdx, globalVertexIdx, buffer.in, handle))
#define READ_OUT(handle) (read_out(idx, buffer.out, handle))
// also write for clarity here also, not for the DSL