diff --git a/src/include/ngspice/gendefs.h b/src/include/ngspice/gendefs.h index 27c89d2e3..ac3a25d6d 100644 --- a/src/include/ngspice/gendefs.h +++ b/src/include/ngspice/gendefs.h @@ -44,6 +44,10 @@ struct GENmodel { /* model structure for a resistor */ GENinstance *GENinstances; /* pointer to list of instances that have this * model */ IFuid GENmodName; /* pointer to character string naming this model */ + +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; /* flag to indicate is the model supports CUDA */ +#endif }; diff --git a/src/spicelib/analysis/CUSPICE/cucktflush.c b/src/spicelib/analysis/CUSPICE/cucktflush.c index fe6bf801f..b72b06f3b 100644 --- a/src/spicelib/analysis/CUSPICE/cucktflush.c +++ b/src/spicelib/analysis/CUSPICE/cucktflush.c @@ -35,16 +35,26 @@ cuCKTflush CKTcircuit *ckt ) { - long unsigned int m, mRHS ; + if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) { + long unsigned int m, mRHS ; - m = (long unsigned int)(ckt->total_n_values + 1) ; // + 1 because of CKTdiagGmin - mRHS = (long unsigned int)ckt->total_n_valuesRHS ; + m = (long unsigned int)(ckt->total_n_values + 1) ; // + 1 because of CKTdiagGmin + mRHS = (long unsigned int)ckt->total_n_valuesRHS ; - /* Clean-up the CKTloadOutput */ - cudaMemset (ckt->d_CKTloadOutput, 0, m * sizeof(double)) ; + /* Clean-up the CKTloadOutput */ + cudaMemset (ckt->d_CKTloadOutput, 0, m * sizeof(double)) ; - /* Clean-up the CKTloadOutputRHS */ - cudaMemset (ckt->d_CKTloadOutputRHS, 0, mRHS * sizeof(double)) ; + /* Clean-up the CKTloadOutputRHS */ + cudaMemset (ckt->d_CKTloadOutputRHS, 0, mRHS * sizeof(double)) ; + } else { + int i, size ; + + size = SMPmatSize (ckt->CKTmatrix) ; + for (i = 0 ; i <= size ; i++) + *(ckt->CKTrhs + i) = 0 ; + + SMPclear (ckt->CKTmatrix) ; + } return (OK) ; } diff --git a/src/spicelib/analysis/CUSPICE/cucktnonconupdate.c b/src/spicelib/analysis/CUSPICE/cucktnonconupdate.c index f56377153..4bdddc9a9 100644 --- a/src/spicelib/analysis/CUSPICE/cucktnonconupdate.c +++ b/src/spicelib/analysis/CUSPICE/cucktnonconupdate.c @@ -47,8 +47,10 @@ CKTcircuit *ckt { cudaError_t status ; - status = cudaMemcpy (ckt->d_CKTnoncon, &(ckt->CKTnoncon), sizeof(int), cudaMemcpyHostToDevice) ; - CUDAMEMCPYCHECK (ckt->d_CKTnoncon, 1, int, status) + if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) { + status = cudaMemcpy (ckt->d_CKTnoncon, &(ckt->CKTnoncon), sizeof(int), cudaMemcpyHostToDevice) ; + CUDAMEMCPYCHECK (ckt->d_CKTnoncon, 1, int, status) + } return (OK) ; } @@ -61,8 +63,10 @@ CKTcircuit *ckt { cudaError_t status ; - status = cudaMemcpy (&(ckt->CKTnoncon), ckt->d_CKTnoncon, sizeof(int), cudaMemcpyDeviceToHost) ; - CUDAMEMCPYCHECK (&(ckt->CKTnoncon), 1, int, status) + if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) { + status = cudaMemcpy (&(ckt->CKTnoncon), ckt->d_CKTnoncon, sizeof(int), cudaMemcpyDeviceToHost) ; + CUDAMEMCPYCHECK (&(ckt->CKTnoncon), 1, int, status) + } return (OK) ; } diff --git a/src/spicelib/analysis/CUSPICE/cucktrhsoldupdate.c b/src/spicelib/analysis/CUSPICE/cucktrhsoldupdate.c index 1e07e0281..78feb1f98 100644 --- a/src/spicelib/analysis/CUSPICE/cucktrhsoldupdate.c +++ b/src/spicelib/analysis/CUSPICE/cucktrhsoldupdate.c @@ -47,8 +47,10 @@ CKTcircuit *ckt { long unsigned int size ; - size = (long unsigned int)(ckt->d_MatrixSize + 1) ; - cudaMemset (ckt->d_CKTrhsOld, 0, size * sizeof(double)) ; + if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) { + size = (long unsigned int)(ckt->d_MatrixSize + 1) ; + cudaMemset (ckt->d_CKTrhsOld, 0, size * sizeof(double)) ; + } return (OK) ; } @@ -62,9 +64,11 @@ CKTcircuit *ckt long unsigned int size ; cudaError_t status ; - size = (long unsigned int)(ckt->d_MatrixSize + 1) ; - status = cudaMemcpy (ckt->d_CKTrhsOld, ckt->CKTrhsOld, size * sizeof(double), cudaMemcpyHostToDevice) ; - CUDAMEMCPYCHECK (ckt->d_CKTrhsOld, size, double, status) + if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) { + size = (long unsigned int)(ckt->d_MatrixSize + 1) ; + status = cudaMemcpy (ckt->d_CKTrhsOld, ckt->CKTrhsOld, size * sizeof(double), cudaMemcpyHostToDevice) ; + CUDAMEMCPYCHECK (ckt->d_CKTrhsOld, size, double, status) + } return (OK) ; } @@ -78,9 +82,11 @@ CKTcircuit *ckt long unsigned int size ; cudaError_t status ; - size = (long unsigned int)(ckt->d_MatrixSize + 1) ; - status = cudaMemcpy (ckt->CKTrhsOld, ckt->d_CKTrhsOld, size * sizeof(double), cudaMemcpyDeviceToHost) ; - CUDAMEMCPYCHECK (ckt->CKTrhsOld, size, double, status) + if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) { + size = (long unsigned int)(ckt->d_MatrixSize + 1) ; + status = cudaMemcpy (ckt->CKTrhsOld, ckt->d_CKTrhsOld, size * sizeof(double), cudaMemcpyDeviceToHost) ; + CUDAMEMCPYCHECK (ckt->CKTrhsOld, size, double, status) + } return (OK) ; } diff --git a/src/spicelib/analysis/CUSPICE/cucktsetup.c b/src/spicelib/analysis/CUSPICE/cucktsetup.c index bd14fc15c..320d5bb47 100644 --- a/src/spicelib/analysis/CUSPICE/cucktsetup.c +++ b/src/spicelib/analysis/CUSPICE/cucktsetup.c @@ -76,88 +76,90 @@ CKTcircuit *ckt size2 = (long unsigned int)ckt->CKTnumStates ; size3 = (long unsigned int)ckt->total_n_timeSteps ; - /* Topology Matrix Handling */ - status = cudaMalloc ((void **)&(ckt->CKTmatrix->d_CKTrhs), (n + 1) * sizeof(double)) ; - CUDAMALLOCCHECK (ckt->CKTmatrix->d_CKTrhs, (n + 1), double, status) + if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) { + /* Topology Matrix Handling */ + status = cudaMalloc ((void **)&(ckt->CKTmatrix->d_CKTrhs), (n + 1) * sizeof(double)) ; + CUDAMALLOCCHECK (ckt->CKTmatrix->d_CKTrhs, (n + 1), double, status) - status = cudaMalloc ((void **)&(ckt->CKTmatrix->d_CKTkluAx), nz * sizeof(double)) ; - CUDAMALLOCCHECK (ckt->CKTmatrix->d_CKTkluAx, nz, double, status) + status = cudaMalloc ((void **)&(ckt->CKTmatrix->d_CKTkluAx), nz * sizeof(double)) ; + CUDAMALLOCCHECK (ckt->CKTmatrix->d_CKTkluAx, nz, double, status) - status = cudaMalloc ((void **)&(ckt->d_CKTloadOutput), m * sizeof(double)) ; - CUDAMALLOCCHECK (ckt->d_CKTloadOutput, m, double, status) + status = cudaMalloc ((void **)&(ckt->d_CKTloadOutput), m * sizeof(double)) ; + CUDAMALLOCCHECK (ckt->d_CKTloadOutput, m, double, status) - status = cudaMalloc ((void **)&(ckt->d_CKTloadOutputRHS), mRHS * sizeof(double)) ; - CUDAMALLOCCHECK (ckt->d_CKTloadOutputRHS, mRHS, double, status) + status = cudaMalloc ((void **)&(ckt->d_CKTloadOutputRHS), mRHS * sizeof(double)) ; + CUDAMALLOCCHECK (ckt->d_CKTloadOutputRHS, mRHS, double, status) - status = cudaMalloc ((void **)&(ckt->d_CKTtopologyMatrixCSRp), (nz + 1) * sizeof(int)) ; - CUDAMALLOCCHECK (ckt->d_CKTtopologyMatrixCSRp, (nz + 1), int, status) + status = cudaMalloc ((void **)&(ckt->d_CKTtopologyMatrixCSRp), (nz + 1) * sizeof(int)) ; + CUDAMALLOCCHECK (ckt->d_CKTtopologyMatrixCSRp, (nz + 1), int, status) - status = cudaMalloc ((void **)&(ckt->d_CKTtopologyMatrixCSRj), TopologyNNZ * sizeof(int)) ; - CUDAMALLOCCHECK (ckt->d_CKTtopologyMatrixCSRj, TopologyNNZ, int, status) + status = cudaMalloc ((void **)&(ckt->d_CKTtopologyMatrixCSRj), TopologyNNZ * sizeof(int)) ; + CUDAMALLOCCHECK (ckt->d_CKTtopologyMatrixCSRj, TopologyNNZ, int, status) - status = cudaMalloc ((void **)&(ckt->d_CKTtopologyMatrixCSRx), TopologyNNZ * sizeof(double)) ; - CUDAMALLOCCHECK (ckt->d_CKTtopologyMatrixCSRx, TopologyNNZ, double, status) + status = cudaMalloc ((void **)&(ckt->d_CKTtopologyMatrixCSRx), TopologyNNZ * sizeof(double)) ; + CUDAMALLOCCHECK (ckt->d_CKTtopologyMatrixCSRx, TopologyNNZ, double, status) - status = cudaMalloc ((void **)&(ckt->d_CKTtopologyMatrixCSRpRHS), ((n + 1) + 1) * sizeof(int)) ; - CUDAMALLOCCHECK (ckt->d_CKTtopologyMatrixCSRpRHS, ((n + 1) + 1), int, status) + status = cudaMalloc ((void **)&(ckt->d_CKTtopologyMatrixCSRpRHS), ((n + 1) + 1) * sizeof(int)) ; + CUDAMALLOCCHECK (ckt->d_CKTtopologyMatrixCSRpRHS, ((n + 1) + 1), int, status) - status = cudaMalloc ((void **)&(ckt->d_CKTtopologyMatrixCSRjRHS), TopologyNNZRHS * sizeof(int)) ; - CUDAMALLOCCHECK (ckt->d_CKTtopologyMatrixCSRjRHS, TopologyNNZRHS, int, status) + status = cudaMalloc ((void **)&(ckt->d_CKTtopologyMatrixCSRjRHS), TopologyNNZRHS * sizeof(int)) ; + CUDAMALLOCCHECK (ckt->d_CKTtopologyMatrixCSRjRHS, TopologyNNZRHS, int, status) - status = cudaMalloc ((void **)&(ckt->d_CKTtopologyMatrixCSRxRHS), TopologyNNZRHS * sizeof(double)) ; - CUDAMALLOCCHECK (ckt->d_CKTtopologyMatrixCSRxRHS, TopologyNNZRHS, double, status) + status = cudaMalloc ((void **)&(ckt->d_CKTtopologyMatrixCSRxRHS), TopologyNNZRHS * sizeof(double)) ; + CUDAMALLOCCHECK (ckt->d_CKTtopologyMatrixCSRxRHS, TopologyNNZRHS, double, status) - cudaMemset (ckt->d_CKTloadOutput + ckt->total_n_values, 0, sizeof(double)) ; //DiagGmin is 0 at the beginning + cudaMemset (ckt->d_CKTloadOutput + ckt->total_n_values, 0, sizeof(double)) ; //DiagGmin is 0 at the beginning - status = cudaMemcpy (ckt->d_CKTtopologyMatrixCSRp, ckt->CKTtopologyMatrixCSRp, (nz + 1) * sizeof(int), cudaMemcpyHostToDevice) ; - CUDAMEMCPYCHECK (ckt->d_CKTtopologyMatrixCSRp, (nz + 1), int, status) + status = cudaMemcpy (ckt->d_CKTtopologyMatrixCSRp, ckt->CKTtopologyMatrixCSRp, (nz + 1) * sizeof(int), cudaMemcpyHostToDevice) ; + CUDAMEMCPYCHECK (ckt->d_CKTtopologyMatrixCSRp, (nz + 1), int, status) - status = cudaMemcpy (ckt->d_CKTtopologyMatrixCSRj, ckt->CKTtopologyMatrixCOOj, TopologyNNZ * sizeof(int), cudaMemcpyHostToDevice) ; - CUDAMEMCPYCHECK (ckt->d_CKTtopologyMatrixCSRj, TopologyNNZ, int, status) + status = cudaMemcpy (ckt->d_CKTtopologyMatrixCSRj, ckt->CKTtopologyMatrixCOOj, TopologyNNZ * sizeof(int), cudaMemcpyHostToDevice) ; + CUDAMEMCPYCHECK (ckt->d_CKTtopologyMatrixCSRj, TopologyNNZ, int, status) - status = cudaMemcpy (ckt->d_CKTtopologyMatrixCSRx, ckt->CKTtopologyMatrixCOOx, TopologyNNZ * sizeof(double), cudaMemcpyHostToDevice) ; - CUDAMEMCPYCHECK (ckt->d_CKTtopologyMatrixCSRx, TopologyNNZ, double, status) + status = cudaMemcpy (ckt->d_CKTtopologyMatrixCSRx, ckt->CKTtopologyMatrixCOOx, TopologyNNZ * sizeof(double), cudaMemcpyHostToDevice) ; + CUDAMEMCPYCHECK (ckt->d_CKTtopologyMatrixCSRx, TopologyNNZ, double, status) - status = cudaMemcpy (ckt->d_CKTtopologyMatrixCSRpRHS, ckt->CKTtopologyMatrixCSRpRHS, ((n + 1) + 1) * sizeof(int), cudaMemcpyHostToDevice) ; - CUDAMEMCPYCHECK (ckt->d_CKTtopologyMatrixCSRpRHS, ((n + 1) + 1), int, status) + status = cudaMemcpy (ckt->d_CKTtopologyMatrixCSRpRHS, ckt->CKTtopologyMatrixCSRpRHS, ((n + 1) + 1) * sizeof(int), cudaMemcpyHostToDevice) ; + CUDAMEMCPYCHECK (ckt->d_CKTtopologyMatrixCSRpRHS, ((n + 1) + 1), int, status) - status = cudaMemcpy (ckt->d_CKTtopologyMatrixCSRjRHS, ckt->CKTtopologyMatrixCOOjRHS, TopologyNNZRHS * sizeof(int), cudaMemcpyHostToDevice) ; - CUDAMEMCPYCHECK (ckt->d_CKTtopologyMatrixCSRjRHS, TopologyNNZRHS, int, status) + status = cudaMemcpy (ckt->d_CKTtopologyMatrixCSRjRHS, ckt->CKTtopologyMatrixCOOjRHS, TopologyNNZRHS * sizeof(int), cudaMemcpyHostToDevice) ; + CUDAMEMCPYCHECK (ckt->d_CKTtopologyMatrixCSRjRHS, TopologyNNZRHS, int, status) - status = cudaMemcpy (ckt->d_CKTtopologyMatrixCSRxRHS, ckt->CKTtopologyMatrixCOOxRHS, TopologyNNZRHS * sizeof(double), cudaMemcpyHostToDevice) ; - CUDAMEMCPYCHECK (ckt->d_CKTtopologyMatrixCSRxRHS, TopologyNNZRHS, double, status) - /* ------------------------ */ + status = cudaMemcpy (ckt->d_CKTtopologyMatrixCSRxRHS, ckt->CKTtopologyMatrixCOOxRHS, TopologyNNZRHS * sizeof(double), cudaMemcpyHostToDevice) ; + CUDAMEMCPYCHECK (ckt->d_CKTtopologyMatrixCSRxRHS, TopologyNNZRHS, double, status) + /* ------------------------ */ - status = cudaMalloc ((void **)&(ckt->d_CKTnoncon), sizeof(int)) ; - CUDAMALLOCCHECK (ckt->d_CKTnoncon, 1, int, status) + status = cudaMalloc ((void **)&(ckt->d_CKTnoncon), sizeof(int)) ; + CUDAMALLOCCHECK (ckt->d_CKTnoncon, 1, int, status) - status = cudaMalloc ((void **)&(ckt->d_CKTrhsOld), size1 * sizeof(double)) ; - CUDAMALLOCCHECK (ckt->d_CKTrhsOld, size1, double, status) + status = cudaMalloc ((void **)&(ckt->d_CKTrhsOld), size1 * sizeof(double)) ; + CUDAMALLOCCHECK (ckt->d_CKTrhsOld, size1, double, status) - for (i = 0 ; i <= MAX (2, ckt->CKTmaxOrder) + 1 ; i++) /* dctran needs 3 states at least */ - { - status = cudaMalloc ((void **)&(ckt->d_CKTstates[i]), size2 * sizeof(double)) ; - CUDAMALLOCCHECK (ckt->d_CKTstates[i], size2, double, status) + for (i = 0 ; i <= MAX (2, ckt->CKTmaxOrder) + 1 ; i++) /* dctran needs 3 states at least */ + { + status = cudaMalloc ((void **)&(ckt->d_CKTstates[i]), size2 * sizeof(double)) ; + CUDAMALLOCCHECK (ckt->d_CKTstates[i], size2, double, status) + } + + + /* Truncation Error */ + status = cudaMalloc ((void **)&(ckt->dD_CKTstates), 8 * sizeof(double *)) ; + CUDAMALLOCCHECK (ckt->dD_CKTstates, 8, double *, status) + + status = cudaMemcpy (ckt->dD_CKTstates, ckt->d_CKTstates, 8 * sizeof(double *), cudaMemcpyHostToDevice) ; + CUDAMEMCPYCHECK (ckt->dD_CKTstates, 8, double *, status) + + status = cudaMalloc ((void **)&(ckt->d_CKTdeltaOld), 7 * sizeof(double)) ; + CUDAMALLOCCHECK (ckt->d_CKTdeltaOld, 7, double, status) + +// ckt->CKTtimeSteps = (double *) malloc (size3 * sizeof(double)) ; + status = cudaMalloc ((void **)&(ckt->d_CKTtimeSteps), size3 * sizeof(double)) ; + CUDAMALLOCCHECK (ckt->d_CKTtimeSteps, size3, double, status) + status = cudaMalloc ((void **)&(ckt->d_CKTtimeStepsOut), size3 * sizeof(double)) ; + CUDAMALLOCCHECK (ckt->d_CKTtimeStepsOut, size3, double, status) } - - /* Truncation Error */ - status = cudaMalloc ((void **)&(ckt->dD_CKTstates), 8 * sizeof(double *)) ; - CUDAMALLOCCHECK (ckt->dD_CKTstates, 8, double *, status) - - status = cudaMemcpy (ckt->dD_CKTstates, ckt->d_CKTstates, 8 * sizeof(double *), cudaMemcpyHostToDevice) ; - CUDAMEMCPYCHECK (ckt->dD_CKTstates, 8, double *, status) - - status = cudaMalloc ((void **)&(ckt->d_CKTdeltaOld), 7 * sizeof(double)) ; - CUDAMALLOCCHECK (ckt->d_CKTdeltaOld, 7, double, status) - -// ckt->CKTtimeSteps = (double *) malloc (size3 * sizeof(double)) ; - status = cudaMalloc ((void **)&(ckt->d_CKTtimeSteps), size3 * sizeof(double)) ; - CUDAMALLOCCHECK (ckt->d_CKTtimeSteps, size3, double, status) - status = cudaMalloc ((void **)&(ckt->d_CKTtimeStepsOut), size3 * sizeof(double)) ; - CUDAMALLOCCHECK (ckt->d_CKTtimeStepsOut, size3, double, status) - return (OK) ; } diff --git a/src/spicelib/analysis/CUSPICE/cucktstatesupdate.c b/src/spicelib/analysis/CUSPICE/cucktstatesupdate.c index 65ba5f6ec..d2d26e0cd 100644 --- a/src/spicelib/analysis/CUSPICE/cucktstatesupdate.c +++ b/src/spicelib/analysis/CUSPICE/cucktstatesupdate.c @@ -28,6 +28,7 @@ #include "ngspice/sperror.h" #include "cuda_runtime_api.h" #include "ngspice/CUSPICE/CUSPICE.h" +#include /* cudaMemcpy MACRO to check it for errors --> CUDAMEMCPYCHECK(name of pointer, dimension, type, status) */ #define CUDAMEMCPYCHECK(a, b, c, d) \ @@ -47,8 +48,10 @@ CKTcircuit *ckt { long unsigned int size ; - size = (long unsigned int)ckt->CKTnumStates ; - cudaMemset (ckt->d_CKTstate0, 0, size * sizeof(double)) ; + if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) { + size = (long unsigned int)ckt->CKTnumStates ; + cudaMemset (ckt->d_CKTstate0, 0, size * sizeof(double)) ; + } return (OK) ; } @@ -62,9 +65,11 @@ CKTcircuit *ckt long unsigned int size ; cudaError_t status ; - size = (long unsigned int)ckt->CKTnumStates ; - status = cudaMemcpy (ckt->d_CKTstate0, ckt->CKTstate0, size * sizeof(double), cudaMemcpyHostToDevice) ; - CUDAMEMCPYCHECK (ckt->d_CKTstate0, size, double, status) + if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) { + size = (long unsigned int)ckt->CKTnumStates ; + status = cudaMemcpy (ckt->d_CKTstate0, ckt->CKTstate0, size * sizeof(double), cudaMemcpyHostToDevice) ; + CUDAMEMCPYCHECK (ckt->d_CKTstate0, size, double, status) + } return (OK) ; } @@ -78,9 +83,11 @@ CKTcircuit *ckt long unsigned int size ; cudaError_t status ; - size = (long unsigned int)ckt->CKTnumStates ; - status = cudaMemcpy (ckt->CKTstate0, ckt->d_CKTstate0, size * sizeof(double), cudaMemcpyDeviceToHost) ; - CUDAMEMCPYCHECK (ckt->CKTstate0, size, double, status) + if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) { + size = (long unsigned int)ckt->CKTnumStates ; + status = cudaMemcpy (ckt->CKTstate0, ckt->d_CKTstate0, size * sizeof(double), cudaMemcpyDeviceToHost) ; + CUDAMEMCPYCHECK (ckt->CKTstate0, size, double, status) + } return (OK) ; } @@ -91,12 +98,16 @@ cuCKTstate01copy CKTcircuit *ckt ) { - long unsigned int size ; - cudaError_t status ; + if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) { + long unsigned int size ; + cudaError_t status ; - size = (long unsigned int)ckt->CKTnumStates ; - status = cudaMemcpy (ckt->d_CKTstate1, ckt->d_CKTstate0, size * sizeof(double), cudaMemcpyDeviceToDevice) ; - CUDAMEMCPYCHECK (ckt->d_CKTstate1, size, double, status) + size = (long unsigned int)ckt->CKTnumStates ; + status = cudaMemcpy (ckt->d_CKTstate1, ckt->d_CKTstate0, size * sizeof(double), cudaMemcpyDeviceToDevice) ; + CUDAMEMCPYCHECK (ckt->d_CKTstate1, size, double, status) + } else { + memcpy (ckt->CKTstate1, ckt->CKTstate0, (size_t) ckt->CKTnumStates * sizeof(double)) ; + } return (OK) ; } @@ -110,11 +121,19 @@ CKTcircuit *ckt int i ; double *temp ; - temp = ckt->d_CKTstates [ckt->CKTmaxOrder + 1] ; - for (i = ckt->CKTmaxOrder ; i >= 0 ; i--) - ckt->d_CKTstates [i + 1] = ckt->d_CKTstates [i] ; + if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) { + temp = ckt->d_CKTstates [ckt->CKTmaxOrder + 1] ; + for (i = ckt->CKTmaxOrder ; i >= 0 ; i--) + ckt->d_CKTstates [i + 1] = ckt->d_CKTstates [i] ; - ckt->d_CKTstates [0] = temp ; + ckt->d_CKTstates [0] = temp ; + } else { + temp = ckt->CKTstates [ckt->CKTmaxOrder + 1] ; + for (i = ckt->CKTmaxOrder ; i >= 0 ; i--) { + ckt->CKTstates [i + 1] = ckt->CKTstates [i] ; + } + ckt->CKTstates [0] = temp ; + } return (OK) ; } @@ -125,16 +144,21 @@ cuCKTstate123copy CKTcircuit *ckt ) { - long unsigned int size ; - cudaError_t status ; + if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) { + long unsigned int size ; + cudaError_t status ; - size = (long unsigned int)ckt->CKTnumStates ; + size = (long unsigned int)ckt->CKTnumStates ; - status = cudaMemcpy (ckt->d_CKTstate2, ckt->d_CKTstate1, size * sizeof(double), cudaMemcpyDeviceToDevice) ; - CUDAMEMCPYCHECK (ckt->d_CKTstate2, size, double, status) + status = cudaMemcpy (ckt->d_CKTstate2, ckt->d_CKTstate1, size * sizeof(double), cudaMemcpyDeviceToDevice) ; + CUDAMEMCPYCHECK (ckt->d_CKTstate2, size, double, status) - status = cudaMemcpy (ckt->d_CKTstate3, ckt->d_CKTstate1, size * sizeof(double), cudaMemcpyDeviceToDevice) ; - CUDAMEMCPYCHECK (ckt->d_CKTstate3, size, double, status) + status = cudaMemcpy (ckt->d_CKTstate3, ckt->d_CKTstate1, size * sizeof(double), cudaMemcpyDeviceToDevice) ; + CUDAMEMCPYCHECK (ckt->d_CKTstate3, size, double, status) + } else { + memcpy (ckt->CKTstate2, ckt->CKTstate1, (size_t) ckt->CKTnumStates * sizeof(double)) ; + memcpy (ckt->CKTstate3, ckt->CKTstate1, (size_t) ckt->CKTnumStates * sizeof(double)) ; + } return (OK) ; } @@ -147,8 +171,10 @@ CKTcircuit *ckt { cudaError_t status ; - status = cudaMemcpy (ckt->d_CKTdeltaOld, ckt->CKTdeltaOld, 7 * sizeof(double), cudaMemcpyHostToDevice) ; - CUDAMEMCPYCHECK (ckt->d_CKTdeltaOld, 7, double, status) + if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) { + status = cudaMemcpy (ckt->d_CKTdeltaOld, ckt->CKTdeltaOld, 7 * sizeof(double), cudaMemcpyHostToDevice) ; + CUDAMEMCPYCHECK (ckt->d_CKTdeltaOld, 7, double, status) + } return (OK) ; } diff --git a/src/spicelib/analysis/CUSPICE/cucktsystem.c b/src/spicelib/analysis/CUSPICE/cucktsystem.c index 447a6d654..b186bab28 100644 --- a/src/spicelib/analysis/CUSPICE/cucktsystem.c +++ b/src/spicelib/analysis/CUSPICE/cucktsystem.c @@ -48,16 +48,18 @@ CKTcircuit *ckt long unsigned int nz, n ; cudaError_t status ; - nz = (long unsigned int)ckt->CKTmatrix->CKTklunz ; - n = (long unsigned int)ckt->CKTmatrix->CKTkluN ; + if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) { + nz = (long unsigned int)ckt->CKTmatrix->CKTklunz ; + n = (long unsigned int)ckt->CKTmatrix->CKTkluN ; - /* Copy back the Matrix */ - status = cudaMemcpy (ckt->CKTmatrix->CKTkluAx, ckt->CKTmatrix->d_CKTkluAx, nz * sizeof(double), cudaMemcpyDeviceToHost) ; - CUDAMEMCPYCHECK (ckt->CKTmatrix->CKTkluAx, nz, double, status) + /* Copy back the Matrix */ + status = cudaMemcpy (ckt->CKTmatrix->CKTkluAx, ckt->CKTmatrix->d_CKTkluAx, nz * sizeof(double), cudaMemcpyDeviceToHost) ; + CUDAMEMCPYCHECK (ckt->CKTmatrix->CKTkluAx, nz, double, status) - /* Copy back the RHS */ - status = cudaMemcpy (ckt->CKTrhs, ckt->CKTmatrix->d_CKTrhs, (n + 1) * sizeof(double), cudaMemcpyDeviceToHost) ; - CUDAMEMCPYCHECK (ckt->CKTrhs, (n + 1), double, status) + /* Copy back the RHS */ + status = cudaMemcpy (ckt->CKTrhs, ckt->CKTmatrix->d_CKTrhs, (n + 1) * sizeof(double), cudaMemcpyDeviceToHost) ; + CUDAMEMCPYCHECK (ckt->CKTrhs, (n + 1), double, status) + } return (OK) ; } diff --git a/src/spicelib/analysis/CUSPICE/cuckttrunc.cu b/src/spicelib/analysis/CUSPICE/cuckttrunc.cu index b215b730f..c3d159968 100644 --- a/src/spicelib/analysis/CUSPICE/cuckttrunc.cu +++ b/src/spicelib/analysis/CUSPICE/cuckttrunc.cu @@ -31,62 +31,66 @@ cuCKTtrunc CKTcircuit *ckt, double timetemp, double *timeStep ) { - long unsigned int size ; - double timetempGPU ; - int thread_x, thread_y, block_x ; + if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) { + long unsigned int size ; + double timetempGPU ; + int thread_x, thread_y, block_x ; - cudaError_t status ; + cudaError_t status ; - /* Determining how many blocks should exist in the kernel */ - thread_x = 1 ; - thread_y = 256 ; - if (ckt->total_n_timeSteps % thread_y != 0) - block_x = (int)((ckt->total_n_timeSteps + thread_y - 1) / thread_y) ; - else - block_x = ckt->total_n_timeSteps / thread_y ; + /* Determining how many blocks should exist in the kernel */ + thread_x = 1 ; + thread_y = 256 ; + if (ckt->total_n_timeSteps % thread_y != 0) + block_x = (int)((ckt->total_n_timeSteps + thread_y - 1) / thread_y) ; + else + block_x = ckt->total_n_timeSteps / thread_y ; - dim3 thread (thread_x, thread_y) ; + dim3 thread (thread_x, thread_y) ; - /* Kernel launch */ - status = cudaGetLastError () ; // clear error status + /* Kernel launch */ + status = cudaGetLastError () ; // clear error status - cuCKTtrunc_kernel <<< block_x, thread, thread_y * sizeof(double) >>> (ckt->d_CKTtimeSteps, ckt->d_CKTtimeStepsOut, ckt->total_n_timeSteps) ; + cuCKTtrunc_kernel <<< block_x, thread, thread_y * sizeof(double) >>> (ckt->d_CKTtimeSteps, ckt->d_CKTtimeStepsOut, ckt->total_n_timeSteps) ; - cudaDeviceSynchronize () ; + cudaDeviceSynchronize () ; - status = cudaGetLastError () ; // check for launch error - if (status != cudaSuccess) - { - fprintf (stderr, "Kernel 1 launch failure in cuCKTtrunc\n\n") ; - return (E_NOMEM) ; - } + status = cudaGetLastError () ; // check for launch error + if (status != cudaSuccess) + { + fprintf (stderr, "Kernel 1 launch failure in cuCKTtrunc\n\n") ; + return (E_NOMEM) ; + } - cuCKTtrunc_kernel <<< 1, thread, thread_y * sizeof(double) >>> (ckt->d_CKTtimeStepsOut, ckt->d_CKTtimeSteps, block_x) ; + cuCKTtrunc_kernel <<< 1, thread, thread_y * sizeof(double) >>> (ckt->d_CKTtimeStepsOut, ckt->d_CKTtimeSteps, block_x) ; - cudaDeviceSynchronize () ; + cudaDeviceSynchronize () ; - status = cudaGetLastError () ; // check for launch error - if (status != cudaSuccess) - { - fprintf (stderr, "Kernel 2 launch failure in cuCKTtrunc\n\n") ; - return (E_NOMEM) ; - } + status = cudaGetLastError () ; // check for launch error + if (status != cudaSuccess) + { + fprintf (stderr, "Kernel 2 launch failure in cuCKTtrunc\n\n") ; + return (E_NOMEM) ; + } - /* Copy back the reduction result */ - size = (long unsigned int)(1) ; - status = cudaMemcpy (&timetempGPU, ckt->d_CKTtimeSteps, size * sizeof(double), cudaMemcpyDeviceToHost) ; - CUDAMEMCPYCHECK (&timetempGPU, size, double, status) + /* Copy back the reduction result */ + size = (long unsigned int)(1) ; + status = cudaMemcpy (&timetempGPU, ckt->d_CKTtimeSteps, size * sizeof(double), cudaMemcpyDeviceToHost) ; + CUDAMEMCPYCHECK (&timetempGPU, size, double, status) - /* Final Comparison */ - if (timetempGPU < timetemp) - { - timetemp = timetempGPU ; - } - if (2 * *timeStep < timetemp) - { - *timeStep = 2 * *timeStep ; + /* Final Comparison */ + if (timetempGPU < timetemp) + { + timetemp = timetempGPU ; + } + if (2 * *timeStep < timetemp) + { + *timeStep = 2 * *timeStep ; + } else { + *timeStep = timetemp ; + } } else { - *timeStep = timetemp ; + *timeStep = MIN (2 * *timeStep, timetemp) ; } return 0 ; diff --git a/src/spicelib/analysis/cktload.c b/src/spicelib/analysis/cktload.c index 8b4e2a697..be3659642 100644 --- a/src/spicelib/analysis/cktload.c +++ b/src/spicelib/analysis/cktload.c @@ -103,8 +103,15 @@ CKTload(CKTcircuit *ckt) return (E_NOMEM) ; #endif + /* Load Sparse Matrix and RHS of all the CUDA supported models */ for (i = 0; i < DEVmaxnum; i++) { + +#ifdef USE_CUSPICE + if (DEVices[i] && DEVices[i]->DEVload && ckt->CKThead[i] && ckt->CKThead[i]->has_cuda) { +#else if (DEVices[i] && DEVices[i]->DEVload && ckt->CKThead[i]) { +#endif + error = DEVices[i]->DEVload (ckt->CKThead[i], ckt); #ifdef USE_CUSPICE @@ -127,50 +134,77 @@ CKTload(CKTcircuit *ckt) } #ifdef USE_CUSPICE - /* Copy the CKTdiagGmin value to the GPU */ - // The real Gmin is needed only when the matrix will reside entirely on the GPU - // Right now, only some models support CUDA, so the matrix is only partially created on the GPU - cudaMemset (ckt->d_CKTloadOutput + ckt->total_n_values, 0, sizeof(double)) ; - //cudaError_t statusMemcpy ; - //statusMemcpy = cudaMemcpy (ckt->d_CKTloadOutput + ckt->total_n_values, &(ckt->CKTdiagGmin), sizeof(double), cudaMemcpyHostToDevice) ; - //CUDAMEMCPYCHECK (ckt->d_CKTloadOutput + ckt->total_n_values, 1, double, statusMemcpy) + int TopologyNNZ, TopologyNNZRHS ; - /* Performing CSRMV for the Sparse Matrix using CUSPARSE */ - cusparseStatus = cusparseDcsrmv ((cusparseHandle_t)(ckt->CKTmatrix->CKTcsrmvHandle), - CUSPARSE_OPERATION_NON_TRANSPOSE, - ckt->CKTmatrix->CKTklunz, ckt->total_n_values + 1, - ckt->total_n_Ptr + ckt->CKTdiagElements, - &alpha, (cusparseMatDescr_t)(ckt->CKTmatrix->CKTcsrmvDescr), - ckt->d_CKTtopologyMatrixCSRx, ckt->d_CKTtopologyMatrixCSRp, - ckt->d_CKTtopologyMatrixCSRj, ckt->d_CKTloadOutput, &beta, - ckt->CKTmatrix->d_CKTkluAx) ; + TopologyNNZ = ckt->total_n_Ptr + ckt->CKTdiagElements ; // + ckt->CKTdiagElements because of CKTdiagGmin + // without the zeroes along the diagonal + TopologyNNZRHS = ckt->total_n_PtrRHS ; - if (cusparseStatus != CUSPARSE_STATUS_SUCCESS) - { - fprintf (stderr, "CUSPARSE MATRIX Call Error\n") ; - return (E_NOMEM) ; + if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) { + /* Copy the CKTdiagGmin value to the GPU */ + // The real Gmin is needed only when the matrix will reside entirely on the GPU + // Right now, only some models support CUDA, so the matrix is only partially created on the GPU + cudaMemset (ckt->d_CKTloadOutput + ckt->total_n_values, 0, sizeof(double)) ; + //cudaError_t statusMemcpy ; + //statusMemcpy = cudaMemcpy (ckt->d_CKTloadOutput + ckt->total_n_values, &(ckt->CKTdiagGmin), sizeof(double), cudaMemcpyHostToDevice) ; + //CUDAMEMCPYCHECK (ckt->d_CKTloadOutput + ckt->total_n_values, 1, double, statusMemcpy) + + /* Performing CSRMV for the Sparse Matrix using CUSPARSE */ + cusparseStatus = cusparseDcsrmv ((cusparseHandle_t)(ckt->CKTmatrix->CKTcsrmvHandle), + CUSPARSE_OPERATION_NON_TRANSPOSE, + ckt->CKTmatrix->CKTklunz, ckt->total_n_values + 1, + ckt->total_n_Ptr + ckt->CKTdiagElements, + &alpha, (cusparseMatDescr_t)(ckt->CKTmatrix->CKTcsrmvDescr), + ckt->d_CKTtopologyMatrixCSRx, ckt->d_CKTtopologyMatrixCSRp, + ckt->d_CKTtopologyMatrixCSRj, ckt->d_CKTloadOutput, &beta, + ckt->CKTmatrix->d_CKTkluAx) ; + + if (cusparseStatus != CUSPARSE_STATUS_SUCCESS) + { + fprintf (stderr, "CUSPARSE MATRIX Call Error\n") ; + return (E_NOMEM) ; + } + + /* Performing CSRMV for the RHS using CUSPARSE */ + cusparseStatus = cusparseDcsrmv ((cusparseHandle_t)(ckt->CKTmatrix->CKTcsrmvHandle), + CUSPARSE_OPERATION_NON_TRANSPOSE, + ckt->CKTmatrix->CKTkluN + 1, ckt->total_n_valuesRHS, ckt->total_n_PtrRHS, + &alpha, (cusparseMatDescr_t)(ckt->CKTmatrix->CKTcsrmvDescr), + ckt->d_CKTtopologyMatrixCSRxRHS, ckt->d_CKTtopologyMatrixCSRpRHS, + ckt->d_CKTtopologyMatrixCSRjRHS, ckt->d_CKTloadOutputRHS, &beta, + ckt->CKTmatrix->d_CKTrhs) ; + + if (cusparseStatus != CUSPARSE_STATUS_SUCCESS) + { + fprintf (stderr, "CUSPARSE RHS Call Error\n") ; + return (E_NOMEM) ; + } + + cudaDeviceSynchronize () ; + + status = cuCKTsystemDtoH (ckt) ; + if (status != 0) + return (E_NOMEM) ; } - /* Performing CSRMV for the RHS using CUSPARSE */ - cusparseStatus = cusparseDcsrmv ((cusparseHandle_t)(ckt->CKTmatrix->CKTcsrmvHandle), - CUSPARSE_OPERATION_NON_TRANSPOSE, - ckt->CKTmatrix->CKTkluN + 1, ckt->total_n_valuesRHS, ckt->total_n_PtrRHS, - &alpha, (cusparseMatDescr_t)(ckt->CKTmatrix->CKTcsrmvDescr), - ckt->d_CKTtopologyMatrixCSRxRHS, ckt->d_CKTtopologyMatrixCSRpRHS, - ckt->d_CKTtopologyMatrixCSRjRHS, ckt->d_CKTloadOutputRHS, &beta, - ckt->CKTmatrix->d_CKTrhs) ; + /* Load Sparse Matrix and RHS of all the CUDA unsupported models */ + for (i = 0; i < DEVmaxnum; i++) { + if (DEVices[i] && DEVices[i]->DEVload && ckt->CKThead[i] && !ckt->CKThead[i]->has_cuda) { + error = DEVices[i]->DEVload (ckt->CKThead[i], ckt); - if (cusparseStatus != CUSPARSE_STATUS_SUCCESS) - { - fprintf (stderr, "CUSPARSE RHS Call Error\n") ; - return (E_NOMEM) ; + if (ckt->CKTnoncon) + ckt->CKTtroubleNode = 0; +#ifdef STEPDEBUG + if (noncon != ckt->CKTnoncon) { + printf("device type %s nonconvergence\n", + DEVices[i]->DEVpublic.name); + noncon = ckt->CKTnoncon; + } +#endif /* STEPDEBUG */ + if (error) return(error); + } } - cudaDeviceSynchronize () ; - - status = cuCKTsystemDtoH (ckt) ; - if (status != 0) - return (E_NOMEM) ; #endif #ifdef XSPICE diff --git a/src/spicelib/analysis/cktsetup.c b/src/spicelib/analysis/cktsetup.c index 8145f6e00..379faea1a 100644 --- a/src/spicelib/analysis/cktsetup.c +++ b/src/spicelib/analysis/cktsetup.c @@ -246,6 +246,8 @@ CKTsetup(CKTcircuit *ckt) ckt->CKTtopologyMatrixCOOxRHS = TMALLOC (double, TopologyNNZRHS) ; + if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) { + /* Topology Matrix Pre-Allocation in CSR format */ ckt->CKTtopologyMatrixCSRp = TMALLOC (int, nz + 1) ; @@ -325,6 +327,7 @@ CKTsetup(CKTcircuit *ckt) ret = Compress (ckt->CKTtopologyMatrixCOOiRHS, ckt->CKTtopologyMatrixCSRpRHS, n + 1, TopologyNNZRHS) ; /* Multiply the Topology Matrix by the M Vector to build the Final CSC Matrix - after the CKTload Call */ + } #endif } else { @@ -337,30 +340,32 @@ CKTsetup(CKTcircuit *ckt) } #ifdef USE_CUSPICE - ckt->d_MatrixSize = SMPmatSize (ckt->CKTmatrix) ; - status = cuCKTsetup (ckt) ; - if (status != 0) - return (E_NOMEM) ; + if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) { + ckt->d_MatrixSize = SMPmatSize (ckt->CKTmatrix) ; + status = cuCKTsetup (ckt) ; + if (status != 0) + return (E_NOMEM) ; - /* CUSPARSE Handle Creation */ - cusparseStatus = cusparseCreate ((cusparseHandle_t *)(&(ckt->CKTmatrix->CKTcsrmvHandle))) ; - if (cusparseStatus != CUSPARSE_STATUS_SUCCESS) - { - fprintf (stderr, "CUSPARSE Handle Setup Error\n") ; - return (E_NOMEM) ; + /* CUSPARSE Handle Creation */ + cusparseStatus = cusparseCreate ((cusparseHandle_t *)(&(ckt->CKTmatrix->CKTcsrmvHandle))) ; + if (cusparseStatus != CUSPARSE_STATUS_SUCCESS) + { + fprintf (stderr, "CUSPARSE Handle Setup Error\n") ; + return (E_NOMEM) ; + } + + /* CUSPARSE Matrix Descriptor Creation */ + cusparseStatus = cusparseCreateMatDescr ((cusparseMatDescr_t *)(&(ckt->CKTmatrix->CKTcsrmvDescr))) ; + if (cusparseStatus != CUSPARSE_STATUS_SUCCESS) + { + fprintf (stderr, "CUSPARSE Matrix Descriptor Setup Error\n") ; + return (E_NOMEM) ; + } + + /* CUSPARSE Matrix Properties Definition */ + cusparseSetMatType ((cusparseMatDescr_t)(ckt->CKTmatrix->CKTcsrmvDescr), CUSPARSE_MATRIX_TYPE_GENERAL) ; + cusparseSetMatIndexBase ((cusparseMatDescr_t)(ckt->CKTmatrix->CKTcsrmvDescr), CUSPARSE_INDEX_BASE_ZERO) ; } - - /* CUSPARSE Matrix Descriptor Creation */ - cusparseStatus = cusparseCreateMatDescr ((cusparseMatDescr_t *)(&(ckt->CKTmatrix->CKTcsrmvDescr))) ; - if (cusparseStatus != CUSPARSE_STATUS_SUCCESS) - { - fprintf (stderr, "CUSPARSE Matrix Descriptor Setup Error\n") ; - return (E_NOMEM) ; - } - - /* CUSPARSE Matrix Properties Definition */ - cusparseSetMatType ((cusparseMatDescr_t)(ckt->CKTmatrix->CKTcsrmvDescr), CUSPARSE_MATRIX_TYPE_GENERAL) ; - cusparseSetMatIndexBase ((cusparseMatDescr_t)(ckt->CKTmatrix->CKTcsrmvDescr), CUSPARSE_INDEX_BASE_ZERO) ; #endif #ifdef WANT_SENSE2 diff --git a/src/spicelib/devices/asrc/asrcdefs.h b/src/spicelib/devices/asrc/asrcdefs.h index 865504307..0652c4669 100644 --- a/src/spicelib/devices/asrc/asrcdefs.h +++ b/src/spicelib/devices/asrc/asrcdefs.h @@ -71,6 +71,12 @@ typedef struct sASRCmodel { /* model structure for a source */ #define ASRCinstances(inst) ((ASRCinstance *)((inst)->gen.GENinstances)) #define ASRCmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + } ASRCmodel; diff --git a/src/spicelib/devices/asrc/asrcset.c b/src/spicelib/devices/asrc/asrcset.c index bff162298..9f49cc5e3 100644 --- a/src/spicelib/devices/asrc/asrcset.c +++ b/src/spicelib/devices/asrc/asrcset.c @@ -32,6 +32,12 @@ ASRCsetup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *states) NG_IGNORE(states); for (; model; model = ASRCnextModel(model)) { + +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + for (here = ASRCinstances(model); here; here=ASRCnextInstance(here)) { if (!here->ASRCtree) diff --git a/src/spicelib/devices/bjt/bjtdefs.h b/src/spicelib/devices/bjt/bjtdefs.h index f3fdd675a..c2bf85025 100644 --- a/src/spicelib/devices/bjt/bjtdefs.h +++ b/src/spicelib/devices/bjt/bjtdefs.h @@ -360,6 +360,12 @@ typedef struct sBJTmodel { /* model structure for a bjt */ #define BJTinstances(inst) ((BJTinstance *)((inst)->gen.GENinstances)) #define BJTmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int BJTtype; int BJTsubs; diff --git a/src/spicelib/devices/bjt/bjtsetup.c b/src/spicelib/devices/bjt/bjtsetup.c index e2cc3d3ba..d3bf16f3c 100644 --- a/src/spicelib/devices/bjt/bjtsetup.c +++ b/src/spicelib/devices/bjt/bjtsetup.c @@ -34,6 +34,11 @@ BJTsetup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *states) /* loop through all the diode models */ for( ; model != NULL; model = BJTnextModel(model)) { +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + if(model->BJTtype != NPN && model->BJTtype != PNP) { model->BJTtype = NPN; } diff --git a/src/spicelib/devices/bsim1/b1set.c b/src/spicelib/devices/bsim1/b1set.c index af910fe7b..d8079f1c9 100644 --- a/src/spicelib/devices/bsim1/b1set.c +++ b/src/spicelib/devices/bsim1/b1set.c @@ -26,8 +26,13 @@ B1setup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, /* loop through all the B1 device models */ for( ; model != NULL; model = B1nextModel(model)) { - -/* Default value Processing for B1 MOSFET Models */ + +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + + /* Default value Processing for B1 MOSFET Models */ if( ! model->B1typeGiven) { model->B1type = NMOS; /* NMOS */ } diff --git a/src/spicelib/devices/bsim1/bsim1def.h b/src/spicelib/devices/bsim1/bsim1def.h index bed659380..d99d3ac9c 100644 --- a/src/spicelib/devices/bsim1/bsim1def.h +++ b/src/spicelib/devices/bsim1/bsim1def.h @@ -349,6 +349,12 @@ typedef struct sBSIM1model { /* model structure for a resistor */ #define B1instances(inst) ((B1instance *)((inst)->gen.GENinstances)) #define B1modName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int B1type; /* device type : 1 = nmos, -1 = pmos */ double B1vfb0; diff --git a/src/spicelib/devices/bsim2/b2set.c b/src/spicelib/devices/bsim2/b2set.c index c868d4811..af9f934ef 100644 --- a/src/spicelib/devices/bsim2/b2set.c +++ b/src/spicelib/devices/bsim2/b2set.c @@ -25,8 +25,13 @@ B2setup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *states) /* loop through all the B2 device models */ for( ; model != NULL; model = B2nextModel(model)) { - -/* Default value Processing for B2 MOSFET Models */ + +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + + /* Default value Processing for B2 MOSFET Models */ if( ! model->B2typeGiven) { model->B2type = NMOS; /* NMOS */ } diff --git a/src/spicelib/devices/bsim2/bsim2def.h b/src/spicelib/devices/bsim2/bsim2def.h index 078f8a62a..178fe2662 100644 --- a/src/spicelib/devices/bsim2/bsim2def.h +++ b/src/spicelib/devices/bsim2/bsim2def.h @@ -272,6 +272,12 @@ typedef struct sBSIM2model { /* model structure for a resistor */ #define B2instances(inst) ((B2instance *)((inst)->gen.GENinstances)) #define B2modName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int B2type; /* device type: 1 = nmos, -1 = pmos */ int pad; diff --git a/src/spicelib/devices/bsim3/b3set.c b/src/spicelib/devices/bsim3/b3set.c index 3480b47ee..d2bb9fb87 100644 --- a/src/spicelib/devices/bsim3/b3set.c +++ b/src/spicelib/devices/bsim3/b3set.c @@ -51,7 +51,13 @@ BSIM3instance **InstArray; /* loop through all the BSIM3 device models */ for( ; model != NULL; model = BSIM3nextModel(model)) { -/* Default value Processing for BSIM3 MOSFET Models */ + +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + + /* Default value Processing for BSIM3 MOSFET Models */ if (!model->BSIM3typeGiven) model->BSIM3type = NMOS; if (!model->BSIM3mobModGiven) diff --git a/src/spicelib/devices/bsim3/bsim3def.h b/src/spicelib/devices/bsim3/bsim3def.h index dbb195994..5342b99fc 100644 --- a/src/spicelib/devices/bsim3/bsim3def.h +++ b/src/spicelib/devices/bsim3/bsim3def.h @@ -445,6 +445,12 @@ typedef struct sBSIM3model #define BSIM3instances(inst) ((BSIM3instance *)((inst)->gen.GENinstances)) #define BSIM3modName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int BSIM3type; int BSIM3mobMod; diff --git a/src/spicelib/devices/bsim3soi_dd/b3soidddef.h b/src/spicelib/devices/bsim3soi_dd/b3soidddef.h index 67cef35e2..f9757add1 100644 --- a/src/spicelib/devices/bsim3soi_dd/b3soidddef.h +++ b/src/spicelib/devices/bsim3soi_dd/b3soidddef.h @@ -704,6 +704,12 @@ typedef struct sB3SOIDDmodel #define B3SOIDDinstances(inst) ((B3SOIDDinstance *)((inst)->gen.GENinstances)) #define B3SOIDDmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int B3SOIDDtype; int B3SOIDDmobMod; diff --git a/src/spicelib/devices/bsim3soi_dd/b3soiddset.c b/src/spicelib/devices/bsim3soi_dd/b3soiddset.c index a9fc99367..d7a9ea27d 100644 --- a/src/spicelib/devices/bsim3soi_dd/b3soiddset.c +++ b/src/spicelib/devices/bsim3soi_dd/b3soiddset.c @@ -49,8 +49,13 @@ IFuid tmpName; /* loop through all the B3SOIDD device models */ for( ; model != NULL; model = B3SOIDDnextModel(model)) { -/* Default value Processing for B3SOIDD MOSFET Models */ +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + + /* Default value Processing for B3SOIDD MOSFET Models */ if (!model->B3SOIDDtypeGiven) model->B3SOIDDtype = NMOS; if (!model->B3SOIDDmobModGiven) diff --git a/src/spicelib/devices/bsim3soi_fd/b3soifddef.h b/src/spicelib/devices/bsim3soi_fd/b3soifddef.h index 6fe56a250..f52cfc059 100644 --- a/src/spicelib/devices/bsim3soi_fd/b3soifddef.h +++ b/src/spicelib/devices/bsim3soi_fd/b3soifddef.h @@ -696,6 +696,12 @@ typedef struct sB3SOIFDmodel #define B3SOIFDinstances(inst) ((B3SOIFDinstance *)((inst)->gen.GENinstances)) #define B3SOIFDmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int B3SOIFDtype; int B3SOIFDmobMod; diff --git a/src/spicelib/devices/bsim3soi_fd/b3soifdset.c b/src/spicelib/devices/bsim3soi_fd/b3soifdset.c index d6ee51654..ca0bf8f5b 100644 --- a/src/spicelib/devices/bsim3soi_fd/b3soifdset.c +++ b/src/spicelib/devices/bsim3soi_fd/b3soifdset.c @@ -49,8 +49,13 @@ IFuid tmpName; /* loop through all the B3SOIFD device models */ for( ; model != NULL; model = B3SOIFDnextModel(model)) { -/* Default value Processing for B3SOIFD MOSFET Models */ +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + + /* Default value Processing for B3SOIFD MOSFET Models */ if (!model->B3SOIFDtypeGiven) model->B3SOIFDtype = NMOS; if (!model->B3SOIFDmobModGiven) diff --git a/src/spicelib/devices/bsim3soi_pd/b3soipddef.h b/src/spicelib/devices/bsim3soi_pd/b3soipddef.h index 779a92c6e..276b177bc 100644 --- a/src/spicelib/devices/bsim3soi_pd/b3soipddef.h +++ b/src/spicelib/devices/bsim3soi_pd/b3soipddef.h @@ -665,6 +665,12 @@ typedef struct sB3SOIPDmodel #define B3SOIPDinstances(inst) ((B3SOIPDinstance *)((inst)->gen.GENinstances)) #define B3SOIPDmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int B3SOIPDtype; int B3SOIPDmobMod; diff --git a/src/spicelib/devices/bsim3soi_pd/b3soipdset.c b/src/spicelib/devices/bsim3soi_pd/b3soipdset.c index 9136e0289..f2bcf1a6e 100644 --- a/src/spicelib/devices/bsim3soi_pd/b3soipdset.c +++ b/src/spicelib/devices/bsim3soi_pd/b3soipdset.c @@ -49,8 +49,13 @@ IFuid tmpName; /* loop through all the B3SOIPD device models */ for( ; model != NULL; model = B3SOIPDnextModel(model)) { -/* Default value Processing for B3SOIPD MOSFET Models */ +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + + /* Default value Processing for B3SOIPD MOSFET Models */ if (!model->B3SOIPDtypeGiven) model->B3SOIPDtype = NMOS; if (!model->B3SOIPDmobModGiven) diff --git a/src/spicelib/devices/bsim3v0/b3v0set.c b/src/spicelib/devices/bsim3v0/b3v0set.c index d518d0e24..29b756f30 100644 --- a/src/spicelib/devices/bsim3v0/b3v0set.c +++ b/src/spicelib/devices/bsim3v0/b3v0set.c @@ -37,7 +37,13 @@ IFuid tmpName; /* loop through all the BSIM3v0 device models */ for( ; model != NULL; model = BSIM3v0nextModel(model)) { -/* Default value Processing for BSIM3v0 MOSFET Models */ + +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + + /* Default value Processing for BSIM3v0 MOSFET Models */ if (!model->BSIM3v0typeGiven) model->BSIM3v0type = NMOS; if (!model->BSIM3v0mobModGiven) diff --git a/src/spicelib/devices/bsim3v0/bsim3v0def.h b/src/spicelib/devices/bsim3v0/bsim3v0def.h index de1b8a4c2..8fe76be05 100644 --- a/src/spicelib/devices/bsim3v0/bsim3v0def.h +++ b/src/spicelib/devices/bsim3v0/bsim3v0def.h @@ -363,6 +363,12 @@ typedef struct sBSIM3v0model #define BSIM3v0instances(inst) ((BSIM3v0instance *)((inst)->gen.GENinstances)) #define BSIM3v0modName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int BSIM3v0type; int BSIM3v0mobMod; diff --git a/src/spicelib/devices/bsim3v1/b3v1set.c b/src/spicelib/devices/bsim3v1/b3v1set.c index 932a1ed47..8cc6a373a 100644 --- a/src/spicelib/devices/bsim3v1/b3v1set.c +++ b/src/spicelib/devices/bsim3v1/b3v1set.c @@ -43,7 +43,13 @@ IFuid tmpName; /* loop through all the BSIM3v1 device models */ for( ; model != NULL; model = BSIM3v1nextModel(model)) { -/* Default value Processing for BSIM3v1 MOSFET Models */ + +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + + /* Default value Processing for BSIM3v1 MOSFET Models */ if (!model->BSIM3v1typeGiven) model->BSIM3v1type = NMOS; if (!model->BSIM3v1mobModGiven) diff --git a/src/spicelib/devices/bsim3v1/bsim3v1def.h b/src/spicelib/devices/bsim3v1/bsim3v1def.h index c37679eed..7f986c94e 100644 --- a/src/spicelib/devices/bsim3v1/bsim3v1def.h +++ b/src/spicelib/devices/bsim3v1/bsim3v1def.h @@ -365,6 +365,12 @@ typedef struct sBSIM3v1model #define BSIM3v1instances(inst) ((BSIM3v1instance *)((inst)->gen.GENinstances)) #define BSIM3v1modName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int BSIM3v1type; int BSIM3v1mobMod; diff --git a/src/spicelib/devices/bsim3v32/b3v32set.c b/src/spicelib/devices/bsim3v32/b3v32set.c index e0c0cd43a..d4d441c9a 100644 --- a/src/spicelib/devices/bsim3v32/b3v32set.c +++ b/src/spicelib/devices/bsim3v32/b3v32set.c @@ -47,7 +47,13 @@ BSIM3v32instance **InstArray; /* loop through all the BSIM3v32 device models */ for( ; model != NULL; model = BSIM3v32nextModel(model)) { -/* Default value Processing for BSIM3v32 MOSFET Models */ + +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + + /* Default value Processing for BSIM3v32 MOSFET Models */ if (!model->BSIM3v32typeGiven) model->BSIM3v32type = NMOS; if (!model->BSIM3v32mobModGiven) diff --git a/src/spicelib/devices/bsim3v32/bsim3v32def.h b/src/spicelib/devices/bsim3v32/bsim3v32def.h index ae9aa7176..4218446d8 100644 --- a/src/spicelib/devices/bsim3v32/bsim3v32def.h +++ b/src/spicelib/devices/bsim3v32/bsim3v32def.h @@ -442,6 +442,12 @@ typedef struct sBSIM3v32model #define BSIM3v32instances(inst) ((BSIM3v32instance *)((inst)->gen.GENinstances)) #define BSIM3v32modName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int BSIM3v32type; int BSIM3v32mobMod; diff --git a/src/spicelib/devices/bsim4/b4set.c b/src/spicelib/devices/bsim4/b4set.c index 01d16d8a4..ac240d868 100644 --- a/src/spicelib/devices/bsim4/b4set.c +++ b/src/spicelib/devices/bsim4/b4set.c @@ -110,7 +110,14 @@ BSIM4instance **InstArray; /* loop through all the BSIM4 device models */ for( ; model != NULL; model = BSIM4nextModel(model)) - { /* process defaults of model parameters */ + { + +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + + /* process defaults of model parameters */ if (!model->BSIM4typeGiven) model->BSIM4type = NMOS; diff --git a/src/spicelib/devices/bsim4/bsim4def.h b/src/spicelib/devices/bsim4/bsim4def.h index 76e22ea5c..211b21137 100644 --- a/src/spicelib/devices/bsim4/bsim4def.h +++ b/src/spicelib/devices/bsim4/bsim4def.h @@ -942,6 +942,12 @@ typedef struct sBSIM4model #define BSIM4instances(inst) ((BSIM4instance *)((inst)->gen.GENinstances)) #define BSIM4modName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int BSIM4type; int BSIM4mobMod; diff --git a/src/spicelib/devices/bsim4v5/b4v5set.c b/src/spicelib/devices/bsim4v5/b4v5set.c index 1f26487e2..1409d8b27 100644 --- a/src/spicelib/devices/bsim4v5/b4v5set.c +++ b/src/spicelib/devices/bsim4v5/b4v5set.c @@ -67,7 +67,14 @@ BSIM4v5instance **InstArray; /* loop through all the BSIM4v5 device models */ for( ; model != NULL; model = BSIM4v5nextModel(model)) - { /* process defaults of model parameters */ + { + +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + + /* process defaults of model parameters */ if (!model->BSIM4v5typeGiven) model->BSIM4v5type = NMOS; diff --git a/src/spicelib/devices/bsim4v5/bsim4v5def.h b/src/spicelib/devices/bsim4v5/bsim4v5def.h index b7d76e037..8b4d0beb9 100644 --- a/src/spicelib/devices/bsim4v5/bsim4v5def.h +++ b/src/spicelib/devices/bsim4v5/bsim4v5def.h @@ -859,6 +859,12 @@ typedef struct sBSIM4v5model #define BSIM4v5instances(inst) ((BSIM4v5instance *)((inst)->gen.GENinstances)) #define BSIM4v5modName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int BSIM4v5type; int BSIM4v5mobMod; diff --git a/src/spicelib/devices/bsim4v6/b4v6set.c b/src/spicelib/devices/bsim4v6/b4v6set.c index 1fa4deee0..a92015edb 100644 --- a/src/spicelib/devices/bsim4v6/b4v6set.c +++ b/src/spicelib/devices/bsim4v6/b4v6set.c @@ -74,7 +74,14 @@ BSIM4v6instance **InstArray; /* loop through all the BSIM4v6 device models */ for( ; model != NULL; model = BSIM4v6nextModel(model)) - { /* process defaults of model parameters */ + { + +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + + /* process defaults of model parameters */ if (!model->BSIM4v6typeGiven) model->BSIM4v6type = NMOS; diff --git a/src/spicelib/devices/bsim4v6/bsim4v6def.h b/src/spicelib/devices/bsim4v6/bsim4v6def.h index abd97f3c2..81640929d 100644 --- a/src/spicelib/devices/bsim4v6/bsim4v6def.h +++ b/src/spicelib/devices/bsim4v6/bsim4v6def.h @@ -880,6 +880,12 @@ typedef struct sBSIM4v6model #define BSIM4v6instances(inst) ((BSIM4v6instance *)((inst)->gen.GENinstances)) #define BSIM4v6modName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int BSIM4v6type; int BSIM4v6mobMod; diff --git a/src/spicelib/devices/bsim4v7/b4v7set.c b/src/spicelib/devices/bsim4v7/b4v7set.c index 91bfe2c9d..67dc01e2b 100644 --- a/src/spicelib/devices/bsim4v7/b4v7set.c +++ b/src/spicelib/devices/bsim4v7/b4v7set.c @@ -2601,6 +2601,9 @@ do { if((here->ptr = SMPmakeElt(matrix,here->first,here->second))==(double *)NUL /* How much instances we have */ model->n_instances = i ; + + /* This model supports CUDA */ + model->has_cuda = 1 ; } /* loop through all the BSIM4v7 models */ diff --git a/src/spicelib/devices/bsim4v7/bsim4v7def.h b/src/spicelib/devices/bsim4v7/bsim4v7def.h index 7423ca786..7b088d7f2 100644 --- a/src/spicelib/devices/bsim4v7/bsim4v7def.h +++ b/src/spicelib/devices/bsim4v7/bsim4v7def.h @@ -1435,6 +1435,12 @@ typedef struct sBSIM4v7model #define BSIM4v7instances(inst) ((BSIM4v7instance *)((inst)->gen.GENinstances)) #define BSIM4v7modName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int BSIM4v7type; int BSIM4v7mobMod; diff --git a/src/spicelib/devices/bsimsoi/b4soidef.h b/src/spicelib/devices/bsimsoi/b4soidef.h index 2a3497120..286cfd19b 100644 --- a/src/spicelib/devices/bsimsoi/b4soidef.h +++ b/src/spicelib/devices/bsimsoi/b4soidef.h @@ -1195,6 +1195,12 @@ typedef struct sB4SOImodel #define B4SOIinstances(inst) ((B4SOIinstance *)((inst)->gen.GENinstances)) #define B4SOImodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int B4SOItype; int B4SOImobMod; diff --git a/src/spicelib/devices/bsimsoi/b4soiset.c b/src/spicelib/devices/bsimsoi/b4soiset.c index 70965f5f2..dd402aec4 100644 --- a/src/spicelib/devices/bsimsoi/b4soiset.c +++ b/src/spicelib/devices/bsimsoi/b4soiset.c @@ -61,8 +61,13 @@ B4SOIinstance **InstArray; /* loop through all the B4SOI device models */ for( ; model != NULL; model = B4SOInextModel(model)) { -/* Default value Processing for B4SOI MOSFET Models */ +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + + /* Default value Processing for B4SOI MOSFET Models */ if (!model->B4SOItypeGiven) model->B4SOItype = NMOS; if (!model->B4SOImobModGiven) diff --git a/src/spicelib/devices/cap/capdefs.h b/src/spicelib/devices/cap/capdefs.h index 604f6d40e..1ca0fd803 100644 --- a/src/spicelib/devices/cap/capdefs.h +++ b/src/spicelib/devices/cap/capdefs.h @@ -120,6 +120,12 @@ typedef struct sCAPmodel { /* model structure for a capacitor */ #define CAPinstances(inst) ((CAPinstance *)((inst)->gen.GENinstances)) #define CAPmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + double CAPtnom; /* temperature at which capacitance measured */ double CAPtempCoeff1; /* linear temperature coefficient */ double CAPtempCoeff2; /* quadratic temperature coefficient */ diff --git a/src/spicelib/devices/cap/capsetup.c b/src/spicelib/devices/cap/capsetup.c index 671e8e4d0..d4ecdab78 100644 --- a/src/spicelib/devices/cap/capsetup.c +++ b/src/spicelib/devices/cap/capsetup.c @@ -136,6 +136,9 @@ do { if((here->ptr = SMPmakeElt(matrix, here->first, here->second)) == NULL){\ /* How much instances we have */ model->n_instances = i ; + + /* This model supports CUDA */ + model->has_cuda = 1 ; } /* loop through all the capacitor models */ diff --git a/src/spicelib/devices/cccs/cccsdefs.h b/src/spicelib/devices/cccs/cccsdefs.h index 37a6de0ce..28540d720 100644 --- a/src/spicelib/devices/cccs/cccsdefs.h +++ b/src/spicelib/devices/cccs/cccsdefs.h @@ -62,6 +62,12 @@ typedef struct sCCCSmodel { /* model structure for a source */ #define CCCSinstances(inst) ((CCCSinstance *)((inst)->gen.GENinstances)) #define CCCSmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + } CCCSmodel; /* device parameters */ diff --git a/src/spicelib/devices/cccs/cccsset.c b/src/spicelib/devices/cccs/cccsset.c index 0c1e7fa92..ffb6d5669 100644 --- a/src/spicelib/devices/cccs/cccsset.c +++ b/src/spicelib/devices/cccs/cccsset.c @@ -29,6 +29,11 @@ CCCSsetup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *states) /* loop through all the voltage source models */ for( ; model != NULL; model = CCCSnextModel(model)) { +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + /* loop through all the instances of the model */ for (here = CCCSinstances(model); here != NULL ; here=CCCSnextInstance(here)) { diff --git a/src/spicelib/devices/ccvs/ccvsdefs.h b/src/spicelib/devices/ccvs/ccvsdefs.h index 9b4acb36e..ac49f4e7f 100644 --- a/src/spicelib/devices/ccvs/ccvsdefs.h +++ b/src/spicelib/devices/ccvs/ccvsdefs.h @@ -69,6 +69,12 @@ typedef struct sCCVSmodel { /* model structure for a CCVsource */ #define CCVSinstances(inst) ((CCVSinstance *)((inst)->gen.GENinstances)) #define CCVSmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + } CCVSmodel; /* device parameters */ diff --git a/src/spicelib/devices/ccvs/ccvsset.c b/src/spicelib/devices/ccvs/ccvsset.c index d3b68dc29..9990ee486 100644 --- a/src/spicelib/devices/ccvs/ccvsset.c +++ b/src/spicelib/devices/ccvs/ccvsset.c @@ -27,6 +27,11 @@ CCVSsetup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *states) /* loop through all the voltage source models */ for( ; model != NULL; model = CCVSnextModel(model)) { +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + /* loop through all the instances of the model */ for (here = CCVSinstances(model); here != NULL ; here=CCVSnextInstance(here)) { diff --git a/src/spicelib/devices/cpl/cpldefs.h b/src/spicelib/devices/cpl/cpldefs.h index 313c6f113..35934d0de 100644 --- a/src/spicelib/devices/cpl/cpldefs.h +++ b/src/spicelib/devices/cpl/cpldefs.h @@ -91,6 +91,12 @@ typedef struct sCPLmodel { /* model structure for a cpl */ #define CPLinstances(inst) ((CPLinstance *)((inst)->gen.GENinstances)) #define CPLmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + double *Rm; int Rm_counter; double *Gm; diff --git a/src/spicelib/devices/cpl/cplsetup.c b/src/spicelib/devices/cpl/cplsetup.c index 6e0328590..7483db3d0 100644 --- a/src/spicelib/devices/cpl/cplsetup.c +++ b/src/spicelib/devices/cpl/cplsetup.c @@ -149,6 +149,11 @@ CPLsetup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *state) /* loop through all the models */ for( ; model != NULL; model = CPLnextModel(model)) { +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + if (!model->Rmgiven) { SPfrontEnd->IFerrorf (ERR_FATAL, "model %s: lossy line series resistance not given", model->CPLmodName); diff --git a/src/spicelib/devices/csw/cswdefs.h b/src/spicelib/devices/csw/cswdefs.h index 722b68ce7..1e211e2e8 100644 --- a/src/spicelib/devices/csw/cswdefs.h +++ b/src/spicelib/devices/csw/cswdefs.h @@ -75,6 +75,12 @@ typedef struct sCSWmodel { /* model structure for a switch */ #define CSWinstances(inst) ((CSWinstance *)((inst)->gen.GENinstances)) #define CSWmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + double CSWonResistance; /* switch "on" resistance */ double CSWoffResistance; /* switch "off" resistance */ double CSWiThreshold; /* switching threshold current */ diff --git a/src/spicelib/devices/csw/cswsetup.c b/src/spicelib/devices/csw/cswsetup.c index edb09ff8c..e0ea66c18 100644 --- a/src/spicelib/devices/csw/cswsetup.c +++ b/src/spicelib/devices/csw/cswsetup.c @@ -24,6 +24,12 @@ CSWsetup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *states) /* loop through all the current source models */ for( ; model != NULL; model = CSWnextModel(model)) { + +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + /* Default Value Processing for Switch Model */ if (!model->CSWthreshGiven) { model->CSWiThreshold = 0; diff --git a/src/spicelib/devices/dio/diodefs.h b/src/spicelib/devices/dio/diodefs.h index 97f7cf65f..d5b077619 100644 --- a/src/spicelib/devices/dio/diodefs.h +++ b/src/spicelib/devices/dio/diodefs.h @@ -185,6 +185,12 @@ typedef struct sDIOmodel { /* model structure for a diode */ #define DIOinstances(inst) ((DIOinstance *)((inst)->gen.GENinstances)) #define DIOmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + unsigned DIOlevelGiven : 1; unsigned DIOsatCurGiven : 1; unsigned DIOsatSWCurGiven : 1; diff --git a/src/spicelib/devices/dio/diosetup.c b/src/spicelib/devices/dio/diosetup.c index e06b13ba1..5017e756f 100644 --- a/src/spicelib/devices/dio/diosetup.c +++ b/src/spicelib/devices/dio/diosetup.c @@ -27,6 +27,11 @@ DIOsetup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *states) /* loop through all the diode models */ for( ; model != NULL; model = DIOnextModel(model)) { +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + if(!model->DIOlevelGiven) { model->DIOlevel = 1; } diff --git a/src/spicelib/devices/hfet1/hfetdefs.h b/src/spicelib/devices/hfet1/hfetdefs.h index 0a08d892b..e9d386b24 100644 --- a/src/spicelib/devices/hfet1/hfetdefs.h +++ b/src/spicelib/devices/hfet1/hfetdefs.h @@ -170,6 +170,12 @@ typedef struct sHFETAmodel { #define HFETAinstances(inst) ((HFETAinstance *)((inst)->gen.GENinstances)) #define HFETAmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int HFETAtype; int HFETAgatemod; diff --git a/src/spicelib/devices/hfet1/hfetsetup.c b/src/spicelib/devices/hfet1/hfetsetup.c index fa29ef325..7ec7991cc 100644 --- a/src/spicelib/devices/hfet1/hfetsetup.c +++ b/src/spicelib/devices/hfet1/hfetsetup.c @@ -30,6 +30,12 @@ HFETAsetup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *states) /* loop through all the diode models */ for( ; model != NULL; model = HFETAnextModel(model)) { + +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + if( (model->HFETAtype != NHFET) && (model->HFETAtype != PHFET) ) { model->HFETAtype = NHFET; } diff --git a/src/spicelib/devices/hfet2/hfet2defs.h b/src/spicelib/devices/hfet2/hfet2defs.h index a41e952e3..474dae3b2 100644 --- a/src/spicelib/devices/hfet2/hfet2defs.h +++ b/src/spicelib/devices/hfet2/hfet2defs.h @@ -115,6 +115,12 @@ typedef struct sHFET2model { #define HFET2instances(inst) ((HFET2instance *)((inst)->gen.GENinstances)) #define HFET2modName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int HFET2type; double HFET2cf; diff --git a/src/spicelib/devices/hfet2/hfet2setup.c b/src/spicelib/devices/hfet2/hfet2setup.c index cb51d3a7c..68d279447 100644 --- a/src/spicelib/devices/hfet2/hfet2setup.c +++ b/src/spicelib/devices/hfet2/hfet2setup.c @@ -21,6 +21,12 @@ int HFET2setup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *state CKTnode *tmp; for( ; model != NULL; model = HFET2nextModel(model)) { + +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + if((TYPE != NHFET) && (TYPE != PHFET) ) TYPE = NHFET; if(!model->HFET2cfGiven) diff --git a/src/spicelib/devices/hisim2/hsm2def.h b/src/spicelib/devices/hisim2/hsm2def.h index 026dda2d0..3cb70a3d8 100644 --- a/src/spicelib/devices/hisim2/hsm2def.h +++ b/src/spicelib/devices/hisim2/hsm2def.h @@ -739,6 +739,12 @@ typedef struct sHSM2model { /* model structure for a resistor */ #define HSM2instances(inst) ((HSM2instance *)((inst)->gen.GENinstances)) #define HSM2modName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int HSM2_type; /* device type: 1 = nmos, -1 = pmos */ int HSM2_level; /* level */ int HSM2_info; /* information */ diff --git a/src/spicelib/devices/hisim2/hsm2set.c b/src/spicelib/devices/hisim2/hsm2set.c index c66a7ec71..0ad5ee6e3 100644 --- a/src/spicelib/devices/hisim2/hsm2set.c +++ b/src/spicelib/devices/hisim2/hsm2set.c @@ -121,6 +121,12 @@ int HSM2setup( /* loop through all the HSM2 device models */ for ( ;model != NULL ;model = HSM2nextModel(model)) { + +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + /* Default value Processing for HSM2 MOSFET Models */ if ( !model->HSM2_type_Given ) model->HSM2_type = NMOS ; diff --git a/src/spicelib/devices/hisimhv1/hsmhvdef.h b/src/spicelib/devices/hisimhv1/hsmhvdef.h index 1e4a4bf0d..5420bb8f7 100644 --- a/src/spicelib/devices/hisimhv1/hsmhvdef.h +++ b/src/spicelib/devices/hisimhv1/hsmhvdef.h @@ -1012,6 +1012,12 @@ typedef struct sHSMHVmodel { /* model structure for a resistor */ #define HSMHVinstances(inst) ((HSMHVinstance *)((inst)->gen.GENinstances)) #define HSMHVmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int HSMHV_type; /* device type: 1 = nmos, -1 = pmos */ int HSMHV_level; /* level */ int HSMHV_info; /* information */ diff --git a/src/spicelib/devices/hisimhv1/hsmhvset.c b/src/spicelib/devices/hisimhv1/hsmhvset.c index 08ba1d1ac..8b6955127 100644 --- a/src/spicelib/devices/hisimhv1/hsmhvset.c +++ b/src/spicelib/devices/hisimhv1/hsmhvset.c @@ -65,6 +65,12 @@ int HSMHVsetup( /* loop through all the HSMHV device models */ for ( ;model != NULL ;model = HSMHVnextModel(model)) { + +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + /* Default value Processing for HVMOS Models */ if ( !model->HSMHV_type_Given ) model->HSMHV_type = NMOS ; diff --git a/src/spicelib/devices/hisimhv2/hsmhv2def.h b/src/spicelib/devices/hisimhv2/hsmhv2def.h index 4fac4f470..aa1760826 100644 --- a/src/spicelib/devices/hisimhv2/hsmhv2def.h +++ b/src/spicelib/devices/hisimhv2/hsmhv2def.h @@ -1137,6 +1137,12 @@ typedef struct sHSMHV2model { /* model structure for a resistor */ #define HSMHV2instances(inst) ((HSMHV2instance *)((inst)->gen.GENinstances)) #define HSMHV2modName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int HSMHV2_type; /* device type: 1 = nmos, -1 = pmos */ int HSMHV2_level; /* level */ int HSMHV2_info; /* information */ diff --git a/src/spicelib/devices/hisimhv2/hsmhv2set.c b/src/spicelib/devices/hisimhv2/hsmhv2set.c index 19d7d6f93..09303d0b2 100644 --- a/src/spicelib/devices/hisimhv2/hsmhv2set.c +++ b/src/spicelib/devices/hisimhv2/hsmhv2set.c @@ -125,6 +125,12 @@ int HSMHV2setup( /* loop through all the HSMHV2 device models */ for ( ;model != NULL ;model = HSMHV2nextModel(model)) { + +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + /* Default value Processing for HVMOS Models */ if ( !model->HSMHV2_type_Given ) model->HSMHV2_type = NMOS ; diff --git a/src/spicelib/devices/ind/inddefs.h b/src/spicelib/devices/ind/inddefs.h index a157dc3ea..ac685e76f 100644 --- a/src/spicelib/devices/ind/inddefs.h +++ b/src/spicelib/devices/ind/inddefs.h @@ -130,6 +130,12 @@ struct sINDmodel { /* model structure for an inductor */ #define INDinstances(inst) ((INDinstance *)((inst)->gen.GENinstances)) #define INDmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + double INDmInd; /* Model inductance */ double INDtnom; /* temperature at which inductance measured */ double INDtempCoeff1; /* first temperature coefficient */ @@ -251,6 +257,12 @@ struct sMUTmodel { /* model structure for a mutual inductor */ #define MUTinstances(inst) ((MUTinstance *)((inst)->gen.GENinstances)) #define MUTmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + +/* --- end of generic struct GENmodel --- */ + #ifdef USE_CUSPICE MUTparamCPUstruct MUTparamCPU; MUTparamGPUstruct MUTparamGPU; diff --git a/src/spicelib/devices/ind/indsetup.c b/src/spicelib/devices/ind/indsetup.c index 9f73b4ea7..232b7211f 100644 --- a/src/spicelib/devices/ind/indsetup.c +++ b/src/spicelib/devices/ind/indsetup.c @@ -123,6 +123,9 @@ do { if((here->ptr = SMPmakeElt(matrix, here->first, here->second)) == NULL){\ /* How much instances we have */ model->n_instances = i ; + + /* This model supports CUDA */ + model->has_cuda = 1 ; } /* loop through all the inductor models */ diff --git a/src/spicelib/devices/ind/mutsetup.c b/src/spicelib/devices/ind/mutsetup.c index 0abf21076..4f1e98ca4 100644 --- a/src/spicelib/devices/ind/mutsetup.c +++ b/src/spicelib/devices/ind/mutsetup.c @@ -79,6 +79,9 @@ MUTsetup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *states) /* How much instances we have */ model->n_instances = i; + + /* This model supports CUDA */ + model->has_cuda = 1 ; } /* loop through all the mutual inductor models */ diff --git a/src/spicelib/devices/jfet/jfetdefs.h b/src/spicelib/devices/jfet/jfetdefs.h index 701dc0375..42281324d 100644 --- a/src/spicelib/devices/jfet/jfetdefs.h +++ b/src/spicelib/devices/jfet/jfetdefs.h @@ -197,6 +197,12 @@ typedef struct sJFETmodel { /* model structure for a jfet */ #define JFETinstances(inst) ((JFETinstance *)((inst)->gen.GENinstances)) #define JFETmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int JFETtype; double JFETthreshold; diff --git a/src/spicelib/devices/jfet/jfetset.c b/src/spicelib/devices/jfet/jfetset.c index d2de3d65f..cb673c0cd 100644 --- a/src/spicelib/devices/jfet/jfetset.c +++ b/src/spicelib/devices/jfet/jfetset.c @@ -29,6 +29,11 @@ JFETsetup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *states) /* loop through all the diode models */ for( ; model != NULL; model = JFETnextModel(model)) { +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + if( (model->JFETtype != NJF) && (model->JFETtype != PJF) ) { model->JFETtype = NJF; } diff --git a/src/spicelib/devices/jfet2/jfet2defs.h b/src/spicelib/devices/jfet2/jfet2defs.h index f8b2a230b..8ed4193e9 100644 --- a/src/spicelib/devices/jfet2/jfet2defs.h +++ b/src/spicelib/devices/jfet2/jfet2defs.h @@ -209,6 +209,12 @@ typedef struct sJFET2model { /* model structure for a jfet */ #define JFET2instances(inst) ((JFET2instance *)((inst)->gen.GENinstances)) #define JFET2modName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int JFET2type; #define PARAM(code,id,flag,ref,default,descrip) double ref; diff --git a/src/spicelib/devices/jfet2/jfet2set.c b/src/spicelib/devices/jfet2/jfet2set.c index 344379320..c47d5c7a2 100644 --- a/src/spicelib/devices/jfet2/jfet2set.c +++ b/src/spicelib/devices/jfet2/jfet2set.c @@ -30,6 +30,11 @@ JFET2setup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *states) /* loop through all the diode models */ for( ; model != NULL; model = JFET2nextModel(model)) { +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + if( (model->JFET2type != NJF) && (model->JFET2type != PJF) ) { model->JFET2type = NJF; } diff --git a/src/spicelib/devices/ltra/ltradefs.h b/src/spicelib/devices/ltra/ltradefs.h index 08e6baeb7..915d91d27 100644 --- a/src/spicelib/devices/ltra/ltradefs.h +++ b/src/spicelib/devices/ltra/ltradefs.h @@ -108,6 +108,12 @@ typedef struct sLTRAmodel { /* model structure for a transmission lines */ #define LTRAinstances(inst) ((LTRAinstance *)((inst)->gen.GENinstances)) #define LTRAmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + double LTRAh1dashFirstVal; /* first needed value of h1dasg at current timepoint */ double LTRAh2FirstVal; /* first needed value of h2 at current diff --git a/src/spicelib/devices/ltra/ltraset.c b/src/spicelib/devices/ltra/ltraset.c index 947dde4fb..3bfd121e2 100644 --- a/src/spicelib/devices/ltra/ltraset.c +++ b/src/spicelib/devices/ltra/ltraset.c @@ -27,6 +27,11 @@ LTRAsetup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *state) /* loop through all the transmission line models */ for (; model != NULL; model = LTRAnextModel(model)) { +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + if (!model->LTRAnlGiven) { model->LTRAnl = .25; } diff --git a/src/spicelib/devices/mes/mesdefs.h b/src/spicelib/devices/mes/mesdefs.h index 24fa8b443..52a7e5ade 100644 --- a/src/spicelib/devices/mes/mesdefs.h +++ b/src/spicelib/devices/mes/mesdefs.h @@ -190,6 +190,12 @@ typedef struct sMESmodel { /* model structure for a mesfet */ #define MESinstances(inst) ((MESinstance *)((inst)->gen.GENinstances)) #define MESmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int MEStype; double MESthreshold; diff --git a/src/spicelib/devices/mes/messetup.c b/src/spicelib/devices/mes/messetup.c index 61f0dac9e..0d2ccc615 100644 --- a/src/spicelib/devices/mes/messetup.c +++ b/src/spicelib/devices/mes/messetup.c @@ -26,6 +26,11 @@ MESsetup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *states) /* loop through all the diode models */ for( ; model != NULL; model = MESnextModel(model)) { +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + if( (model->MEStype != NMF) && (model->MEStype != PMF) ) { model->MEStype = NMF; } diff --git a/src/spicelib/devices/mesa/mesadefs.h b/src/spicelib/devices/mesa/mesadefs.h index 50bdcece4..12cc4b62d 100644 --- a/src/spicelib/devices/mesa/mesadefs.h +++ b/src/spicelib/devices/mesa/mesadefs.h @@ -256,6 +256,12 @@ typedef struct sMESAmodel { /* model structure for a MESAfet */ #define MESAinstances(inst) ((MESAinstance *)((inst)->gen.GENinstances)) #define MESAmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int MESAtype; double MESAthreshold; diff --git a/src/spicelib/devices/mesa/mesasetup.c b/src/spicelib/devices/mesa/mesasetup.c index 95c7a05ec..775af4b19 100644 --- a/src/spicelib/devices/mesa/mesasetup.c +++ b/src/spicelib/devices/mesa/mesasetup.c @@ -28,6 +28,12 @@ MESAsetup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *states) /* loop through all the diode models */ for( ; model != NULL; model = MESAnextModel(model)) { + +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + if( (model->MESAtype != NMF) ) { fprintf(stderr, "Only nmf model type supported, set to nmf\n"); model->MESAtype = NMF; diff --git a/src/spicelib/devices/mos1/mos1defs.h b/src/spicelib/devices/mos1/mos1defs.h index 8117bf655..b8d185b93 100644 --- a/src/spicelib/devices/mos1/mos1defs.h +++ b/src/spicelib/devices/mos1/mos1defs.h @@ -347,6 +347,12 @@ typedef struct sMOS1model { /* model structure for a resistor */ #define MOS1instances(inst) ((MOS1instance *)((inst)->gen.GENinstances)) #define MOS1modName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int MOS1type; /* device type : 1 = nmos, -1 = pmos */ double MOS1tnom; /* temperature at which parameters measured */ double MOS1latDiff; diff --git a/src/spicelib/devices/mos1/mos1set.c b/src/spicelib/devices/mos1/mos1set.c index a4483de3a..32727ed73 100644 --- a/src/spicelib/devices/mos1/mos1set.c +++ b/src/spicelib/devices/mos1/mos1set.c @@ -27,6 +27,11 @@ MOS1setup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, /* loop through all the MOS1 device models */ for( ; model != NULL; model = MOS1nextModel(model)) { +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + if(!model->MOS1typeGiven) { model->MOS1type = NMOS; } diff --git a/src/spicelib/devices/mos2/mos2defs.h b/src/spicelib/devices/mos2/mos2defs.h index 407ac3026..2483db120 100644 --- a/src/spicelib/devices/mos2/mos2defs.h +++ b/src/spicelib/devices/mos2/mos2defs.h @@ -354,6 +354,12 @@ typedef struct sMOS2model { /* model structure for a resistor */ #define MOS2instances(inst) ((MOS2instance *)((inst)->gen.GENinstances)) #define MOS2modName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int MOS2type; /* device type : 1 = nmos, -1 = pmos */ int MOS2gateType; diff --git a/src/spicelib/devices/mos2/mos2set.c b/src/spicelib/devices/mos2/mos2set.c index b12eb1a98..c138765f2 100644 --- a/src/spicelib/devices/mos2/mos2set.c +++ b/src/spicelib/devices/mos2/mos2set.c @@ -26,6 +26,11 @@ MOS2setup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *states) /* loop through all the MOS2 device models */ for( ; model != NULL; model = MOS2nextModel(model)) { +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + if(!model->MOS2typeGiven) { model->MOS2type = NMOS; } diff --git a/src/spicelib/devices/mos3/mos3defs.h b/src/spicelib/devices/mos3/mos3defs.h index 268a9b158..15bc5a45e 100644 --- a/src/spicelib/devices/mos3/mos3defs.h +++ b/src/spicelib/devices/mos3/mos3defs.h @@ -352,6 +352,12 @@ typedef struct sMOS3model { /* model structure for a resistor */ #define MOS3instances(inst) ((MOS3instance *)((inst)->gen.GENinstances)) #define MOS3modName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int MOS3type; /* device type : 1 = nmos, -1 = pmos */ double MOS3tnom; /* temperature at which parameters measured */ double MOS3latDiff; diff --git a/src/spicelib/devices/mos3/mos3set.c b/src/spicelib/devices/mos3/mos3set.c index 4432ba461..0b0093d6c 100644 --- a/src/spicelib/devices/mos3/mos3set.c +++ b/src/spicelib/devices/mos3/mos3set.c @@ -30,6 +30,11 @@ MOS3setup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *states) /* loop through all the MOS3 device models */ for( ; model != NULL; model = MOS3nextModel(model)) { +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + /* perform model defaulting */ if(!model->MOS3typeGiven) { model->MOS3type = NMOS; diff --git a/src/spicelib/devices/mos6/mos6defs.h b/src/spicelib/devices/mos6/mos6defs.h index 580a3b2af..116f8e718 100644 --- a/src/spicelib/devices/mos6/mos6defs.h +++ b/src/spicelib/devices/mos6/mos6defs.h @@ -279,6 +279,12 @@ typedef struct sMOS6model { /* model structure for a resistor */ #define MOS6instances(inst) ((MOS6instance *)((inst)->gen.GENinstances)) #define MOS6modName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int MOS6type; /* device type : 1 = nmos, -1 = pmos */ double MOS6tnom; /* temperature at which parameters measured */ double MOS6latDiff; diff --git a/src/spicelib/devices/mos6/mos6set.c b/src/spicelib/devices/mos6/mos6set.c index a46840630..68758730a 100644 --- a/src/spicelib/devices/mos6/mos6set.c +++ b/src/spicelib/devices/mos6/mos6set.c @@ -27,6 +27,11 @@ MOS6setup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, /* loop through all the MOS6 device models */ for( ; model != NULL; model = MOS6nextModel(model)) { +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + if(!model->MOS6typeGiven) { model->MOS6type = NMOS; } diff --git a/src/spicelib/devices/mos9/mos9defs.h b/src/spicelib/devices/mos9/mos9defs.h index 6f673df91..3cce08f9d 100644 --- a/src/spicelib/devices/mos9/mos9defs.h +++ b/src/spicelib/devices/mos9/mos9defs.h @@ -354,6 +354,12 @@ typedef struct sMOS9model { /* model structure for a resistor */ #define MOS9instances(inst) ((MOS9instance *)((inst)->gen.GENinstances)) #define MOS9modName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int MOS9type; /* device type : 1 = nmos, -1 = pmos */ double MOS9tnom; /* temperature at which parameters measured */ double MOS9latDiff; diff --git a/src/spicelib/devices/mos9/mos9set.c b/src/spicelib/devices/mos9/mos9set.c index 1dddbd7e7..d0d354dc1 100644 --- a/src/spicelib/devices/mos9/mos9set.c +++ b/src/spicelib/devices/mos9/mos9set.c @@ -30,6 +30,11 @@ MOS9setup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *states) /* loop through all the MOS9 device models */ for( ; model != NULL; model = MOS9nextModel(model)) { +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + /* perform model defaulting */ if(!model->MOS9typeGiven) { model->MOS9type = NMOS; diff --git a/src/spicelib/devices/res/resdefs.h b/src/spicelib/devices/res/resdefs.h index 327031cbc..33a83e480 100644 --- a/src/spicelib/devices/res/resdefs.h +++ b/src/spicelib/devices/res/resdefs.h @@ -148,6 +148,12 @@ typedef struct sRESmodel { /* model structure for a resistor */ #define RESinstances(inst) ((RESinstance *)((inst)->gen.GENinstances)) #define RESmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + double REStnom; /* temperature at which resistance measured */ double REStempCoeff1; /* first temperature coefficient of resistors */ double REStempCoeff2; /* second temperature coefficient of resistors */ diff --git a/src/spicelib/devices/res/ressetup.c b/src/spicelib/devices/res/ressetup.c index 6dfae725c..337b3a1a8 100644 --- a/src/spicelib/devices/res/ressetup.c +++ b/src/spicelib/devices/res/ressetup.c @@ -94,6 +94,9 @@ do { if((here->ptr = SMPmakeElt(matrix, here->first, here->second)) == NULL){\ /* How much instances we have */ model->n_instances = i ; + + /* This model supports CUDA */ + model->has_cuda = 1 ; } /* loop through all the resistor models */ diff --git a/src/spicelib/devices/soi3/soi3defs.h b/src/spicelib/devices/soi3/soi3defs.h index 465597cf8..5c3dc7426 100644 --- a/src/spicelib/devices/soi3/soi3defs.h +++ b/src/spicelib/devices/soi3/soi3defs.h @@ -482,6 +482,12 @@ typedef struct sSOI3model { /* model structure for an SOI3 MOSFET */ #define SOI3instances(inst) ((SOI3instance *)((inst)->gen.GENinstances)) #define SOI3modName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int SOI3type; /* device type : 1 = nsoi, -1 = psoi */ double SOI3tnom; /* temperature at which parameters measured */ double SOI3latDiff; diff --git a/src/spicelib/devices/soi3/soi3set.c b/src/spicelib/devices/soi3/soi3set.c index 11eade08a..9d45f78af 100644 --- a/src/spicelib/devices/soi3/soi3set.c +++ b/src/spicelib/devices/soi3/soi3set.c @@ -55,6 +55,11 @@ SOI3setup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *states) /* loop through all the SOI3 device models */ for( ; model != NULL; model = SOI3nextModel(model)) { +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + if(!model->SOI3typeGiven) { model->SOI3type = NSOI3; } diff --git a/src/spicelib/devices/sw/swdefs.h b/src/spicelib/devices/sw/swdefs.h index b22f12e6d..f67f822b9 100644 --- a/src/spicelib/devices/sw/swdefs.h +++ b/src/spicelib/devices/sw/swdefs.h @@ -74,6 +74,12 @@ typedef struct sSWmodel { /* model structure for a switch */ #define SWinstances(inst) ((SWinstance *)((inst)->gen.GENinstances)) #define SWmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + double SWonResistance; /* switch "on" resistance */ double SWoffResistance; /* switch "off" resistance */ double SWvThreshold; /* switching threshold voltage */ diff --git a/src/spicelib/devices/sw/swsetup.c b/src/spicelib/devices/sw/swsetup.c index 11eea0421..64ada8e09 100644 --- a/src/spicelib/devices/sw/swsetup.c +++ b/src/spicelib/devices/sw/swsetup.c @@ -24,6 +24,12 @@ SWsetup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *states) /* loop through all the current source models */ for( ; model != NULL; model = SWnextModel(model)) { + +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + /* Default Value Processing for Switch Model */ if (!model->SWthreshGiven) { model->SWvThreshold = 0; diff --git a/src/spicelib/devices/tra/tradefs.h b/src/spicelib/devices/tra/tradefs.h index 7e1df5453..76f096040 100644 --- a/src/spicelib/devices/tra/tradefs.h +++ b/src/spicelib/devices/tra/tradefs.h @@ -123,6 +123,12 @@ typedef struct sTRAmodel { /* model structure for a transmission lines */ #define TRAinstances(inst) ((TRAinstance *)((inst)->gen.GENinstances)) #define TRAmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + } TRAmodel; /* device parameters */ diff --git a/src/spicelib/devices/tra/trasetup.c b/src/spicelib/devices/tra/trasetup.c index 8ad68e8ca..afccfff13 100644 --- a/src/spicelib/devices/tra/trasetup.c +++ b/src/spicelib/devices/tra/trasetup.c @@ -30,6 +30,11 @@ TRAsetup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *state) /* loop through all the transmission line models */ for( ; model != NULL; model = TRAnextModel(model)) { +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + /* loop through all the instances of the model */ for (here = TRAinstances(model); here != NULL ; here=TRAnextInstance(here)) { diff --git a/src/spicelib/devices/txl/txldefs.h b/src/spicelib/devices/txl/txldefs.h index 290074c48..b0821d037 100644 --- a/src/spicelib/devices/txl/txldefs.h +++ b/src/spicelib/devices/txl/txldefs.h @@ -81,6 +81,12 @@ typedef struct sTXLmodel { /* model structure for a txl */ #define TXLinstances(inst) ((TXLinstance *)((inst)->gen.GENinstances)) #define TXLmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + double R; double L; double G; diff --git a/src/spicelib/devices/txl/txlsetup.c b/src/spicelib/devices/txl/txlsetup.c index bd2a5ef2a..62b53e024 100644 --- a/src/spicelib/devices/txl/txlsetup.c +++ b/src/spicelib/devices/txl/txlsetup.c @@ -87,6 +87,11 @@ TXLsetup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit*ckt, int *state) /* loop through all the models */ for( ; model != NULL; model = TXLnextModel(model)) { +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + if (!model->Rgiven) { SPfrontEnd->IFerrorf (ERR_FATAL, "model %s: lossy line series resistance not given", model->TXLmodName); diff --git a/src/spicelib/devices/urc/urcdefs.h b/src/spicelib/devices/urc/urcdefs.h index 7f96c1716..48c1452cf 100644 --- a/src/spicelib/devices/urc/urcdefs.h +++ b/src/spicelib/devices/urc/urcdefs.h @@ -48,6 +48,12 @@ typedef struct sURCmodel { /* model structure for a resistor */ #define URCinstances(inst) ((URCinstance *)((inst)->gen.GENinstances)) #define URCmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + double URCk; /* propagation constant for URC */ double URCfmax; /* max frequence of interest */ double URCrPerL; /* resistance per unit length */ diff --git a/src/spicelib/devices/urc/urcsetup.c b/src/spicelib/devices/urc/urcsetup.c index 0a7adbc30..14341cc9b 100644 --- a/src/spicelib/devices/urc/urcsetup.c +++ b/src/spicelib/devices/urc/urcsetup.c @@ -62,6 +62,12 @@ URCsetup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *state) dtype = CKTtypelook("Diode"); /* loop through all the URC models */ for( ; model != NULL; model = URCnextModel(model)) { + +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + if(!model->URCkGiven) model->URCk = 1.5; if(!model->URCfmaxGiven) diff --git a/src/spicelib/devices/vbic/vbicdefs.h b/src/spicelib/devices/vbic/vbicdefs.h index 1dd356607..adf0c683b 100644 --- a/src/spicelib/devices/vbic/vbicdefs.h +++ b/src/spicelib/devices/vbic/vbicdefs.h @@ -414,6 +414,12 @@ typedef struct sVBICmodel { /* model structure for a vbic */ #define VBICinstances(inst) ((VBICinstance *)((inst)->gen.GENinstances)) #define VBICmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + int VBICtype; double VBICtnom; diff --git a/src/spicelib/devices/vbic/vbicsetup.c b/src/spicelib/devices/vbic/vbicsetup.c index 47907538b..afd4c8a17 100644 --- a/src/spicelib/devices/vbic/vbicsetup.c +++ b/src/spicelib/devices/vbic/vbicsetup.c @@ -35,6 +35,11 @@ VBICsetup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *states) /* loop through all the transistor models */ for( ; model != NULL; model = VBICnextModel(model)) { +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + if(model->VBICtype != NPN && model->VBICtype != PNP) { model->VBICtype = NPN; } diff --git a/src/spicelib/devices/vccs/vccsdefs.h b/src/spicelib/devices/vccs/vccsdefs.h index a97ef7c59..80fd5adc9 100644 --- a/src/spicelib/devices/vccs/vccsdefs.h +++ b/src/spicelib/devices/vccs/vccsdefs.h @@ -72,6 +72,12 @@ typedef struct sVCCSmodel { /* model structure for a source */ #define VCCSinstances(inst) ((VCCSinstance *)((inst)->gen.GENinstances)) #define VCCSmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + } VCCSmodel; /* device parameters */ diff --git a/src/spicelib/devices/vccs/vccsset.c b/src/spicelib/devices/vccs/vccsset.c index e956708ca..05b924c63 100644 --- a/src/spicelib/devices/vccs/vccsset.c +++ b/src/spicelib/devices/vccs/vccsset.c @@ -30,6 +30,11 @@ VCCSsetup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *states) /* loop through all the current source models */ for( ; model != NULL; model = VCCSnextModel(model)) { +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + /* loop through all the instances of the model */ for (here = VCCSinstances(model); here != NULL ; here=VCCSnextInstance(here)) { diff --git a/src/spicelib/devices/vcvs/vcvsdefs.h b/src/spicelib/devices/vcvs/vcvsdefs.h index a89fc79ca..110008bf8 100644 --- a/src/spicelib/devices/vcvs/vcvsdefs.h +++ b/src/spicelib/devices/vcvs/vcvsdefs.h @@ -78,6 +78,12 @@ typedef struct sVCVSmodel { /* model structure for a source */ #define VCVSinstances(inst) ((VCVSinstance *)((inst)->gen.GENinstances)) #define VCVSmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + } VCVSmodel; /* device parameters */ diff --git a/src/spicelib/devices/vcvs/vcvsset.c b/src/spicelib/devices/vcvs/vcvsset.c index d3196b4a4..88945f646 100644 --- a/src/spicelib/devices/vcvs/vcvsset.c +++ b/src/spicelib/devices/vcvs/vcvsset.c @@ -28,6 +28,11 @@ VCVSsetup(SMPmatrix *matrix, GENmodel *inModel, CKTcircuit *ckt, int *states) /* loop through all the voltage source models */ for( ; model != NULL; model = VCVSnextModel(model)) { +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + /* loop through all the instances of the model */ for (here = VCVSinstances(model); here != NULL ; here=VCVSnextInstance(here)) { diff --git a/src/spicelib/devices/vsrc/vsrcdefs.h b/src/spicelib/devices/vsrc/vsrcdefs.h index 773526296..d1195736a 100644 --- a/src/spicelib/devices/vsrc/vsrcdefs.h +++ b/src/spicelib/devices/vsrc/vsrcdefs.h @@ -142,6 +142,12 @@ typedef struct sVSRCmodel { #define VSRCinstances(inst) ((VSRCinstance *)((inst)->gen.GENinstances)) #define VSRCmodName gen.GENmodName +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + + /* --- end of generic struct GENmodel --- */ + #ifdef USE_CUSPICE VSRCparamCPUstruct VSRCparamCPU ; VSRCparamGPUstruct VSRCparamGPU ; diff --git a/src/spicelib/devices/vsrc/vsrcset.c b/src/spicelib/devices/vsrc/vsrcset.c index b595230a7..57e910b28 100644 --- a/src/spicelib/devices/vsrc/vsrcset.c +++ b/src/spicelib/devices/vsrc/vsrcset.c @@ -80,6 +80,9 @@ do { if((here->ptr = SMPmakeElt(matrix, here->first, here->second)) == NULL){\ /* How much instances we have */ model->n_instances = i ; + + /* This model supports CUDA */ + model->has_cuda = 1 ; } /* loop through all the voltage source models */