Added functions acSetVertexBuffer, acNodeSetVertexBuffer, and acDeviceSetVertexBuffer for setting device memory directly to some constant

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"

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]);

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);

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,