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

Finetuning some error checks

Browse Source
This commit is contained in:
jpekkila
2019-10-07 20:40:32 +03:00
parent ea4438f331
commit 08f155cbec
2 changed files with 17 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files

18
src/core/device.cu
View File

@@ -351,9 +351,15 @@ acDeviceAutoOptimize(const Device device)
cudaEventCreate(&tstart); cudaEventCreate(&tstart);
cudaEventCreate(&tstop); cudaEventCreate(&tstop);
cudaEventRecord(tstart); // ---------------------------------------- Timing start // #ifdef AC_dt
acDeviceLoadScalarUniform(device, STREAM_DEFAULT, AC_dt, FLT_EPSILON); acDeviceLoadScalarUniform(device, STREAM_DEFAULT, AC_dt, FLT_EPSILON);
/*#else
ERROR("FATAL ERROR: acDeviceAutoOptimize() or
acDeviceIntegrateSubstep() was " "called, but AC_dt was not defined. Either define
it or call the generated " "device function acDeviceKernel_<kernel name> which does
not require the " "timestep to be defined.\n"); #endif*/
cudaEventRecord(tstart); // ---------------------------------------- Timing start
for (int i = 0; i < num_iterations; ++i) for (int i = 0; i < num_iterations; ++i)
solve<2><<<bpg, tpb>>>(start, end, device->vba); solve<2><<<bpg, tpb>>>(start, end, device->vba);
@@ -666,7 +672,15 @@ acDeviceIntegrateSubstep(const Device device, const Stream stream, const int ste
(unsigned int)ceil(n.y / AcReal(tpb.y)), // (unsigned int)ceil(n.y / AcReal(tpb.y)), //
(unsigned int)ceil(n.z / AcReal(tpb.z))); (unsigned int)ceil(n.z / AcReal(tpb.z)));
//#ifdef AC_dt
acDeviceLoadScalarUniform(device, stream, AC_dt, dt); acDeviceLoadScalarUniform(device, stream, AC_dt, dt);
/*#else
(void)dt;
ERROR("FATAL ERROR: acDeviceAutoOptimize() or acDeviceIntegrateSubstep() was "
"called, but AC_dt was not defined. Either define it or call the generated "
"device function acDeviceKernel_<kernel name> which does not require the "
"timestep to be defined.\n");
#endif*/
if (step_number == 0) if (step_number == 0)
solve<0><<<bpg, tpb, 0, device->streams[stream]>>>(start, end, device->vba); solve<0><<<bpg, tpb, 0, device->streams[stream]>>>(start, end, device->vba);
else if (step_number == 1) else if (step_number == 1)

3
src/core/kernels/integration.cuh
View File

@@ -29,8 +29,7 @@
#include <assert.h> #include <assert.h>
static_assert(NUM_VTXBUF_HANDLES > 0, static_assert(NUM_VTXBUF_HANDLES > 0, "ERROR: At least one uniform ScalarField must be declared.");
"ERROR: NUM_VTXBUF_HANDLES (uniform ScalarFields) must be > 0");
static __device__ constexpr int static __device__ constexpr int
IDX(const int i) IDX(const int i)