From 65ff4f46e1f0c0163a7c0a753724a41803a42a7b Mon Sep 17 00:00:00 2001 From: rlar Date: Sun, 29 Apr 2018 11:17:14 +0200 Subject: [PATCH] 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 --- .../devices/bsim4v7/CUSPICE/cubsim4v7getic.c | 2 +- .../devices/bsim4v7/CUSPICE/cubsim4v7load.cu | 8 ++-- .../devices/bsim4v7/CUSPICE/cubsim4v7setup.c | 2 +- .../devices/bsim4v7/CUSPICE/cubsim4v7temp.c | 2 +- .../devices/bsim4v7/CUSPICE/cubsim4v7trunc.cu | 8 ++-- src/spicelib/devices/bsim4v7/b4v7getic.c | 9 +---- src/spicelib/devices/bsim4v7/b4v7set.c | 30 ++++---------- src/spicelib/devices/bsim4v7/bsim4v7def.h | 2 - src/spicelib/devices/cap/CUSPICE/cucapgetic.c | 2 +- src/spicelib/devices/cap/CUSPICE/cucapload.cu | 8 ++-- src/spicelib/devices/cap/CUSPICE/cucapsetup.c | 2 +- src/spicelib/devices/cap/CUSPICE/cucaptemp.c | 2 +- .../devices/cap/CUSPICE/cucaptrunc.cu | 8 ++-- src/spicelib/devices/cap/capdefs.h | 2 - src/spicelib/devices/cap/capsetup.c | 40 ++++++------------- src/spicelib/devices/cap/captemp.c | 9 +---- src/spicelib/devices/ind/CUSPICE/cuindload.cu | 8 ++-- src/spicelib/devices/ind/CUSPICE/cuindsetup.c | 2 +- src/spicelib/devices/ind/CUSPICE/cuindtemp.c | 2 +- .../devices/ind/CUSPICE/cuindtrunc.cu | 8 ++-- src/spicelib/devices/ind/CUSPICE/cumutload.cu | 8 ++-- src/spicelib/devices/ind/CUSPICE/cumutsetup.c | 2 +- src/spicelib/devices/ind/CUSPICE/cumuttemp.c | 2 +- src/spicelib/devices/ind/inddefs.h | 4 -- src/spicelib/devices/ind/indsetup.c | 40 ++++++------------- src/spicelib/devices/ind/indtemp.c | 9 +---- src/spicelib/devices/ind/mutsetup.c | 31 +++++--------- src/spicelib/devices/ind/muttemp.c | 9 +---- .../devices/isrc/CUSPICE/cuisrcload.cu | 8 ++-- .../devices/isrc/CUSPICE/cuisrcsetup.c | 2 +- .../devices/isrc/CUSPICE/cuisrctemp.c | 2 +- src/spicelib/devices/isrc/isrcdefs.h | 2 - src/spicelib/devices/isrc/isrcsetup.c | 8 ++-- src/spicelib/devices/isrc/isrctemp.c | 10 ++--- .../devices/vsrc/CUSPICE/cuvsrcload.cu | 8 ++-- .../devices/vsrc/CUSPICE/cuvsrcsetup.c | 2 +- .../devices/vsrc/CUSPICE/cuvsrctemp.c | 2 +- src/spicelib/devices/vsrc/vsrcdefs.h | 2 - src/spicelib/devices/vsrc/vsrcset.c | 34 +++++----------- src/spicelib/devices/vsrc/vsrctemp.c | 10 ++--- 40 files changed, 122 insertions(+), 229 deletions(-) diff --git a/src/spicelib/devices/bsim4v7/CUSPICE/cubsim4v7getic.c b/src/spicelib/devices/bsim4v7/CUSPICE/cubsim4v7getic.c index 14cbe506f..8d2fb16f9 100644 --- a/src/spicelib/devices/bsim4v7/CUSPICE/cubsim4v7getic.c +++ b/src/spicelib/devices/bsim4v7/CUSPICE/cubsim4v7getic.c @@ -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) diff --git a/src/spicelib/devices/bsim4v7/CUSPICE/cubsim4v7load.cu b/src/spicelib/devices/bsim4v7/CUSPICE/cubsim4v7load.cu index ef93ee18d..a4ae2b86f 100644 --- a/src/spicelib/devices/bsim4v7/CUSPICE/cubsim4v7load.cu +++ b/src/spicelib/devices/bsim4v7/CUSPICE/cubsim4v7load.cu @@ -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, diff --git a/src/spicelib/devices/bsim4v7/CUSPICE/cubsim4v7setup.c b/src/spicelib/devices/bsim4v7/CUSPICE/cubsim4v7setup.c index 6f110bbe8..eeb7351cb 100644 --- a/src/spicelib/devices/bsim4v7/CUSPICE/cubsim4v7setup.c +++ b/src/spicelib/devices/bsim4v7/CUSPICE/cubsim4v7setup.c @@ -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)) ; diff --git a/src/spicelib/devices/bsim4v7/CUSPICE/cubsim4v7temp.c b/src/spicelib/devices/bsim4v7/CUSPICE/cubsim4v7temp.c index 0c07a9a22..05bd20c61 100644 --- a/src/spicelib/devices/bsim4v7/CUSPICE/cubsim4v7temp.c +++ b/src/spicelib/devices/bsim4v7/CUSPICE/cubsim4v7temp.c @@ -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 *)) ; diff --git a/src/spicelib/devices/bsim4v7/CUSPICE/cubsim4v7trunc.cu b/src/spicelib/devices/bsim4v7/CUSPICE/cubsim4v7trunc.cu index 5e39b5bc4..b701c33cb 100644 --- a/src/spicelib/devices/bsim4v7/CUSPICE/cubsim4v7trunc.cu +++ b/src/spicelib/devices/bsim4v7/CUSPICE/cubsim4v7trunc.cu @@ -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, diff --git a/src/spicelib/devices/bsim4v7/b4v7getic.c b/src/spicelib/devices/bsim4v7/b4v7getic.c index fce7c3ba1..907e1bcd6 100644 --- a/src/spicelib/devices/bsim4v7/b4v7getic.c +++ b/src/spicelib/devices/bsim4v7/b4v7getic.c @@ -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 } diff --git a/src/spicelib/devices/bsim4v7/b4v7set.c b/src/spicelib/devices/bsim4v7/b4v7set.c index a462a1e16..3df934a37 100644 --- a/src/spicelib/devices/bsim4v7/b4v7set.c +++ b/src/spicelib/devices/bsim4v7/b4v7set.c @@ -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 ; diff --git a/src/spicelib/devices/bsim4v7/bsim4v7def.h b/src/spicelib/devices/bsim4v7/bsim4v7def.h index 7423ca786..e2ca56979 100644 --- a/src/spicelib/devices/bsim4v7/bsim4v7def.h +++ b/src/spicelib/devices/bsim4v7/bsim4v7def.h @@ -3351,8 +3351,6 @@ typedef struct sBSIM4v7model int n_timeSteps; int *PositionVector_timeSteps; int *d_PositionVector_timeSteps; - - int n_instances; #endif } BSIM4v7model; diff --git a/src/spicelib/devices/cap/CUSPICE/cucapgetic.c b/src/spicelib/devices/cap/CUSPICE/cucapgetic.c index 62b75f566..736092072 100644 --- a/src/spicelib/devices/cap/CUSPICE/cucapgetic.c +++ b/src/spicelib/devices/cap/CUSPICE/cucapgetic.c @@ -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) diff --git a/src/spicelib/devices/cap/CUSPICE/cucapload.cu b/src/spicelib/devices/cap/CUSPICE/cucapload.cu index cf4de91d3..d549e5704 100644 --- a/src/spicelib/devices/cap/CUSPICE/cucapload.cu +++ b/src/spicelib/devices/cap/CUSPICE/cucapload.cu @@ -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) ; diff --git a/src/spicelib/devices/cap/CUSPICE/cucapsetup.c b/src/spicelib/devices/cap/CUSPICE/cucapsetup.c index 21d182b82..2c76e3169 100644 --- a/src/spicelib/devices/cap/CUSPICE/cucapsetup.c +++ b/src/spicelib/devices/cap/CUSPICE/cucapsetup.c @@ -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)) ; diff --git a/src/spicelib/devices/cap/CUSPICE/cucaptemp.c b/src/spicelib/devices/cap/CUSPICE/cucaptemp.c index 2adf40bd6..d63468641 100644 --- a/src/spicelib/devices/cap/CUSPICE/cucaptemp.c +++ b/src/spicelib/devices/cap/CUSPICE/cucaptemp.c @@ -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) ; diff --git a/src/spicelib/devices/cap/CUSPICE/cucaptrunc.cu b/src/spicelib/devices/cap/CUSPICE/cucaptrunc.cu index fa09659e5..469a8c35d 100644 --- a/src/spicelib/devices/cap/CUSPICE/cucaptrunc.cu +++ b/src/spicelib/devices/cap/CUSPICE/cucaptrunc.cu @@ -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, diff --git a/src/spicelib/devices/cap/capdefs.h b/src/spicelib/devices/cap/capdefs.h index 052eace7a..620f39009 100644 --- a/src/spicelib/devices/cap/capdefs.h +++ b/src/spicelib/devices/cap/capdefs.h @@ -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; diff --git a/src/spicelib/devices/cap/capsetup.c b/src/spicelib/devices/cap/capsetup.c index c343a5aa0..e6cb971a4 100644 --- a/src/spicelib/devices/cap/capsetup.c +++ b/src/spicelib/devices/cap/capsetup.c @@ -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 ; } diff --git a/src/spicelib/devices/cap/captemp.c b/src/spicelib/devices/cap/captemp.c index b3cdd38c0..ca483a62d 100644 --- a/src/spicelib/devices/cap/captemp.c +++ b/src/spicelib/devices/cap/captemp.c @@ -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 } diff --git a/src/spicelib/devices/ind/CUSPICE/cuindload.cu b/src/spicelib/devices/ind/CUSPICE/cuindload.cu index e5c699113..bb6aaa23f 100644 --- a/src/spicelib/devices/ind/CUSPICE/cuindload.cu +++ b/src/spicelib/devices/ind/CUSPICE/cuindload.cu @@ -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) ; diff --git a/src/spicelib/devices/ind/CUSPICE/cuindsetup.c b/src/spicelib/devices/ind/CUSPICE/cuindsetup.c index e60972073..eb4ff2ca7 100644 --- a/src/spicelib/devices/ind/CUSPICE/cuindsetup.c +++ b/src/spicelib/devices/ind/CUSPICE/cuindsetup.c @@ -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)) ; diff --git a/src/spicelib/devices/ind/CUSPICE/cuindtemp.c b/src/spicelib/devices/ind/CUSPICE/cuindtemp.c index 05d67a972..5544438f7 100644 --- a/src/spicelib/devices/ind/CUSPICE/cuindtemp.c +++ b/src/spicelib/devices/ind/CUSPICE/cuindtemp.c @@ -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) ; diff --git a/src/spicelib/devices/ind/CUSPICE/cuindtrunc.cu b/src/spicelib/devices/ind/CUSPICE/cuindtrunc.cu index 3f7f91b43..32890e0da 100644 --- a/src/spicelib/devices/ind/CUSPICE/cuindtrunc.cu +++ b/src/spicelib/devices/ind/CUSPICE/cuindtrunc.cu @@ -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, diff --git a/src/spicelib/devices/ind/CUSPICE/cumutload.cu b/src/spicelib/devices/ind/CUSPICE/cumutload.cu index 56405c08e..f68653497 100644 --- a/src/spicelib/devices/ind/CUSPICE/cumutload.cu +++ b/src/spicelib/devices/ind/CUSPICE/cumutload.cu @@ -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) ; diff --git a/src/spicelib/devices/ind/CUSPICE/cumutsetup.c b/src/spicelib/devices/ind/CUSPICE/cumutsetup.c index b5f39fb6e..0ddd0d514 100644 --- a/src/spicelib/devices/ind/CUSPICE/cumutsetup.c +++ b/src/spicelib/devices/ind/CUSPICE/cumutsetup.c @@ -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)) ; diff --git a/src/spicelib/devices/ind/CUSPICE/cumuttemp.c b/src/spicelib/devices/ind/CUSPICE/cumuttemp.c index eb61d2e64..a39ef649f 100644 --- a/src/spicelib/devices/ind/CUSPICE/cumuttemp.c +++ b/src/spicelib/devices/ind/CUSPICE/cumuttemp.c @@ -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) ; diff --git a/src/spicelib/devices/ind/inddefs.h b/src/spicelib/devices/ind/inddefs.h index 4e7f43d5f..db48eb47c 100644 --- a/src/spicelib/devices/ind/inddefs.h +++ b/src/spicelib/devices/ind/inddefs.h @@ -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 diff --git a/src/spicelib/devices/ind/indsetup.c b/src/spicelib/devices/ind/indsetup.c index b5ef543ea..6c711ec75 100644 --- a/src/spicelib/devices/ind/indsetup.c +++ b/src/spicelib/devices/ind/indsetup.c @@ -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 ; } diff --git a/src/spicelib/devices/ind/indtemp.c b/src/spicelib/devices/ind/indtemp.c index 7eeba855c..f1f4a107a 100644 --- a/src/spicelib/devices/ind/indtemp.c +++ b/src/spicelib/devices/ind/indtemp.c @@ -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 } diff --git a/src/spicelib/devices/ind/mutsetup.c b/src/spicelib/devices/ind/mutsetup.c index d66162ddf..1e95ca1ec 100644 --- a/src/spicelib/devices/ind/mutsetup.c +++ b/src/spicelib/devices/ind/mutsetup.c @@ -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; diff --git a/src/spicelib/devices/ind/muttemp.c b/src/spicelib/devices/ind/muttemp.c index 27ca67551..e6eae929d 100644 --- a/src/spicelib/devices/ind/muttemp.c +++ b/src/spicelib/devices/ind/muttemp.c @@ -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) { diff --git a/src/spicelib/devices/isrc/CUSPICE/cuisrcload.cu b/src/spicelib/devices/isrc/CUSPICE/cuisrcload.cu index fde1bcf97..639d7b819 100644 --- a/src/spicelib/devices/isrc/CUSPICE/cuisrcload.cu +++ b/src/spicelib/devices/isrc/CUSPICE/cuisrcload.cu @@ -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 () ; diff --git a/src/spicelib/devices/isrc/CUSPICE/cuisrcsetup.c b/src/spicelib/devices/isrc/CUSPICE/cuisrcsetup.c index ce6450c7f..8c2da4f23 100644 --- a/src/spicelib/devices/isrc/CUSPICE/cuisrcsetup.c +++ b/src/spicelib/devices/isrc/CUSPICE/cuisrcsetup.c @@ -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)) ; diff --git a/src/spicelib/devices/isrc/CUSPICE/cuisrctemp.c b/src/spicelib/devices/isrc/CUSPICE/cuisrctemp.c index dbeaf06da..48087eff1 100644 --- a/src/spicelib/devices/isrc/CUSPICE/cuisrctemp.c +++ b/src/spicelib/devices/isrc/CUSPICE/cuisrctemp.c @@ -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 ; diff --git a/src/spicelib/devices/isrc/isrcdefs.h b/src/spicelib/devices/isrc/isrcdefs.h index 32ac5b4a2..3f9f8affb 100644 --- a/src/spicelib/devices/isrc/isrcdefs.h +++ b/src/spicelib/devices/isrc/isrcdefs.h @@ -129,8 +129,6 @@ typedef struct sISRCmodel { int n_PtrRHS ; int *PositionVectorRHS ; int *d_PositionVectorRHS ; - - int n_instances ; #endif } ISRCmodel; diff --git a/src/spicelib/devices/isrc/isrcsetup.c b/src/spicelib/devices/isrc/isrcsetup.c index e1a001d45..a181a67a0 100644 --- a/src/spicelib/devices/isrc/isrcsetup.c +++ b/src/spicelib/devices/isrc/isrcsetup.c @@ -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 ; } diff --git a/src/spicelib/devices/isrc/isrctemp.c b/src/spicelib/devices/isrc/isrctemp.c index db9a03837..16e054b52 100644 --- a/src/spicelib/devices/isrc/isrctemp.c +++ b/src/spicelib/devices/isrc/isrctemp.c @@ -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 } diff --git a/src/spicelib/devices/vsrc/CUSPICE/cuvsrcload.cu b/src/spicelib/devices/vsrc/CUSPICE/cuvsrcload.cu index 7a2f88b0a..587b79c2b 100644 --- a/src/spicelib/devices/vsrc/CUSPICE/cuvsrcload.cu +++ b/src/spicelib/devices/vsrc/CUSPICE/cuvsrcload.cu @@ -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) ; diff --git a/src/spicelib/devices/vsrc/CUSPICE/cuvsrcsetup.c b/src/spicelib/devices/vsrc/CUSPICE/cuvsrcsetup.c index 39d6c558a..b8fdf4765 100644 --- a/src/spicelib/devices/vsrc/CUSPICE/cuvsrcsetup.c +++ b/src/spicelib/devices/vsrc/CUSPICE/cuvsrcsetup.c @@ -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)) ; diff --git a/src/spicelib/devices/vsrc/CUSPICE/cuvsrctemp.c b/src/spicelib/devices/vsrc/CUSPICE/cuvsrctemp.c index 2de620cf5..7ec78ed52 100644 --- a/src/spicelib/devices/vsrc/CUSPICE/cuvsrctemp.c +++ b/src/spicelib/devices/vsrc/CUSPICE/cuvsrctemp.c @@ -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 ; diff --git a/src/spicelib/devices/vsrc/vsrcdefs.h b/src/spicelib/devices/vsrc/vsrcdefs.h index c7714d3db..1c2d7b06a 100644 --- a/src/spicelib/devices/vsrc/vsrcdefs.h +++ b/src/spicelib/devices/vsrc/vsrcdefs.h @@ -157,8 +157,6 @@ typedef struct sVSRCmodel { int n_PtrRHS ; int *PositionVectorRHS ; int *d_PositionVectorRHS ; - - int n_instances ; #endif } VSRCmodel; diff --git a/src/spicelib/devices/vsrc/vsrcset.c b/src/spicelib/devices/vsrc/vsrcset.c index 835d6bd75..5d7d353ce 100644 --- a/src/spicelib/devices/vsrc/vsrcset.c +++ b/src/spicelib/devices/vsrc/vsrcset.c @@ -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 ; } diff --git a/src/spicelib/devices/vsrc/vsrctemp.c b/src/spicelib/devices/vsrc/vsrctemp.c index b6d20d74f..0088aed56 100644 --- a/src/spicelib/devices/vsrc/vsrctemp.c +++ b/src/spicelib/devices/vsrc/vsrctemp.c @@ -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 }