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

Added functions for loading int, int3, scalar and vector constants to the device layer (acDeviceLoad...Constant)

Browse Source
This commit is contained in:
jpekkila
2019-08-19 15:28:16 +03:00
parent 41805dcb68
commit 787363226b
3 changed files with 50 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files

16
include/astaroth_device.h
View File

@@ -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,

37
src/core/device.cu
View File

@@ -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)
{

2
src/core/node.cu
View File

@@ -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;
}