From 6090572fe0dafd69c8bbb3c61830a375da6169b5 Mon Sep 17 00:00:00 2001 From: Francesco Lannutti Date: Thu, 19 Apr 2018 00:53:55 +0200 Subject: [PATCH] Fix the case when there is no RHS computed by device models --- src/spicelib/analysis/CUSPICE/cucktflush.c | 2 +- .../analysis/CUSPICE/cucktnonconupdate.c | 4 +- .../analysis/CUSPICE/cucktrhsoldupdate.c | 6 +- src/spicelib/analysis/CUSPICE/cucktsetup.c | 60 +++---- .../analysis/CUSPICE/cucktstatesupdate.c | 14 +- src/spicelib/analysis/CUSPICE/cucktsystem.c | 25 ++- src/spicelib/analysis/CUSPICE/cuckttrunc.cu | 2 +- src/spicelib/analysis/cktload.c | 57 ++++--- src/spicelib/analysis/cktsetup.c | 152 +++++++++--------- src/spicelib/devices/res/CUSPICE/curesfree.c | 27 ---- src/spicelib/devices/res/CUSPICE/curesload.cu | 23 +-- src/spicelib/devices/res/CUSPICE/curessetup.c | 38 +---- src/spicelib/devices/res/CUSPICE/curestemp.c | 26 +-- src/spicelib/devices/res/resdefs.h | 44 ++--- src/spicelib/devices/res/restemp.c | 8 - 15 files changed, 191 insertions(+), 297 deletions(-) diff --git a/src/spicelib/analysis/CUSPICE/cucktflush.c b/src/spicelib/analysis/CUSPICE/cucktflush.c index b72b06f3b..7dc1c8f9c 100644 --- a/src/spicelib/analysis/CUSPICE/cucktflush.c +++ b/src/spicelib/analysis/CUSPICE/cucktflush.c @@ -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 diff --git a/src/spicelib/analysis/CUSPICE/cucktnonconupdate.c b/src/spicelib/analysis/CUSPICE/cucktnonconupdate.c index 4bdddc9a9..f6f9dded5 100644 --- a/src/spicelib/analysis/CUSPICE/cucktnonconupdate.c +++ b/src/spicelib/analysis/CUSPICE/cucktnonconupdate.c @@ -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) } diff --git a/src/spicelib/analysis/CUSPICE/cucktrhsoldupdate.c b/src/spicelib/analysis/CUSPICE/cucktrhsoldupdate.c index 78feb1f98..bb31c70b7 100644 --- a/src/spicelib/analysis/CUSPICE/cucktrhsoldupdate.c +++ b/src/spicelib/analysis/CUSPICE/cucktrhsoldupdate.c @@ -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) diff --git a/src/spicelib/analysis/CUSPICE/cucktsetup.c b/src/spicelib/analysis/CUSPICE/cucktsetup.c index 320d5bb47..b152a1032 100644 --- a/src/spicelib/analysis/CUSPICE/cucktsetup.c +++ b/src/spicelib/analysis/CUSPICE/cucktsetup.c @@ -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)) ; diff --git a/src/spicelib/analysis/CUSPICE/cucktstatesupdate.c b/src/spicelib/analysis/CUSPICE/cucktstatesupdate.c index d2d26e0cd..7cec763b4 100644 --- a/src/spicelib/analysis/CUSPICE/cucktstatesupdate.c +++ b/src/spicelib/analysis/CUSPICE/cucktstatesupdate.c @@ -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) } diff --git a/src/spicelib/analysis/CUSPICE/cucktsystem.c b/src/spicelib/analysis/CUSPICE/cucktsystem.c index b186bab28..f97713433 100644 --- a/src/spicelib/analysis/CUSPICE/cucktsystem.c +++ b/src/spicelib/analysis/CUSPICE/cucktsystem.c @@ -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) ; diff --git a/src/spicelib/analysis/CUSPICE/cuckttrunc.cu b/src/spicelib/analysis/CUSPICE/cuckttrunc.cu index c3d159968..079b0a042 100644 --- a/src/spicelib/analysis/CUSPICE/cuckttrunc.cu +++ b/src/spicelib/analysis/CUSPICE/cuckttrunc.cu @@ -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 ; diff --git a/src/spicelib/analysis/cktload.c b/src/spicelib/analysis/cktload.c index be3659642..ee04aebdb 100644 --- a/src/spicelib/analysis/cktload.c +++ b/src/spicelib/analysis/cktload.c @@ -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 () ; diff --git a/src/spicelib/analysis/cktsetup.c b/src/spicelib/analysis/cktsetup.c index 379faea1a..9082e4168 100644 --- a/src/spicelib/analysis/cktsetup.c +++ b/src/spicelib/analysis/cktsetup.c @@ -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) diff --git a/src/spicelib/devices/res/CUSPICE/curesfree.c b/src/spicelib/devices/res/CUSPICE/curesfree.c index 6efa9fbc7..d3488907d 100644 --- a/src/spicelib/devices/res/CUSPICE/curesfree.c +++ b/src/spicelib/devices/res/CUSPICE/curesfree.c @@ -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) ; diff --git a/src/spicelib/devices/res/CUSPICE/curesload.cu b/src/spicelib/devices/res/CUSPICE/curesload.cu index ebeac3a4c..072526a28 100644 --- a/src/spicelib/devices/res/CUSPICE/curesload.cu +++ b/src/spicelib/devices/res/CUSPICE/curesload.cu @@ -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] ; } } diff --git a/src/spicelib/devices/res/CUSPICE/curessetup.c b/src/spicelib/devices/res/CUSPICE/curessetup.c index 444389bad..a348f45d7 100644 --- a/src/spicelib/devices/res/CUSPICE/curessetup.c +++ b/src/spicelib/devices/res/CUSPICE/curessetup.c @@ -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) diff --git a/src/spicelib/devices/res/CUSPICE/curestemp.c b/src/spicelib/devices/res/CUSPICE/curestemp.c index b996764ee..5f23321d3 100644 --- a/src/spicelib/devices/res/CUSPICE/curestemp.c +++ b/src/spicelib/devices/res/CUSPICE/curestemp.c @@ -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) diff --git a/src/spicelib/devices/res/resdefs.h b/src/spicelib/devices/res/resdefs.h index 327031cbc..81ed42b67 100644 --- a/src/spicelib/devices/res/resdefs.h +++ b/src/spicelib/devices/res/resdefs.h @@ -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 diff --git a/src/spicelib/devices/res/restemp.c b/src/spicelib/devices/res/restemp.c index 0f9116557..14df9177b 100644 --- a/src/spicelib/devices/res/restemp.c +++ b/src/spicelib/devices/res/restemp.c @@ -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