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

Fixed some stupid logic when loading uniforms

Browse Source
This commit is contained in:
jpekkila
2020-08-21 20:23:50 +03:00
parent 764e4dda69
commit 9f676a6d5d
1 changed files with 5 additions and 13 deletions
Download Patch File Download Diff File Expand all files Collapse all files

18
src/core/kernels/kernels.cu
View File

@@ -105,7 +105,8 @@ acDeviceLoadScalarUniform(const Device device, const Stream stream, const AcReal
{ {
cudaSetDevice(device->id); cudaSetDevice(device->id);
if (param < 0 || param >= NUM_REAL_PARAMS) { if (param < 0 || param >= NUM_REAL_PARAMS) {
fprintf(stderr, "WARNING: invalid AcRealParam %d\n", param); fprintf(stderr, "WARNING: invalid AcRealParam %d. Skipping.\n", param);
return AC_FAILURE;
} }
if (!is_valid(value)) { if (!is_valid(value)) {
@@ -114,9 +115,6 @@ acDeviceLoadScalarUniform(const Device device, const Stream stream, const AcReal
return AC_FAILURE; return AC_FAILURE;
} }
if (param >= NUM_REAL_PARAMS)
return AC_FAILURE;
const size_t offset = (size_t)&d_mesh_info.real_params[param] - (size_t)&d_mesh_info; const size_t offset = (size_t)&d_mesh_info.real_params[param] - (size_t)&d_mesh_info;
ERRCHK_CUDA(cudaMemcpyToSymbolAsync(d_mesh_info, &value, sizeof(value), offset, ERRCHK_CUDA(cudaMemcpyToSymbolAsync(d_mesh_info, &value, sizeof(value), offset,
cudaMemcpyHostToDevice, device->streams[stream])); cudaMemcpyHostToDevice, device->streams[stream]));
@@ -130,6 +128,7 @@ acDeviceLoadVectorUniform(const Device device, const Stream stream, const AcReal
cudaSetDevice(device->id); cudaSetDevice(device->id);
if (param < 0 || param >= NUM_REAL3_PARAMS) { if (param < 0 || param >= NUM_REAL3_PARAMS) {
fprintf(stderr, "WARNING: invalid AcReal3Param %d\n", param); fprintf(stderr, "WARNING: invalid AcReal3Param %d\n", param);
return AC_FAILURE;
} }
if (!is_valid(value)) { if (!is_valid(value)) {
@@ -139,9 +138,6 @@ acDeviceLoadVectorUniform(const Device device, const Stream stream, const AcReal
return AC_FAILURE; return AC_FAILURE;
} }
if (param >= NUM_REAL3_PARAMS || !NUM_REAL3_PARAMS)
return AC_FAILURE;
const size_t offset = (size_t)&d_mesh_info.real3_params[param] - (size_t)&d_mesh_info; 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, ERRCHK_CUDA(cudaMemcpyToSymbolAsync(d_mesh_info, &value, sizeof(value), offset,
cudaMemcpyHostToDevice, device->streams[stream])); cudaMemcpyHostToDevice, device->streams[stream]));
@@ -155,6 +151,7 @@ acDeviceLoadIntUniform(const Device device, const Stream stream, const AcIntPara
cudaSetDevice(device->id); cudaSetDevice(device->id);
if (param < 0 || param >= NUM_INT_PARAMS) { if (param < 0 || param >= NUM_INT_PARAMS) {
fprintf(stderr, "WARNING: invalid AcIntParam %d\n", param); fprintf(stderr, "WARNING: invalid AcIntParam %d\n", param);
return AC_FAILURE;
} }
if (!is_valid(value)) { if (!is_valid(value)) {
@@ -163,9 +160,6 @@ acDeviceLoadIntUniform(const Device device, const Stream stream, const AcIntPara
return AC_FAILURE; return AC_FAILURE;
} }
if (param >= NUM_INT_PARAMS)
return AC_FAILURE;
const size_t offset = (size_t)&d_mesh_info.int_params[param] - (size_t)&d_mesh_info; 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, ERRCHK_CUDA(cudaMemcpyToSymbolAsync(d_mesh_info, &value, sizeof(value), offset,
cudaMemcpyHostToDevice, device->streams[stream])); cudaMemcpyHostToDevice, device->streams[stream]));
@@ -179,6 +173,7 @@ acDeviceLoadInt3Uniform(const Device device, const Stream stream, const AcInt3Pa
cudaSetDevice(device->id); cudaSetDevice(device->id);
if (param < 0 || param >= NUM_INT3_PARAMS) { if (param < 0 || param >= NUM_INT3_PARAMS) {
fprintf(stderr, "WARNING: invalid AcInt3Param %d\n", param); fprintf(stderr, "WARNING: invalid AcInt3Param %d\n", param);
return AC_FAILURE;
} }
if (!is_valid(value.x) || !is_valid(value.y) || !is_valid(value.z)) { if (!is_valid(value.x) || !is_valid(value.y) || !is_valid(value.z)) {
@@ -189,9 +184,6 @@ acDeviceLoadInt3Uniform(const Device device, const Stream stream, const AcInt3Pa
return AC_FAILURE; return AC_FAILURE;
} }
if (param >= NUM_INT3_PARAMS)
return AC_FAILURE;
const size_t offset = (size_t)&d_mesh_info.int3_params[param] - (size_t)&d_mesh_info; 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, ERRCHK_CUDA(cudaMemcpyToSymbolAsync(d_mesh_info, &value, sizeof(value), offset,
cudaMemcpyHostToDevice, device->streams[stream])); cudaMemcpyHostToDevice, device->streams[stream]));