luke
/
ngspice
mirror of https://git.code.sf.net/p/ngspice/ngspice
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

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

This commit is contained in:
Francesco Lannutti 2018-04-28 12:29:15 +02:00
parent fcb5163a95
commit 958e209866
9 changed files with 62 additions and 38 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
src/include/ngspice/gendefs.h
View File

@ -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
};

8
src/spicelib/devices/res/CUSPICE/curesload.cu
View File

@ -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 () ;

2
src/spicelib/devices/res/CUSPICE/curessetup.c
View File

@ -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)) ;

2
src/spicelib/devices/res/CUSPICE/curestemp.c
View File

@ -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) ;

7
src/spicelib/devices/res/resdefs.h
View File

@ -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;

18
src/spicelib/devices/res/resparam.c
View File

@ -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);
}

32
src/spicelib/devices/res/ressetup.c
View File

@ -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

13
src/spicelib/devices/res/restemp.c
View File

@ -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

12
src/spicelib/parser/inp2r.c
View File

@ -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