make res changes more generic

src/spicelib/analysis/dctrcurv.c

@@ -486,8 +486,8 @@ DCtrCurv(CKTcircuit *ckt, int restart)
 
 #ifdef USE_CUSPICE
             RESmodel *model = RESmodPtr(here) ;
-            if (model->RESinitCUDA) {
+            if (model->gen.GENinitCUDA) {
-                model->RESparamCPU.RESconductArray[here->REScudaIndex] = here->RESconduct;
+                model->RESparamCPU.RESconductArray[here->gen.GENcudaIndex] = here->RESconduct;
                 status = cuREStemp ((GENmodel *)model);
                 if (status != 0)
                     return E_NOMEM;

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->RESnInstances % thread_y != 0)
+        if (model->gen.GENnInstances % thread_y != 0)
-            block_x = (int)((model->RESnInstances + thread_y - 1) / thread_y) ;
+            block_x = (int)((model->gen.GENnInstances + thread_y - 1) / thread_y) ;
         else
-            block_x = model->RESnInstances / thread_y ;
+            block_x = model->gen.GENnInstances / 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->RESnInstances,
+        cuRESload_kernel <<< block_x, thread >>> (model->RESparamGPU, ckt->d_CKTrhsOld, model->gen.GENnInstances,
                                                   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->RESnInstances ;
+    size = (long unsigned int) model->gen.GENnInstances ;
 
     /* 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->RESnInstances ;
+    size = (long unsigned int) model->gen.GENnInstances ;
 
     /* DOUBLE */
     status = cudaMemcpy (model->RESparamGPU.d_RESconductArray, model->RESparamCPU.RESconductArray, size * sizeof(double), cudaMemcpyHostToDevice) ;

src/spicelib/devices/res/resdefs.h

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

src/spicelib/devices/res/resparam.c

@@ -96,8 +96,8 @@ RESparam(int param, IFvalue *value, GENinstance *inst, IFvalue *select)
     RESmodel *model ;
 
     model = RESmodPtr(here) ;
-    if (model->RESinitCUDA) {
+    if (model->gen.GENinitCUDA) {
-        model->RESparamCPU.RESconductArray[here->REScudaIndex] = here->RESconduct;
+        model->RESparamCPU.RESconductArray[here->gen.GENcudaIndex] = here->RESconduct;
         status = cuREStemp ((GENmodel *)model);
         if (status != 0)
             return E_NOMEM;

src/spicelib/devices/res/ressetup.c

@@ -108,7 +108,7 @@ do { if((here->ptr = SMPmakeElt(matrix, here->first, here->second)) == NULL){\
                 j++ ;
         }
 
-        model->n_values = model->RESnInstances ;
+        model->n_values = model->gen.GENnInstances ;
         ckt->total_n_values += model->n_values ;
 
         model->n_Ptr = j ;
@@ -116,9 +116,9 @@ do { if((here->ptr = SMPmakeElt(matrix, here->first, here->second)) == NULL){\
 
 
         /* Position Vector assignment */
-        model->PositionVector = TMALLOC (int, model->RESnInstances) ;
+        model->PositionVector = TMALLOC (int, model->gen.GENnInstances) ;
 
-        for (j = 0 ; j < model->RESnInstances ; j++)
+        for (j = 0 ; j < model->gen.GENnInstances ; j++)
             model->PositionVector [j] = model->offset + j ;
     }
 
@@ -129,7 +129,7 @@ do { if((here->ptr = SMPmakeElt(matrix, here->first, here->second)) == NULL){\
         if (status != 0)
             return (E_NOMEM) ;
 
-        model->RESinitCUDA = 1 ;
+        model->gen.GENinitCUDA = 1 ;
     }
 #endif
 

src/spicelib/devices/res/restemp.c

@@ -49,9 +49,10 @@ REStemp(GENmodel *inModel, CKTcircuit *ckt)
             RESupdate_conduct(here, TRUE);
 
 #ifdef USE_CUSPICE
-            model->RESparamCPU.RESposNodeArray[here->REScudaIndex] = here->RESposNode;
+            int i = here->gen.GENcudaIndex;
-            model->RESparamCPU.RESnegNodeArray[here->REScudaIndex] = here->RESnegNode;
+            model->RESparamCPU.RESposNodeArray[i] = here->RESposNode;
-            model->RESparamCPU.RESconductArray[here->REScudaIndex] = here->RESconduct;
+            model->RESparamCPU.RESnegNodeArray[i] = here->RESnegNode;
+            model->RESparamCPU.RESconductArray[i] = here->RESconduct;
 #endif
 
         }