diff --git a/acc/src/code_generator.c b/acc/src/code_generator.c index 2810f5b..93ca42f 100644 --- a/acc/src/code_generator.c +++ b/acc/src/code_generator.c @@ -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;\ }\ diff --git a/src/core/kernels/kernels.cuh b/src/core/kernels/kernels.cuh index 52b48f6..c604205 100644 --- a/src/core/kernels/kernels.cuh +++ b/src/core/kernels/kernels.cuh @@ -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