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

Made globalVertexIdx available during preprocessing. NOTE: potentially dangerous. globalVertexIdx should never be used for reading data from the vertex buffers.

Browse Source
This commit is contained in:
jpekkila
2019-08-05 15:03:02 +03:00
parent 62100b1140
commit f3de2fa03c
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 // Preprocessed parameter boilerplate
if (node->type == NODE_TYPE_QUALIFIER && node->token == PREPROCESSED) if (node->type == NODE_TYPE_QUALIFIER && node->token == PREPROCESSED)
inside_preprocessed = true; 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) if (inside_preprocessed && node->type == NODE_FUNCTION_PARAMETER_DECLARATION)
printf("%s ", preprocessed_parameter_boilerplate); printf("%s ", preprocessed_parameter_boilerplate);
// BOILERPLATE END//////////////////////////////////////////////////////// // BOILERPLATE END////////////////////////////////////////////////////////
@@ -491,6 +492,7 @@ generate_preprocessed_structures(void)
// FILLING THE DATA STRUCT // FILLING THE DATA STRUCT
printf("static __device__ __forceinline__ AcRealData\ printf("static __device__ __forceinline__ AcRealData\
read_data(const int3 vertexIdx,\ read_data(const int3 vertexIdx,\
const int3 globalVertexIdx,\
AcReal* __restrict__ buf[], const int handle)\ AcReal* __restrict__ buf[], const int handle)\
{\n\ {\n\
%sData data;\n", %sData data;\n",
@@ -498,7 +500,7 @@ generate_preprocessed_structures(void)
for (int i = 0; i < num_symbols; ++i) { for (int i = 0; i < num_symbols; ++i) {
if (symbol_table[i].type_qualifier == PREPROCESSED) 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); symbol_table[i].identifier, symbol_table[i].identifier);
} }
printf("return data;\n"); printf("return data;\n");
@@ -526,13 +528,14 @@ generate_preprocessed_structures(void)
\ \
static __device__ __forceinline__ AcReal3Data\ static __device__ __forceinline__ AcReal3Data\
read_data(const int3 vertexIdx,\ read_data(const int3 vertexIdx,\
const int3 globalVertexIdx,\
AcReal* __restrict__ buf[], const int3& handle)\ AcReal* __restrict__ buf[], const int3& handle)\
{\ {\
AcReal3Data data;\ AcReal3Data data;\
\ \
data.x = read_data(vertexIdx, buf, handle.x);\ data.x = read_data(vertexIdx, globalVertexIdx, buf, handle.x);\
data.y = read_data(vertexIdx, buf, handle.y);\ data.y = read_data(vertexIdx, globalVertexIdx, buf, handle.y);\
data.z = read_data(vertexIdx, buf, handle.z);\ data.z = read_data(vertexIdx, globalVertexIdx, buf, handle.z);\
\ \
return data;\ 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 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)) #define READ_OUT(handle) (read_out(idx, buffer.out, handle))
// also write for clarity here also, not for the DSL // also write for clarity here also, not for the DSL