From 787363226b5510c06990ee5ac66939f7f09c9304 Mon Sep 17 00:00:00 2001 From: jpekkila Date: Mon, 19 Aug 2019 15:28:16 +0300 Subject: [PATCH] Added functions for loading int, int3, scalar and vector constants to the device layer (acDeviceLoad...Constant) --- include/astaroth_device.h | 16 ++++++++++++++-- src/core/device.cu | 37 +++++++++++++++++++++++++++++++++++-- src/core/node.cu | 2 +- 3 files changed, 50 insertions(+), 5 deletions(-) diff --git a/include/astaroth_device.h b/include/astaroth_device.h index 95cbb83..1d7926f 100644 --- a/include/astaroth_device.h +++ b/include/astaroth_device.h @@ -46,8 +46,20 @@ AcResult acDeviceSynchronizeStream(const Device device, const Stream stream); AcResult acDeviceSwapBuffers(const Device device); /** */ -AcResult acDeviceLoadConstant(const Device device, const Stream stream, const AcRealParam param, - const AcReal value); +AcResult acDeviceLoadScalarConstant(const Device device, const Stream stream, + const AcRealParam param, const AcReal value); + +/** */ +AcResult acDeviceLoadVectorConstant(const Device device, const Stream stream, + const AcReal3Param param, const AcReal3 value); + +/** */ +AcResult acDeviceLoadIntConstant(const Device device, const Stream stream, const AcIntParam param, + const int value); + +/** */ +AcResult acDeviceLoadInt3Constant(const Device device, const Stream stream, const AcInt3Param param, + const int3 value); /** */ AcResult acDeviceLoadMeshInfo(const Device device, const Stream stream, diff --git a/src/core/device.cu b/src/core/device.cu index b50478a..8b89a78 100644 --- a/src/core/device.cu +++ b/src/core/device.cu @@ -356,8 +356,8 @@ acDeviceSwapBuffers(const Device device) } AcResult -acDeviceLoadConstant(const Device device, const Stream stream, const AcRealParam param, - const AcReal value) +acDeviceLoadScalarConstant(const Device device, const Stream stream, const AcRealParam param, + const AcReal value) { cudaSetDevice(device->id); const size_t offset = (size_t)&d_mesh_info.real_params[param] - (size_t)&d_mesh_info; @@ -366,6 +366,39 @@ acDeviceLoadConstant(const Device device, const Stream stream, const AcRealParam return AC_SUCCESS; } +AcResult +acDeviceLoadVectorConstant(const Device device, const Stream stream, const AcReal3Param param, + const AcReal3 value) +{ + cudaSetDevice(device->id); + const size_t offset = (size_t)&d_mesh_info.real3_params[param] - (size_t)&d_mesh_info; + ERRCHK_CUDA(cudaMemcpyToSymbolAsync(d_mesh_info, &value, sizeof(value), offset, + cudaMemcpyHostToDevice, device->streams[stream])); + return AC_SUCCESS; +} + +AcResult +acDeviceLoadIntConstant(const Device device, const Stream stream, const AcIntParam param, + const int value) +{ + cudaSetDevice(device->id); + const size_t offset = (size_t)&d_mesh_info.int_params[param] - (size_t)&d_mesh_info; + ERRCHK_CUDA(cudaMemcpyToSymbolAsync(d_mesh_info, &value, sizeof(value), offset, + cudaMemcpyHostToDevice, device->streams[stream])); + return AC_SUCCESS; +} + +AcResult +acDeviceLoadInt3Constant(const Device device, const Stream stream, const AcInt3Param param, + const int3 value) +{ + cudaSetDevice(device->id); + const size_t offset = (size_t)&d_mesh_info.int3_params[param] - (size_t)&d_mesh_info; + ERRCHK_CUDA(cudaMemcpyToSymbolAsync(d_mesh_info, &value, sizeof(value), offset, + cudaMemcpyHostToDevice, device->streams[stream])); + return AC_SUCCESS; +} + AcResult acDeviceLoadMeshInfo(const Device device, const Stream stream, const AcMeshInfo device_config) { diff --git a/src/core/node.cu b/src/core/node.cu index 6fa80f2..9950dd4 100644 --- a/src/core/node.cu +++ b/src/core/node.cu @@ -429,7 +429,7 @@ acNodeLoadConstant(const Node node, const Stream stream, const AcRealParam param acNodeSynchronizeStream(node, stream); // #pragma omp parallel for for (int i = 0; i < node->num_devices; ++i) { - acDeviceLoadConstant(node->devices[i], stream, param, value); + acDeviceLoadScalarConstant(node->devices[i], stream, param, value); } return AC_SUCCESS; }