Fix concept proposal for DEVsetparam needed for 'alter' and 'resume' command support

src/include/ngspice/gendefs.h

@@ -22,6 +22,10 @@ struct GENinstance {
     IFuid GENname;  /* pointer to character string naming this instance */
     int GENstate;   /* state index number */
 
+#ifdef USE_CUSPICE
+    int GENcudaIndex ; /* device index for CUDA */
+#endif
+
     /* The actual device instance structs have to place their node elements
      *   right after the the end of struct GENinstance
      *   where they will be accessed by generic GENnode()[]
@@ -47,6 +51,8 @@ struct GENmodel {       /* model structure for a resistor */
 
 #ifdef USE_CUSPICE
     unsigned int has_cuda:1 ;   /* flag to indicate is the model supports CUDA */
+    unsigned int GENinitCUDA:1 ; /* flag to initialize CUDA data */
+    int GENnInstances ; /* number of instances for CUDA */
 #endif
 };
 

src/spicelib/devices/res/CUSPICE/curesload.cu

@@ -47,17 +47,17 @@ GENmodel *inModel, CKTcircuit *ckt
         /* Determining how many blocks should exist in the kernel */
         thread_x = 1 ;
         thread_y = 256 ;
-        if (model->n_instances % thread_y != 0)
+        if (model->RESnInstances % thread_y != 0)
-            block_x = (int)((model->n_instances + thread_y - 1) / thread_y) ;
+            block_x = (int)((model->RESnInstances + thread_y - 1) / thread_y) ;
         else
-            block_x = model->n_instances / thread_y ;
+            block_x = model->RESnInstances / thread_y ;
 
         dim3 thread (thread_x, thread_y) ;
 
         /* Kernel launch */
         status = cudaGetLastError () ; // clear error status
 
-        cuRESload_kernel <<< block_x, thread >>> (model->RESparamGPU, ckt->d_CKTrhsOld, model->n_instances,
+        cuRESload_kernel <<< block_x, thread >>> (model->RESparamGPU, ckt->d_CKTrhsOld, model->RESnInstances,
                                                   model->d_PositionVector, ckt->d_CKTloadOutput) ;
 
         cudaDeviceSynchronize () ;

src/spicelib/devices/res/CUSPICE/curessetup.c

@@ -58,7 +58,7 @@ GENmodel *inModel
     cudaError_t status ;
     RESmodel *model = (RESmodel *)inModel ;
 
-    size = (long unsigned int) model->n_instances;
+    size = (long unsigned int) model->RESnInstances ;
 
     /* Space Allocation to GPU */
     status = cudaMalloc ((void **)&(model->d_PositionVector), size * sizeof(int)) ;

src/spicelib/devices/res/CUSPICE/curestemp.c

@@ -48,7 +48,7 @@ GENmodel *inModel
     cudaError_t status ;
     RESmodel *model = (RESmodel *)inModel ;
 
-    size = (long unsigned int) model->n_instances;
+    size = (long unsigned int) model->RESnInstances ;
 
     /* DOUBLE */
     status = cudaMemcpy (model->RESparamGPU.d_RESconductArray, model->RESparamCPU.RESconductArray, size * sizeof(double), cudaMemcpyHostToDevice) ;

src/spicelib/devices/res/resdefs.h

@@ -95,6 +95,11 @@ typedef struct sRESinstance {
     BindElement *RESnegPosBinding ;
 #endif
 
+#ifdef USE_CUSPICE
+    #define REScudaIndex gen.GENcudaIndex
+    #define RESnInstances gen.GENnInstances
+#endif
+
 } RESinstance ;
 
 #ifdef USE_CUSPICE
@@ -173,7 +178,7 @@ typedef struct sRESmodel {       /* model structure for a resistor */
     int *PositionVector ;
     int *d_PositionVector ;
 
-    int n_instances ;
+    #define RESinitCUDA gen.GENinitCUDA
 #endif
 
 } RESmodel;

src/spicelib/devices/res/resparam.c

@@ -12,6 +12,10 @@ Modified: Apr 2000 - Paolo Nenzi
 #include "ngspice/missing_math.h"
 #include "ngspice/fteext.h"
 
+#ifdef USE_CUSPICE
+#include "ngspice/CUSPICE/CUSPICE.h"
+#endif
+
 int
 RESparam(int param, IFvalue *value, GENinstance *inst, IFvalue *select)
 {
@@ -86,5 +90,19 @@ RESparam(int param, IFvalue *value, GENinstance *inst, IFvalue *select)
         return(E_BADPARM);
     }
     RESupdate_conduct(here, FALSE);
+
+#ifdef USE_CUSPICE
+    int status ;
+    RESmodel *model ;
+
+    model = RESmodPtr(here) ;
+    if (model->RESinitCUDA) {
+        model->RESparamCPU.RESconductArray[here->REScudaIndex] = here->RESconduct;
+        status = cuREStemp ((GENmodel *)model);
+        if (status != 0)
+            return E_NOMEM;
+    }
+#endif
+
     return(OK);
 }

src/spicelib/devices/res/ressetup.c

@@ -80,28 +80,14 @@ do { if((here->ptr = SMPmakeElt(matrix, here->first, here->second)) == NULL){\
     }
 
 #ifdef USE_CUSPICE
-    int i, j, status ;
+    int j, status ;
-
-    /* Counting the instances */
-    for (model = (RESmodel *)inModel ; model != NULL ; model = RESnextModel(model))
-    {
-        i = 0 ;
-
-        for (here = RESinstances(model); here != NULL ; here = RESnextInstance(here))
-        {
-            i++ ;
-        }
-
-        /* How much instances we have */
-        model->n_instances = i ;
-
-        /* This model supports CUDA */
-        model->gen.has_cuda = 1 ;
-    }
 
     /*  loop through all the resistor models */
     for (model = (RESmodel *)inModel ; model != NULL ; model = RESnextModel(model))
     {
+        /* This model supports CUDA */
+        model->gen.has_cuda = 1 ;
+
         model->offset = ckt->total_n_values ;
 
         j = 0 ;
@@ -122,7 +108,7 @@ do { if((here->ptr = SMPmakeElt(matrix, here->first, here->second)) == NULL){\
                 j++ ;
         }
 
-        model->n_values = model->n_instances ;
+        model->n_values = model->RESnInstances ;
         ckt->total_n_values += model->n_values ;
 
         model->n_Ptr = j ;
@@ -130,18 +116,20 @@ do { if((here->ptr = SMPmakeElt(matrix, here->first, here->second)) == NULL){\
 
 
         /* Position Vector assignment */
-        model->PositionVector = TMALLOC (int, model->n_instances) ;
+        model->PositionVector = TMALLOC (int, model->RESnInstances) ;
 
-        for (j = 0 ; j < model->n_instances ; j++)
+        for (j = 0 ; j < model->RESnInstances ; j++)
             model->PositionVector [j] = model->offset + j ;
     }
 
-    /*  loop through all the resistor models */
+    /* loop through all the resistor models */
     for (model = (RESmodel *)inModel ; model != NULL ; model = RESnextModel(model))
     {
         status = cuRESsetup ((GENmodel *)model) ;
         if (status != 0)
             return (E_NOMEM) ;
+
+        model->RESinitCUDA = 1 ;
     }
 #endif
 

src/spicelib/devices/res/restemp.c

@@ -30,10 +30,6 @@ REStemp(GENmodel *inModel, CKTcircuit *ckt)
     /*  loop through all the resistor models */
     for( ; model != NULL; model = RESnextModel(model)) {
 
-#ifdef USE_CUSPICE
-        int i = 0;
-#endif
-
         /* loop through all the instances of the model */
         for (here = RESinstances(model); here != NULL ;
                 here=RESnextInstance(here)) {
@@ -53,12 +49,11 @@ REStemp(GENmodel *inModel, CKTcircuit *ckt)
             RESupdate_conduct(here, TRUE);
 
 #ifdef USE_CUSPICE
-            model->RESparamCPU.RESposNodeArray[i] = here->RESposNode;
+            model->RESparamCPU.RESposNodeArray[here->REScudaIndex] = here->RESposNode;
-            model->RESparamCPU.RESnegNodeArray[i] = here->RESnegNode;
+            model->RESparamCPU.RESnegNodeArray[here->REScudaIndex] = here->RESnegNode;
-            model->RESparamCPU.RESconductArray[i] = here->RESconduct;
+            model->RESparamCPU.RESconductArray[here->REScudaIndex] = here->RESconduct;
-
-            i++;
 #endif
+
         }
 
 #ifdef USE_CUSPICE

src/spicelib/parser/inp2r.c

@@ -176,8 +176,20 @@ void INP2R(CKTcircuit *ckt, INPtables * tab, struct card *current)
           /* create default R model */
           IFnewUid(ckt, &uid, NULL, "R", UID_MODEL, NULL);
           IFC(newModel, (ckt, type, &(tab->defRmod), uid));
+
+#ifdef USE_CUSPICE
+          tab->defRmod->GENnInstances = 0 ;
+          tab->defRmod->GENinitCUDA = 0 ;
+#endif
+
       }
       IFC(newInstance, (ckt, tab->defRmod, &fast, name));
+
+#ifdef USE_CUSPICE
+      fast->GENcudaIndex = tab->defRmod->GENnInstances ;
+      tab->defRmod->GENnInstances++ ;
+#endif
+
       if (error1 == 1) {		/* was a r=val construction */
         val = INPevaluate(&line, &error1, 1);	/* [<val>] */
 #ifdef TRACE