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

use generic gen.GENcudaIndex ... 

carefull, this commit is quick and dirty,
  created in a half automated way, low confidence.

This is missing in this commit
   model->gen.GENinitCUDA
This commit is contained in:
rlar 2018-04-29 11:17:14 +02:00
parent ec2c710ecc
commit 65ff4f46e1
40 changed files with 122 additions and 229 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
src/spicelib/devices/bsim4v7/CUSPICE/cubsim4v7getic.c
View File

@ -48,7 +48,7 @@ GENmodel *inModel
cudaError_t status ;
BSIM4v7model *model = (BSIM4v7model *)inModel ;
size = (long unsigned int) model->n_instances;
size = (long unsigned int) model->gen.GENnInstances;
status = cudaMemcpy (model->BSIM4v7paramGPU.d_BSIM4v7icVDSArray, model->BSIM4v7paramCPU.BSIM4v7icVDSArray, size * sizeof(double), cudaMemcpyHostToDevice) ;
CUDAMEMCPYCHECK (model->BSIM4v7paramGPU.d_BSIM4v7icVDSArray, size, double, status)

8
src/spicelib/devices/bsim4v7/CUSPICE/cubsim4v7load.cu
View File

@ -274,10 +274,10 @@ 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->gen.GENnInstances % thread_y != 0)
block_x = (int)((model->gen.GENnInstances + thread_y - 1) / thread_y) ;
else
block_x = model->n_instances / thread_y ;
block_x = model->gen.GENnInstances / thread_y ;
dim3 thread (thread_x, thread_y) ;
@ -290,7 +290,7 @@ GENmodel *inModel, CKTcircuit *ckt
ckt->CKTdelta, ckt->CKTdeltaOld [1], ckt->CKTreltol,
ckt->CKTvoltTol, ckt->CKTabstol, ckt->CKTag [0], ckt->CKTag [1],
ckt->CKTgmin, ckt->CKTbypass, ckt->CKTmode, ckt->CKTorder,
model->n_instances, ckt->CKTtemp, ckt->d_CKTnoncon, CONSTvt0,
model->gen.GENnInstances, ckt->CKTtemp, ckt->d_CKTnoncon, CONSTvt0,
/* Model */
model->BSIM4v7type, model->BSIM4v7rdsMod, model->BSIM4v7igcMod,
model->BSIM4v7igbMod, model->BSIM4v7vcrit, model->BSIM4v7vtm,

2
src/spicelib/devices/bsim4v7/CUSPICE/cubsim4v7setup.c
View File

@ -58,7 +58,7 @@ GENmodel *inModel
cudaError_t status ;
BSIM4v7model *model = (BSIM4v7model *)inModel ;
size = (long unsigned int) model->n_instances;
size = (long unsigned int) model->gen.GENnInstances;
/* Space Allocation to GPU */
status = cudaMalloc ((void **)&(model->d_PositionVector), size * sizeof(int)) ;

2
src/spicelib/devices/bsim4v7/CUSPICE/cubsim4v7temp.c
View File

@ -60,7 +60,7 @@ GENmodel *inModel
BSIM4v7model *model = (BSIM4v7model *)inModel ;
BSIM4v7instance *here ;
size = (long unsigned int) model->n_instances;
size = (long unsigned int) model->gen.GENnInstances;
/* Special case here->d_pParam */
model->pParamHost = (struct bsim4SizeDependParam **) malloc (size * sizeof(struct bsim4SizeDependParam *)) ;

8
src/spicelib/devices/bsim4v7/CUSPICE/cubsim4v7trunc.cu
View File

@ -31,17 +31,17 @@ GENmodel *inModel, CKTcircuit *ckt, double *timeStep
/* 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->gen.GENnInstances % thread_y != 0)
block_x = (int)((model->gen.GENnInstances + thread_y - 1) / thread_y) ;
else
block_x = model->n_instances / thread_y ;
block_x = model->gen.GENnInstances / thread_y ;
dim3 thread (thread_x, thread_y) ;
/* Kernel launch */
status = cudaGetLastError () ; // clear error status
cuBSIM4v7trunc_kernel <<< block_x, thread >>> (model->BSIM4v7paramGPU, model->n_instances,
cuBSIM4v7trunc_kernel <<< block_x, thread >>> (model->BSIM4v7paramGPU, model->gen.GENnInstances,
ckt->dD_CKTstates, ckt->d_CKTdeltaOld,
ckt->CKTdelta, ckt->CKTorder, ckt->CKTintegrateMethod,
ckt->CKTabstol, ckt->CKTreltol, ckt->CKTchgtol, ckt->CKTtrtol,

9
src/spicelib/devices/bsim4v7/b4v7getic.c
View File

@ -29,15 +29,11 @@ BSIM4v7model *model = (BSIM4v7model*)inModel;
BSIM4v7instance *here;
#ifdef USE_CUSPICE
int i, status;
int status;
#endif
for (; model; model = BSIM4v7nextModel(model)) {
#ifdef USE_CUSPICE
i = 0;
#endif
for (here = BSIM4v7instances(model); here; here = BSIM4v7nextInstance(here))
{
if (!here->BSIM4v7icVDSGiven)
@ -54,11 +50,10 @@ BSIM4v7instance *here;
}
#ifdef USE_CUSPICE
int i = here->gen.GENcudaIndex;
model->BSIM4v7paramCPU.BSIM4v7icVDSArray [i] = here->BSIM4v7icVDS;
model->BSIM4v7paramCPU.BSIM4v7icVGSArray [i] = here->BSIM4v7icVGS;
model->BSIM4v7paramCPU.BSIM4v7icVBSArray [i] = here->BSIM4v7icVBS;
i++;
#endif
}

30
src/spicelib/devices/bsim4v7/b4v7set.c
View File

@ -2587,36 +2587,22 @@ do { if((here->ptr = SMPmakeElt(matrix,here->first,here->second))==(double *)NUL
}
#ifdef USE_CUSPICE
int i, j, jRHS, l, lRHS, lTimeSteps, status ;
/* Counting the instances */
for (model = (BSIM4v7model *)inModel ; model != NULL ; model = BSIM4v7nextModel(model))
{
i = 0 ;
for (here = BSIM4v7instances(model); here != NULL ; here = BSIM4v7nextInstance(here))
{
i++ ;
}
/* How much instances we have */
model->n_instances = i ;
/* This model supports CUDA */
model->gen.has_cuda = 1 ;
}
int j, jRHS, l, lRHS, lTimeSteps, status ;
/* loop through all the BSIM4v7 models */
for (model = (BSIM4v7model *)inModel ; model != NULL ; model = BSIM4v7nextModel(model))
{
/* This model supports CUDA */
model->gen.has_cuda = 1 ;
/* Position Vector Allocation */
model->PositionVector = TMALLOC (int, model->n_instances) ;
model->PositionVector = TMALLOC (int, model->gen.GENnInstances) ;
/* Position Vector Allocation for the RHS */
model->PositionVectorRHS = TMALLOC (int, model->n_instances) ;
model->PositionVectorRHS = TMALLOC (int, model->gen.GENnInstances) ;
/* Position Vector Allocation for timeSteps */
model->PositionVector_timeSteps = TMALLOC (int, model->n_instances) ;
model->PositionVector_timeSteps = TMALLOC (int, model->gen.GENnInstances) ;
model->offset = ckt->total_n_values ;
@ -2625,7 +2611,7 @@ do { if((here->ptr = SMPmakeElt(matrix,here->first,here->second))==(double *)NUL
model->offset_timeSteps = ckt->total_n_timeSteps ;
i = 0 ;
int i = 0 ;
j = 0 ;
jRHS = 0 ;
l = 0 ;

2
src/spicelib/devices/bsim4v7/bsim4v7def.h
View File

@ -3351,8 +3351,6 @@ typedef struct sBSIM4v7model
int n_timeSteps;
int *PositionVector_timeSteps;
int *d_PositionVector_timeSteps;
int n_instances;
#endif
} BSIM4v7model;

2
src/spicelib/devices/cap/CUSPICE/cucapgetic.c
View File

@ -48,7 +48,7 @@ GENmodel *inModel
cudaError_t status ;
CAPmodel *model = (CAPmodel *)inModel ;
size = (long unsigned int) model->n_instances;
size = (long unsigned int) model->gen.GENnInstances;
status = cudaMemcpy (model->CAPparamGPU.d_CAPinitCondArray, model->CAPparamCPU.CAPinitCondArray, size * sizeof(double), cudaMemcpyHostToDevice) ;
CUDAMEMCPYCHECK(model->CAPparamGPU.d_CAPinitCondArray, size, double, status)

8
src/spicelib/devices/cap/CUSPICE/cucapload.cu
View File

@ -56,10 +56,10 @@ 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->gen.GENnInstances % thread_y != 0)
block_x = (int)((model->gen.GENnInstances + thread_y - 1) / thread_y) ;
else
block_x = model->n_instances / thread_y ;
block_x = model->gen.GENnInstances / thread_y ;
dim3 thread (thread_x, thread_y) ;
@ -68,7 +68,7 @@ GENmodel *inModel, CKTcircuit *ckt
cuCAPload_kernel <<< block_x, thread >>> (model->CAPparamGPU, ckt->d_CKTrhsOld, ckt->d_CKTstate0,
ckt->d_CKTstate1, ckt->CKTmode, ckt->CKTag [0], ckt->CKTag [1],
ckt->CKTorder, model->n_instances, cond1,
ckt->CKTorder, model->gen.GENnInstances, cond1,
model->d_PositionVector, ckt->d_CKTloadOutput,
model->d_PositionVectorRHS, ckt->d_CKTloadOutputRHS) ;

2
src/spicelib/devices/cap/CUSPICE/cucapsetup.c
View File

@ -58,7 +58,7 @@ GENmodel *inModel
cudaError_t status ;
CAPmodel *model = (CAPmodel *)inModel ;
size = (long unsigned int) model->n_instances;
size = (long unsigned int) model->gen.GENnInstances;
/* Space Allocation to GPU */
status = cudaMalloc ((void **)&(model->d_PositionVector), size * sizeof(int)) ;

2
src/spicelib/devices/cap/CUSPICE/cucaptemp.c
View File

@ -48,7 +48,7 @@ GENmodel *inModel
cudaError_t status ;
CAPmodel *model = (CAPmodel *)inModel ;
size = (long unsigned int) model->n_instances;
size = (long unsigned int) model->gen.GENnInstances;
/* DOUBLE */
status = cudaMemcpy (model->CAPparamGPU.d_CAPcapacArray, model->CAPparamCPU.CAPcapacArray, size * sizeof(double), cudaMemcpyHostToDevice) ;

8
src/spicelib/devices/cap/CUSPICE/cucaptrunc.cu
View File

@ -31,17 +31,17 @@ GENmodel *inModel, CKTcircuit *ckt, double *timeStep
/* 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->gen.GENnInstances % thread_y != 0)
block_x = (int)((model->gen.GENnInstances + thread_y - 1) / thread_y) ;
else
block_x = model->n_instances / thread_y ;
block_x = model->gen.GENnInstances / thread_y ;
dim3 thread (thread_x, thread_y) ;
/* Kernel launch */
status = cudaGetLastError () ; // clear error status
cuCAPtrunc_kernel <<< block_x, thread >>> (model->CAPparamGPU, model->n_instances,
cuCAPtrunc_kernel <<< block_x, thread >>> (model->CAPparamGPU, model->gen.GENnInstances,
ckt->dD_CKTstates, ckt->d_CKTdeltaOld,
ckt->CKTdelta, ckt->CKTorder, ckt->CKTintegrateMethod,
ckt->CKTabstol, ckt->CKTreltol, ckt->CKTchgtol, ckt->CKTtrtol,

2
src/spicelib/devices/cap/capdefs.h
View File

@ -174,8 +174,6 @@ typedef struct sCAPmodel { /* model structure for a capacitor */
int n_timeSteps ;
int *PositionVector_timeSteps ;
int *d_PositionVector_timeSteps ;
int n_instances ;
#endif
} CAPmodel;

40
src/spicelib/devices/cap/capsetup.c
View File

@ -122,28 +122,14 @@ do { if((here->ptr = SMPmakeElt(matrix, here->first, here->second)) == NULL){\
}
#ifdef USE_CUSPICE
int i, j, k, status ;
/* Counting the instances */
for (model = (CAPmodel *)inModel ; model != NULL ; model = CAPnextModel(model))
{
i = 0 ;
for (here = CAPinstances(model); here != NULL ; here = CAPnextInstance(here))
{
i++ ;
}
/* How much instances we have */
model->n_instances = i ;
/* This model supports CUDA */
model->gen.has_cuda = 1 ;
}
int j, k, status ;
/* loop through all the capacitor models */
for (model = (CAPmodel *)inModel ; model != NULL ; model = CAPnextModel(model))
{
/* This model supports CUDA */
model->gen.has_cuda = 1 ;
model->offset = ckt->total_n_values ;
model->offsetRHS = ckt->total_n_valuesRHS ;
@ -174,13 +160,13 @@ do { if((here->ptr = SMPmakeElt(matrix, here->first, here->second)) == NULL){\
k++ ;
}
model->n_values = model->n_instances;
model->n_values = model->gen.GENnInstances;
ckt->total_n_values += model->n_values ;
model->n_Ptr = j ;
ckt->total_n_Ptr += model->n_Ptr ;
model->n_valuesRHS = model->n_instances;
model->n_valuesRHS = model->gen.GENnInstances;
ckt->total_n_valuesRHS += model->n_valuesRHS ;
model->n_PtrRHS = k ;
@ -188,27 +174,27 @@ 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->gen.GENnInstances) ;
for (j = 0 ; j < model->n_instances; j++)
for (j = 0 ; j < model->gen.GENnInstances; j++)
model->PositionVector [j] = model->offset + j ;
/* Position Vector assignment for the RHS */
model->PositionVectorRHS = TMALLOC (int, model->n_instances) ;
model->PositionVectorRHS = TMALLOC (int, model->gen.GENnInstances) ;
for (j = 0 ; j < model->n_instances; j++)
for (j = 0 ; j < model->gen.GENnInstances; j++)
model->PositionVectorRHS [j] = model->offsetRHS + j ;
/* Position Vector for timeSteps */
model->offset_timeSteps = ckt->total_n_timeSteps ;
model->n_timeSteps = model->n_instances;
model->n_timeSteps = model->gen.GENnInstances;
ckt->total_n_timeSteps += model->n_timeSteps ;
/* Position Vector assignment for timeSteps */
model->PositionVector_timeSteps = TMALLOC (int, model->n_instances) ;
model->PositionVector_timeSteps = TMALLOC (int, model->gen.GENnInstances) ;
for (j = 0 ; j < model->n_instances; j++)
for (j = 0 ; j < model->gen.GENnInstances; j++)
model->PositionVector_timeSteps [j] = model->offset_timeSteps + j ;
}

9
src/spicelib/devices/cap/captemp.c
View File

@ -32,16 +32,12 @@ CAPtemp(GENmodel *inModel, CKTcircuit *ckt)
double tc1, tc2;
#ifdef USE_CUSPICE
int i, status ;
int status ;
#endif
/* loop through all the capacitor models */
for( ; model != NULL; model = CAPnextModel(model)) {
#ifdef USE_CUSPICE
i = 0 ;
#endif
/* loop through all the instances of the model */
for (here = CAPinstances(model); here != NULL ;
here=CAPnextInstance(here)) {
@ -97,13 +93,12 @@ CAPtemp(GENmodel *inModel, CKTcircuit *ckt)
here->CAPcapac = here->CAPcapac * factor * here->CAPscale;
#ifdef USE_CUSPICE
int i = here->gen.GENcudaIndex;
model->CAPparamCPU.CAPcapacArray[i] = here->CAPcapac ;
model->CAPparamCPU.CAPmArray[i] = here->CAPm ;
model->CAPparamCPU.CAPposNodeArray[i] = here->CAPposNode ;
model->CAPparamCPU.CAPnegNodeArray[i] = here->CAPnegNode ;
model->CAPparamCPU.CAPstateArray[i] = here->CAPstate ;
i++ ;
#endif
}

8
src/spicelib/devices/ind/CUSPICE/cuindload.cu
View File

@ -68,10 +68,10 @@ 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 ;
if (model->gen.GENnInstances % thread_y != 0)
block_x = (int)(model->gen.GENnInstances / thread_y) + 1 ;
else
block_x = model->n_instances / thread_y ;
block_x = model->gen.GENnInstances / thread_y ;
dim3 thread (thread_x, thread_y) ;
@ -80,7 +80,7 @@ GENmodel *inModel, CKTcircuit *ckt
cuINDload_kernel <<< block_x, thread >>> (model->INDparamGPU, ckt->d_CKTrhsOld, ckt->d_CKTstate0,
ckt->d_CKTstate1, ckt->CKTmode, ckt->CKTag [0], ckt->CKTag [1],
ckt->CKTorder, model->n_instances,
ckt->CKTorder, model->gen.GENnInstances,
model->d_PositionVector, ckt->d_CKTloadOutput,
model->d_PositionVectorRHS, ckt->d_CKTloadOutputRHS) ;

2
src/spicelib/devices/ind/CUSPICE/cuindsetup.c
View File

@ -58,7 +58,7 @@ GENmodel *inModel
cudaError_t status ;
INDmodel *model = (INDmodel *)inModel ;
size = (long unsigned int) model->n_instances;
size = (long unsigned int) model->gen.GENnInstances;
/* Space Allocation to GPU */
status = cudaMalloc ((void **)&(model->d_PositionVector), size * sizeof(int)) ;

2
src/spicelib/devices/ind/CUSPICE/cuindtemp.c
View File

@ -48,7 +48,7 @@ GENmodel *inModel
cudaError_t status ;
INDmodel *model = (INDmodel *)inModel ;
size = (long unsigned int) model->n_instances;
size = (long unsigned int) model->gen.GENnInstances;
/* DOUBLE */
status = cudaMemcpy (model->INDparamGPU.d_INDinitCondArray, model->INDparamCPU.INDinitCondArray, size * sizeof(double), cudaMemcpyHostToDevice) ;

8
src/spicelib/devices/ind/CUSPICE/cuindtrunc.cu
View File

@ -31,17 +31,17 @@ GENmodel *inModel, CKTcircuit *ckt, double *timeStep
/* 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->gen.GENnInstances % thread_y != 0)
block_x = (int)((model->gen.GENnInstances + thread_y - 1) / thread_y) ;
else
block_x = model->n_instances / thread_y ;
block_x = model->gen.GENnInstances / thread_y ;
dim3 thread (thread_x, thread_y) ;
/* Kernel launch */
status = cudaGetLastError () ; // clear error status
cuINDtrunc_kernel <<< block_x, thread >>> (model->INDparamGPU, model->n_instances,
cuINDtrunc_kernel <<< block_x, thread >>> (model->INDparamGPU, model->gen.GENnInstances,
ckt->dD_CKTstates, ckt->d_CKTdeltaOld,
ckt->CKTdelta, ckt->CKTorder, ckt->CKTintegrateMethod,
ckt->CKTabstol, ckt->CKTreltol, ckt->CKTchgtol, ckt->CKTtrtol,

8
src/spicelib/devices/ind/CUSPICE/cumutload.cu
View File

@ -68,10 +68,10 @@ 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 ;
if (model->gen.GENnInstances % thread_y != 0)
block_x = (int)(model->gen.GENnInstances / thread_y) + 1 ;
else
block_x = model->n_instances / thread_y ;
block_x = model->gen.GENnInstances / thread_y ;
dim3 thread (thread_x, thread_y) ;
@ -80,7 +80,7 @@ GENmodel *inModel, CKTcircuit *ckt
cuMUTload_kernel <<< block_x, thread >>> (model->MUTparamGPU, ckt->d_CKTrhsOld, ckt->d_CKTstate0,
ckt->d_CKTstate1, ckt->CKTmode, ckt->CKTag [0], ckt->CKTag [1],
ckt->CKTorder, model->n_instances,
ckt->CKTorder, model->gen.GENnInstances,
model->d_PositionVector, ckt->d_CKTloadOutput,
model->d_PositionVectorRHS, ckt->d_CKTloadOutputRHS) ;

2
src/spicelib/devices/ind/CUSPICE/cumutsetup.c
View File

@ -58,7 +58,7 @@ GENmodel *inModel
cudaError_t status ;
MUTmodel *model = (MUTmodel *)inModel ;
size = (long unsigned int) model->n_instances;
size = (long unsigned int) model->gen.GENnInstances;
/* Space Allocation to GPU */
status = cudaMalloc ((void **)&(model->d_PositionVector), size * sizeof(int)) ;

2
src/spicelib/devices/ind/CUSPICE/cumuttemp.c
View File

@ -48,7 +48,7 @@ GENmodel *inModel
cudaError_t status ;
MUTmodel *model = (MUTmodel *)inModel ;
size = (long unsigned int) model->n_instances;
size = (long unsigned int) model->gen.GENnInstances;
/* DOUBLE */
status = cudaMemcpy (model->MUTparamGPU.d_MUTfactorArray, model->MUTparamCPU.MUTfactorArray, size * sizeof(double), cudaMemcpyHostToDevice) ;

4
src/spicelib/devices/ind/inddefs.h
View File

@ -175,8 +175,6 @@ struct sINDmodel { /* model structure for an inductor */
int n_timeSteps;
int *PositionVector_timeSteps;
int *d_PositionVector_timeSteps;
int n_instances;
#endif
};
@ -268,8 +266,6 @@ struct sMUTmodel { /* model structure for a mutual inductor */
int *PositionVectorRHS;
int *d_PositionVectorRHS;
int n_instances;
/* PARTICULAR SITUATION */
int n_instancesRHS;
#endif

40
src/spicelib/devices/ind/indsetup.c
View File

@ -109,28 +109,14 @@ do { if((here->ptr = SMPmakeElt(matrix, here->first, here->second)) == NULL){\
}
#ifdef USE_CUSPICE
int i, j, k, status ;
/* Counting the instances */
for (model = (INDmodel *)inModel ; model != NULL ; model = INDnextModel(model))
{
i = 0 ;
for (here = INDinstances(model); here != NULL ; here = INDnextInstance(here))
{
i++ ;
}
/* How much instances we have */
model->n_instances = i ;
/* This model supports CUDA */
model->gen.has_cuda = 1 ;
}
int j, k, status ;
/* loop through all the inductor models */
for (model = (INDmodel *)inModel ; model != NULL ; model = INDnextModel(model))
{
/* This model supports CUDA */
model->gen.has_cuda = 1 ;
model->offset = ckt->total_n_values ;
model->offsetRHS = ckt->total_n_valuesRHS ;
@ -162,13 +148,13 @@ do { if((here->ptr = SMPmakeElt(matrix, here->first, here->second)) == NULL){\
}
/* 2 Different Values for Every Instance */
model->n_values = 2 * model->n_instances;
model->n_values = 2 * model->gen.GENnInstances;
ckt->total_n_values += model->n_values ;
model->n_Ptr = j ;
ckt->total_n_Ptr += model->n_Ptr ;
model->n_valuesRHS = model->n_instances;
model->n_valuesRHS = model->gen.GENnInstances;
ckt->total_n_valuesRHS += model->n_valuesRHS ;
model->n_PtrRHS = k ;
@ -176,30 +162,30 @@ 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->gen.GENnInstances) ;
for (j = 0 ; j < model->n_instances; j++)
for (j = 0 ; j < model->gen.GENnInstances; j++)
{
/* 2 Different Values for Every Instance */
model->PositionVector [j] = model->offset + 2 * j ;
}
/* Position Vector assignment for the RHS */
model->PositionVectorRHS = TMALLOC (int, model->n_instances) ;
model->PositionVectorRHS = TMALLOC (int, model->gen.GENnInstances) ;
for (j = 0 ; j < model->n_instances; j++)
for (j = 0 ; j < model->gen.GENnInstances; j++)
model->PositionVectorRHS [j] = model->offsetRHS + j ;
/* Position Vector for timeSteps */
model->offset_timeSteps = ckt->total_n_timeSteps ;
model->n_timeSteps = model->n_instances;
model->n_timeSteps = model->gen.GENnInstances;
ckt->total_n_timeSteps += model->n_timeSteps ;
/* Position Vector assignment for timeSteps */
model->PositionVector_timeSteps = TMALLOC (int, model->n_instances) ;
model->PositionVector_timeSteps = TMALLOC (int, model->gen.GENnInstances) ;
for (j = 0 ; j < model->n_instances; j++)
for (j = 0 ; j < model->gen.GENnInstances; j++)
model->PositionVector_timeSteps [j] = model->offset_timeSteps + j ;
}

9
src/spicelib/devices/ind/indtemp.c
View File

@ -26,16 +26,12 @@ INDtemp(GENmodel *inModel, CKTcircuit *ckt)
double tc1, tc2;
#ifdef USE_CUSPICE
int i, status ;
int status ;
#endif
/* loop through all the inductor models */
for( ; model != NULL; model = INDnextModel(model)) {
#ifdef USE_CUSPICE
i = 0 ;
#endif
/* loop through all the instances of the model */
for (here = INDinstances(model); here != NULL ;
here=INDnextInstance(here)) {
@ -81,12 +77,11 @@ INDtemp(GENmodel *inModel, CKTcircuit *ckt)
here->INDinduct = here->INDinduct * factor * here->INDscale;
#ifdef USE_CUSPICE
int i = here->gen.GENcudaIndex;
model->INDparamCPU.INDinitCondArray[i] = here->INDinitCond ;
model->INDparamCPU.INDinductArray[i] = here->INDinduct ;
model->INDparamCPU.INDbrEqArray[i] = here->INDbrEq ;
model->INDparamCPU.INDstateArray[i] = here->INDstate ;
i++ ;
#endif
}

31
src/spicelib/devices/ind/mutsetup.c
View File

@ -66,26 +66,15 @@ MUTsetup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *states)
}
#ifdef USE_CUSPICE
int i, j, status;
int j, status;
INDmodel *indmodel;
INDinstance *indhere;
/* Counting the instances */
for (model = (MUTmodel *)inModel; model; model = MUTnextModel(model)) {
i = 0;
for (here = MUTinstances(model); here; here = MUTnextInstance(here))
i++;
/* How much instances we have */
model->n_instances = i;
/* This model supports CUDA */
model->gen.has_cuda = 1 ;
}
/* loop through all the mutual inductor models */
for (model = (MUTmodel *)inModel; model; model = MUTnextModel(model)) {
/* This model supports CUDA */
model->gen.has_cuda = 1 ;
model->offset = ckt->total_n_values;
j = 0;
@ -100,7 +89,7 @@ MUTsetup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *states)
j++;
}
model->n_values = model->n_instances;
model->n_values = model->gen.GENnInstances;
ckt->total_n_values += model->n_values;
model->n_Ptr = j;
@ -108,16 +97,16 @@ MUTsetup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *states)
/* Position Vector assignment */
model->PositionVector = TMALLOC(int, model->n_instances);
model->PositionVector = TMALLOC(int, model->gen.GENnInstances);
for (j = 0; j < model->n_instances; j++)
for (j = 0; j < model->gen.GENnInstances; j++)
model->PositionVector [j] = model->offset + j;
/* PARTICULAR SITUATION */
/* Pick up the IND model from one of the two IND instances */
indmodel = INDmodPtr(MUTinstances(model)->MUTind1);
model->n_instancesRHS = indmodel->n_instances;
model->n_instancesRHS = indmodel->gen.GENnInstances;
/* Position Vector assignment for the RHS */
model->PositionVectorRHS = TMALLOC(int, model->n_instancesRHS);
@ -132,9 +121,9 @@ MUTsetup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *states)
/* InstanceID storing for every MUT instance */
model->MUTparamCPU.MUTinstanceIND1Array = TMALLOC(int,
model->n_instances);
model->gen.GENnInstances);
model->MUTparamCPU.MUTinstanceIND2Array = TMALLOC(int,
model->n_instances);
model->gen.GENnInstances);
j = 0;
for (here = MUTinstances(model); here; here = MUTnextInstance(here)) {
model->MUTparamCPU.MUTinstanceIND1Array [j] = here->MUTind1->instanceID;

9
src/spicelib/devices/ind/muttemp.c
View File

@ -44,17 +44,13 @@ MUTtemp(GENmodel *inModel, CKTcircuit *ckt)
struct INDsystem *first_system = NULL;
#ifdef USE_CUSPICE
int i, status;
int status;
#endif
NG_IGNORE(ckt);
for (; model; model = MUTnextModel(model)) {
#ifdef USE_CUSPICE
i = 0;
#endif
for (here = MUTinstances(model); here; here = MUTnextInstance(here)) {
/* Value Processing for mutual inductors */
@ -68,13 +64,12 @@ MUTtemp(GENmodel *inModel, CKTcircuit *ckt)
here->MUTfactor = here->MUTcoupling * sqrt(fabs(ind1 * ind2));
#ifdef USE_CUSPICE
int i = here->gen.GENcudaIndex;
model->MUTparamCPU.MUTfactorArray[i] = here->MUTfactor;
model->MUTparamCPU.MUTflux1Array[i] = here->MUTind1->INDflux;
model->MUTparamCPU.MUTflux2Array[i] = here->MUTind2->INDflux;
model->MUTparamCPU.MUTbrEq1Array[i] = here->MUTind1->INDbrEq;
model->MUTparamCPU.MUTbrEq2Array[i] = here->MUTind2->INDbrEq;
i++;
#endif
if (ckt->CKTindverbosity > 0) {

8
src/spicelib/devices/isrc/CUSPICE/cuisrcload.cu
View File

@ -72,10 +72,10 @@ 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->gen.GENnInstances % thread_y != 0)
block_x = (int)((model->gen.GENnInstances + thread_y - 1) / thread_y) ;
else
block_x = model->n_instances / thread_y ;
block_x = model->gen.GENnInstances / thread_y ;
dim3 thread (thread_x, thread_y) ;
@ -84,7 +84,7 @@ GENmodel *inModel, CKTcircuit *ckt
cuISRCload_kernel <<< block_x, thread >>> (model->ISRCparamGPU, ckt->CKTmode, ckt->CKTtime,
ckt->CKTstep, ckt->CKTfinalTime, ckt->CKTsrcFact,
model->n_instances, model->d_PositionVectorRHS,
model->gen.GENnInstances, model->d_PositionVectorRHS,
ckt->d_CKTloadOutputRHS) ;
cudaDeviceSynchronize () ;

2
src/spicelib/devices/isrc/CUSPICE/cuisrcsetup.c
View File

@ -60,7 +60,7 @@ GENmodel *inModel
ISRCmodel *model = (ISRCmodel *)inModel ;
ISRCinstance *here ;
size1 = (long unsigned int) model->n_instances;
size1 = (long unsigned int) model->gen.GENnInstances;
/* Space Allocation to GPU */
status = cudaMalloc ((void **)&(model->d_PositionVectorRHS), size1 * sizeof(int)) ;

2
src/spicelib/devices/isrc/CUSPICE/cuisrctemp.c
View File

@ -50,7 +50,7 @@ GENmodel *inModel
ISRCmodel *model = (ISRCmodel *)inModel ;
ISRCinstance *here ;
size1 = (long unsigned int) model->n_instances;
size1 = (long unsigned int) model->gen.GENnInstances;
i = 0 ;

2
src/spicelib/devices/isrc/isrcdefs.h
View File

@ -129,8 +129,6 @@ typedef struct sISRCmodel {
int n_PtrRHS ;
int *PositionVectorRHS ;
int *d_PositionVectorRHS ;
int n_instances ;
#endif
} ISRCmodel;

8
src/spicelib/devices/isrc/isrcsetup.c
View File

@ -34,7 +34,7 @@ ISRCsetup (SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *state)
}
/* How many instances we have */
model->n_instances = i ;
model->gen.GENnInstances = i ;
/* This model supports CUDA */
model->gen.has_cuda = 1 ;
@ -58,7 +58,7 @@ ISRCsetup (SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *state)
j++ ;
}
model->n_valuesRHS = model->n_instances;
model->n_valuesRHS = model->gen.GENnInstances;
ckt->total_n_valuesRHS += model->n_valuesRHS ;
model->n_PtrRHS = j ;
@ -66,9 +66,9 @@ ISRCsetup (SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *state)
/* Position Vector assignment for the RHS */
model->PositionVectorRHS = TMALLOC (int, model->n_instances) ;
model->PositionVectorRHS = TMALLOC (int, model->gen.GENnInstances) ;
for (j = 0 ; j < model->n_instances; j++)
for (j = 0 ; j < model->gen.GENnInstances; j++)
model->PositionVectorRHS [j] = model->offsetRHS + j ;
}

10
src/spicelib/devices/isrc/isrctemp.c
View File

@ -27,16 +27,12 @@ ISRCtemp(GENmodel *inModel, CKTcircuit *ckt)
NG_IGNORE(ckt);
#ifdef USE_CUSPICE
int i, j, status ;
int j, status ;
#endif
/* loop through all the voltage source models */
for( ; model != NULL; model = ISRCnextModel(model)) {
#ifdef USE_CUSPICE
i = 0 ;
#endif
/* loop through all the instances of the model */
for (here = ISRCinstances(model); here != NULL ;
here=ISRCnextInstance(here)) {
@ -66,6 +62,8 @@ ISRCtemp(GENmodel *inModel, CKTcircuit *ckt)
here->ISRCacImag = here->ISRCacMag * sin(radians);
#ifdef USE_CUSPICE
int i = here->gen.GENcudaIndex;
for (j = 0 ; j < here->n_coeffs ; j++)
{
model->ISRCparamCPU.ISRCcoeffsArrayHost [i] [j] = here->ISRCcoeffs [j] ;
@ -75,8 +73,6 @@ ISRCtemp(GENmodel *inModel, CKTcircuit *ckt)
model->ISRCparamCPU.ISRCdcGivenArray[i] = here->ISRCdcGiven ;
model->ISRCparamCPU.ISRCfunctionTypeArray[i] = here->ISRCfunctionType ;
model->ISRCparamCPU.ISRCfunctionOrderArray[i] = here->ISRCfunctionOrder ;
i++ ;
#endif
}

8
src/spicelib/devices/vsrc/CUSPICE/cuvsrcload.cu
View File

@ -120,10 +120,10 @@ 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->gen.GENnInstances % thread_y != 0)
block_x = (int)((model->gen.GENnInstances + thread_y - 1) / thread_y) ;
else
block_x = model->n_instances / thread_y ;
block_x = model->gen.GENnInstances / thread_y ;
dim3 thread (thread_x, thread_y) ;
@ -132,7 +132,7 @@ GENmodel *inModel, CKTcircuit *ckt
cuVSRCload_kernel <<< block_x, thread >>> (model->VSRCparamGPU, ckt->CKTmode, ckt->CKTtime,
ckt->CKTstep, ckt->CKTfinalTime, ckt->CKTsrcFact,
model->n_instances, model->d_PositionVector,
model->gen.GENnInstances, model->d_PositionVector,
ckt->d_CKTloadOutput, model->d_PositionVectorRHS,
ckt->d_CKTloadOutputRHS) ;

2
src/spicelib/devices/vsrc/CUSPICE/cuvsrcsetup.c
View File

@ -60,7 +60,7 @@ GENmodel *inModel
VSRCmodel *model = (VSRCmodel *)inModel ;
VSRCinstance *here ;
size1 = (long unsigned int) model->n_instances;
size1 = (long unsigned int) model->gen.GENnInstances;
/* Space Allocation to GPU */
status = cudaMalloc ((void **)&(model->d_PositionVector), size1 * sizeof(int)) ;

2
src/spicelib/devices/vsrc/CUSPICE/cuvsrctemp.c
View File

@ -50,7 +50,7 @@ GENmodel *inModel
VSRCmodel *model = (VSRCmodel *)inModel ;
VSRCinstance *here ;
size1 = (long unsigned int) model->n_instances;
size1 = (long unsigned int) model->gen.GENnInstances;
i = 0 ;

2
src/spicelib/devices/vsrc/vsrcdefs.h
View File

@ -157,8 +157,6 @@ typedef struct sVSRCmodel {
int n_PtrRHS ;
int *PositionVectorRHS ;
int *d_PositionVectorRHS ;
int n_instances ;
#endif
} VSRCmodel;

34
src/spicelib/devices/vsrc/vsrcset.c
View File

@ -66,28 +66,14 @@ do { if((here->ptr = SMPmakeElt(matrix, here->first, here->second)) == NULL){\
}
#ifdef USE_CUSPICE
int i, j, k, status ;
/* Counting the instances */
for (model = (VSRCmodel *)inModel ; model != NULL ; model = VSRCnextModel(model))
{
i = 0 ;
for (here = VSRCinstances(model); here != NULL ; here = VSRCnextInstance(here))
{
i++ ;
}
/* How much instances we have */
model->n_instances = i ;
/* This model supports CUDA */
model->gen.has_cuda = 1 ;
}
int j, k, status ;
/* loop through all the voltage source models */
for (model = (VSRCmodel *)inModel ; model != NULL ; model = VSRCnextModel(model))
{
/* This model supports CUDA */
model->gen.has_cuda = 1 ;
model->offset = ckt->total_n_values ;
model->offsetRHS = ckt->total_n_valuesRHS ;
@ -115,13 +101,13 @@ do { if((here->ptr = SMPmakeElt(matrix, here->first, here->second)) == NULL){\
k++ ;
}
model->n_values = model->n_instances ;
model->n_values = model->gen.GENnInstances ;
ckt->total_n_values += model->n_values ;
model->n_Ptr = j ;
ckt->total_n_Ptr += model->n_Ptr ;
model->n_valuesRHS = model->n_instances ;
model->n_valuesRHS = model->gen.GENnInstances ;
ckt->total_n_valuesRHS += model->n_valuesRHS ;
model->n_PtrRHS = k ;
@ -129,15 +115,15 @@ 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->gen.GENnInstances) ;
for (j = 0 ; j < model->n_instances ; j++)
for (j = 0 ; j < model->gen.GENnInstances ; j++)
model->PositionVector [j] = model->offset + j ;
/* Position Vector assignment for the RHS */
model->PositionVectorRHS = TMALLOC (int, model->n_instances) ;
model->PositionVectorRHS = TMALLOC (int, model->gen.GENnInstances) ;
for (j = 0 ; j < model->n_instances ; j++)
for (j = 0 ; j < model->gen.GENnInstances ; j++)
model->PositionVectorRHS [j] = model->offsetRHS + j ;
}

10
src/spicelib/devices/vsrc/vsrctemp.c
View File

@ -27,16 +27,12 @@ VSRCtemp(GENmodel *inModel, CKTcircuit *ckt)
NG_IGNORE(ckt);
#ifdef USE_CUSPICE
int i, j, status ;
int j, status ;
#endif
/* loop through all the voltage source models */
for( ; model != NULL; model = VSRCnextModel(model)) {
#ifdef USE_CUSPICE
i = 0 ;
#endif
/* loop through all the instances of the model */
for (here = VSRCinstances(model); here != NULL ;
here=VSRCnextInstance(here)) {
@ -64,6 +60,8 @@ VSRCtemp(GENmodel *inModel, CKTcircuit *ckt)
here->VSRCacImag = here->VSRCacMag * sin(radians);
#ifdef USE_CUSPICE
int i = here->gen.GENcudaIndex;
for (j = 0 ; j < here->n_coeffs ; j++)
{
model->VSRCparamCPU.VSRCcoeffsArrayHost [i] [j] = here->VSRCcoeffs [j] ;
@ -76,8 +74,6 @@ VSRCtemp(GENmodel *inModel, CKTcircuit *ckt)
model->VSRCparamCPU.VSRCfunctionOrderArray[i] = here->VSRCfunctionOrder ;
model->VSRCparamCPU.VSRCrGivenArray[i] = here->VSRCrGiven ;
model->VSRCparamCPU.VSRCrBreakptArray[i] = here->VSRCrBreakpt ;
i++ ;
#endif
}