diff --git a/src/spicelib/analysis/dctrcurv.c b/src/spicelib/analysis/dctrcurv.c index 9eae8bfe0..24a610fb7 100644 --- a/src/spicelib/analysis/dctrcurv.c +++ b/src/spicelib/analysis/dctrcurv.c @@ -486,8 +486,8 @@ DCtrCurv(CKTcircuit *ckt, int restart) #ifdef USE_CUSPICE RESmodel *model = RESmodPtr(here) ; - if (model->RESinitCUDA) { - model->RESparamCPU.RESconductArray[here->REScudaIndex] = here->RESconduct; + if (model->gen.GENinitCUDA) { + model->RESparamCPU.RESconductArray[here->gen.GENcudaIndex] = here->RESconduct; status = cuREStemp ((GENmodel *)model); if (status != 0) return E_NOMEM; diff --git a/src/spicelib/devices/res/CUSPICE/curesload.cu b/src/spicelib/devices/res/CUSPICE/curesload.cu index 4ff523b62..ba20d5db3 100644 --- a/src/spicelib/devices/res/CUSPICE/curesload.cu +++ b/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) - block_x = (int)((model->RESnInstances + thread_y - 1) / thread_y) ; + if (model->gen.GENnInstances % thread_y != 0) + 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 () ; diff --git a/src/spicelib/devices/res/CUSPICE/curessetup.c b/src/spicelib/devices/res/CUSPICE/curessetup.c index 6bad9d323..5f6d74071 100644 --- a/src/spicelib/devices/res/CUSPICE/curessetup.c +++ b/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)) ; diff --git a/src/spicelib/devices/res/CUSPICE/curestemp.c b/src/spicelib/devices/res/CUSPICE/curestemp.c index 17a313ac1..209732e83 100644 --- a/src/spicelib/devices/res/CUSPICE/curestemp.c +++ b/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) ; diff --git a/src/spicelib/devices/res/resdefs.h b/src/spicelib/devices/res/resdefs.h index 325b873ff..c5778fb00 100644 --- a/src/spicelib/devices/res/resdefs.h +++ b/src/spicelib/devices/res/resdefs.h @@ -93,11 +93,6 @@ typedef struct sRESinstance { BindElement *RESnegPosBinding ; #endif -#ifdef USE_CUSPICE - #define REScudaIndex gen.GENcudaIndex - #define RESnInstances gen.GENnInstances -#endif - } RESinstance ; #ifdef USE_CUSPICE @@ -175,8 +170,6 @@ typedef struct sRESmodel { /* model structure for a resistor */ int n_Ptr ; int *PositionVector ; int *d_PositionVector ; - - #define RESinitCUDA gen.GENinitCUDA #endif } RESmodel; diff --git a/src/spicelib/devices/res/resparam.c b/src/spicelib/devices/res/resparam.c index 2ab559160..41b0115a4 100644 --- a/src/spicelib/devices/res/resparam.c +++ b/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) { - model->RESparamCPU.RESconductArray[here->REScudaIndex] = here->RESconduct; + if (model->gen.GENinitCUDA) { + model->RESparamCPU.RESconductArray[here->gen.GENcudaIndex] = here->RESconduct; status = cuREStemp ((GENmodel *)model); if (status != 0) return E_NOMEM; diff --git a/src/spicelib/devices/res/ressetup.c b/src/spicelib/devices/res/ressetup.c index cebe8ecde..f9c779f5b 100644 --- a/src/spicelib/devices/res/ressetup.c +++ b/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 diff --git a/src/spicelib/devices/res/restemp.c b/src/spicelib/devices/res/restemp.c index 5ce6e82bd..cb15d9769 100644 --- a/src/spicelib/devices/res/restemp.c +++ b/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; - model->RESparamCPU.RESnegNodeArray[here->REScudaIndex] = here->RESnegNode; - model->RESparamCPU.RESconductArray[here->REScudaIndex] = here->RESconduct; + int i = here->gen.GENcudaIndex; + model->RESparamCPU.RESposNodeArray[i] = here->RESposNode; + model->RESparamCPU.RESnegNodeArray[i] = here->RESnegNode; + model->RESparamCPU.RESconductArray[i] = here->RESconduct; #endif }