From 958e209866eaccb7386dab33caa5c1b12a87f29b Mon Sep 17 00:00:00 2001 From: Francesco Lannutti Date: Sat, 28 Apr 2018 12:29:15 +0200 Subject: [PATCH] Fix concept proposal for DEVsetparam needed for 'alter' and 'resume' command support --- src/include/ngspice/gendefs.h | 6 ++++ src/spicelib/devices/res/CUSPICE/curesload.cu | 8 ++--- src/spicelib/devices/res/CUSPICE/curessetup.c | 2 +- src/spicelib/devices/res/CUSPICE/curestemp.c | 2 +- src/spicelib/devices/res/resdefs.h | 7 +++- src/spicelib/devices/res/resparam.c | 18 +++++++++++ src/spicelib/devices/res/ressetup.c | 32 ++++++------------- src/spicelib/devices/res/restemp.c | 13 +++----- src/spicelib/parser/inp2r.c | 12 +++++++ 9 files changed, 62 insertions(+), 38 deletions(-) diff --git a/src/include/ngspice/gendefs.h b/src/include/ngspice/gendefs.h index ac3a25d6d..ff4720df7 100644 --- a/src/include/ngspice/gendefs.h +++ b/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 }; diff --git a/src/spicelib/devices/res/CUSPICE/curesload.cu b/src/spicelib/devices/res/CUSPICE/curesload.cu index 1170d3338..4ff523b62 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->n_instances % thread_y != 0) - block_x = (int)((model->n_instances + thread_y - 1) / thread_y) ; + if (model->RESnInstances % thread_y != 0) + 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 () ; diff --git a/src/spicelib/devices/res/CUSPICE/curessetup.c b/src/spicelib/devices/res/CUSPICE/curessetup.c index ddcebf5cb..6bad9d323 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->n_instances; + size = (long unsigned int) model->RESnInstances ; /* 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 957cf732a..17a313ac1 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->n_instances; + size = (long unsigned int) model->RESnInstances ; /* 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 6d394f8d2..325b873ff 100644 --- a/src/spicelib/devices/res/resdefs.h +++ b/src/spicelib/devices/res/resdefs.h @@ -93,6 +93,11 @@ typedef struct sRESinstance { BindElement *RESnegPosBinding ; #endif +#ifdef USE_CUSPICE + #define REScudaIndex gen.GENcudaIndex + #define RESnInstances gen.GENnInstances +#endif + } RESinstance ; #ifdef USE_CUSPICE @@ -171,7 +176,7 @@ typedef struct sRESmodel { /* model structure for a resistor */ int *PositionVector ; int *d_PositionVector ; - int n_instances ; + #define RESinitCUDA gen.GENinitCUDA #endif } RESmodel; diff --git a/src/spicelib/devices/res/resparam.c b/src/spicelib/devices/res/resparam.c index 1385a71b9..2ab559160 100644 --- a/src/spicelib/devices/res/resparam.c +++ b/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); } diff --git a/src/spicelib/devices/res/ressetup.c b/src/spicelib/devices/res/ressetup.c index d1fa24273..cebe8ecde 100644 --- a/src/spicelib/devices/res/ressetup.c +++ b/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 ; - - /* 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 ; - } + int j, status ; /* 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 diff --git a/src/spicelib/devices/res/restemp.c b/src/spicelib/devices/res/restemp.c index 57656a77e..5ce6e82bd 100644 --- a/src/spicelib/devices/res/restemp.c +++ b/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.RESnegNodeArray[i] = here->RESnegNode; - model->RESparamCPU.RESconductArray[i] = here->RESconduct; - - i++; + model->RESparamCPU.RESposNodeArray[here->REScudaIndex] = here->RESposNode; + model->RESparamCPU.RESnegNodeArray[here->REScudaIndex] = here->RESnegNode; + model->RESparamCPU.RESconductArray[here->REScudaIndex] = here->RESconduct; #endif + } #ifdef USE_CUSPICE diff --git a/src/spicelib/parser/inp2r.c b/src/spicelib/parser/inp2r.c index 1f5bf28b2..ceb4e43f6 100644 --- a/src/spicelib/parser/inp2r.c +++ b/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); /* [] */ #ifdef TRACE