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 a4b9168863
commit a2dcdbc972
107 changed files with 874 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
};

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

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

@ -71,6 +71,12 @@ typedef struct sASRCmodel { /* model structure for a source */
#define ASRCinstances(inst) ((ASRCinstance *)((inst)->gen.GENinstances))
#define ASRCmodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
} ASRCmodel;

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 = ASRCnextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
for (here = ASRCinstances(model); here; here=ASRCnextInstance(here)) {
if (!here->ASRCtree)

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

@ -360,6 +360,12 @@ typedef struct sBJTmodel { /* model structure for a bjt */
#define BJTinstances(inst) ((BJTinstance *)((inst)->gen.GENinstances))
#define BJTmodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int BJTtype;
int BJTsubs;

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 = BJTnextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
if(model->BJTtype != NPN && model->BJTtype != PNP) {
model->BJTtype = NPN;
}

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 = B1nextModel(model)) {
/* Default value Processing for B1 MOSFET Models */
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
/* Default value Processing for B1 MOSFET Models */
if( ! model->B1typeGiven) {
model->B1type = NMOS; /* NMOS */
}

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

@ -349,6 +349,12 @@ typedef struct sBSIM1model { /* model structure for a resistor */
#define B1instances(inst) ((B1instance *)((inst)->gen.GENinstances))
#define B1modName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int B1type; /* device type : 1 = nmos, -1 = pmos */
double B1vfb0;

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 = B2nextModel(model)) {
/* Default value Processing for B2 MOSFET Models */
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
/* Default value Processing for B2 MOSFET Models */
if( ! model->B2typeGiven) {
model->B2type = NMOS; /* NMOS */
}

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

@ -272,6 +272,12 @@ typedef struct sBSIM2model { /* model structure for a resistor */
#define B2instances(inst) ((B2instance *)((inst)->gen.GENinstances))
#define B2modName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int B2type; /* device type: 1 = nmos, -1 = pmos */
int pad;

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

@ -51,7 +51,13 @@ BSIM3instance **InstArray;
/* loop through all the BSIM3 device models */
for( ; model != NULL; model = BSIM3nextModel(model))
{
/* Default value Processing for BSIM3 MOSFET Models */
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
/* Default value Processing for BSIM3 MOSFET Models */
if (!model->BSIM3typeGiven)
model->BSIM3type = NMOS;
if (!model->BSIM3mobModGiven)

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

@ -445,6 +445,12 @@ typedef struct sBSIM3model
#define BSIM3instances(inst) ((BSIM3instance *)((inst)->gen.GENinstances))
#define BSIM3modName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int BSIM3type;
int BSIM3mobMod;

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

@ -704,6 +704,12 @@ typedef struct sB3SOIDDmodel
#define B3SOIDDinstances(inst) ((B3SOIDDinstance *)((inst)->gen.GENinstances))
#define B3SOIDDmodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int B3SOIDDtype;
int B3SOIDDmobMod;

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 = B3SOIDDnextModel(model))
{
/* Default value Processing for B3SOIDD MOSFET Models */
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
/* Default value Processing for B3SOIDD MOSFET Models */
if (!model->B3SOIDDtypeGiven)
model->B3SOIDDtype = NMOS;
if (!model->B3SOIDDmobModGiven)

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

@ -696,6 +696,12 @@ typedef struct sB3SOIFDmodel
#define B3SOIFDinstances(inst) ((B3SOIFDinstance *)((inst)->gen.GENinstances))
#define B3SOIFDmodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int B3SOIFDtype;
int B3SOIFDmobMod;

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 = B3SOIFDnextModel(model))
{
/* Default value Processing for B3SOIFD MOSFET Models */
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
/* Default value Processing for B3SOIFD MOSFET Models */
if (!model->B3SOIFDtypeGiven)
model->B3SOIFDtype = NMOS;
if (!model->B3SOIFDmobModGiven)

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

@ -665,6 +665,12 @@ typedef struct sB3SOIPDmodel
#define B3SOIPDinstances(inst) ((B3SOIPDinstance *)((inst)->gen.GENinstances))
#define B3SOIPDmodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int B3SOIPDtype;
int B3SOIPDmobMod;

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 = B3SOIPDnextModel(model))
{
/* Default value Processing for B3SOIPD MOSFET Models */
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
/* Default value Processing for B3SOIPD MOSFET Models */
if (!model->B3SOIPDtypeGiven)
model->B3SOIPDtype = NMOS;
if (!model->B3SOIPDmobModGiven)

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 = BSIM3v0nextModel(model))
{
/* Default value Processing for BSIM3v0 MOSFET Models */
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
/* Default value Processing for BSIM3v0 MOSFET Models */
if (!model->BSIM3v0typeGiven)
model->BSIM3v0type = NMOS;
if (!model->BSIM3v0mobModGiven)

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

@ -363,6 +363,12 @@ typedef struct sBSIM3v0model
#define BSIM3v0instances(inst) ((BSIM3v0instance *)((inst)->gen.GENinstances))
#define BSIM3v0modName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int BSIM3v0type;
int BSIM3v0mobMod;

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 = BSIM3v1nextModel(model))
{
/* Default value Processing for BSIM3v1 MOSFET Models */
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
/* Default value Processing for BSIM3v1 MOSFET Models */
if (!model->BSIM3v1typeGiven)
model->BSIM3v1type = NMOS;
if (!model->BSIM3v1mobModGiven)

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

@ -365,6 +365,12 @@ typedef struct sBSIM3v1model
#define BSIM3v1instances(inst) ((BSIM3v1instance *)((inst)->gen.GENinstances))
#define BSIM3v1modName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int BSIM3v1type;
int BSIM3v1mobMod;

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

@ -47,7 +47,13 @@ BSIM3v32instance **InstArray;
/* loop through all the BSIM3v32 device models */
for( ; model != NULL; model = BSIM3v32nextModel(model))
{
/* Default value Processing for BSIM3v32 MOSFET Models */
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
/* Default value Processing for BSIM3v32 MOSFET Models */
if (!model->BSIM3v32typeGiven)
model->BSIM3v32type = NMOS;
if (!model->BSIM3v32mobModGiven)

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

@ -442,6 +442,12 @@ typedef struct sBSIM3v32model
#define BSIM3v32instances(inst) ((BSIM3v32instance *)((inst)->gen.GENinstances))
#define BSIM3v32modName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int BSIM3v32type;
int BSIM3v32mobMod;

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 = BSIM4nextModel(model))
{ /* process defaults of model parameters */
{
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
/* process defaults of model parameters */
if (!model->BSIM4typeGiven)
model->BSIM4type = NMOS;

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

@ -942,6 +942,12 @@ typedef struct sBSIM4model
#define BSIM4instances(inst) ((BSIM4instance *)((inst)->gen.GENinstances))
#define BSIM4modName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int BSIM4type;
int BSIM4mobMod;

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 = BSIM4v5nextModel(model))
{ /* process defaults of model parameters */
{
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
/* process defaults of model parameters */
if (!model->BSIM4v5typeGiven)
model->BSIM4v5type = NMOS;

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

@ -859,6 +859,12 @@ typedef struct sBSIM4v5model
#define BSIM4v5instances(inst) ((BSIM4v5instance *)((inst)->gen.GENinstances))
#define BSIM4v5modName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int BSIM4v5type;
int BSIM4v5mobMod;

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 = BSIM4v6nextModel(model))
{ /* process defaults of model parameters */
{
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
/* process defaults of model parameters */
if (!model->BSIM4v6typeGiven)
model->BSIM4v6type = NMOS;

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

@ -880,6 +880,12 @@ typedef struct sBSIM4v6model
#define BSIM4v6instances(inst) ((BSIM4v6instance *)((inst)->gen.GENinstances))
#define BSIM4v6modName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int BSIM4v6type;
int BSIM4v6mobMod;

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

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

@ -1435,6 +1435,12 @@ typedef struct sBSIM4v7model
#define BSIM4v7instances(inst) ((BSIM4v7instance *)((inst)->gen.GENinstances))
#define BSIM4v7modName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int BSIM4v7type;
int BSIM4v7mobMod;

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

@ -1195,6 +1195,12 @@ typedef struct sB4SOImodel
#define B4SOIinstances(inst) ((B4SOIinstance *)((inst)->gen.GENinstances))
#define B4SOImodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int B4SOItype;
int B4SOImobMod;

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 = B4SOInextModel(model))
{
/* Default value Processing for B4SOI MOSFET Models */
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
/* Default value Processing for B4SOI MOSFET Models */
if (!model->B4SOItypeGiven)
model->B4SOItype = NMOS;
if (!model->B4SOImobModGiven)

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

@ -120,6 +120,12 @@ typedef struct sCAPmodel { /* model structure for a capacitor */
#define CAPinstances(inst) ((CAPinstance *)((inst)->gen.GENinstances))
#define CAPmodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
double CAPtnom; /* temperature at which capacitance measured */
double CAPtempCoeff1; /* linear temperature coefficient */
double CAPtempCoeff2; /* quadratic temperature coefficient */

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

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

@ -62,6 +62,12 @@ typedef struct sCCCSmodel { /* model structure for a source */
#define CCCSinstances(inst) ((CCCSinstance *)((inst)->gen.GENinstances))
#define CCCSmodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
} CCCSmodel;
/* device parameters */

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 = CCCSnextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
/* loop through all the instances of the model */
for (here = CCCSinstances(model); here != NULL ;
here=CCCSnextInstance(here)) {

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

@ -69,6 +69,12 @@ typedef struct sCCVSmodel { /* model structure for a CCVsource */
#define CCVSinstances(inst) ((CCVSinstance *)((inst)->gen.GENinstances))
#define CCVSmodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
} CCVSmodel;
/* device parameters */

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 = CCVSnextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
/* loop through all the instances of the model */
for (here = CCVSinstances(model); here != NULL ;
here=CCVSnextInstance(here)) {

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

@ -91,6 +91,12 @@ typedef struct sCPLmodel { /* model structure for a cpl */
#define CPLinstances(inst) ((CPLinstance *)((inst)->gen.GENinstances))
#define CPLmodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
double *Rm;
int Rm_counter;
double *Gm;

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 = CPLnextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
if (!model->Rmgiven) {
SPfrontEnd->IFerrorf (ERR_FATAL,
"model %s: lossy line series resistance not given", model->CPLmodName);

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

@ -75,6 +75,12 @@ typedef struct sCSWmodel { /* model structure for a switch */
#define CSWinstances(inst) ((CSWinstance *)((inst)->gen.GENinstances))
#define CSWmodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
double CSWonResistance; /* switch "on" resistance */
double CSWoffResistance; /* switch "off" resistance */
double CSWiThreshold; /* switching threshold current */

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 = CSWnextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
/* Default Value Processing for Switch Model */
if (!model->CSWthreshGiven) {
model->CSWiThreshold = 0;

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

@ -185,6 +185,12 @@ typedef struct sDIOmodel { /* model structure for a diode */
#define DIOinstances(inst) ((DIOinstance *)((inst)->gen.GENinstances))
#define DIOmodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
unsigned DIOlevelGiven : 1;
unsigned DIOsatCurGiven : 1;
unsigned DIOsatSWCurGiven : 1;

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 = DIOnextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
if(!model->DIOlevelGiven) {
model->DIOlevel = 1;
}

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

@ -170,6 +170,12 @@ typedef struct sHFETAmodel {
#define HFETAinstances(inst) ((HFETAinstance *)((inst)->gen.GENinstances))
#define HFETAmodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int HFETAtype;
int HFETAgatemod;

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 = HFETAnextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
if( (model->HFETAtype != NHFET) && (model->HFETAtype != PHFET) ) {
model->HFETAtype = NHFET;
}

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

@ -115,6 +115,12 @@ typedef struct sHFET2model {
#define HFET2instances(inst) ((HFET2instance *)((inst)->gen.GENinstances))
#define HFET2modName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int HFET2type;
double HFET2cf;

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 = HFET2nextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
if((TYPE != NHFET) && (TYPE != PHFET) )
TYPE = NHFET;
if(!model->HFET2cfGiven)

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

@ -739,6 +739,12 @@ typedef struct sHSM2model { /* model structure for a resistor */
#define HSM2instances(inst) ((HSM2instance *)((inst)->gen.GENinstances))
#define HSM2modName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int HSM2_type; /* device type: 1 = nmos, -1 = pmos */
int HSM2_level; /* level */
int HSM2_info; /* information */

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 = HSM2nextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
/* Default value Processing for HSM2 MOSFET Models */
if ( !model->HSM2_type_Given )
model->HSM2_type = NMOS ;

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

@ -1012,6 +1012,12 @@ typedef struct sHSMHVmodel { /* model structure for a resistor */
#define HSMHVinstances(inst) ((HSMHVinstance *)((inst)->gen.GENinstances))
#define HSMHVmodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int HSMHV_type; /* device type: 1 = nmos, -1 = pmos */
int HSMHV_level; /* level */
int HSMHV_info; /* information */

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 = HSMHVnextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
/* Default value Processing for HVMOS Models */
if ( !model->HSMHV_type_Given )
model->HSMHV_type = NMOS ;

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

@ -1137,6 +1137,12 @@ typedef struct sHSMHV2model { /* model structure for a resistor */
#define HSMHV2instances(inst) ((HSMHV2instance *)((inst)->gen.GENinstances))
#define HSMHV2modName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int HSMHV2_type; /* device type: 1 = nmos, -1 = pmos */
int HSMHV2_level; /* level */
int HSMHV2_info; /* information */

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 = HSMHV2nextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
/* Default value Processing for HVMOS Models */
if ( !model->HSMHV2_type_Given )
model->HSMHV2_type = NMOS ;

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

@ -130,6 +130,12 @@ struct sINDmodel { /* model structure for an inductor */
#define INDinstances(inst) ((INDinstance *)((inst)->gen.GENinstances))
#define INDmodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
double INDmInd; /* Model inductance */
double INDtnom; /* temperature at which inductance measured */
double INDtempCoeff1; /* first temperature coefficient */
@ -251,6 +257,12 @@ struct sMUTmodel { /* model structure for a mutual inductor */
#define MUTinstances(inst) ((MUTinstance *)((inst)->gen.GENinstances))
#define MUTmodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
#ifdef USE_CUSPICE
MUTparamCPUstruct MUTparamCPU;
MUTparamGPUstruct MUTparamGPU;

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

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

@ -197,6 +197,12 @@ typedef struct sJFETmodel { /* model structure for a jfet */
#define JFETinstances(inst) ((JFETinstance *)((inst)->gen.GENinstances))
#define JFETmodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int JFETtype;
double JFETthreshold;

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 = JFETnextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
if( (model->JFETtype != NJF) && (model->JFETtype != PJF) ) {
model->JFETtype = NJF;
}

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

@ -209,6 +209,12 @@ typedef struct sJFET2model { /* model structure for a jfet */
#define JFET2instances(inst) ((JFET2instance *)((inst)->gen.GENinstances))
#define JFET2modName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int JFET2type;
#define PARAM(code,id,flag,ref,default,descrip) double ref;

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 = JFET2nextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
if( (model->JFET2type != NJF) && (model->JFET2type != PJF) ) {
model->JFET2type = NJF;
}

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

@ -108,6 +108,12 @@ typedef struct sLTRAmodel { /* model structure for a transmission lines */
#define LTRAinstances(inst) ((LTRAinstance *)((inst)->gen.GENinstances))
#define LTRAmodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
double LTRAh1dashFirstVal; /* first needed value of h1dasg at
current timepoint */
double LTRAh2FirstVal; /* first needed value of h2 at current

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 = LTRAnextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
if (!model->LTRAnlGiven) {
model->LTRAnl = .25;
}

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

@ -190,6 +190,12 @@ typedef struct sMESmodel { /* model structure for a mesfet */
#define MESinstances(inst) ((MESinstance *)((inst)->gen.GENinstances))
#define MESmodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int MEStype;
double MESthreshold;

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 = MESnextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
if( (model->MEStype != NMF) && (model->MEStype != PMF) ) {
model->MEStype = NMF;
}

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

@ -256,6 +256,12 @@ typedef struct sMESAmodel { /* model structure for a MESAfet */
#define MESAinstances(inst) ((MESAinstance *)((inst)->gen.GENinstances))
#define MESAmodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int MESAtype;
double MESAthreshold;

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 = MESAnextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
if( (model->MESAtype != NMF) ) {
fprintf(stderr, "Only nmf model type supported, set to nmf\n");
model->MESAtype = NMF;

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

@ -347,6 +347,12 @@ typedef struct sMOS1model { /* model structure for a resistor */
#define MOS1instances(inst) ((MOS1instance *)((inst)->gen.GENinstances))
#define MOS1modName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int MOS1type; /* device type : 1 = nmos, -1 = pmos */
double MOS1tnom; /* temperature at which parameters measured */
double MOS1latDiff;

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 = MOS1nextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
if(!model->MOS1typeGiven) {
model->MOS1type = NMOS;
}

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

@ -354,6 +354,12 @@ typedef struct sMOS2model { /* model structure for a resistor */
#define MOS2instances(inst) ((MOS2instance *)((inst)->gen.GENinstances))
#define MOS2modName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int MOS2type; /* device type : 1 = nmos, -1 = pmos */
int MOS2gateType;

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 = MOS2nextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
if(!model->MOS2typeGiven) {
model->MOS2type = NMOS;
}

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

@ -352,6 +352,12 @@ typedef struct sMOS3model { /* model structure for a resistor */
#define MOS3instances(inst) ((MOS3instance *)((inst)->gen.GENinstances))
#define MOS3modName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int MOS3type; /* device type : 1 = nmos, -1 = pmos */
double MOS3tnom; /* temperature at which parameters measured */
double MOS3latDiff;

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 = MOS3nextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
/* perform model defaulting */
if(!model->MOS3typeGiven) {
model->MOS3type = NMOS;

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

@ -279,6 +279,12 @@ typedef struct sMOS6model { /* model structure for a resistor */
#define MOS6instances(inst) ((MOS6instance *)((inst)->gen.GENinstances))
#define MOS6modName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int MOS6type; /* device type : 1 = nmos, -1 = pmos */
double MOS6tnom; /* temperature at which parameters measured */
double MOS6latDiff;

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 = MOS6nextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
if(!model->MOS6typeGiven) {
model->MOS6type = NMOS;
}

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

@ -354,6 +354,12 @@ typedef struct sMOS9model { /* model structure for a resistor */
#define MOS9instances(inst) ((MOS9instance *)((inst)->gen.GENinstances))
#define MOS9modName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int MOS9type; /* device type : 1 = nmos, -1 = pmos */
double MOS9tnom; /* temperature at which parameters measured */
double MOS9latDiff;

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 = MOS9nextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
/* perform model defaulting */
if(!model->MOS9typeGiven) {
model->MOS9type = NMOS;

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

@ -148,6 +148,12 @@ typedef struct sRESmodel { /* model structure for a resistor */
#define RESinstances(inst) ((RESinstance *)((inst)->gen.GENinstances))
#define RESmodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
double REStnom; /* temperature at which resistance measured */
double REStempCoeff1; /* first temperature coefficient of resistors */
double REStempCoeff2; /* second temperature coefficient of resistors */

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

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

@ -482,6 +482,12 @@ typedef struct sSOI3model { /* model structure for an SOI3 MOSFET */
#define SOI3instances(inst) ((SOI3instance *)((inst)->gen.GENinstances))
#define SOI3modName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int SOI3type; /* device type : 1 = nsoi, -1 = psoi */
double SOI3tnom; /* temperature at which parameters measured */
double SOI3latDiff;

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 = SOI3nextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
if(!model->SOI3typeGiven) {
model->SOI3type = NSOI3;
}

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

@ -74,6 +74,12 @@ typedef struct sSWmodel { /* model structure for a switch */
#define SWinstances(inst) ((SWinstance *)((inst)->gen.GENinstances))
#define SWmodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
double SWonResistance; /* switch "on" resistance */
double SWoffResistance; /* switch "off" resistance */
double SWvThreshold; /* switching threshold voltage */

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 = SWnextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
/* Default Value Processing for Switch Model */
if (!model->SWthreshGiven) {
model->SWvThreshold = 0;

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

@ -123,6 +123,12 @@ typedef struct sTRAmodel { /* model structure for a transmission lines */
#define TRAinstances(inst) ((TRAinstance *)((inst)->gen.GENinstances))
#define TRAmodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
} TRAmodel;
/* device parameters */

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 = TRAnextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
/* loop through all the instances of the model */
for (here = TRAinstances(model); here != NULL ;
here=TRAnextInstance(here)) {

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

@ -81,6 +81,12 @@ typedef struct sTXLmodel { /* model structure for a txl */
#define TXLinstances(inst) ((TXLinstance *)((inst)->gen.GENinstances))
#define TXLmodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
double R;
double L;
double G;

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 = TXLnextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
if (!model->Rgiven) {
SPfrontEnd->IFerrorf (ERR_FATAL,
"model %s: lossy line series resistance not given", model->TXLmodName);

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

@ -48,6 +48,12 @@ typedef struct sURCmodel { /* model structure for a resistor */
#define URCinstances(inst) ((URCinstance *)((inst)->gen.GENinstances))
#define URCmodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
double URCk; /* propagation constant for URC */
double URCfmax; /* max frequence of interest */
double URCrPerL; /* resistance per unit length */

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 = URCnextModel(model)) {
#ifdef USE_CUSPICE
/* This model doesn't support CUDA */
model->has_cuda = 0 ;
#endif
if(!model->URCkGiven)
model->URCk = 1.5;
if(!model->URCfmaxGiven)

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

@ -414,6 +414,12 @@ typedef struct sVBICmodel { /* model structure for a vbic */
#define VBICinstances(inst) ((VBICinstance *)((inst)->gen.GENinstances))
#define VBICmodName gen.GENmodName
#ifdef USE_CUSPICE
unsigned int has_cuda:1 ;
#endif
/* --- end of generic struct GENmodel --- */
int VBICtype;
double VBICtnom;

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