diff --git a/src/include/ngspice/gendefs.h b/src/include/ngspice/gendefs.h index dae4e1fe6..2e5c26003 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 }; #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 6ee35ee34..9b73a9798 100644 --- a/src/spicelib/devices/asrc/asrcdefs.h +++ b/src/spicelib/devices/asrc/asrcdefs.h @@ -67,6 +67,10 @@ typedef struct sASRCmodel { /* model structure for a source */ * that have this model */ IFuid ASRCmodName; /* pointer to character string naming this model */ +#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 a5c63cdf2..1132d550a 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 = model->ASRCnextModel) { + +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + for (here = model->ASRCinstances; here; here=here->ASRCnextInstance) { if (!here->ASRCtree) diff --git a/src/spicelib/devices/bjt/bjtdefs.h b/src/spicelib/devices/bjt/bjtdefs.h index ceb3225cd..28b614a51 100644 --- a/src/spicelib/devices/bjt/bjtdefs.h +++ b/src/spicelib/devices/bjt/bjtdefs.h @@ -357,6 +357,10 @@ typedef struct sBJTmodel { /* model structure for a bjt */ * that have this model */ IFuid BJTmodName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int BJTtype; diff --git a/src/spicelib/devices/bjt/bjtsetup.c b/src/spicelib/devices/bjt/bjtsetup.c index 239914838..d1bdda425 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 = model->BJTnextModel ) { +#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 58d1f2e41..989c86c88 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 = model->B1nextModel ) { - -/* 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 7da460bb5..872bbbd41 100644 --- a/src/spicelib/devices/bsim1/bsim1def.h +++ b/src/spicelib/devices/bsim1/bsim1def.h @@ -346,6 +346,10 @@ typedef struct sBSIM1model { /* model structure for a resistor */ * that have this model */ IFuid B1modName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int B1type; /* device type : 1 = nmos, -1 = pmos */ diff --git a/src/spicelib/devices/bsim2/b2set.c b/src/spicelib/devices/bsim2/b2set.c index 9f3895234..3070deca5 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 = model->B2nextModel ) { - -/* 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 cec4c35b8..38f3b1d0a 100644 --- a/src/spicelib/devices/bsim2/bsim2def.h +++ b/src/spicelib/devices/bsim2/bsim2def.h @@ -269,6 +269,10 @@ typedef struct sBSIM2model { /* model structure for a resistor */ * that have this model */ IFuid B2modName; /* pointer to the name of this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int B2type; /* device type: 1 = nmos, -1 = pmos */ diff --git a/src/spicelib/devices/bsim3/b3set.c b/src/spicelib/devices/bsim3/b3set.c index 2b104bb59..dd18f0453 100644 --- a/src/spicelib/devices/bsim3/b3set.c +++ b/src/spicelib/devices/bsim3/b3set.c @@ -50,7 +50,13 @@ BSIM3instance **InstArray; /* loop through all the BSIM3 device models */ for( ; model != NULL; model = model->BSIM3nextModel ) { -/* 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 520b76892..3e9e3891d 100644 --- a/src/spicelib/devices/bsim3/bsim3def.h +++ b/src/spicelib/devices/bsim3/bsim3def.h @@ -438,6 +438,10 @@ typedef struct sBSIM3model BSIM3instance *BSIM3instances; IFuid BSIM3modName; +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int BSIM3type; diff --git a/src/spicelib/devices/bsim3soi_dd/b3soidddef.h b/src/spicelib/devices/bsim3soi_dd/b3soidddef.h index 71b7117f1..a2b975cdb 100644 --- a/src/spicelib/devices/bsim3soi_dd/b3soidddef.h +++ b/src/spicelib/devices/bsim3soi_dd/b3soidddef.h @@ -698,6 +698,10 @@ typedef struct sB3SOIDDmodel B3SOIDDinstance *B3SOIDDinstances; IFuid B3SOIDDmodName; +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int B3SOIDDtype; diff --git a/src/spicelib/devices/bsim3soi_dd/b3soiddset.c b/src/spicelib/devices/bsim3soi_dd/b3soiddset.c index 069117c63..2cfc51008 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 = model->B3SOIDDnextModel ) { -/* 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 b6aae9042..0591e8699 100644 --- a/src/spicelib/devices/bsim3soi_fd/b3soifddef.h +++ b/src/spicelib/devices/bsim3soi_fd/b3soifddef.h @@ -690,6 +690,10 @@ typedef struct sB3SOIFDmodel B3SOIFDinstance *B3SOIFDinstances; IFuid B3SOIFDmodName; +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int B3SOIFDtype; diff --git a/src/spicelib/devices/bsim3soi_fd/b3soifdset.c b/src/spicelib/devices/bsim3soi_fd/b3soifdset.c index feaa79db6..2a90266ad 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 = model->B3SOIFDnextModel ) { -/* 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 70ac112ee..81a8f483e 100644 --- a/src/spicelib/devices/bsim3soi_pd/b3soipddef.h +++ b/src/spicelib/devices/bsim3soi_pd/b3soipddef.h @@ -659,6 +659,10 @@ typedef struct sB3SOIPDmodel B3SOIPDinstance *B3SOIPDinstances; IFuid B3SOIPDmodName; +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int B3SOIPDtype; diff --git a/src/spicelib/devices/bsim3soi_pd/b3soipdset.c b/src/spicelib/devices/bsim3soi_pd/b3soipdset.c index 674d2de97..dc66cccb1 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 = model->B3SOIPDnextModel ) { -/* 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 b0574396d..dbfb975ee 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 = model->BSIM3v0nextModel ) { -/* 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 c14456d69..bf0a08a3b 100644 --- a/src/spicelib/devices/bsim3v0/bsim3v0def.h +++ b/src/spicelib/devices/bsim3v0/bsim3v0def.h @@ -357,6 +357,10 @@ typedef struct sBSIM3v0model BSIM3v0instance *BSIM3v0instances; IFuid BSIM3v0modName; +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int BSIM3v0type; diff --git a/src/spicelib/devices/bsim3v1/b3v1set.c b/src/spicelib/devices/bsim3v1/b3v1set.c index a1818084e..65429b6ce 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 = model->BSIM3v1nextModel ) { -/* 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 6356068f2..82527a807 100644 --- a/src/spicelib/devices/bsim3v1/bsim3v1def.h +++ b/src/spicelib/devices/bsim3v1/bsim3v1def.h @@ -359,6 +359,10 @@ typedef struct sBSIM3v1model BSIM3v1instance *BSIM3v1instances; IFuid BSIM3v1modName; +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int BSIM3v1type; diff --git a/src/spicelib/devices/bsim3v32/b3v32set.c b/src/spicelib/devices/bsim3v32/b3v32set.c index 76bcda994..6008c0168 100644 --- a/src/spicelib/devices/bsim3v32/b3v32set.c +++ b/src/spicelib/devices/bsim3v32/b3v32set.c @@ -46,7 +46,13 @@ BSIM3v32instance **InstArray; /* loop through all the BSIM3v32 device models */ for( ; model != NULL; model = model->BSIM3v32nextModel ) { -/* 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 da4cf27ca..afc2ea67b 100644 --- a/src/spicelib/devices/bsim3v32/bsim3v32def.h +++ b/src/spicelib/devices/bsim3v32/bsim3v32def.h @@ -435,6 +435,10 @@ typedef struct sBSIM3v32model BSIM3v32instance *BSIM3v32instances; IFuid BSIM3v32modName; +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int BSIM3v32type; diff --git a/src/spicelib/devices/bsim4/b4set.c b/src/spicelib/devices/bsim4/b4set.c index 7139ef8fc..b1b8565c5 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 = model->BSIM4nextModel ) - { /* 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 1868c768b..84b2b815f 100644 --- a/src/spicelib/devices/bsim4/bsim4def.h +++ b/src/spicelib/devices/bsim4/bsim4def.h @@ -935,6 +935,10 @@ typedef struct sBSIM4model BSIM4instance *BSIM4instances; IFuid BSIM4modName; +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int BSIM4type; diff --git a/src/spicelib/devices/bsim4v5/b4v5set.c b/src/spicelib/devices/bsim4v5/b4v5set.c index 6da70c71f..aec7fbb92 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 = model->BSIM4v5nextModel ) - { /* 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 f71d72afc..14edc5df5 100644 --- a/src/spicelib/devices/bsim4v5/bsim4v5def.h +++ b/src/spicelib/devices/bsim4v5/bsim4v5def.h @@ -852,6 +852,10 @@ typedef struct sBSIM4v5model BSIM4v5instance *BSIM4v5instances; IFuid BSIM4v5modName; +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int BSIM4v5type; diff --git a/src/spicelib/devices/bsim4v6/b4v6set.c b/src/spicelib/devices/bsim4v6/b4v6set.c index 3d71a4c8a..82e7a0bcc 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 = model->BSIM4v6nextModel ) - { /* 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 0f6a3c1f1..b436a2808 100644 --- a/src/spicelib/devices/bsim4v6/bsim4v6def.h +++ b/src/spicelib/devices/bsim4v6/bsim4v6def.h @@ -873,6 +873,10 @@ typedef struct sBSIM4v6model BSIM4v6instance *BSIM4v6instances; IFuid BSIM4v6modName; +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int BSIM4v6type; diff --git a/src/spicelib/devices/bsim4v7/b4v7set.c b/src/spicelib/devices/bsim4v7/b4v7set.c index 95ee31ca9..12adca309 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 adba486cf..b46dee399 100644 --- a/src/spicelib/devices/bsim4v7/bsim4v7def.h +++ b/src/spicelib/devices/bsim4v7/bsim4v7def.h @@ -1428,6 +1428,10 @@ typedef struct sBSIM4v7model BSIM4v7instance *BSIM4v7instances; IFuid BSIM4v7modName; +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int BSIM4v7type; diff --git a/src/spicelib/devices/bsimsoi/b4soidef.h b/src/spicelib/devices/bsimsoi/b4soidef.h index eb7ee506d..8902c6af1 100644 --- a/src/spicelib/devices/bsimsoi/b4soidef.h +++ b/src/spicelib/devices/bsimsoi/b4soidef.h @@ -1189,6 +1189,10 @@ typedef struct sB4SOImodel B4SOIinstance *B4SOIinstances; IFuid B4SOImodName; +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int B4SOItype; diff --git a/src/spicelib/devices/bsimsoi/b4soiset.c b/src/spicelib/devices/bsimsoi/b4soiset.c index 78c36a02c..1db437c67 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 = model->B4SOInextModel ) { -/* 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 a3dfa117d..9a4715c76 100644 --- a/src/spicelib/devices/cap/capdefs.h +++ b/src/spicelib/devices/cap/capdefs.h @@ -116,6 +116,10 @@ typedef struct sCAPmodel { /* model structure for a capacitor */ * model */ IFuid CAPmodName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ double CAPtnom; /* temperature at which capacitance measured */ diff --git a/src/spicelib/devices/cap/capsetup.c b/src/spicelib/devices/cap/capsetup.c index fdef6e7ec..39d2fbdd9 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 e1ebb3b12..9f2624358 100644 --- a/src/spicelib/devices/cccs/cccsdefs.h +++ b/src/spicelib/devices/cccs/cccsdefs.h @@ -59,6 +59,10 @@ typedef struct sCCCSmodel { /* model structure for a source */ that have this model */ IFuid CCCSmodName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ } CCCSmodel; diff --git a/src/spicelib/devices/cccs/cccsset.c b/src/spicelib/devices/cccs/cccsset.c index be3dab93f..e67373a69 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 = model->CCCSnextModel ) { +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + /* loop through all the instances of the model */ for (here = model->CCCSinstances; here != NULL ; here=here->CCCSnextInstance) { diff --git a/src/spicelib/devices/ccvs/ccvsdefs.h b/src/spicelib/devices/ccvs/ccvsdefs.h index d6856be4f..93456f998 100644 --- a/src/spicelib/devices/ccvs/ccvsdefs.h +++ b/src/spicelib/devices/ccvs/ccvsdefs.h @@ -66,6 +66,10 @@ typedef struct sCCVSmodel { /* model structure for a CCVsource */ that have this model */ IFuid CCVSmodName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ } CCVSmodel; diff --git a/src/spicelib/devices/ccvs/ccvsset.c b/src/spicelib/devices/ccvs/ccvsset.c index 5521312b3..4024a8f36 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 = model->CCVSnextModel ) { +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + /* loop through all the instances of the model */ for (here = model->CCVSinstances; here != NULL ; here=here->CCVSnextInstance) { diff --git a/src/spicelib/devices/cpl/cpldefs.h b/src/spicelib/devices/cpl/cpldefs.h index b2bcc39f1..766e0221b 100644 --- a/src/spicelib/devices/cpl/cpldefs.h +++ b/src/spicelib/devices/cpl/cpldefs.h @@ -89,6 +89,10 @@ typedef struct sCPLmodel { /* model structure for a cpl */ * model */ IFuid CPLmodName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ double *Rm; diff --git a/src/spicelib/devices/cpl/cplsetup.c b/src/spicelib/devices/cpl/cplsetup.c index c913bff06..77a9362b9 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 = model->CPLnextModel ) { +#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 9d77356ee..e70178566 100644 --- a/src/spicelib/devices/csw/cswdefs.h +++ b/src/spicelib/devices/csw/cswdefs.h @@ -72,6 +72,10 @@ typedef struct sCSWmodel { /* model structure for a switch */ * model */ IFuid CSWmodName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ double CSWonResistance; /* switch "on" resistance */ diff --git a/src/spicelib/devices/csw/cswsetup.c b/src/spicelib/devices/csw/cswsetup.c index d334276c1..8efd24a56 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 = model->CSWnextModel ) { + +#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 7d543c67b..4ea6b97a3 100644 --- a/src/spicelib/devices/dio/diodefs.h +++ b/src/spicelib/devices/dio/diodefs.h @@ -181,6 +181,10 @@ typedef struct sDIOmodel { /* model structure for a diode */ * that have this model */ IFuid DIOmodName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ unsigned DIOlevelGiven : 1; diff --git a/src/spicelib/devices/dio/diosetup.c b/src/spicelib/devices/dio/diosetup.c index 3ef214e74..a5547cf37 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 = model->DIOnextModel ) { +#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 fccad2c88..8144a6d6b 100644 --- a/src/spicelib/devices/hfet1/hfetdefs.h +++ b/src/spicelib/devices/hfet1/hfetdefs.h @@ -164,6 +164,10 @@ typedef struct sHFETAmodel { HFETAinstance *HFETAinstances; IFuid HFETAmodName; +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int HFETAtype; diff --git a/src/spicelib/devices/hfet1/hfetsetup.c b/src/spicelib/devices/hfet1/hfetsetup.c index 845f67296..e12c6f101 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 = model->HFETAnextModel ) { + +#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 819881ebf..609ad33c2 100644 --- a/src/spicelib/devices/hfet2/hfet2defs.h +++ b/src/spicelib/devices/hfet2/hfet2defs.h @@ -108,6 +108,10 @@ typedef struct sHFET2model { struct sHFET2model *HFET2nextModel; HFET2instance * HFET2instances; +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ IFuid HFET2modName; diff --git a/src/spicelib/devices/hfet2/hfet2setup.c b/src/spicelib/devices/hfet2/hfet2setup.c index dc267a3f1..be044309d 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 = model->HFET2nextModel ) { + +#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 a5f752017..76d760161 100644 --- a/src/spicelib/devices/hisim2/hsm2def.h +++ b/src/spicelib/devices/hisim2/hsm2def.h @@ -736,6 +736,10 @@ typedef struct sHSM2model { /* model structure for a resistor */ that have this model */ IFuid HSM2modName; /* pointer to the name of this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int HSM2_type; /* device type: 1 = nmos, -1 = pmos */ diff --git a/src/spicelib/devices/hisim2/hsm2set.c b/src/spicelib/devices/hisim2/hsm2set.c index d9f346241..45a78d02d 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 = model->HSM2nextModel ) { + +#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 f0072b3eb..691e56f3b 100644 --- a/src/spicelib/devices/hisimhv1/hsmhvdef.h +++ b/src/spicelib/devices/hisimhv1/hsmhvdef.h @@ -1009,6 +1009,10 @@ typedef struct sHSMHVmodel { /* model structure for a resistor */ that have this model */ IFuid HSMHVmodName; /* pointer to the name of this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int HSMHV_type; /* device type: 1 = nmos, -1 = pmos */ diff --git a/src/spicelib/devices/hisimhv1/hsmhvset.c b/src/spicelib/devices/hisimhv1/hsmhvset.c index 0dc67ae9b..b35d71480 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 = model->HSMHVnextModel ) { + +#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 5dcba20bf..f797715aa 100644 --- a/src/spicelib/devices/hisimhv2/hsmhv2def.h +++ b/src/spicelib/devices/hisimhv2/hsmhv2def.h @@ -1134,6 +1134,10 @@ typedef struct sHSMHV2model { /* model structure for a resistor */ that have this model */ IFuid HSMHV2modName; /* pointer to the name of this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int HSMHV2_type; /* device type: 1 = nmos, -1 = pmos */ diff --git a/src/spicelib/devices/hisimhv2/hsmhv2set.c b/src/spicelib/devices/hisimhv2/hsmhv2set.c index a434f0b5e..9b03cdfcc 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 = model->HSMHV2nextModel ) { + +#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 62af208f9..9108d4374 100644 --- a/src/spicelib/devices/ind/inddefs.h +++ b/src/spicelib/devices/ind/inddefs.h @@ -126,6 +126,10 @@ struct sINDmodel { /* model structure for an inductor */ * model */ IFuid INDmodName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ double INDmInd; /* Model inductance */ @@ -244,7 +248,12 @@ struct sMUTmodel { /* model structure for a mutual inductor */ * model */ IFuid MUTmodName; /* pointer to character string naming this model */ +#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 72d88df9b..085bcc0ec 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 9cc094a40..d42d2f294 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 4740825cc..d37349a37 100644 --- a/src/spicelib/devices/jfet/jfetdefs.h +++ b/src/spicelib/devices/jfet/jfetdefs.h @@ -193,6 +193,10 @@ typedef struct sJFETmodel { /* model structure for a jfet */ * that have this model */ IFuid JFETmodName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int JFETtype; diff --git a/src/spicelib/devices/jfet/jfetset.c b/src/spicelib/devices/jfet/jfetset.c index 1ee580749..0000324d3 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 = model->JFETnextModel ) { +#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 dc99a4fd6..0503dcde7 100644 --- a/src/spicelib/devices/jfet2/jfet2defs.h +++ b/src/spicelib/devices/jfet2/jfet2defs.h @@ -205,6 +205,10 @@ typedef struct sJFET2model { /* model structure for a jfet */ * that have this model */ IFuid JFET2modName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int JFET2type; diff --git a/src/spicelib/devices/jfet2/jfet2set.c b/src/spicelib/devices/jfet2/jfet2set.c index 9ba95678c..f80981c67 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 = model->JFET2nextModel ) { +#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 9ba680c3d..dbaa4e4e3 100644 --- a/src/spicelib/devices/ltra/ltradefs.h +++ b/src/spicelib/devices/ltra/ltradefs.h @@ -104,6 +104,10 @@ typedef struct sLTRAmodel { /* model structure for a transmission lines */ * model */ IFuid LTRAmodName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ double LTRAh1dashFirstVal; /* first needed value of h1dasg at diff --git a/src/spicelib/devices/ltra/ltraset.c b/src/spicelib/devices/ltra/ltraset.c index 30f21b3a1..af77cf4bd 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 = model->LTRAnextModel) { +#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 96e6c4967..7a8815714 100644 --- a/src/spicelib/devices/mes/mesdefs.h +++ b/src/spicelib/devices/mes/mesdefs.h @@ -186,6 +186,10 @@ typedef struct sMESmodel { /* model structure for a mesfet */ * that have this model */ IFuid MESmodName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int MEStype; diff --git a/src/spicelib/devices/mes/messetup.c b/src/spicelib/devices/mes/messetup.c index e7c885361..1ac8dd888 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 = model->MESnextModel ) { +#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 cb6f84710..1628ea41b 100644 --- a/src/spicelib/devices/mesa/mesadefs.h +++ b/src/spicelib/devices/mesa/mesadefs.h @@ -253,6 +253,10 @@ typedef struct sMESAmodel { /* model structure for a MESAfet */ * that have this model */ IFuid MESAmodName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int MESAtype; diff --git a/src/spicelib/devices/mesa/mesasetup.c b/src/spicelib/devices/mesa/mesasetup.c index b11ebac5b..2907bad2c 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 = model->MESAnextModel ) { + +#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 f0eae205a..9824770fd 100644 --- a/src/spicelib/devices/mos1/mos1defs.h +++ b/src/spicelib/devices/mos1/mos1defs.h @@ -344,6 +344,10 @@ typedef struct sMOS1model { /* model structure for a resistor */ * that have this model */ IFuid MOS1modName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int MOS1type; /* device type : 1 = nmos, -1 = pmos */ diff --git a/src/spicelib/devices/mos1/mos1set.c b/src/spicelib/devices/mos1/mos1set.c index 128c6769e..4491b3fc5 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 = model->MOS1nextModel ) { +#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 8e1c5bfae..c01e1ee59 100644 --- a/src/spicelib/devices/mos2/mos2defs.h +++ b/src/spicelib/devices/mos2/mos2defs.h @@ -350,6 +350,10 @@ typedef struct sMOS2model { /* model structure for a resistor */ * that have this model */ IFuid MOS2modName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int MOS2type; /* device type : 1 = nmos, -1 = pmos */ diff --git a/src/spicelib/devices/mos2/mos2set.c b/src/spicelib/devices/mos2/mos2set.c index 5ed8d5ce8..652567691 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 = model->MOS2nextModel ) { +#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 7f5db234b..6e9d08be2 100644 --- a/src/spicelib/devices/mos3/mos3defs.h +++ b/src/spicelib/devices/mos3/mos3defs.h @@ -348,6 +348,10 @@ typedef struct sMOS3model { /* model structure for a resistor */ * that have this model */ IFuid MOS3modName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int MOS3type; /* device type : 1 = nmos, -1 = pmos */ diff --git a/src/spicelib/devices/mos3/mos3set.c b/src/spicelib/devices/mos3/mos3set.c index 1e8cac082..a63a27bca 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 = model->MOS3nextModel ) { +#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 b448a17f7..1de132617 100644 --- a/src/spicelib/devices/mos6/mos6defs.h +++ b/src/spicelib/devices/mos6/mos6defs.h @@ -275,6 +275,10 @@ typedef struct sMOS6model { /* model structure for a resistor */ * that have this model */ IFuid MOS6modName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int MOS6type; /* device type : 1 = nmos, -1 = pmos */ diff --git a/src/spicelib/devices/mos6/mos6set.c b/src/spicelib/devices/mos6/mos6set.c index 5e8a913b3..a8cbb2fbb 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 = model->MOS6nextModel ) { +#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 66e3753f8..2b8313f96 100644 --- a/src/spicelib/devices/mos9/mos9defs.h +++ b/src/spicelib/devices/mos9/mos9defs.h @@ -350,6 +350,10 @@ typedef struct sMOS9model { /* model structure for a resistor */ * that have this model */ IFuid MOS9modName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int MOS9type; /* device type : 1 = nmos, -1 = pmos */ diff --git a/src/spicelib/devices/mos9/mos9set.c b/src/spicelib/devices/mos9/mos9set.c index 78d4abe47..19d92f34b 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 = model->MOS9nextModel ) { +#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 2ba794eb3..6f48691d5 100644 --- a/src/spicelib/devices/res/resdefs.h +++ b/src/spicelib/devices/res/resdefs.h @@ -145,6 +145,10 @@ typedef struct sRESmodel { /* model structure for a resistor */ * model */ IFuid RESmodName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ double REStnom; /* temperature at which resistance measured */ diff --git a/src/spicelib/devices/res/ressetup.c b/src/spicelib/devices/res/ressetup.c index f4c3fe5f7..50db1db6f 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 82359b7d4..27d33afcb 100644 --- a/src/spicelib/devices/soi3/soi3defs.h +++ b/src/spicelib/devices/soi3/soi3defs.h @@ -482,6 +482,10 @@ typedef struct sSOI3model { /* model structure for an SOI3 MOSFET */ * that have this model */ IFuid SOI3modName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int SOI3type; /* device type : 1 = nsoi, -1 = psoi */ diff --git a/src/spicelib/devices/soi3/soi3set.c b/src/spicelib/devices/soi3/soi3set.c index fa526700d..81dff1f5f 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 = model->SOI3nextModel ) { +#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 126d4e119..6a14544f6 100644 --- a/src/spicelib/devices/sw/swdefs.h +++ b/src/spicelib/devices/sw/swdefs.h @@ -71,6 +71,10 @@ typedef struct sSWmodel { /* model structure for a switch */ * model */ IFuid SWmodName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ double SWonResistance; /* switch "on" resistance */ diff --git a/src/spicelib/devices/sw/swsetup.c b/src/spicelib/devices/sw/swsetup.c index 7ba4a4905..ae3235a05 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 = model->SWnextModel ) { + +#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 a8f12b050..ff94d6bc4 100644 --- a/src/spicelib/devices/tra/tradefs.h +++ b/src/spicelib/devices/tra/tradefs.h @@ -119,6 +119,10 @@ typedef struct sTRAmodel { /* model structure for a transmission lines */ * model */ IFuid TRAmodName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ } TRAmodel; diff --git a/src/spicelib/devices/tra/trasetup.c b/src/spicelib/devices/tra/trasetup.c index bdf83a243..b23dca9ef 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 = model->TRAnextModel ) { +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + /* loop through all the instances of the model */ for (here = model->TRAinstances; here != NULL ; here=here->TRAnextInstance) { diff --git a/src/spicelib/devices/txl/txldefs.h b/src/spicelib/devices/txl/txldefs.h index 59dcf0e44..8eebebb01 100644 --- a/src/spicelib/devices/txl/txldefs.h +++ b/src/spicelib/devices/txl/txldefs.h @@ -79,6 +79,10 @@ typedef struct sTXLmodel { /* model structure for a txl */ * model */ IFuid TXLmodName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ double R; diff --git a/src/spicelib/devices/txl/txlsetup.c b/src/spicelib/devices/txl/txlsetup.c index 0eea24240..936b11228 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 = model->TXLnextModel ) { +#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 4adbb88c4..18230c651 100644 --- a/src/spicelib/devices/urc/urcdefs.h +++ b/src/spicelib/devices/urc/urcdefs.h @@ -44,6 +44,10 @@ typedef struct sURCmodel { /* model structure for a resistor */ * model */ IFuid URCmodName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ double URCk; /* propagation constant for URC */ diff --git a/src/spicelib/devices/urc/urcsetup.c b/src/spicelib/devices/urc/urcsetup.c index 2df4ef01f..31beeefbd 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 = model->URCnextModel ) { + +#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 ebeb65ecc..5e8ad6074 100644 --- a/src/spicelib/devices/vbic/vbicdefs.h +++ b/src/spicelib/devices/vbic/vbicdefs.h @@ -412,6 +412,10 @@ typedef struct sVBICmodel { /* model structure for a vbic */ IFuid VBICmodName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ int VBICtype; diff --git a/src/spicelib/devices/vbic/vbicsetup.c b/src/spicelib/devices/vbic/vbicsetup.c index 7c483a8aa..641555103 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 = model->VBICnextModel ) { +#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 7bc5b5244..7c35e2221 100644 --- a/src/spicelib/devices/vccs/vccsdefs.h +++ b/src/spicelib/devices/vccs/vccsdefs.h @@ -69,6 +69,10 @@ typedef struct sVCCSmodel { /* model structure for a source */ * that have this model */ IFuid VCCSmodName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ } VCCSmodel; diff --git a/src/spicelib/devices/vccs/vccsset.c b/src/spicelib/devices/vccs/vccsset.c index 374b2e926..5cb0b57e8 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 = model->VCCSnextModel ) { +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + /* loop through all the instances of the model */ for (here = model->VCCSinstances; here != NULL ; here=here->VCCSnextInstance) { diff --git a/src/spicelib/devices/vcvs/vcvsdefs.h b/src/spicelib/devices/vcvs/vcvsdefs.h index f8f815337..c2512f1f8 100644 --- a/src/spicelib/devices/vcvs/vcvsdefs.h +++ b/src/spicelib/devices/vcvs/vcvsdefs.h @@ -75,6 +75,10 @@ typedef struct sVCVSmodel { /* model structure for a source */ * that have this model */ IFuid VCVSmodName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ } VCVSmodel; diff --git a/src/spicelib/devices/vcvs/vcvsset.c b/src/spicelib/devices/vcvs/vcvsset.c index 981f8af79..ae5dccd4d 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 = model->VCVSnextModel ) { +#ifdef USE_CUSPICE + /* This model doesn't support CUDA */ + model->has_cuda = 0 ; +#endif + /* loop through all the instances of the model */ for (here = model->VCVSinstances; here != NULL ; here=here->VCVSnextInstance) { diff --git a/src/spicelib/devices/vsrc/vsrcdefs.h b/src/spicelib/devices/vsrc/vsrcdefs.h index f8b693cdb..ebf11d8f6 100644 --- a/src/spicelib/devices/vsrc/vsrcdefs.h +++ b/src/spicelib/devices/vsrc/vsrcdefs.h @@ -139,6 +139,10 @@ typedef struct sVSRCmodel { * that have this model */ IFuid VSRCmodName; /* pointer to character string naming this model */ +#ifdef USE_CUSPICE + unsigned int has_cuda:1 ; +#endif + /* --- end of generic struct GENmodel --- */ #ifdef USE_CUSPICE diff --git a/src/spicelib/devices/vsrc/vsrcset.c b/src/spicelib/devices/vsrc/vsrcset.c index bf894c1eb..20cc120ac 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 */