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

Added the support for CUDA and non-CUDA models within the same netlist in CUSPICE 

Added KLU support for '.ic' and '.nodeset' instructions, needed also by CUSPICE
This commit is contained in:
Francesco Lannutti 2017-09-19 21:19:23 +02:00 committed by rlar
parent 5c766ed162
commit 0d4c76fd5d
107 changed files with 777 additions and 230 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
src/include/ngspice/gendefs.h
View File

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

24
src/spicelib/analysis/CUSPICE/cucktflush.c
View File

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

12
src/spicelib/analysis/CUSPICE/cucktnonconupdate.c
View File

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

22
src/spicelib/analysis/CUSPICE/cucktrhsoldupdate.c
View File

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

122
src/spicelib/analysis/CUSPICE/cucktsetup.c
View File

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

78
src/spicelib/analysis/CUSPICE/cucktstatesupdate.c
View File

@ -28,6 +28,7 @@
#include "ngspice/sperror.h"
#include "cuda_runtime_api.h"
#include "ngspice/CUSPICE/CUSPICE.h"
#include <string.h>
/* 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) ;
}

18
src/spicelib/analysis/CUSPICE/cucktsystem.c
View File

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

90
src/spicelib/analysis/CUSPICE/cuckttrunc.cu
View File

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

108
src/spicelib/analysis/cktload.c
View File

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

49
src/spicelib/analysis/cktsetup.c
View File

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

4
src/spicelib/devices/asrc/asrcdefs.h
View File

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

6
src/spicelib/devices/asrc/asrcset.c
View File

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

4
src/spicelib/devices/bjt/bjtdefs.h
View File

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

5
src/spicelib/devices/bjt/bjtsetup.c
View File

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

9
src/spicelib/devices/bsim1/b1set.c
View File

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

4
src/spicelib/devices/bsim1/bsim1def.h
View File

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

9
src/spicelib/devices/bsim2/b2set.c
View File

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

4
src/spicelib/devices/bsim2/bsim2def.h
View File

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

8
src/spicelib/devices/bsim3/b3set.c
View File

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

4
src/spicelib/devices/bsim3/bsim3def.h
View File

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

4
src/spicelib/devices/bsim3soi_dd/b3soidddef.h
View File

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

7
src/spicelib/devices/bsim3soi_dd/b3soiddset.c
View File

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

4
src/spicelib/devices/bsim3soi_fd/b3soifddef.h
View File

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

7
src/spicelib/devices/bsim3soi_fd/b3soifdset.c
View File

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

4
src/spicelib/devices/bsim3soi_pd/b3soipddef.h
View File

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

7
src/spicelib/devices/bsim3soi_pd/b3soipdset.c
View File

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

8
src/spicelib/devices/bsim3v0/b3v0set.c
View File

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

4
src/spicelib/devices/bsim3v0/bsim3v0def.h
View File

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

8
src/spicelib/devices/bsim3v1/b3v1set.c
View File

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

4
src/spicelib/devices/bsim3v1/bsim3v1def.h
View File

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

8
src/spicelib/devices/bsim3v32/b3v32set.c
View File

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

4
src/spicelib/devices/bsim3v32/bsim3v32def.h
View File

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

9
src/spicelib/devices/bsim4/b4set.c
View File

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

4
src/spicelib/devices/bsim4/bsim4def.h
View File

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

9
src/spicelib/devices/bsim4v5/b4v5set.c
View File

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

4
src/spicelib/devices/bsim4v5/bsim4v5def.h
View File

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

9
src/spicelib/devices/bsim4v6/b4v6set.c
View File

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

4
src/spicelib/devices/bsim4v6/bsim4v6def.h
View File

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

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

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

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

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

4
src/spicelib/devices/bsimsoi/b4soidef.h
View File

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

7
src/spicelib/devices/bsimsoi/b4soiset.c
View File

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

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

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

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

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

4
src/spicelib/devices/cccs/cccsdefs.h
View File

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

5
src/spicelib/devices/cccs/cccsset.c
View File

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

4
src/spicelib/devices/ccvs/ccvsdefs.h
View File

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

5
src/spicelib/devices/ccvs/ccvsset.c
View File

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

4
src/spicelib/devices/cpl/cpldefs.h
View File

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

5
src/spicelib/devices/cpl/cplsetup.c
View File

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

4
src/spicelib/devices/csw/cswdefs.h
View File

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

6
src/spicelib/devices/csw/cswsetup.c
View File

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

4
src/spicelib/devices/dio/diodefs.h
View File

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

5
src/spicelib/devices/dio/diosetup.c
View File

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

4
src/spicelib/devices/hfet1/hfetdefs.h
View File

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

6
src/spicelib/devices/hfet1/hfetsetup.c
View File

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

4
src/spicelib/devices/hfet2/hfet2defs.h
View File

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

6
src/spicelib/devices/hfet2/hfet2setup.c
View File

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

4
src/spicelib/devices/hisim2/hsm2def.h
View File

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

6
src/spicelib/devices/hisim2/hsm2set.c
View File

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

4
src/spicelib/devices/hisimhv1/hsmhvdef.h
View File

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

6
src/spicelib/devices/hisimhv1/hsmhvset.c
View File

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

4
src/spicelib/devices/hisimhv2/hsmhv2def.h
View File

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

6
src/spicelib/devices/hisimhv2/hsmhv2set.c
View File

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

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

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

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

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

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

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

4
src/spicelib/devices/jfet/jfetdefs.h
View File

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

5
src/spicelib/devices/jfet/jfetset.c
View File

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

4
src/spicelib/devices/jfet2/jfet2defs.h
View File

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

5
src/spicelib/devices/jfet2/jfet2set.c
View File

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

4
src/spicelib/devices/ltra/ltradefs.h
View File

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

5
src/spicelib/devices/ltra/ltraset.c
View File

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

4
src/spicelib/devices/mes/mesdefs.h
View File

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

5
src/spicelib/devices/mes/messetup.c
View File

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

4
src/spicelib/devices/mesa/mesadefs.h
View File

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

6
src/spicelib/devices/mesa/mesasetup.c
View File

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

4
src/spicelib/devices/mos1/mos1defs.h
View File

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

5
src/spicelib/devices/mos1/mos1set.c
View File

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

4
src/spicelib/devices/mos2/mos2defs.h
View File

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

5
src/spicelib/devices/mos2/mos2set.c
View File

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

4
src/spicelib/devices/mos3/mos3defs.h
View File

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

5
src/spicelib/devices/mos3/mos3set.c
View File

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

4
src/spicelib/devices/mos6/mos6defs.h
View File

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

5
src/spicelib/devices/mos6/mos6set.c
View File

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

4
src/spicelib/devices/mos9/mos9defs.h
View File

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

5
src/spicelib/devices/mos9/mos9set.c
View File

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

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

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

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

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

4
src/spicelib/devices/soi3/soi3defs.h
View File

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

5
src/spicelib/devices/soi3/soi3set.c
View File

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

4
src/spicelib/devices/sw/swdefs.h
View File

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

6
src/spicelib/devices/sw/swsetup.c
View File

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

4
src/spicelib/devices/tra/tradefs.h
View File

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

5
src/spicelib/devices/tra/trasetup.c
View File

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

4
src/spicelib/devices/txl/txldefs.h
View File

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

5
src/spicelib/devices/txl/txlsetup.c
View File

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

4
src/spicelib/devices/urc/urcdefs.h
View File

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

6
src/spicelib/devices/urc/urcsetup.c
View File

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

4
src/spicelib/devices/vbic/vbicdefs.h
View File

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

Some files were not shown because too many files have changed in this diff Show More