From 095f8630976c5429db0ac345a52ffe4712ab2be2 Mon Sep 17 00:00:00 2001 From: jpekkila Date: Tue, 24 Nov 2020 21:29:14 +0200 Subject: [PATCH 1/3] Added functions acSetVertexBuffer, acNodeSetVertexBuffer, and acDeviceSetVertexBuffer for setting device memory directly to some constant --- include/astaroth.h | 17 +++++++++++++---- src/core/astaroth.cc | 14 ++++++++++---- src/core/device.cc | 30 +++++++++++++++++++++++++----- src/core/node.cc | 12 ++++++++++++ 4 files changed, 60 insertions(+), 13 deletions(-) diff --git a/include/astaroth.h b/include/astaroth.h index 07b4d61..4a6fd03 100644 --- a/include/astaroth.h +++ b/include/astaroth.h @@ -225,6 +225,9 @@ AcResult acLoadDeviceConstant(const AcRealParam param, const AcReal value); /** Loads an AcMesh to the devices visible to the caller */ AcResult acLoad(const AcMesh host_mesh); +/** Sets the whole mesh to some value */ +AcResult acSetVertexBuffer(const VertexBufferHandle handle, const AcReal value); + /** Stores the AcMesh distributed among the devices visible to the caller back to the host*/ AcResult acStore(AcMesh* host_mesh); @@ -432,6 +435,9 @@ AcResult acNodeLoadVertexBuffer(const Node node, const Stream stream, const AcMe /** */ AcResult acNodeLoadMesh(const Node node, const Stream stream, const AcMesh host_mesh); +/** */ +AcResult acNodeSetVertexBuffer(const Node node, const Stream stream, const VertexBufferHandle handle, const AcReal value); + /** Deprecated ? */ AcResult acNodeStoreVertexBufferWithOffset(const Node node, const Stream stream, const VertexBufferHandle vtxbuf_handle, const int3 src, @@ -554,6 +560,9 @@ AcResult acDeviceLoadVertexBuffer(const Device device, const Stream stream, cons /** */ AcResult acDeviceLoadMesh(const Device device, const Stream stream, const AcMesh host_mesh); +/** */ +AcResult acDeviceSetVertexBuffer(const Device device, const Stream stream, const VertexBufferHandle handle, const AcReal value); + /** */ AcResult acDeviceStoreVertexBufferWithOffset(const Device device, const Stream stream, const VertexBufferHandle vtxbuf_handle, const int3 src, @@ -631,16 +640,16 @@ AcResult acDeviceRunMPITest(void); * ============================================================================= */ /** Updates the built-in parameters based on nx, ny and nz */ -AcResult acUpdateBuiltinParams(AcMeshInfo* config); +AcResult acHostUpdateBuiltinParams(AcMeshInfo* config); /** Creates a mesh stored in host memory */ -AcResult acMeshCreate(const AcMeshInfo mesh_info, AcMesh* mesh); +AcResult acHostMeshCreate(const AcMeshInfo mesh_info, AcMesh* mesh); /** Randomizes a host mesh */ -AcResult acMeshRandomize(AcMesh* mesh); +AcResult acHostMeshRandomize(AcMesh* mesh); /** Destroys a mesh stored in host memory */ -AcResult acMeshDestroy(AcMesh* mesh); +AcResult acHostMeshDestroy(AcMesh* mesh); #ifdef __cplusplus } // extern "C" diff --git a/src/core/astaroth.cc b/src/core/astaroth.cc index 2123224..f5e97ce 100644 --- a/src/core/astaroth.cc +++ b/src/core/astaroth.cc @@ -74,6 +74,12 @@ acLoad(const AcMesh host_mesh) return acNodeLoadMesh(nodes[0], STREAM_DEFAULT, host_mesh); } +AcResult +acSetVertexBuffer(const VertexBufferHandle handle, const AcReal value) +{ + return acNodeSetVertexBuffer(nodes[0], STREAM_DEFAULT, handle, value); +} + AcResult acStore(AcMesh* host_mesh) { @@ -182,7 +188,7 @@ acGetNode(void) } AcResult -acUpdateBuiltinParams(AcMeshInfo* config) +acHostUpdateBuiltinParams(AcMeshInfo* config) { config->int_params[AC_mx] = config->int_params[AC_nx] + STENCIL_ORDER; ///////////// PAD TEST @@ -221,7 +227,7 @@ acUpdateBuiltinParams(AcMeshInfo* config) } AcResult -acMeshCreate(const AcMeshInfo info, AcMesh* mesh) +acHostMeshCreate(const AcMeshInfo info, AcMesh* mesh) { mesh->info = info; @@ -241,7 +247,7 @@ randf(void) } AcResult -acMeshRandomize(AcMesh* mesh) +acHostMeshRandomize(AcMesh* mesh) { const int n = acVertexBufferSize(mesh->info); for (int w = 0; w < NUM_VTXBUF_HANDLES; ++w) @@ -252,7 +258,7 @@ acMeshRandomize(AcMesh* mesh) } AcResult -acMeshDestroy(AcMesh* mesh) +acHostMeshDestroy(AcMesh* mesh) { for (int w = 0; w < NUM_VTXBUF_HANDLES; ++w) free(mesh->vertex_buffer[w]); diff --git a/src/core/device.cc b/src/core/device.cc index c1bb43f..123dee2 100644 --- a/src/core/device.cc +++ b/src/core/device.cc @@ -294,6 +294,26 @@ acDeviceLoadMesh(const Device device, const Stream stream, const AcMesh host_mes return AC_SUCCESS; } +AcResult +acDeviceSetVertexBuffer(const Device device, const Stream stream, const VertexBufferHandle handle, const AcReal value) +{ + acDeviceSynchronizeStream(device, stream); + + const size_t count = acVertexBufferSize(device->local_config); + AcReal* data = (AcReal*) malloc(sizeof(AcReal) * count); + ERRCHK_ALWAYS(data); + + for (size_t i = 0; i < count; ++i) + data[i] = value; + + // Set both in and out for safety (not strictly needed) + ERRCHK_CUDA_ALWAYS(cudaMemcpyAsync(device->vba.in[handle], data, sizeof(data[0]) * count, cudaMemcpyHostToDevice, device->streams[stream])); + ERRCHK_CUDA_ALWAYS(cudaMemcpyAsync(device->vba.out[handle], data, sizeof(data[0]) * count, cudaMemcpyHostToDevice, device->streams[stream])); + + free(data); + return AC_SUCCESS; +} + AcResult acDeviceStoreVertexBufferWithOffset(const Device device, const Stream stream, const VertexBufferHandle vtxbuf_handle, const int3 src, @@ -1269,10 +1289,10 @@ acGridRandomize(void) ERRCHK(grid.initialized); AcMesh host; - acMeshCreate(grid.submesh.info, &host); + acHostMeshCreate(grid.submesh.info, &host); acMeshRandomize(&host); acDeviceLoadMesh(grid.device, STREAM_DEFAULT, host); - acMeshDestroy(&host); + acHostMeshDestroy(&host); return AC_SUCCESS; } @@ -1320,7 +1340,7 @@ acGridInit(const AcMeshInfo info) }; submesh_info.int3_params[AC_multigpu_offset] = pid3d * (int3){submesh_nx, submesh_ny, submesh_nz}; - acUpdateBuiltinParams(&submesh_info); + acHostUpdateBuiltinParams(&submesh_info); // GPU alloc int devices_per_node = -1; @@ -1331,7 +1351,7 @@ acGridInit(const AcMeshInfo info) // CPU alloc AcMesh submesh; - acMeshCreate(submesh_info, &submesh); + acHostMeshCreate(submesh_info, &submesh); // Setup the global grid structure grid.device = device; @@ -1380,7 +1400,7 @@ acGridQuit(void) grid.initialized = false; grid.decomposition = (uint3_64){0, 0, 0}; - acMeshDestroy(&grid.submesh); + acHostMeshDestroy(&grid.submesh); acDeviceDestroy(grid.device); acGridSynchronizeStream(STREAM_ALL); diff --git a/src/core/node.cc b/src/core/node.cc index fdc4f49..14542e1 100644 --- a/src/core/node.cc +++ b/src/core/node.cc @@ -536,6 +536,18 @@ acNodeLoadMesh(const Node node, const Stream stream, const AcMesh host_mesh) return AC_SUCCESS; } +AcResult +acNodeSetVertexBuffer(const Node node, const Stream stream, const VertexBufferHandle handle, const AcReal value) +{ + acNodeSynchronizeStream(node, stream); + + for (int i = 0; i < node->num_devices; ++i) + acDeviceSetVertexBuffer(node->devices[i], stream, handle, value); + + acNodeSynchronizeStream(node, stream); // For safety + return AC_SUCCESS; +} + AcResult acNodeStoreVertexBufferWithOffset(const Node node, const Stream stream, const VertexBufferHandle vtxbuf_handle, const int3 src, From bcacc357d35fb32e56f6727a88e77b9b7fc9871a Mon Sep 17 00:00:00 2001 From: jpekkila Date: Tue, 24 Nov 2020 21:32:43 +0200 Subject: [PATCH 2/3] Now all host functions start with acHost to avoid confusion on whether the function operates on host or device memory --- acc/src/code_generator.c | 2 +- include/astaroth_utils.h | 14 +++++++------- samples/benchmark/main.cc | 14 +++++++------- samples/cpptest/main.cc | 12 ++++++------ samples/ctest/main.c | 13 +++++++------ samples/fortrantest/main.f90 | 2 +- samples/mpitest/main.cc | 22 +++++++++++----------- samples/standalone_mpi/host_memory.cc | 18 +++++++++--------- samples/standalone_mpi/main.cc | 24 ++++++++++++------------ src/core/astaroth_fortran.cc | 4 ++-- src/core/astaroth_fortran.h | 2 +- src/utils/config_loader.c | 2 +- src/utils/memory.c | 12 ++++++------ src/utils/modelreduce.c | 4 ++-- src/utils/modelsolver.c | 10 +++++----- 15 files changed, 78 insertions(+), 77 deletions(-) diff --git a/acc/src/code_generator.c b/acc/src/code_generator.c index b1de527..7b6ae2c 100644 --- a/acc/src/code_generator.c +++ b/acc/src/code_generator.c @@ -606,7 +606,7 @@ generate_headers(void) ! -*-f90-*- (for emacs) vim:set filetype=fortran: (for vim) ! Utils (see astaroth_fortran.cc for definitions) -external acupdatebuiltinparams +external achostupdatebuiltinparams external acgetdevicecount ! Device interface (see astaroth_fortran.cc for definitions) diff --git a/include/astaroth_utils.h b/include/astaroth_utils.h index 4289c1b..b456b49 100644 --- a/include/astaroth_utils.h +++ b/include/astaroth_utils.h @@ -45,25 +45,25 @@ typedef struct { AcResult acLoadConfig(const char* config_path, AcMeshInfo* config); /** */ -AcResult acVertexBufferSet(const VertexBufferHandle handle, const AcReal value, AcMesh* mesh); +AcResult acHostVertexBufferSet(const VertexBufferHandle handle, const AcReal value, AcMesh* mesh); /** */ -AcResult acMeshSet(const AcReal value, AcMesh* mesh); +AcResult acHostMeshSet(const AcReal value, AcMesh* mesh); /** */ -AcResult acMeshApplyPeriodicBounds(AcMesh* mesh); +AcResult acHostMeshApplyPeriodicBounds(AcMesh* mesh); /** */ -AcResult acMeshClear(AcMesh* mesh); +AcResult acHostMeshClear(AcMesh* mesh); /** */ -AcResult acModelIntegrateStep(AcMesh mesh, const AcReal dt); +AcResult acHostIntegrateStep(AcMesh mesh, const AcReal dt); /** */ -AcReal acModelReduceScal(const AcMesh mesh, const ReductionType rtype, const VertexBufferHandle a); +AcReal acHostReduceScal(const AcMesh mesh, const ReductionType rtype, const VertexBufferHandle a); /** */ -AcReal acModelReduceVec(const AcMesh mesh, const ReductionType rtype, const VertexBufferHandle a, +AcReal acHostReduceVec(const AcMesh mesh, const ReductionType rtype, const VertexBufferHandle a, const VertexBufferHandle b, const VertexBufferHandle c); Error acGetError(const AcReal model, const AcReal candidate); diff --git a/samples/benchmark/main.cc b/samples/benchmark/main.cc index 76223bb..861d2bd 100644 --- a/samples/benchmark/main.cc +++ b/samples/benchmark/main.cc @@ -98,7 +98,7 @@ main(int argc, char** argv) info.int_params[AC_nx] = nx; info.int_params[AC_ny] = ny; info.int_params[AC_nz] = nz; - acUpdateBuiltinParams(&info); + acHostUpdateBuiltinParams(&info); printf("Benchmark mesh dimensions: (%d, %d, %d)\n", nx, ny, nz); } else { @@ -118,8 +118,8 @@ main(int argc, char** argv) /* AcMesh model, candidate; if (pid == 0) { - acMeshCreate(info, &model); - acMeshCreate(info, &candidate); + acHostMeshCreate(info, &model); + acHostMeshCreate(info, &candidate); acMeshRandomize(&model); acMeshRandomize(&candidate); }*/ @@ -130,7 +130,7 @@ main(int argc, char** argv) /* AcMesh model; - acMeshCreate(info, &model); + acHostMeshCreate(info, &model); acMeshRandomize(&model); acGridLoadMesh(STREAM_DEFAULT, model); */ @@ -145,12 +145,12 @@ main(int argc, char** argv) // Verify if (pid == 0) { - acModelIntegrateStep(model, FLT_EPSILON); + acHostIntegrateStep(model, FLT_EPSILON); acMeshApplyPeriodicBounds(&model); AcResult retval = acVerifyMesh(model, candidate); - acMeshDestroy(&model); - acMeshDestroy(&candidate); + acHostMeshDestroy(&model); + acHostMeshDestroy(&candidate); if (retval != AC_SUCCESS) { fprintf(stderr, "Failures found, benchmark invalid. Skipping\n"); diff --git a/samples/cpptest/main.cc b/samples/cpptest/main.cc index 3dde831..99f37e6 100644 --- a/samples/cpptest/main.cc +++ b/samples/cpptest/main.cc @@ -31,12 +31,12 @@ main(void) // Alloc AcMesh model, candidate; - acMeshCreate(info, &model); - acMeshCreate(info, &candidate); + acHostMeshCreate(info, &model); + acHostMeshCreate(info, &candidate); // Init - acMeshRandomize(&model); - acMeshApplyPeriodicBounds(&model); + acHostMeshRandomize(&model); + acHostMeshApplyPeriodicBounds(&model); // Verify that the mesh was loaded and stored correctly acInit(info); @@ -55,8 +55,8 @@ main(void) // Destroy acQuit(); - acMeshDestroy(&model); - acMeshDestroy(&candidate); + acHostMeshDestroy(&model); + acHostMeshDestroy(&candidate); puts("cpptest complete."); return EXIT_SUCCESS; diff --git a/samples/ctest/main.c b/samples/ctest/main.c index b340154..751798b 100644 --- a/samples/ctest/main.c +++ b/samples/ctest/main.c @@ -30,12 +30,12 @@ main(void) // Alloc AcMesh model, candidate; - acMeshCreate(info, &model); - acMeshCreate(info, &candidate); + acHostMeshCreate(info, &model); + acHostMeshCreate(info, &candidate); // Init - acMeshRandomize(&model); - acMeshApplyPeriodicBounds(&model); + acHostMeshRandomize(&model); + acHostMeshApplyPeriodicBounds(&model); // Verify that the mesh was loaded and stored correctly acInit(info); @@ -46,6 +46,7 @@ main(void) // Attempt to integrate and check max and min printf("Integrating... "); acIntegrate(FLT_EPSILON); + printf("Done.\nVTXBUF ranges after one integration step:\n"); for (size_t i = 0; i < NUM_VTXBUF_HANDLES; ++i) printf("\t%-15s... [%.3g, %.3g]\n", vtxbuf_names[i], // @@ -54,8 +55,8 @@ main(void) // Destroy acQuit(); - acMeshDestroy(&model); - acMeshDestroy(&candidate); + acHostMeshDestroy(&model); + acHostMeshDestroy(&candidate); puts("ctest complete."); return EXIT_SUCCESS; diff --git a/samples/fortrantest/main.f90 b/samples/fortrantest/main.f90 index 1eba03f..7e09a49 100644 --- a/samples/fortrantest/main.f90 +++ b/samples/fortrantest/main.f90 @@ -14,7 +14,7 @@ program pc info%int_params(AC_nx + 1) = 128 info%int_params(AC_ny + 1) = 128 info%int_params(AC_nz + 1) = 128 - call acupdatebuiltinparams(info) + call achostupdatebuiltinparams(info) call acdevicecreate(0, info, device) call acdeviceprintinfo(device) diff --git a/samples/mpitest/main.cc b/samples/mpitest/main.cc index f3ad473..e690352 100644 --- a/samples/mpitest/main.cc +++ b/samples/mpitest/main.cc @@ -47,8 +47,8 @@ main(void) AcMesh model, candidate; if (pid == 0) { - acMeshCreate(info, &model); - acMeshCreate(info, &candidate); + acHostMeshCreate(info, &model); + acHostMeshCreate(info, &candidate); acMeshRandomize(&model); acMeshRandomize(&candidate); } @@ -73,7 +73,7 @@ main(void) acGridPeriodicBoundconds(STREAM_DEFAULT); acGridStoreMesh(STREAM_DEFAULT, &candidate); if (pid == 0) { - acModelIntegrateStep(model, FLT_EPSILON); + acHostIntegrateStep(model, FLT_EPSILON); acMeshApplyPeriodicBounds(&model); const AcResult res = acVerifyMesh("Integration", model, candidate); ERRCHK_ALWAYS(res == AC_SUCCESS); @@ -93,10 +93,10 @@ main(void) AcReal candval; acGridReduceScal(STREAM_DEFAULT, (ReductionType)i, v0, &candval); if (pid == 0) { - const AcReal modelval = acModelReduceScal(model, (ReductionType)i, v0); + const AcReal modelval = acHostReduceScal(model, (ReductionType)i, v0); Error error = acGetError(modelval, candval); - error.maximum_magnitude = acModelReduceScal(model, RTYPE_MAX, v0); - error.minimum_magnitude = acModelReduceScal(model, RTYPE_MIN, v0); + error.maximum_magnitude = acHostReduceScal(model, RTYPE_MAX, v0); + error.minimum_magnitude = acHostReduceScal(model, RTYPE_MIN, v0); ERRCHK_ALWAYS(acEvalError(rtype_names[i], error)); } } @@ -114,17 +114,17 @@ main(void) AcReal candval; acGridReduceVec(STREAM_DEFAULT, (ReductionType)i, v0, v1, v2, &candval); if (pid == 0) { - const AcReal modelval = acModelReduceVec(model, (ReductionType)i, v0, v1, v2); + const AcReal modelval = acHostReduceVec(model, (ReductionType)i, v0, v1, v2); Error error = acGetError(modelval, candval); - error.maximum_magnitude = acModelReduceVec(model, RTYPE_MAX, v0, v1, v2); - error.minimum_magnitude = acModelReduceVec(model, RTYPE_MIN, v0, v1, v1); + error.maximum_magnitude = acHostReduceVec(model, RTYPE_MAX, v0, v1, v2); + error.minimum_magnitude = acHostReduceVec(model, RTYPE_MIN, v0, v1, v1); ERRCHK_ALWAYS(acEvalError(rtype_names[i], error)); } } if (pid == 0) { - acMeshDestroy(&model); - acMeshDestroy(&candidate); + acHostMeshDestroy(&model); + acHostMeshDestroy(&candidate); } acGridQuit(); diff --git a/samples/standalone_mpi/host_memory.cc b/samples/standalone_mpi/host_memory.cc index 1e809df..a2d20d5 100644 --- a/samples/standalone_mpi/host_memory.cc +++ b/samples/standalone_mpi/host_memory.cc @@ -576,7 +576,7 @@ acmesh_init_to(const InitType& init_type, AcMesh* mesh) switch (init_type) { case INIT_TYPE_RANDOM: { - acMeshClear(mesh); + acHostMeshClear(mesh); const AcReal range = AcReal(0.01); for (int w = 0; w < NUM_VTXBUF_HANDLES; ++w) for (int i = 0; i < n; ++i) @@ -585,14 +585,14 @@ acmesh_init_to(const InitType& init_type, AcMesh* mesh) break; } case INIT_TYPE_GAUSSIAN_RADIAL_EXPL: - acMeshClear(mesh); - acVertexBufferSet(VTXBUF_LNRHO, mesh->info.real_params[AC_ampl_lnrho], mesh); + acHostMeshClear(mesh); + acHostVertexBufferSet(VTXBUF_LNRHO, mesh->info.real_params[AC_ampl_lnrho], mesh); // acmesh_init_to(INIT_TYPE_RANDOM, mesh); gaussian_radial_explosion(mesh); break; case INIT_TYPE_XWAVE: - acMeshClear(mesh); + acHostMeshClear(mesh); acmesh_init_to(INIT_TYPE_RANDOM, mesh); for (int k = 0; k < mz; k++) { for (int j = 0; j < my; j++) { @@ -605,24 +605,24 @@ acmesh_init_to(const InitType& init_type, AcMesh* mesh) } break; case INIT_TYPE_SIMPLE_CORE: - acMeshClear(mesh); + acHostMeshClear(mesh); simple_uniform_core(mesh); break; case INIT_TYPE_VEDGE: - acMeshClear(mesh); + acHostMeshClear(mesh); inflow_vedge_freefall(mesh); break; case INIT_TYPE_VEDGEX: - acMeshClear(mesh); + acHostMeshClear(mesh); inflow_freefall_x(mesh); break; case INIT_TYPE_RAYLEIGH_TAYLOR: - acMeshClear(mesh); + acHostMeshClear(mesh); inflow_freefall_x(mesh); lnrho_step(mesh); break; case INIT_TYPE_ABC_FLOW: { - acMeshClear(mesh); + acHostMeshClear(mesh); acmesh_init_to(INIT_TYPE_RANDOM, mesh); for (int k = nz_min; k < nz_max; k++) { for (int j = ny_min; j < ny_max; j++) { diff --git a/samples/standalone_mpi/main.cc b/samples/standalone_mpi/main.cc index c5bb51a..23080af 100644 --- a/samples/standalone_mpi/main.cc +++ b/samples/standalone_mpi/main.cc @@ -199,9 +199,9 @@ print_diagnostics_host(const AcMesh mesh, const int step, const AcReal dt, const const int max_name_width = 16; // Calculate rms, min and max from the velocity vector field - buf_max = acModelReduceVec(mesh, RTYPE_MAX, VTXBUF_UUX, VTXBUF_UUY, VTXBUF_UUZ); - buf_min = acModelReduceVec(mesh, RTYPE_MIN, VTXBUF_UUX, VTXBUF_UUY, VTXBUF_UUZ); - buf_rms = acModelReduceVec(mesh, RTYPE_RMS, VTXBUF_UUX, VTXBUF_UUY, VTXBUF_UUZ); + buf_max = acHostReduceVec(mesh, RTYPE_MAX, VTXBUF_UUX, VTXBUF_UUY, VTXBUF_UUZ); + buf_min = acHostReduceVec(mesh, RTYPE_MIN, VTXBUF_UUX, VTXBUF_UUY, VTXBUF_UUZ); + buf_rms = acHostReduceVec(mesh, RTYPE_RMS, VTXBUF_UUX, VTXBUF_UUY, VTXBUF_UUZ); // MV: The ordering in the earlier version was wrong in terms of variable // MV: name and its diagnostics. @@ -213,9 +213,9 @@ print_diagnostics_host(const AcMesh mesh, const int step, const AcReal dt, const // Calculate rms, min and max from the variables as scalars for (int i = 0; i < NUM_VTXBUF_HANDLES; ++i) { - buf_max = acModelReduceScal(mesh, RTYPE_MAX, VertexBufferHandle(i)); - buf_min = acModelReduceScal(mesh, RTYPE_MIN, VertexBufferHandle(i)); - buf_rms = acModelReduceScal(mesh, RTYPE_RMS, VertexBufferHandle(i)); + buf_max = acHostReduceScal(mesh, RTYPE_MAX, VertexBufferHandle(i)); + buf_min = acHostReduceScal(mesh, RTYPE_MIN, VertexBufferHandle(i)); + buf_rms = acHostReduceScal(mesh, RTYPE_RMS, VertexBufferHandle(i)); printf(" %*s: min %.3e,\trms %.3e,\tmax %.3e\n", max_name_width, vtxbuf_names[i], double(buf_min), double(buf_rms), double(buf_max)); @@ -330,7 +330,7 @@ main(int argc, char** argv) AcMesh mesh; if (pid == 0) { - acMeshCreate(info, &mesh); + acHostMeshCreate(info, &mesh); acmesh_init_to(INIT_TYPE_GAUSSIAN_RADIAL_EXPL, &mesh); } acGridInit(info); @@ -362,7 +362,7 @@ main(int argc, char** argv) } } if (pid == 0) - acMeshDestroy(&mesh); + acHostMeshDestroy(&mesh); acGridQuit(); /////////////// Simple example END @@ -376,7 +376,7 @@ main(int argc, char** argv) if (argc == 3 && (!strcmp(argv[1], "-c") || !strcmp(argv[1], "--config"))) { acLoadConfig(argv[2], &info); load_config(argv[2], &info); - acUpdateBuiltinParams(&info); + acHostUpdateBuiltinParams(&info); } else { printf("Usage: ./ac_run\n"); @@ -388,7 +388,7 @@ main(int argc, char** argv) else { acLoadConfig(AC_DEFAULT_CONFIG, &info); load_config(AC_DEFAULT_CONFIG, &info); - acUpdateBuiltinParams(&info); + acHostUpdateBuiltinParams(&info); } const int start_step = info.int_params[AC_start_step]; @@ -406,7 +406,7 @@ main(int argc, char** argv) AcMesh mesh; ///////////////////////////////// PROC 0 BLOCK START /////////////////////////////////////////// if (pid == 0) { - acMeshCreate(info, &mesh); + acHostMeshCreate(info, &mesh); // TODO: This need to be possible to define in astaroth.conf acmesh_init_to(INIT_TYPE_GAUSSIAN_RADIAL_EXPL, &mesh); // acmesh_init_to(INIT_TYPE_SIMPLE_CORE, mesh); //Initial condition for a collapse test @@ -564,7 +564,7 @@ main(int argc, char** argv) acGridQuit(); if (pid == 0) - acMeshDestroy(&mesh); + acHostMeshDestroy(&mesh); fclose(diag_file); #endif diff --git a/src/core/astaroth_fortran.cc b/src/core/astaroth_fortran.cc index 18aa749..03f3749 100644 --- a/src/core/astaroth_fortran.cc +++ b/src/core/astaroth_fortran.cc @@ -7,9 +7,9 @@ * Utils */ void -acupdatebuiltinparams_(AcMeshInfo* info) +achostupdatebuiltinparams_(AcMeshInfo* info) { - acUpdateBuiltinParams(info); + acHostUpdateBuiltinParams(info); } void diff --git a/src/core/astaroth_fortran.h b/src/core/astaroth_fortran.h index fa0ba3a..9e760c2 100644 --- a/src/core/astaroth_fortran.h +++ b/src/core/astaroth_fortran.h @@ -8,7 +8,7 @@ extern "C" { /** * Utils */ -void acupdatebuiltinparams_(AcMeshInfo* info); +void achostupdatebuiltinparams_(AcMeshInfo* info); void acgetdevicecount_(int* count); diff --git a/src/utils/config_loader.c b/src/utils/config_loader.c index 5799157..305e35a 100644 --- a/src/utils/config_loader.c +++ b/src/utils/config_loader.c @@ -89,7 +89,7 @@ acLoadConfig(const char* config_path, AcMeshInfo* config) memset(config, (uint8_t)0xFF, sizeof(*config)); parse_config(config_path, config); - acUpdateBuiltinParams(config); + acHostUpdateBuiltinParams(config); #if AC_VERBOSE printf("###############################################################\n"); printf("Config dimensions loaded:\n"); diff --git a/src/utils/memory.c b/src/utils/memory.c index f52fa12..de36fac 100644 --- a/src/utils/memory.c +++ b/src/utils/memory.c @@ -21,7 +21,7 @@ #include "errchk.h" AcResult -acVertexBufferSet(const VertexBufferHandle handle, const AcReal value, AcMesh* mesh) +acHostVertexBufferSet(const VertexBufferHandle handle, const AcReal value, AcMesh* mesh) { const int n = acVertexBufferSize(mesh->info); for (int i = 0; i < n; ++i) @@ -30,16 +30,16 @@ acVertexBufferSet(const VertexBufferHandle handle, const AcReal value, AcMesh* m return AC_SUCCESS; } AcResult -acMeshSet(const AcReal value, AcMesh* mesh) +acHostMeshSet(const AcReal value, AcMesh* mesh) { for (int w = 0; w < NUM_VTXBUF_HANDLES; ++w) - acVertexBufferSet(w, value, mesh); + acHostVertexBufferSet(w, value, mesh); return AC_SUCCESS; } AcResult -acMeshApplyPeriodicBounds(AcMesh* mesh) +acHostMeshApplyPeriodicBounds(AcMesh* mesh) { const AcMeshInfo info = mesh->info; for (int w = 0; w < NUM_VTXBUF_HANDLES; ++w) { @@ -105,7 +105,7 @@ acMeshApplyPeriodicBounds(AcMesh* mesh) } AcResult -acMeshClear(AcMesh* mesh) +acHostMeshClear(AcMesh* mesh) { - return acMeshSet(0, mesh); + return acHostMeshSet(0, mesh); } diff --git a/src/utils/modelreduce.c b/src/utils/modelreduce.c index d95fc8c..012617b 100644 --- a/src/utils/modelreduce.c +++ b/src/utils/modelreduce.c @@ -74,7 +74,7 @@ exp_squared_vec(const AcReal a, const AcReal b, const AcReal c) { return exp_squ // clang-format on AcReal -acModelReduceScal(const AcMesh mesh, const ReductionType rtype, const VertexBufferHandle a) +acHostReduceScal(const AcMesh mesh, const ReductionType rtype, const VertexBufferHandle a) { ReduceInitialScalFunc reduce_initial; ReduceFunc reduce; @@ -139,7 +139,7 @@ acModelReduceScal(const AcMesh mesh, const ReductionType rtype, const VertexBuff } AcReal -acModelReduceVec(const AcMesh mesh, const ReductionType rtype, const VertexBufferHandle a, +acHostReduceVec(const AcMesh mesh, const ReductionType rtype, const VertexBufferHandle a, const VertexBufferHandle b, const VertexBufferHandle c) { // AcReal (*reduce_initial)(AcReal, AcReal, AcReal); diff --git a/src/utils/modelsolver.c b/src/utils/modelsolver.c index 3347d75..3bed567 100644 --- a/src/utils/modelsolver.c +++ b/src/utils/modelsolver.c @@ -30,7 +30,7 @@ #include #include "errchk.h" -#include "memory.h" // acMeshCreate, acMeshDestroy, acMeshApplyPeriodicBounds +#include "memory.h" // acHostMeshCreate, acHostMeshDestroy, acMeshApplyPeriodicBounds // Standalone flags #define LDENSITY (1) @@ -985,7 +985,7 @@ checkConfiguration(const AcMeshInfo info) } AcResult -acModelIntegrateStep(AcMesh mesh, const AcReal dt) +acHostIntegrateStep(AcMesh mesh, const AcReal dt) { mesh_info = &(mesh.info); @@ -998,7 +998,7 @@ acModelIntegrateStep(AcMesh mesh, const AcReal dt) checkConfiguration(*mesh_info); AcMesh intermediate_mesh; - acMeshCreate(mesh.info, &intermediate_mesh); + acHostMeshCreate(mesh.info, &intermediate_mesh); const int nx_min = getInt(AC_nx_min); const int nx_max = getInt(AC_nx_max); @@ -1012,7 +1012,7 @@ acModelIntegrateStep(AcMesh mesh, const AcReal dt) for (int step_number = 0; step_number < 3; ++step_number) { // Boundconds - acMeshApplyPeriodicBounds(&mesh); + acHostMeshApplyPeriodicBounds(&mesh); // Alpha step // #pragma omp parallel for @@ -1035,7 +1035,7 @@ acModelIntegrateStep(AcMesh mesh, const AcReal dt) } } - acMeshDestroy(&intermediate_mesh); + acHostMeshDestroy(&intermediate_mesh); mesh_info = NULL; return AC_SUCCESS; } From b6bb53a75cfd333b9cefbf076938d0d71c9860b5 Mon Sep 17 00:00:00 2001 From: jpekkila Date: Tue, 24 Nov 2020 21:39:44 +0200 Subject: [PATCH 3/3] Missed some renamings --- samples/benchmark/main.cc | 8 ++++---- samples/mpitest/main.cc | 12 ++++++------ samples/standalone_mpi/main.cc | 2 +- src/core/device.cc | 2 +- src/utils/modelsolver.c | 2 +- 5 files changed, 13 insertions(+), 13 deletions(-) diff --git a/samples/benchmark/main.cc b/samples/benchmark/main.cc index 861d2bd..1a6c366 100644 --- a/samples/benchmark/main.cc +++ b/samples/benchmark/main.cc @@ -120,8 +120,8 @@ main(int argc, char** argv) if (pid == 0) { acHostMeshCreate(info, &model); acHostMeshCreate(info, &candidate); - acMeshRandomize(&model); - acMeshRandomize(&candidate); + acHostMeshRandomize(&model); + acHostMeshRandomize(&candidate); }*/ // GPU alloc & compute @@ -131,7 +131,7 @@ main(int argc, char** argv) /* AcMesh model; acHostMeshCreate(info, &model); - acMeshRandomize(&model); + acHostMeshRandomize(&model); acGridLoadMesh(STREAM_DEFAULT, model); */ @@ -146,7 +146,7 @@ main(int argc, char** argv) // Verify if (pid == 0) { acHostIntegrateStep(model, FLT_EPSILON); - acMeshApplyPeriodicBounds(&model); + acHostMeshApplyPeriodicBounds(&model); AcResult retval = acVerifyMesh(model, candidate); acHostMeshDestroy(&model); diff --git a/samples/mpitest/main.cc b/samples/mpitest/main.cc index e690352..2ef4c0e 100644 --- a/samples/mpitest/main.cc +++ b/samples/mpitest/main.cc @@ -49,8 +49,8 @@ main(void) if (pid == 0) { acHostMeshCreate(info, &model); acHostMeshCreate(info, &candidate); - acMeshRandomize(&model); - acMeshRandomize(&candidate); + acHostMeshRandomize(&model); + acHostMeshRandomize(&candidate); } // GPU alloc & compute @@ -61,10 +61,10 @@ main(void) acGridPeriodicBoundconds(STREAM_DEFAULT); acGridStoreMesh(STREAM_DEFAULT, &candidate); if (pid == 0) { - acMeshApplyPeriodicBounds(&model); + acHostMeshApplyPeriodicBounds(&model); const AcResult res = acVerifyMesh("Boundconds", model, candidate); ERRCHK_ALWAYS(res == AC_SUCCESS); - acMeshRandomize(&model); + acHostMeshRandomize(&model); } // Integration @@ -74,10 +74,10 @@ main(void) acGridStoreMesh(STREAM_DEFAULT, &candidate); if (pid == 0) { acHostIntegrateStep(model, FLT_EPSILON); - acMeshApplyPeriodicBounds(&model); + acHostMeshApplyPeriodicBounds(&model); const AcResult res = acVerifyMesh("Integration", model, candidate); ERRCHK_ALWAYS(res == AC_SUCCESS); - acMeshRandomize(&model); + acHostMeshRandomize(&model); } // Scalar reductions diff --git a/samples/standalone_mpi/main.cc b/samples/standalone_mpi/main.cc index 23080af..897e595 100644 --- a/samples/standalone_mpi/main.cc +++ b/samples/standalone_mpi/main.cc @@ -445,7 +445,7 @@ main(int argc, char** argv) #endif } - acMeshApplyPeriodicBounds(&mesh); + acHostMeshApplyPeriodicBounds(&mesh); if (start_step == 0) { save_mesh(mesh, 0, t_step); } diff --git a/src/core/device.cc b/src/core/device.cc index 123dee2..d4fdafe 100644 --- a/src/core/device.cc +++ b/src/core/device.cc @@ -1290,7 +1290,7 @@ acGridRandomize(void) AcMesh host; acHostMeshCreate(grid.submesh.info, &host); - acMeshRandomize(&host); + acHostMeshRandomize(&host); acDeviceLoadMesh(grid.device, STREAM_DEFAULT, host); acHostMeshDestroy(&host); diff --git a/src/utils/modelsolver.c b/src/utils/modelsolver.c index 3bed567..0a20e3b 100644 --- a/src/utils/modelsolver.c +++ b/src/utils/modelsolver.c @@ -30,7 +30,7 @@ #include #include "errchk.h" -#include "memory.h" // acHostMeshCreate, acHostMeshDestroy, acMeshApplyPeriodicBounds +#include "memory.h" // acHostMeshCreate, acHostMeshDestroy, acHostMeshApplyPeriodicBounds // Standalone flags #define LDENSITY (1)