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

Merged in host-layer-renaming-2020-11-24 (pull request #17)

Browse Source 
Renaming host layer functions and introducing acSetVertexBuffer

Approved-by: Miikka Väisälä <mvaisala@asiaa.sinica.edu.tw>
This commit is contained in:
jpekkila
2020-11-25 03:37:40 +00:00
committed by Miikka Väisälä
parent 5a29543727 5df695b4b1
commit 7878811820
19 changed files with 150 additions and 102 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
acc/src/code_generator.c
View File

@@ -606,7 +606,7 @@ generate_headers(void)
! -*-f90-*- (for emacs) vim:set filetype=fortran: (for vim) ! -*-f90-*- (for emacs) vim:set filetype=fortran: (for vim)
! Utils (see astaroth_fortran.cc for definitions) ! Utils (see astaroth_fortran.cc for definitions)
external acupdatebuiltinparams external achostupdatebuiltinparams
external acgetdevicecount external acgetdevicecount
! Device interface (see astaroth_fortran.cc for definitions) ! Device interface (see astaroth_fortran.cc for definitions)

17
include/astaroth.h
View File

@@ -225,6 +225,9 @@ AcResult acLoadDeviceConstant(const AcRealParam param, const AcReal value);
/** Loads an AcMesh to the devices visible to the caller */ /** Loads an AcMesh to the devices visible to the caller */
AcResult acLoad(const AcMesh host_mesh); 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*/ /** Stores the AcMesh distributed among the devices visible to the caller back to the host*/
AcResult acStore(AcMesh* host_mesh); 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 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 ? */ /** Deprecated ? */
AcResult acNodeStoreVertexBufferWithOffset(const Node node, const Stream stream, AcResult acNodeStoreVertexBufferWithOffset(const Node node, const Stream stream,
const VertexBufferHandle vtxbuf_handle, const int3 src, 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 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, AcResult acDeviceStoreVertexBufferWithOffset(const Device device, const Stream stream,
const VertexBufferHandle vtxbuf_handle, const int3 src, 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 */ /** 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 */ /** 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 */ /** Randomizes a host mesh */
AcResult acMeshRandomize(AcMesh* mesh); AcResult acHostMeshRandomize(AcMesh* mesh);
/** Destroys a mesh stored in host memory */ /** Destroys a mesh stored in host memory */
AcResult acMeshDestroy(AcMesh* mesh); AcResult acHostMeshDestroy(AcMesh* mesh);
#ifdef __cplusplus #ifdef __cplusplus
} // extern "C" } // extern "C"

14
include/astaroth_utils.h
View File

@@ -45,25 +45,25 @@ typedef struct {
AcResult acLoadConfig(const char* config_path, AcMeshInfo* config); 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); const VertexBufferHandle b, const VertexBufferHandle c);
Error acGetError(const AcReal model, const AcReal candidate); Error acGetError(const AcReal model, const AcReal candidate);

22
samples/benchmark/main.cc
View File

@@ -98,7 +98,7 @@ main(int argc, char** argv)
info.int_params[AC_nx] = nx; info.int_params[AC_nx] = nx;
info.int_params[AC_ny] = ny; info.int_params[AC_ny] = ny;
info.int_params[AC_nz] = nz; info.int_params[AC_nz] = nz;
acUpdateBuiltinParams(&info); acHostUpdateBuiltinParams(&info);
printf("Benchmark mesh dimensions: (%d, %d, %d)\n", nx, ny, nz); printf("Benchmark mesh dimensions: (%d, %d, %d)\n", nx, ny, nz);
} }
else { else {
@@ -118,10 +118,10 @@ main(int argc, char** argv)
/* /*
AcMesh model, candidate; AcMesh model, candidate;
if (pid == 0) { if (pid == 0) {
acMeshCreate(info, &model); acHostMeshCreate(info, &model);
acMeshCreate(info, &candidate); acHostMeshCreate(info, &candidate);
acMeshRandomize(&model); acHostMeshRandomize(&model);
acMeshRandomize(&candidate); acHostMeshRandomize(&candidate);
}*/ }*/
// GPU alloc & compute // GPU alloc & compute
@@ -130,8 +130,8 @@ main(int argc, char** argv)
/* /*
AcMesh model; AcMesh model;
acMeshCreate(info, &model); acHostMeshCreate(info, &model);
acMeshRandomize(&model); acHostMeshRandomize(&model);
acGridLoadMesh(STREAM_DEFAULT, model); acGridLoadMesh(STREAM_DEFAULT, model);
*/ */
@@ -145,12 +145,12 @@ main(int argc, char** argv)
// Verify // Verify
if (pid == 0) { if (pid == 0) {
acModelIntegrateStep(model, FLT_EPSILON); acHostIntegrateStep(model, FLT_EPSILON);
acMeshApplyPeriodicBounds(&model); acHostMeshApplyPeriodicBounds(&model);
AcResult retval = acVerifyMesh(model, candidate); AcResult retval = acVerifyMesh(model, candidate);
acMeshDestroy(&model); acHostMeshDestroy(&model);
acMeshDestroy(&candidate); acHostMeshDestroy(&candidate);
if (retval != AC_SUCCESS) { if (retval != AC_SUCCESS) {
fprintf(stderr, "Failures found, benchmark invalid. Skipping\n"); fprintf(stderr, "Failures found, benchmark invalid. Skipping\n");

12
samples/cpptest/main.cc
View File

@@ -31,12 +31,12 @@ main(void)
// Alloc // Alloc
AcMesh model, candidate; AcMesh model, candidate;
acMeshCreate(info, &model); acHostMeshCreate(info, &model);
acMeshCreate(info, &candidate); acHostMeshCreate(info, &candidate);
// Init // Init
acMeshRandomize(&model); acHostMeshRandomize(&model);
acMeshApplyPeriodicBounds(&model); acHostMeshApplyPeriodicBounds(&model);
// Verify that the mesh was loaded and stored correctly // Verify that the mesh was loaded and stored correctly
acInit(info); acInit(info);
@@ -55,8 +55,8 @@ main(void)
// Destroy // Destroy
acQuit(); acQuit();
acMeshDestroy(&model); acHostMeshDestroy(&model);
acMeshDestroy(&candidate); acHostMeshDestroy(&candidate);
puts("cpptest complete."); puts("cpptest complete.");
return EXIT_SUCCESS; return EXIT_SUCCESS;

13
samples/ctest/main.c
View File

@@ -30,12 +30,12 @@ main(void)
// Alloc // Alloc
AcMesh model, candidate; AcMesh model, candidate;
acMeshCreate(info, &model); acHostMeshCreate(info, &model);
acMeshCreate(info, &candidate); acHostMeshCreate(info, &candidate);
// Init // Init
acMeshRandomize(&model); acHostMeshRandomize(&model);
acMeshApplyPeriodicBounds(&model); acHostMeshApplyPeriodicBounds(&model);
// Verify that the mesh was loaded and stored correctly // Verify that the mesh was loaded and stored correctly
acInit(info); acInit(info);
@@ -46,6 +46,7 @@ main(void)
// Attempt to integrate and check max and min // Attempt to integrate and check max and min
printf("Integrating... "); printf("Integrating... ");
acIntegrate(FLT_EPSILON); acIntegrate(FLT_EPSILON);
printf("Done.\nVTXBUF ranges after one integration step:\n"); printf("Done.\nVTXBUF ranges after one integration step:\n");
for (size_t i = 0; i < NUM_VTXBUF_HANDLES; ++i) for (size_t i = 0; i < NUM_VTXBUF_HANDLES; ++i)
printf("\t%-15s... [%.3g, %.3g]\n", vtxbuf_names[i], // printf("\t%-15s... [%.3g, %.3g]\n", vtxbuf_names[i], //
@@ -54,8 +55,8 @@ main(void)
// Destroy // Destroy
acQuit(); acQuit();
acMeshDestroy(&model); acHostMeshDestroy(&model);
acMeshDestroy(&candidate); acHostMeshDestroy(&candidate);
puts("ctest complete."); puts("ctest complete.");
return EXIT_SUCCESS; return EXIT_SUCCESS;

2
samples/fortrantest/main.f90
View File

@@ -14,7 +14,7 @@ program pc
info%int_params(AC_nx + 1) = 128 info%int_params(AC_nx + 1) = 128
info%int_params(AC_ny + 1) = 128 info%int_params(AC_ny + 1) = 128
info%int_params(AC_nz + 1) = 128 info%int_params(AC_nz + 1) = 128
call acupdatebuiltinparams(info) call achostupdatebuiltinparams(info)
call acdevicecreate(0, info, device) call acdevicecreate(0, info, device)
call acdeviceprintinfo(device) call acdeviceprintinfo(device)

34
samples/mpitest/main.cc
View File

@@ -47,10 +47,10 @@ main(void)
AcMesh model, candidate; AcMesh model, candidate;
if (pid == 0) { if (pid == 0) {
acMeshCreate(info, &model); acHostMeshCreate(info, &model);
acMeshCreate(info, &candidate); acHostMeshCreate(info, &candidate);
acMeshRandomize(&model); acHostMeshRandomize(&model);
acMeshRandomize(&candidate); acHostMeshRandomize(&candidate);
} }
// GPU alloc & compute // GPU alloc & compute
@@ -61,10 +61,10 @@ main(void)
acGridPeriodicBoundconds(STREAM_DEFAULT); acGridPeriodicBoundconds(STREAM_DEFAULT);
acGridStoreMesh(STREAM_DEFAULT, &candidate); acGridStoreMesh(STREAM_DEFAULT, &candidate);
if (pid == 0) { if (pid == 0) {
acMeshApplyPeriodicBounds(&model); acHostMeshApplyPeriodicBounds(&model);
const AcResult res = acVerifyMesh("Boundconds", model, candidate); const AcResult res = acVerifyMesh("Boundconds", model, candidate);
ERRCHK_ALWAYS(res == AC_SUCCESS); ERRCHK_ALWAYS(res == AC_SUCCESS);
acMeshRandomize(&model); acHostMeshRandomize(&model);
} }
// Integration // Integration
@@ -73,11 +73,11 @@ main(void)
acGridPeriodicBoundconds(STREAM_DEFAULT); acGridPeriodicBoundconds(STREAM_DEFAULT);
acGridStoreMesh(STREAM_DEFAULT, &candidate); acGridStoreMesh(STREAM_DEFAULT, &candidate);
if (pid == 0) { if (pid == 0) {
acModelIntegrateStep(model, FLT_EPSILON); acHostIntegrateStep(model, FLT_EPSILON);
acMeshApplyPeriodicBounds(&model); acHostMeshApplyPeriodicBounds(&model);
const AcResult res = acVerifyMesh("Integration", model, candidate); const AcResult res = acVerifyMesh("Integration", model, candidate);
ERRCHK_ALWAYS(res == AC_SUCCESS); ERRCHK_ALWAYS(res == AC_SUCCESS);
acMeshRandomize(&model); acHostMeshRandomize(&model);
} }
// Scalar reductions // Scalar reductions
@@ -93,10 +93,10 @@ main(void)
AcReal candval; AcReal candval;
acGridReduceScal(STREAM_DEFAULT, (ReductionType)i, v0, &candval); acGridReduceScal(STREAM_DEFAULT, (ReductionType)i, v0, &candval);
if (pid == 0) { 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 error = acGetError(modelval, candval);
error.maximum_magnitude = acModelReduceScal(model, RTYPE_MAX, v0); error.maximum_magnitude = acHostReduceScal(model, RTYPE_MAX, v0);
error.minimum_magnitude = acModelReduceScal(model, RTYPE_MIN, v0); error.minimum_magnitude = acHostReduceScal(model, RTYPE_MIN, v0);
ERRCHK_ALWAYS(acEvalError(rtype_names[i], error)); ERRCHK_ALWAYS(acEvalError(rtype_names[i], error));
} }
} }
@@ -114,17 +114,17 @@ main(void)
AcReal candval; AcReal candval;
acGridReduceVec(STREAM_DEFAULT, (ReductionType)i, v0, v1, v2, &candval); acGridReduceVec(STREAM_DEFAULT, (ReductionType)i, v0, v1, v2, &candval);
if (pid == 0) { 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 error = acGetError(modelval, candval);
error.maximum_magnitude = acModelReduceVec(model, RTYPE_MAX, v0, v1, v2); error.maximum_magnitude = acHostReduceVec(model, RTYPE_MAX, v0, v1, v2);
error.minimum_magnitude = acModelReduceVec(model, RTYPE_MIN, v0, v1, v1); error.minimum_magnitude = acHostReduceVec(model, RTYPE_MIN, v0, v1, v1);
ERRCHK_ALWAYS(acEvalError(rtype_names[i], error)); ERRCHK_ALWAYS(acEvalError(rtype_names[i], error));
} }
} }
if (pid == 0) { if (pid == 0) {
acMeshDestroy(&model); acHostMeshDestroy(&model);
acMeshDestroy(&candidate); acHostMeshDestroy(&candidate);
} }
acGridQuit(); acGridQuit();

18
samples/standalone_mpi/host_memory.cc
View File

@@ -576,7 +576,7 @@ acmesh_init_to(const InitType& init_type, AcMesh* mesh)
switch (init_type) { switch (init_type) {
case INIT_TYPE_RANDOM: { case INIT_TYPE_RANDOM: {
acMeshClear(mesh); acHostMeshClear(mesh);
const AcReal range = AcReal(0.01); const AcReal range = AcReal(0.01);
for (int w = 0; w < NUM_VTXBUF_HANDLES; ++w) for (int w = 0; w < NUM_VTXBUF_HANDLES; ++w)
for (int i = 0; i < n; ++i) for (int i = 0; i < n; ++i)
@@ -585,14 +585,14 @@ acmesh_init_to(const InitType& init_type, AcMesh* mesh)
break; break;
} }
case INIT_TYPE_GAUSSIAN_RADIAL_EXPL: case INIT_TYPE_GAUSSIAN_RADIAL_EXPL:
acMeshClear(mesh); acHostMeshClear(mesh);
acVertexBufferSet(VTXBUF_LNRHO, mesh->info.real_params[AC_ampl_lnrho], mesh); acHostVertexBufferSet(VTXBUF_LNRHO, mesh->info.real_params[AC_ampl_lnrho], mesh);
// acmesh_init_to(INIT_TYPE_RANDOM, mesh); // acmesh_init_to(INIT_TYPE_RANDOM, mesh);
gaussian_radial_explosion(mesh); gaussian_radial_explosion(mesh);
break; break;
case INIT_TYPE_XWAVE: case INIT_TYPE_XWAVE:
acMeshClear(mesh); acHostMeshClear(mesh);
acmesh_init_to(INIT_TYPE_RANDOM, mesh); acmesh_init_to(INIT_TYPE_RANDOM, mesh);
for (int k = 0; k < mz; k++) { for (int k = 0; k < mz; k++) {
for (int j = 0; j < my; j++) { for (int j = 0; j < my; j++) {
@@ -605,24 +605,24 @@ acmesh_init_to(const InitType& init_type, AcMesh* mesh)
} }
break; break;
case INIT_TYPE_SIMPLE_CORE: case INIT_TYPE_SIMPLE_CORE:
acMeshClear(mesh); acHostMeshClear(mesh);
simple_uniform_core(mesh); simple_uniform_core(mesh);
break; break;
case INIT_TYPE_VEDGE: case INIT_TYPE_VEDGE:
acMeshClear(mesh); acHostMeshClear(mesh);
inflow_vedge_freefall(mesh); inflow_vedge_freefall(mesh);
break; break;
case INIT_TYPE_VEDGEX: case INIT_TYPE_VEDGEX:
acMeshClear(mesh); acHostMeshClear(mesh);
inflow_freefall_x(mesh); inflow_freefall_x(mesh);
break; break;
case INIT_TYPE_RAYLEIGH_TAYLOR: case INIT_TYPE_RAYLEIGH_TAYLOR:
acMeshClear(mesh); acHostMeshClear(mesh);
inflow_freefall_x(mesh); inflow_freefall_x(mesh);
lnrho_step(mesh); lnrho_step(mesh);
break; break;
case INIT_TYPE_ABC_FLOW: { case INIT_TYPE_ABC_FLOW: {
acMeshClear(mesh); acHostMeshClear(mesh);
acmesh_init_to(INIT_TYPE_RANDOM, mesh); acmesh_init_to(INIT_TYPE_RANDOM, mesh);
for (int k = nz_min; k < nz_max; k++) { for (int k = nz_min; k < nz_max; k++) {
for (int j = ny_min; j < ny_max; j++) { for (int j = ny_min; j < ny_max; j++) {

26
samples/standalone_mpi/main.cc
View File

@@ -199,9 +199,9 @@ print_diagnostics_host(const AcMesh mesh, const int step, const AcReal dt, const
const int max_name_width = 16; const int max_name_width = 16;
// Calculate rms, min and max from the velocity vector field // Calculate rms, min and max from the velocity vector field
buf_max = acModelReduceVec(mesh, RTYPE_MAX, VTXBUF_UUX, VTXBUF_UUY, VTXBUF_UUZ); buf_max = acHostReduceVec(mesh, RTYPE_MAX, VTXBUF_UUX, VTXBUF_UUY, VTXBUF_UUZ);
buf_min = acModelReduceVec(mesh, RTYPE_MIN, VTXBUF_UUX, VTXBUF_UUY, VTXBUF_UUZ); buf_min = acHostReduceVec(mesh, RTYPE_MIN, VTXBUF_UUX, VTXBUF_UUY, VTXBUF_UUZ);
buf_rms = acModelReduceVec(mesh, RTYPE_RMS, 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: The ordering in the earlier version was wrong in terms of variable
// MV: name and its diagnostics. // 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 // Calculate rms, min and max from the variables as scalars
for (int i = 0; i < NUM_VTXBUF_HANDLES; ++i) { for (int i = 0; i < NUM_VTXBUF_HANDLES; ++i) {
buf_max = acModelReduceScal(mesh, RTYPE_MAX, VertexBufferHandle(i)); buf_max = acHostReduceScal(mesh, RTYPE_MAX, VertexBufferHandle(i));
buf_min = acModelReduceScal(mesh, RTYPE_MIN, VertexBufferHandle(i)); buf_min = acHostReduceScal(mesh, RTYPE_MIN, VertexBufferHandle(i));
buf_rms = acModelReduceScal(mesh, RTYPE_RMS, 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], printf(" %*s: min %.3e,\trms %.3e,\tmax %.3e\n", max_name_width, vtxbuf_names[i],
double(buf_min), double(buf_rms), double(buf_max)); double(buf_min), double(buf_rms), double(buf_max));
@@ -330,7 +330,7 @@ main(int argc, char** argv)
AcMesh mesh; AcMesh mesh;
if (pid == 0) { if (pid == 0) {
acMeshCreate(info, &mesh); acHostMeshCreate(info, &mesh);
acmesh_init_to(INIT_TYPE_GAUSSIAN_RADIAL_EXPL, &mesh); acmesh_init_to(INIT_TYPE_GAUSSIAN_RADIAL_EXPL, &mesh);
} }
acGridInit(info); acGridInit(info);
@@ -362,7 +362,7 @@ main(int argc, char** argv)
} }
} }
if (pid == 0) if (pid == 0)
acMeshDestroy(&mesh); acHostMeshDestroy(&mesh);
acGridQuit(); acGridQuit();
/////////////// Simple example END /////////////// Simple example END
@@ -376,7 +376,7 @@ main(int argc, char** argv)
if (argc == 3 && (!strcmp(argv[1], "-c") || !strcmp(argv[1], "--config"))) { if (argc == 3 && (!strcmp(argv[1], "-c") || !strcmp(argv[1], "--config"))) {
acLoadConfig(argv[2], &info); acLoadConfig(argv[2], &info);
load_config(argv[2], &info); load_config(argv[2], &info);
acUpdateBuiltinParams(&info); acHostUpdateBuiltinParams(&info);
} }
else { else {
printf("Usage: ./ac_run\n"); printf("Usage: ./ac_run\n");
@@ -388,7 +388,7 @@ main(int argc, char** argv)
else { else {
acLoadConfig(AC_DEFAULT_CONFIG, &info); acLoadConfig(AC_DEFAULT_CONFIG, &info);
load_config(AC_DEFAULT_CONFIG, &info); load_config(AC_DEFAULT_CONFIG, &info);
acUpdateBuiltinParams(&info); acHostUpdateBuiltinParams(&info);
} }
const int start_step = info.int_params[AC_start_step]; const int start_step = info.int_params[AC_start_step];
@@ -406,7 +406,7 @@ main(int argc, char** argv)
AcMesh mesh; AcMesh mesh;
///////////////////////////////// PROC 0 BLOCK START /////////////////////////////////////////// ///////////////////////////////// PROC 0 BLOCK START ///////////////////////////////////////////
if (pid == 0) { if (pid == 0) {
acMeshCreate(info, &mesh); acHostMeshCreate(info, &mesh);
// TODO: This need to be possible to define in astaroth.conf // 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_GAUSSIAN_RADIAL_EXPL, &mesh);
// acmesh_init_to(INIT_TYPE_SIMPLE_CORE, mesh); //Initial condition for a collapse test // acmesh_init_to(INIT_TYPE_SIMPLE_CORE, mesh); //Initial condition for a collapse test
@@ -445,7 +445,7 @@ main(int argc, char** argv)
#endif #endif
} }
acMeshApplyPeriodicBounds(&mesh); acHostMeshApplyPeriodicBounds(&mesh);
if (start_step == 0) { if (start_step == 0) {
save_mesh(mesh, 0, t_step); save_mesh(mesh, 0, t_step);
} }
@@ -564,7 +564,7 @@ main(int argc, char** argv)
acGridQuit(); acGridQuit();
if (pid == 0) if (pid == 0)
acMeshDestroy(&mesh); acHostMeshDestroy(&mesh);
fclose(diag_file); fclose(diag_file);
#endif #endif

14
src/core/astaroth.cc
View File

@@ -74,6 +74,12 @@ acLoad(const AcMesh host_mesh)
return acNodeLoadMesh(nodes[0], STREAM_DEFAULT, 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 AcResult
acStore(AcMesh* host_mesh) acStore(AcMesh* host_mesh)
{ {
@@ -182,7 +188,7 @@ acGetNode(void)
} }
AcResult AcResult
acUpdateBuiltinParams(AcMeshInfo* config) acHostUpdateBuiltinParams(AcMeshInfo* config)
{ {
config->int_params[AC_mx] = config->int_params[AC_nx] + STENCIL_ORDER; config->int_params[AC_mx] = config->int_params[AC_nx] + STENCIL_ORDER;
///////////// PAD TEST ///////////// PAD TEST
@@ -221,7 +227,7 @@ acUpdateBuiltinParams(AcMeshInfo* config)
} }
AcResult AcResult
acMeshCreate(const AcMeshInfo info, AcMesh* mesh) acHostMeshCreate(const AcMeshInfo info, AcMesh* mesh)
{ {
mesh->info = info; mesh->info = info;
@@ -241,7 +247,7 @@ randf(void)
} }
AcResult AcResult
acMeshRandomize(AcMesh* mesh) acHostMeshRandomize(AcMesh* mesh)
{ {
const int n = acVertexBufferSize(mesh->info); const int n = acVertexBufferSize(mesh->info);
for (int w = 0; w < NUM_VTXBUF_HANDLES; ++w) for (int w = 0; w < NUM_VTXBUF_HANDLES; ++w)
@@ -252,7 +258,7 @@ acMeshRandomize(AcMesh* mesh)
} }
AcResult AcResult
acMeshDestroy(AcMesh* mesh) acHostMeshDestroy(AcMesh* mesh)
{ {
for (int w = 0; w < NUM_VTXBUF_HANDLES; ++w) for (int w = 0; w < NUM_VTXBUF_HANDLES; ++w)
free(mesh->vertex_buffer[w]); free(mesh->vertex_buffer[w]);

4
src/core/astaroth_fortran.cc
View File

@@ -7,9 +7,9 @@
* Utils * Utils
*/ */
void void
acupdatebuiltinparams_(AcMeshInfo* info) achostupdatebuiltinparams_(AcMeshInfo* info)
{ {
acUpdateBuiltinParams(info); acHostUpdateBuiltinParams(info);
} }
void void

2
src/core/astaroth_fortran.h
View File

@@ -8,7 +8,7 @@ extern "C" {
/** /**
* Utils * Utils
*/ */
void acupdatebuiltinparams_(AcMeshInfo* info); void achostupdatebuiltinparams_(AcMeshInfo* info);
void acgetdevicecount_(int* count); void acgetdevicecount_(int* count);

32
src/core/device.cc
View File

@@ -294,6 +294,26 @@ acDeviceLoadMesh(const Device device, const Stream stream, const AcMesh host_mes
return AC_SUCCESS; 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 AcResult
acDeviceStoreVertexBufferWithOffset(const Device device, const Stream stream, acDeviceStoreVertexBufferWithOffset(const Device device, const Stream stream,
const VertexBufferHandle vtxbuf_handle, const int3 src, const VertexBufferHandle vtxbuf_handle, const int3 src,
@@ -1269,10 +1289,10 @@ acGridRandomize(void)
ERRCHK(grid.initialized); ERRCHK(grid.initialized);
AcMesh host; AcMesh host;
acMeshCreate(grid.submesh.info, &host); acHostMeshCreate(grid.submesh.info, &host);
acMeshRandomize(&host); acHostMeshRandomize(&host);
acDeviceLoadMesh(grid.device, STREAM_DEFAULT, host); acDeviceLoadMesh(grid.device, STREAM_DEFAULT, host);
acMeshDestroy(&host); acHostMeshDestroy(&host);
return AC_SUCCESS; return AC_SUCCESS;
} }
@@ -1320,7 +1340,7 @@ acGridInit(const AcMeshInfo info)
}; };
submesh_info.int3_params[AC_multigpu_offset] = pid3d * submesh_info.int3_params[AC_multigpu_offset] = pid3d *
(int3){submesh_nx, submesh_ny, submesh_nz}; (int3){submesh_nx, submesh_ny, submesh_nz};
acUpdateBuiltinParams(&submesh_info); acHostUpdateBuiltinParams(&submesh_info);
// GPU alloc // GPU alloc
int devices_per_node = -1; int devices_per_node = -1;
@@ -1331,7 +1351,7 @@ acGridInit(const AcMeshInfo info)
// CPU alloc // CPU alloc
AcMesh submesh; AcMesh submesh;
acMeshCreate(submesh_info, &submesh); acHostMeshCreate(submesh_info, &submesh);
// Setup the global grid structure // Setup the global grid structure
grid.device = device; grid.device = device;
@@ -1380,7 +1400,7 @@ acGridQuit(void)
grid.initialized = false; grid.initialized = false;
grid.decomposition = (uint3_64){0, 0, 0}; grid.decomposition = (uint3_64){0, 0, 0};
acMeshDestroy(&grid.submesh); acHostMeshDestroy(&grid.submesh);
acDeviceDestroy(grid.device); acDeviceDestroy(grid.device);
acGridSynchronizeStream(STREAM_ALL); acGridSynchronizeStream(STREAM_ALL);

12
src/core/node.cc
View File

@@ -536,6 +536,18 @@ acNodeLoadMesh(const Node node, const Stream stream, const AcMesh host_mesh)
return AC_SUCCESS; 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 AcResult
acNodeStoreVertexBufferWithOffset(const Node node, const Stream stream, acNodeStoreVertexBufferWithOffset(const Node node, const Stream stream,
const VertexBufferHandle vtxbuf_handle, const int3 src, const VertexBufferHandle vtxbuf_handle, const int3 src,

2
src/utils/config_loader.c
View File

@@ -89,7 +89,7 @@ acLoadConfig(const char* config_path, AcMeshInfo* config)
memset(config, (uint8_t)0xFF, sizeof(*config)); memset(config, (uint8_t)0xFF, sizeof(*config));
parse_config(config_path, config); parse_config(config_path, config);
acUpdateBuiltinParams(config); acHostUpdateBuiltinParams(config);
#if AC_VERBOSE #if AC_VERBOSE
printf("###############################################################\n"); printf("###############################################################\n");
printf("Config dimensions loaded:\n"); printf("Config dimensions loaded:\n");

12
src/utils/memory.c
View File

@@ -21,7 +21,7 @@
#include "errchk.h" #include "errchk.h"
AcResult 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); const int n = acVertexBufferSize(mesh->info);
for (int i = 0; i < n; ++i) for (int i = 0; i < n; ++i)
@@ -30,16 +30,16 @@ acVertexBufferSet(const VertexBufferHandle handle, const AcReal value, AcMesh* m
return AC_SUCCESS; return AC_SUCCESS;
} }
AcResult AcResult
acMeshSet(const AcReal value, AcMesh* mesh) acHostMeshSet(const AcReal value, AcMesh* mesh)
{ {
for (int w = 0; w < NUM_VTXBUF_HANDLES; ++w) for (int w = 0; w < NUM_VTXBUF_HANDLES; ++w)
acVertexBufferSet(w, value, mesh); acHostVertexBufferSet(w, value, mesh);
return AC_SUCCESS; return AC_SUCCESS;
} }
AcResult AcResult
acMeshApplyPeriodicBounds(AcMesh* mesh) acHostMeshApplyPeriodicBounds(AcMesh* mesh)
{ {
const AcMeshInfo info = mesh->info; const AcMeshInfo info = mesh->info;
for (int w = 0; w < NUM_VTXBUF_HANDLES; ++w) { for (int w = 0; w < NUM_VTXBUF_HANDLES; ++w) {
@@ -105,7 +105,7 @@ acMeshApplyPeriodicBounds(AcMesh* mesh)
} }
AcResult AcResult
acMeshClear(AcMesh* mesh) acHostMeshClear(AcMesh* mesh)
{ {
return acMeshSet(0, mesh); return acHostMeshSet(0, mesh);
} }

4
src/utils/modelreduce.c
View File

@@ -74,7 +74,7 @@ exp_squared_vec(const AcReal a, const AcReal b, const AcReal c) { return exp_squ
// clang-format on // clang-format on
AcReal AcReal
acModelReduceScal(const AcMesh mesh, const ReductionType rtype, const VertexBufferHandle a) acHostReduceScal(const AcMesh mesh, const ReductionType rtype, const VertexBufferHandle a)
{ {
ReduceInitialScalFunc reduce_initial; ReduceInitialScalFunc reduce_initial;
ReduceFunc reduce; ReduceFunc reduce;
@@ -139,7 +139,7 @@ acModelReduceScal(const AcMesh mesh, const ReductionType rtype, const VertexBuff
} }
AcReal 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) const VertexBufferHandle b, const VertexBufferHandle c)
{ {
// AcReal (*reduce_initial)(AcReal, AcReal, AcReal); // AcReal (*reduce_initial)(AcReal, AcReal, AcReal);

10
src/utils/modelsolver.c
View File

@@ -30,7 +30,7 @@
#include <stdbool.h> #include <stdbool.h>
#include "errchk.h" #include "errchk.h"
#include "memory.h" // acMeshCreate, acMeshDestroy, acMeshApplyPeriodicBounds #include "memory.h" // acHostMeshCreate, acHostMeshDestroy, acHostMeshApplyPeriodicBounds
// Standalone flags // Standalone flags
#define LDENSITY (1) #define LDENSITY (1)
@@ -985,7 +985,7 @@ checkConfiguration(const AcMeshInfo info)
} }
AcResult AcResult
acModelIntegrateStep(AcMesh mesh, const AcReal dt) acHostIntegrateStep(AcMesh mesh, const AcReal dt)
{ {
mesh_info = &(mesh.info); mesh_info = &(mesh.info);
@@ -998,7 +998,7 @@ acModelIntegrateStep(AcMesh mesh, const AcReal dt)
checkConfiguration(*mesh_info); checkConfiguration(*mesh_info);
AcMesh intermediate_mesh; AcMesh intermediate_mesh;
acMeshCreate(mesh.info, &intermediate_mesh); acHostMeshCreate(mesh.info, &intermediate_mesh);
const int nx_min = getInt(AC_nx_min); const int nx_min = getInt(AC_nx_min);
const int nx_max = getInt(AC_nx_max); 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) { for (int step_number = 0; step_number < 3; ++step_number) {
// Boundconds // Boundconds
acMeshApplyPeriodicBounds(&mesh); acHostMeshApplyPeriodicBounds(&mesh);
// Alpha step // Alpha step
// #pragma omp parallel for // #pragma omp parallel for
@@ -1035,7 +1035,7 @@ acModelIntegrateStep(AcMesh mesh, const AcReal dt)
} }
} }
acMeshDestroy(&intermediate_mesh); acHostMeshDestroy(&intermediate_mesh);
mesh_info = NULL; mesh_info = NULL;
return AC_SUCCESS; return AC_SUCCESS;
} }