From 095f8630976c5429db0ac345a52ffe4712ab2be2 Mon Sep 17 00:00:00 2001 From: jpekkila Date: Tue, 24 Nov 2020 21:29:14 +0200 Subject: [PATCH] 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,