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

make res changes more generic

This commit is contained in:
rlar 2018-04-28 22:58:00 +02:00
parent 0b23419e0f
commit cff3e65ab2
8 changed files with 18 additions and 24 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
src/spicelib/analysis/dctrcurv.c
View File

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

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

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->RESnInstances ;
size = (long unsigned int) model->gen.GENnInstances ;
/* 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->RESnInstances ;
size = (long unsigned int) model->gen.GENnInstances ;
/* DOUBLE */
status = cudaMemcpy (model->RESparamGPU.d_RESconductArray, model->RESparamCPU.RESconductArray, size * sizeof(double), cudaMemcpyHostToDevice) ;

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

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

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

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

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

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

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

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