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

Fix the case when there is no RHS computed by device models

This commit is contained in:
Francesco Lannutti 2018-04-19 00:53:55 +02:00
parent 750cc1e7fd
commit 6090572fe0
15 changed files with 191 additions and 297 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -35,7 +35,7 @@ cuCKTflush
CKTcircuit *ckt
)
{
if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) {
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

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

@ -47,7 +47,7 @@ CKTcircuit *ckt
{
cudaError_t status ;
if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) {
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)
}
@ -63,7 +63,7 @@ CKTcircuit *ckt
{
cudaError_t status ;
if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) {
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)
}

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

@ -47,7 +47,7 @@ CKTcircuit *ckt
{
long unsigned int size ;
if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) {
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)) ;
}
@ -64,7 +64,7 @@ CKTcircuit *ckt
long unsigned int size ;
cudaError_t status ;
if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) {
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)
@ -82,7 +82,7 @@ CKTcircuit *ckt
long unsigned int size ;
cudaError_t status ;
if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) {
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)

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

@ -76,7 +76,7 @@ CKTcircuit *ckt
size2 = (long unsigned int)ckt->CKTnumStates ;
size3 = (long unsigned int)ckt->total_n_timeSteps ;
if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) {
if (TopologyNNZ > 0 || TopologyNNZRHS > 0) {
/* Topology Matrix Handling */
status = cudaMalloc ((void **)&(ckt->CKTmatrix->d_CKTrhs), (n + 1) * sizeof(double)) ;
CUDAMALLOCCHECK (ckt->CKTmatrix->d_CKTrhs, (n + 1), double, status)
@ -90,45 +90,51 @@ CKTcircuit *ckt
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)
if (TopologyNNZ > 0) {
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)
if (TopologyNNZRHS > 0) {
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_CKTtopologyMatrixCSRxRHS), TopologyNNZRHS * sizeof(double)) ;
CUDAMALLOCCHECK (ckt->d_CKTtopologyMatrixCSRxRHS, TopologyNNZRHS, double, 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)
}
cudaMemset (ckt->d_CKTloadOutput + ckt->total_n_values, 0, sizeof(double)) ; //DiagGmin is 0 at the beginning
if (TopologyNNZ > 0) {
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)
if (TopologyNNZRHS > 0) {
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)) ;

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

@ -48,7 +48,7 @@ CKTcircuit *ckt
{
long unsigned int size ;
if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) {
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)) ;
}
@ -65,7 +65,7 @@ CKTcircuit *ckt
long unsigned int size ;
cudaError_t status ;
if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) {
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)
@ -83,7 +83,7 @@ CKTcircuit *ckt
long unsigned int size ;
cudaError_t status ;
if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) {
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)
@ -98,7 +98,7 @@ cuCKTstate01copy
CKTcircuit *ckt
)
{
if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) {
if (ckt->total_n_Ptr > 0 || ckt->total_n_PtrRHS > 0) {
long unsigned int size ;
cudaError_t status ;
@ -121,7 +121,7 @@ CKTcircuit *ckt
int i ;
double *temp ;
if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) {
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] ;
@ -144,7 +144,7 @@ cuCKTstate123copy
CKTcircuit *ckt
)
{
if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) {
if (ckt->total_n_Ptr > 0 || ckt->total_n_PtrRHS > 0) {
long unsigned int size ;
cudaError_t status ;
@ -171,7 +171,7 @@ CKTcircuit *ckt
{
cudaError_t status ;
if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) {
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)
}

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

@ -45,20 +45,29 @@ cuCKTsystemDtoH
CKTcircuit *ckt
)
{
long unsigned int nz, n ;
long unsigned int i, nz, n ;
cudaError_t status ;
if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) {
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)
if (ckt->total_n_Ptr > 0) {
/* 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)
if (ckt->total_n_PtrRHS > 0) {
/* Copy back the RHS */
status = cudaMemcpy (ckt->CKTrhs, ckt->CKTmatrix->d_CKTrhs, (n + 1) * sizeof(double), cudaMemcpyDeviceToHost) ;
CUDAMEMCPYCHECK (ckt->CKTrhs, (n + 1), double, status)
} else {
/* RHS is empty */
for (i = 0 ; i < n + 1 ; i++) {
ckt->CKTrhs [i] = 0 ;
}
}
}
return (OK) ;

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

@ -31,7 +31,7 @@ cuCKTtrunc
CKTcircuit *ckt, double timetemp, double *timeStep
)
{
if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) {
if (ckt->total_n_Ptr > 0 || ckt->total_n_PtrRHS > 0) {
long unsigned int size ;
double timetempGPU ;
int thread_x, thread_y, block_x ;

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

@ -140,7 +140,7 @@ CKTload(CKTcircuit *ckt)
// without the zeroes along the diagonal
TopologyNNZRHS = ckt->total_n_PtrRHS ;
if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) {
if (TopologyNNZ > 0 || TopologyNNZRHS > 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
@ -149,35 +149,40 @@ CKTload(CKTcircuit *ckt)
//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 (TopologyNNZ > 0) {
/* 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,
TopologyNNZ,
&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) ;
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 (TopologyNNZRHS > 0) {
/* 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,
TopologyNNZRHS,
&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) ;
if (cusparseStatus != CUSPARSE_STATUS_SUCCESS)
{
fprintf (stderr, "CUSPARSE RHS Call Error\n") ;
return (E_NOMEM) ;
}
}
cudaDeviceSynchronize () ;

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

@ -232,101 +232,105 @@ CKTsetup(CKTcircuit *ckt)
ckt->CKTdiagElements = j ;
/* Topology Matrix Pre-Allocation in COO format */
TopologyNNZ = ckt->total_n_Ptr + ckt->CKTdiagElements ; // + ckt->CKTdiagElements because of CKTdiagGmin
// without the zeroes along the diagonal
ckt->CKTtopologyMatrixCOOi = TMALLOC (int, TopologyNNZ) ;
ckt->CKTtopologyMatrixCOOj = TMALLOC (int, TopologyNNZ) ;
ckt->CKTtopologyMatrixCOOx = TMALLOC (double, TopologyNNZ) ;
/* Topology Matrix for the RHS Pre-Allocation in COO format */
TopologyNNZRHS = ckt->total_n_PtrRHS ;
ckt->CKTtopologyMatrixCOOiRHS = TMALLOC (int, TopologyNNZRHS) ;
ckt->CKTtopologyMatrixCOOjRHS = TMALLOC (int, TopologyNNZRHS) ;
ckt->CKTtopologyMatrixCOOxRHS = TMALLOC (double, TopologyNNZRHS) ;
if (TopologyNNZ > 0 || TopologyNNZRHS > 0) {
/* Topology Matrix Construction & Topology Matrix for the RHS Construction */
if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) {
if (TopologyNNZ > 0) {
/* Topology Matrix Pre-Allocation in COO format */
ckt->CKTtopologyMatrixCOOi = TMALLOC (int, TopologyNNZ) ;
ckt->CKTtopologyMatrixCOOj = TMALLOC (int, TopologyNNZ) ;
ckt->CKTtopologyMatrixCOOx = TMALLOC (double, TopologyNNZ) ;
/* Topology Matrix Pre-Allocation in CSR format */
ckt->CKTtopologyMatrixCSRp = TMALLOC (int, nz + 1) ;
/* Topology Matrix for the RHS Pre-Allocation in CSR format */
ckt->CKTtopologyMatrixCSRpRHS = TMALLOC (int, (n + 1) + 1) ;
/* Topology Matrix Construction & Topology Matrix for the RHS Construction */
u = 0 ;
uRHS = 0 ;
for (i = 0 ; i < DEVmaxnum ; i++)
if (DEVices [i] && DEVices [i]->DEVtopology && ckt->CKThead [i])
DEVices [i]->DEVtopology (ckt->CKThead [i], ckt, &u, &uRHS) ;
/* CKTdiagGmin Contribute Addition to the Topology Matrix */
k = u ;
for (j = 0 ; j < n ; j++)
{
if (ckt->CKTmatrix->CKTdiag_CSC [j] >= ckt->CKTmatrix->CKTkluAx)
{
ckt->CKTtopologyMatrixCOOi [k] = (int)(ckt->CKTmatrix->CKTdiag_CSC [j] - ckt->CKTmatrix->CKTkluAx) ;
ckt->CKTtopologyMatrixCOOj [k] = ckt->total_n_values ;
ckt->CKTtopologyMatrixCOOx [k] = 1 ;
k++ ;
u = 0 ;
uRHS = 0 ;
for (i = 0 ; i < DEVmaxnum ; i++)
if (DEVices [i] && DEVices [i]->DEVtopology && ckt->CKThead [i])
DEVices [i]->DEVtopology (ckt->CKThead [i], ckt, &u, &uRHS) ;
}
}
/* Copy the Topology Matrix to the GPU in COO format */
if (ckt->CKTdiagElements > 0) {
/* CKTdiagGmin Contribute Addition to the Topology Matrix */
k = u ;
for (j = 0 ; j < n ; j++)
{
if (ckt->CKTmatrix->CKTdiag_CSC [j] >= ckt->CKTmatrix->CKTkluAx)
{
ckt->CKTtopologyMatrixCOOi [k] = (int)(ckt->CKTmatrix->CKTdiag_CSC [j] - ckt->CKTmatrix->CKTkluAx) ;
ckt->CKTtopologyMatrixCOOj [k] = ckt->total_n_values ;
ckt->CKTtopologyMatrixCOOx [k] = 1 ;
k++ ;
}
}
}
/* Copy the Topology Matrix to the GPU in COO format */
/* COO format to CSR format Conversion using Quick Sort */
if (TopologyNNZ > 0) {
/* COO format to CSR format Conversion using Quick Sort */
Element *TopologyStruct ;
TopologyStruct = TMALLOC (Element, TopologyNNZ) ;
Element *TopologyStruct ;
TopologyStruct = TMALLOC (Element, TopologyNNZ) ;
for (i = 0 ; i < TopologyNNZ ; i++)
{
TopologyStruct [i].row = ckt->CKTtopologyMatrixCOOi [i] ;
TopologyStruct [i].col = ckt->CKTtopologyMatrixCOOj [i] ;
TopologyStruct [i].val = ckt->CKTtopologyMatrixCOOx [i] ;
}
for (i = 0 ; i < TopologyNNZ ; i++)
{
TopologyStruct [i].row = ckt->CKTtopologyMatrixCOOi [i] ;
TopologyStruct [i].col = ckt->CKTtopologyMatrixCOOj [i] ;
TopologyStruct [i].val = ckt->CKTtopologyMatrixCOOx [i] ;
}
qsort (TopologyStruct, (size_t)TopologyNNZ, sizeof(Element), Compare) ;
qsort (TopologyStruct, (size_t)TopologyNNZ, sizeof(Element), Compare) ;
for (i = 0 ; i < TopologyNNZ ; i++)
{
ckt->CKTtopologyMatrixCOOi [i] = TopologyStruct [i].row ;
ckt->CKTtopologyMatrixCOOj [i] = TopologyStruct [i].col ;
ckt->CKTtopologyMatrixCOOx [i] = TopologyStruct [i].val ;
}
for (i = 0 ; i < TopologyNNZ ; i++)
{
ckt->CKTtopologyMatrixCOOi [i] = TopologyStruct [i].row ;
ckt->CKTtopologyMatrixCOOj [i] = TopologyStruct [i].col ;
ckt->CKTtopologyMatrixCOOx [i] = TopologyStruct [i].val ;
}
ret = Compress (ckt->CKTtopologyMatrixCOOi, ckt->CKTtopologyMatrixCSRp, nz, TopologyNNZ) ;
/* Topology Matrix Pre-Allocation in CSR format */
ckt->CKTtopologyMatrixCSRp = TMALLOC (int, nz + 1) ;
/* COO format to CSR format Conversion for the RHS using Quick Sort */
ret = Compress (ckt->CKTtopologyMatrixCOOi, ckt->CKTtopologyMatrixCSRp, nz, TopologyNNZ) ;
}
Element *TopologyStructRHS ;
TopologyStructRHS = TMALLOC (Element, TopologyNNZRHS) ;
if (TopologyNNZRHS > 0) {
/* Topology Matrix for the RHS Pre-Allocation in COO format */
ckt->CKTtopologyMatrixCOOiRHS = TMALLOC (int, TopologyNNZRHS) ;
ckt->CKTtopologyMatrixCOOjRHS = TMALLOC (int, TopologyNNZRHS) ;
ckt->CKTtopologyMatrixCOOxRHS = TMALLOC (double, TopologyNNZRHS) ;
for (i = 0 ; i < TopologyNNZRHS ; i++)
{
TopologyStructRHS [i].row = ckt->CKTtopologyMatrixCOOiRHS [i] ;
TopologyStructRHS [i].col = ckt->CKTtopologyMatrixCOOjRHS [i] ;
TopologyStructRHS [i].val = ckt->CKTtopologyMatrixCOOxRHS [i] ;
}
/* COO format to CSR format Conversion for the RHS using Quick Sort */
qsort (TopologyStructRHS, (size_t)TopologyNNZRHS, sizeof(Element), Compare) ;
Element *TopologyStructRHS ;
TopologyStructRHS = TMALLOC (Element, TopologyNNZRHS) ;
for (i = 0 ; i < TopologyNNZRHS ; i++)
{
ckt->CKTtopologyMatrixCOOiRHS [i] = TopologyStructRHS [i].row ;
ckt->CKTtopologyMatrixCOOjRHS [i] = TopologyStructRHS [i].col ;
ckt->CKTtopologyMatrixCOOxRHS [i] = TopologyStructRHS [i].val ;
}
for (i = 0 ; i < TopologyNNZRHS ; i++)
{
TopologyStructRHS [i].row = ckt->CKTtopologyMatrixCOOiRHS [i] ;
TopologyStructRHS [i].col = ckt->CKTtopologyMatrixCOOjRHS [i] ;
TopologyStructRHS [i].val = ckt->CKTtopologyMatrixCOOxRHS [i] ;
}
ret = Compress (ckt->CKTtopologyMatrixCOOiRHS, ckt->CKTtopologyMatrixCSRpRHS, n + 1, TopologyNNZRHS) ;
qsort (TopologyStructRHS, (size_t)TopologyNNZRHS, sizeof(Element), Compare) ;
/* Multiply the Topology Matrix by the M Vector to build the Final CSC Matrix - after the CKTload Call */
for (i = 0 ; i < TopologyNNZRHS ; i++)
{
ckt->CKTtopologyMatrixCOOiRHS [i] = TopologyStructRHS [i].row ;
ckt->CKTtopologyMatrixCOOjRHS [i] = TopologyStructRHS [i].col ;
ckt->CKTtopologyMatrixCOOxRHS [i] = TopologyStructRHS [i].val ;
}
/* Topology Matrix for the RHS Pre-Allocation in CSR format */
ckt->CKTtopologyMatrixCSRpRHS = TMALLOC (int, (n + 1) + 1) ;
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
@ -340,7 +344,7 @@ CKTsetup(CKTcircuit *ckt)
}
#ifdef USE_CUSPICE
if (ckt->total_n_Ptr > 0 && ckt->total_n_PtrRHS > 0) {
if (ckt->total_n_Ptr > 0 || ckt->total_n_PtrRHS > 0) {
ckt->d_MatrixSize = SMPmatSize (ckt->CKTmatrix) ;
status = cuCKTsetup (ckt) ;
if (status != 0)

27
src/spicelib/devices/res/CUSPICE/curesfree.c
View File

@ -39,40 +39,13 @@ GENmodel *inModel
for ( ; model != NULL ; model = RESnextModel(model))
{
/* DOUBLE */
free (model->RESparamCPU.REStc1Array) ;
cudaFree (model->RESparamGPU.d_REStc1Array) ;
free (model->RESparamCPU.REStc2Array) ;
cudaFree (model->RESparamGPU.d_REStc2Array) ;
free (model->RESparamCPU.RESmArray) ;
cudaFree (model->RESparamGPU.d_RESmArray) ;
free (model->RESparamCPU.RESconductArray) ;
cudaFree (model->RESparamGPU.d_RESconductArray) ;
free (model->RESparamCPU.REStempArray) ;
cudaFree (model->RESparamGPU.d_REStempArray) ;
free (model->RESparamCPU.RESdtempArray) ;
cudaFree (model->RESparamGPU.d_RESdtempArray) ;
free (model->RESparamCPU.REScurrentArray) ;
cudaFree (model->RESparamGPU.d_REScurrentArray) ;
free (model->RESparamCPU.RESgValueArray) ;
cudaFree (model->RESparamGPU.d_RESgValueArray) ;
/* INT */
free (model->RESparamCPU.REStc1GivenArray) ;
cudaFree (model->RESparamGPU.d_REStc1GivenArray) ;
free (model->RESparamCPU.REStc2GivenArray) ;
cudaFree (model->RESparamGPU.d_REStc2GivenArray) ;
free (model->RESparamCPU.RESmGivenArray) ;
cudaFree (model->RESparamGPU.d_RESmGivenArray) ;
free (model->RESparamCPU.RESposNodeArray) ;
cudaFree (model->RESparamGPU.d_RESposNodeArray) ;

23
src/spicelib/devices/res/CUSPICE/curesload.cu
View File

@ -78,11 +78,9 @@ __global__
void
cuRESload_kernel
(
RESparamGPUstruct RESentry, double *CKTrhsOld, int n_instances, int *d_PositionVector, double * d_CKTloadOutput
RESparamGPUstruct RESentry, double *CKTrhsOld, int n_instances, int *d_PositionVector, double *d_CKTloadOutput
)
{
double m, difference, factor ;
int instance_ID ;
instance_ID = threadIdx.y + blockDim.y * blockIdx.x ;
@ -90,26 +88,11 @@ RESparamGPUstruct RESentry, double *CKTrhsOld, int n_instances, int *d_PositionV
{
if (threadIdx.x == 0)
{
if (!(RESentry.d_REStc1GivenArray [instance_ID]))
RESentry.d_REStc1Array [instance_ID] = 0.0 ;
if (!(RESentry.d_REStc2GivenArray [instance_ID]))
RESentry.d_REStc2Array [instance_ID] = 0.0 ;
if (!(RESentry.d_RESmGivenArray [instance_ID]))
RESentry.d_RESmArray [instance_ID] = 1.0 ;
RESentry.d_REScurrentArray [instance_ID] = (CKTrhsOld [RESentry.d_RESposNodeArray [instance_ID]] -
CKTrhsOld [RESentry.d_RESnegNodeArray [instance_ID]]) *
RESentry.d_RESconductArray [instance_ID] ;
difference = (RESentry.d_REStempArray [instance_ID] + RESentry.d_RESdtempArray [instance_ID]) - 300.15 ;
factor = 1.0 + (RESentry.d_REStc1Array [instance_ID]) * difference +
(RESentry.d_REStc2Array [instance_ID]) * difference * difference ;
m = (RESentry.d_RESmArray [instance_ID]) / factor ;
d_CKTloadOutput [d_PositionVector [instance_ID]] = m * RESentry.d_RESconductArray [instance_ID] ;
d_CKTloadOutput [d_PositionVector [instance_ID]] = RESentry.d_RESconductArray [instance_ID] ;
}
}

38
src/spicelib/devices/res/CUSPICE/curessetup.c
View File

@ -58,7 +58,7 @@ GENmodel *inModel
cudaError_t status ;
RESmodel *model = (RESmodel *)inModel ;
size = (long unsigned int) model->n_instances;
size = (long unsigned int) model->n_instances ;
/* Space Allocation to GPU */
status = cudaMalloc ((void **)&(model->d_PositionVector), size * sizeof(int)) ;
@ -68,51 +68,15 @@ GENmodel *inModel
CUDAMEMCPYCHECK (model->d_PositionVector, size, int, status)
/* DOUBLE */
model->RESparamCPU.REStc1Array = (double *) malloc (size * sizeof(double)) ;
status = cudaMalloc ((void **)&(model->RESparamGPU.d_REStc1Array), size * sizeof(double)) ;
CUDAMALLOCCHECK (model->RESparamGPU.d_REStc1Array, size, double, status)
model->RESparamCPU.REStc2Array = (double *) malloc (size * sizeof(double)) ;
status = cudaMalloc ((void **)&(model->RESparamGPU.d_REStc2Array), size * sizeof(double)) ;
CUDAMALLOCCHECK (model->RESparamGPU.d_REStc2Array, size, double, status)
model->RESparamCPU.RESmArray = (double *) malloc (size * sizeof(double)) ;
status = cudaMalloc ((void **)&(model->RESparamGPU.d_RESmArray), size * sizeof(double)) ;
CUDAMALLOCCHECK (model->RESparamGPU.d_RESmArray, size, double, status)
model->RESparamCPU.RESconductArray = (double *) malloc (size * sizeof(double)) ;
status = cudaMalloc ((void **)&(model->RESparamGPU.d_RESconductArray), size * sizeof(double)) ;
CUDAMALLOCCHECK (model->RESparamGPU.d_RESconductArray, size, double, status)
model->RESparamCPU.REStempArray = (double *) malloc (size * sizeof(double)) ;
status = cudaMalloc ((void **)&(model->RESparamGPU.d_REStempArray), size * sizeof(double)) ;
CUDAMALLOCCHECK (model->RESparamGPU.d_REStempArray, size, double, status)
model->RESparamCPU.RESdtempArray = (double *) malloc (size * sizeof(double)) ;
status = cudaMalloc ((void **)&(model->RESparamGPU.d_RESdtempArray), size * sizeof(double)) ;
CUDAMALLOCCHECK (model->RESparamGPU.d_RESdtempArray, size, double, status)
model->RESparamCPU.REScurrentArray = (double *) malloc (size * sizeof(double)) ;
status = cudaMalloc ((void **)&(model->RESparamGPU.d_REScurrentArray), size * sizeof(double)) ;
CUDAMALLOCCHECK (model->RESparamGPU.d_REScurrentArray, size, double, status)
model->RESparamCPU.RESgValueArray = (double *) malloc (size * sizeof(double)) ;
status = cudaMalloc ((void **)&(model->RESparamGPU.d_RESgValueArray), size * sizeof(double)) ;
CUDAMALLOCCHECK (model->RESparamGPU.d_RESgValueArray, size, double, status)
/* INT */
model->RESparamCPU.REStc1GivenArray = (int *) malloc (size * sizeof(int)) ;
status = cudaMalloc ((void **)&(model->RESparamGPU.d_REStc1GivenArray), size * sizeof(int)) ;
CUDAMALLOCCHECK (model->RESparamGPU.d_REStc1GivenArray, size, int, status)
model->RESparamCPU.REStc2GivenArray = (int *) malloc (size * sizeof(int)) ;
status = cudaMalloc ((void **)&(model->RESparamGPU.d_REStc2GivenArray), size * sizeof(int)) ;
CUDAMALLOCCHECK (model->RESparamGPU.d_REStc2GivenArray, size, int, status)
model->RESparamCPU.RESmGivenArray = (int *) malloc (size * sizeof(int)) ;
status = cudaMalloc ((void **)&(model->RESparamGPU.d_RESmGivenArray), size * sizeof(int)) ;
CUDAMALLOCCHECK (model->RESparamGPU.d_RESmGivenArray, size, int, status)
model->RESparamCPU.RESposNodeArray = (int *) malloc (size * sizeof(int)) ;
status = cudaMalloc ((void **)&(model->RESparamGPU.d_RESposNodeArray), size * sizeof(int)) ;
CUDAMALLOCCHECK (model->RESparamGPU.d_RESposNodeArray, size, int, status)

26
src/spicelib/devices/res/CUSPICE/curestemp.c
View File

@ -48,37 +48,13 @@ GENmodel *inModel
cudaError_t status ;
RESmodel *model = (RESmodel *)inModel ;
size = (long unsigned int) model->n_instances;
size = (long unsigned int) model->n_instances ;
/* DOUBLE */
status = cudaMemcpy (model->RESparamGPU.d_REStc1Array, model->RESparamCPU.REStc1Array, size * sizeof(double), cudaMemcpyHostToDevice) ;
CUDAMEMCPYCHECK(model->RESparamGPU.d_REStc1Array, size, double, status)
status = cudaMemcpy (model->RESparamGPU.d_REStc2Array, model->RESparamCPU.REStc2Array, size * sizeof(double), cudaMemcpyHostToDevice) ;
CUDAMEMCPYCHECK(model->RESparamGPU.d_REStc2Array, size, double, status)
status = cudaMemcpy (model->RESparamGPU.d_RESmArray, model->RESparamCPU.RESmArray, size * sizeof(double), cudaMemcpyHostToDevice) ;
CUDAMEMCPYCHECK(model->RESparamGPU.d_RESmArray, size, double, status)
status = cudaMemcpy (model->RESparamGPU.d_RESconductArray, model->RESparamCPU.RESconductArray, size * sizeof(double), cudaMemcpyHostToDevice) ;
CUDAMEMCPYCHECK(model->RESparamGPU.d_RESconductArray, size, double, status)
status = cudaMemcpy (model->RESparamGPU.d_REStempArray, model->RESparamCPU.REStempArray, size * sizeof(double), cudaMemcpyHostToDevice) ;
CUDAMEMCPYCHECK(model->RESparamGPU.d_REStempArray, size, double, status)
status = cudaMemcpy (model->RESparamGPU.d_RESdtempArray, model->RESparamCPU.RESdtempArray, size * sizeof(double), cudaMemcpyHostToDevice) ;
CUDAMEMCPYCHECK(model->RESparamGPU.d_RESdtempArray, size, double, status)
/* INT */
status = cudaMemcpy (model->RESparamGPU.d_REStc1GivenArray, model->RESparamCPU.REStc1GivenArray, size * sizeof(int), cudaMemcpyHostToDevice) ;
CUDAMEMCPYCHECK(model->RESparamGPU.d_REStc1GivenArray, size, int, status)
status = cudaMemcpy (model->RESparamGPU.d_REStc2GivenArray, model->RESparamCPU.REStc2GivenArray, size * sizeof(int), cudaMemcpyHostToDevice) ;
CUDAMEMCPYCHECK(model->RESparamGPU.d_REStc2GivenArray, size, int, status)
status = cudaMemcpy (model->RESparamGPU.d_RESmGivenArray, model->RESparamCPU.RESmGivenArray, size * sizeof(int), cudaMemcpyHostToDevice) ;
CUDAMEMCPYCHECK(model->RESparamGPU.d_RESmGivenArray, size, int, status)
status = cudaMemcpy (model->RESparamGPU.d_RESposNodeArray, model->RESparamCPU.RESposNodeArray, size * sizeof(int), cudaMemcpyHostToDevice) ;
CUDAMEMCPYCHECK(model->RESparamGPU.d_RESposNodeArray, size, int, status)

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

@ -99,41 +99,23 @@ typedef struct sRESinstance {
#ifdef USE_CUSPICE
typedef struct sRESparamCPUstruct {
double *REScpuPointersD [8] ;
#define REStc1Array REScpuPointersD[0]
#define REStc2Array REScpuPointersD[1]
#define RESmArray REScpuPointersD[2]
#define RESconductArray REScpuPointersD[3]
#define REStempArray REScpuPointersD[4]
#define RESdtempArray REScpuPointersD[5]
#define REScurrentArray REScpuPointersD[6]
#define RESgValueArray REScpuPointersD[7]
double *REScpuPointersD [2] ;
#define RESconductArray REScpuPointersD[0]
#define REScurrentArray REScpuPointersD[1]
int *REScpuPointersI [5] ;
#define REStc1GivenArray REScpuPointersI[0]
#define REStc2GivenArray REScpuPointersI[1]
#define RESmGivenArray REScpuPointersI[2]
#define RESposNodeArray REScpuPointersI[3]
#define RESnegNodeArray REScpuPointersI[4]
int *REScpuPointersI [2] ;
#define RESposNodeArray REScpuPointersI[0]
#define RESnegNodeArray REScpuPointersI[1]
} RESparamCPUstruct ;
typedef struct sRESparamGPUstruct {
double *REScudaPointersD [8] ;
#define d_REStc1Array REScudaPointersD[0]
#define d_REStc2Array REScudaPointersD[1]
#define d_RESmArray REScudaPointersD[2]
#define d_RESconductArray REScudaPointersD[3]
#define d_REStempArray REScudaPointersD[4]
#define d_RESdtempArray REScudaPointersD[5]
#define d_REScurrentArray REScudaPointersD[6]
#define d_RESgValueArray REScudaPointersD[7]
int *REScudaPointersI [5] ;
#define d_REStc1GivenArray REScudaPointersI[0]
#define d_REStc2GivenArray REScudaPointersI[1]
#define d_RESmGivenArray REScudaPointersI[2]
#define d_RESposNodeArray REScudaPointersI[3]
#define d_RESnegNodeArray REScudaPointersI[4]
double *REScudaPointersD [2] ;
#define d_RESconductArray REScudaPointersD[0]
#define d_REScurrentArray REScudaPointersD[1]
int *REScudaPointersI [2] ;
#define d_RESposNodeArray REScudaPointersI[0]
#define d_RESnegNodeArray REScudaPointersI[1]
} RESparamGPUstruct ;
#endif

8
src/spicelib/devices/res/restemp.c
View File

@ -53,17 +53,9 @@ REStemp(GENmodel *inModel, CKTcircuit *ckt)
RESupdate_conduct(here, TRUE);
#ifdef USE_CUSPICE
model->RESparamCPU.REStc1GivenArray[i] = here->REStc1Given;
model->RESparamCPU.REStc2GivenArray[i] = here->REStc2Given;
model->RESparamCPU.RESmGivenArray[i] = here->RESmGiven;
model->RESparamCPU.REStc1Array[i] = here->REStc1;
model->RESparamCPU.REStc2Array[i] = here->REStc2;
model->RESparamCPU.RESmArray[i] = here->RESm;
model->RESparamCPU.RESposNodeArray[i] = here->RESposNode;
model->RESparamCPU.RESnegNodeArray[i] = here->RESnegNode;
model->RESparamCPU.RESconductArray[i] = here->RESconduct;
model->RESparamCPU.REStempArray[i] = here->REStemp;
model->RESparamCPU.RESdtempArray[i] = here->RESdtemp;
i++;
#endif