From abed2083e1182b94aeb00aca9f38f10ad70e9024 Mon Sep 17 00:00:00 2001
From: Francesco Lannutti <nicolati@users.sourceforge.net>
Date: Wed, 16 Sep 2015 19:54:26 +0200
Subject: [PATCH] Fixed a bug in CKTtrunc reduction when there is nothing to
 reduce

---
 src/spicelib/analysis/CUSPICE/cuckttrunc.cu |  9 +++++++-
 src/spicelib/analysis/ckttrunc.c            | 23 ++++++++++++++++-----
 2 files changed, 26 insertions(+), 6 deletions(-)

diff --git a/src/spicelib/analysis/CUSPICE/cuckttrunc.cu b/src/spicelib/analysis/CUSPICE/cuckttrunc.cu
index f93eb7667..b215b730f 100644
--- a/src/spicelib/analysis/CUSPICE/cuckttrunc.cu
+++ b/src/spicelib/analysis/CUSPICE/cuckttrunc.cu
@@ -54,6 +54,13 @@ CKTcircuit *ckt, double timetemp, double *timeStep
 
     cudaDeviceSynchronize () ;
 
+    status = cudaGetLastError () ; // check for launch error
+    if (status != cudaSuccess)
+    {
+        fprintf (stderr, "Kernel 1 launch failure in cuCKTtrunc\n\n") ;
+        return (E_NOMEM) ;
+    }
+
     cuCKTtrunc_kernel <<< 1, thread, thread_y * sizeof(double) >>> (ckt->d_CKTtimeStepsOut, ckt->d_CKTtimeSteps, block_x) ;
 
     cudaDeviceSynchronize () ;
@@ -61,7 +68,7 @@ CKTcircuit *ckt, double timetemp, double *timeStep
     status = cudaGetLastError () ; // check for launch error
     if (status != cudaSuccess)
     {
-        fprintf (stderr, "Kernel launch failure in cuCKTtrunc\n\n") ;
+        fprintf (stderr, "Kernel 2 launch failure in cuCKTtrunc\n\n") ;
         return (E_NOMEM) ;
     }
 
diff --git a/src/spicelib/analysis/ckttrunc.c b/src/spicelib/analysis/ckttrunc.c
index bd2a75c40..c78adcc03 100644
--- a/src/spicelib/analysis/ckttrunc.c
+++ b/src/spicelib/analysis/ckttrunc.c
@@ -33,8 +33,16 @@ CKTtrunc (CKTcircuit *ckt, double *timeStep)
     double startTime ;
     int error = OK ;
 
+#ifdef USE_CUSPICE
+    int doReduction, status ;
+#endif
+
     startTime = SPfrontEnd->IFseconds () ;
 
+#ifdef USE_CUSPICE
+    doReduction = 0 ;
+#endif
+
     timetemp = HUGE ;
     for (i = 0 ; i < DEVmaxnum ; i++)
     {
@@ -45,6 +53,10 @@ CKTtrunc (CKTcircuit *ckt, double *timeStep)
             debugtemp = timetemp ;
 #endif /* STEPDEBUG */
 
+#ifdef USE_CUSPICE
+            doReduction = 1 ;
+#endif
+
 	    error = DEVices[i]->DEVtrunc (ckt->CKThead[i], ckt, &timetemp) ;
 	    if (error)
             {
@@ -63,11 +75,12 @@ CKTtrunc (CKTcircuit *ckt, double *timeStep)
     }
 
 #ifdef USE_CUSPICE
-    int status ;
-
-    status = cuCKTtrunc (ckt, HUGE, timeStep) ;
-    if (status != 0)
-        return (E_NOMEM) ;
+    if (doReduction)
+    {
+        status = cuCKTtrunc (ckt, HUGE, timeStep) ;
+        if (status != 0)
+            return (E_NOMEM) ;
+    }
 #else
     *timeStep = MIN (2 * *timeStep, timetemp) ;
 #endif