From 08f155cbecacf0fdf45934b73070b9d895fe6641 Mon Sep 17 00:00:00 2001 From: jpekkila Date: Mon, 7 Oct 2019 20:40:32 +0300 Subject: [PATCH] Finetuning some error checks --- src/core/device.cu | 18 ++++++++++++++++-- src/core/kernels/integration.cuh | 3 +-- 2 files changed, 17 insertions(+), 4 deletions(-) diff --git a/src/core/device.cu b/src/core/device.cu index 17e87f7..3dcdd24 100644 --- a/src/core/device.cu +++ b/src/core/device.cu @@ -351,9 +351,15 @@ acDeviceAutoOptimize(const Device device) cudaEventCreate(&tstart); cudaEventCreate(&tstop); - cudaEventRecord(tstart); // ---------------------------------------- Timing start - + // #ifdef AC_dt 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_ which does + not require the " "timestep to be defined.\n"); #endif*/ + + cudaEventRecord(tstart); // ---------------------------------------- Timing start for (int i = 0; i < num_iterations; ++i) solve<2><<>>(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.z / AcReal(tpb.z))); + //#ifdef AC_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_ which does not require the " + "timestep to be defined.\n"); + #endif*/ if (step_number == 0) solve<0><<streams[stream]>>>(start, end, device->vba); else if (step_number == 1) diff --git a/src/core/kernels/integration.cuh b/src/core/kernels/integration.cuh index 6524903..695bb2a 100644 --- a/src/core/kernels/integration.cuh +++ b/src/core/kernels/integration.cuh @@ -29,8 +29,7 @@ #include -static_assert(NUM_VTXBUF_HANDLES > 0, - "ERROR: NUM_VTXBUF_HANDLES (uniform ScalarFields) must be > 0"); +static_assert(NUM_VTXBUF_HANDLES > 0, "ERROR: At least one uniform ScalarField must be declared."); static __device__ constexpr int IDX(const int i)