From a4a3da9dfdc328585de3e23a06db2908751de882 Mon Sep 17 00:00:00 2001 From: Francesco Lannutti Date: Tue, 19 Sep 2017 21:19:23 +0200 Subject: [PATCH] whitespace, indentation for the next commit --- src/spicelib/analysis/CUSPICE/cucktflush.c | 14 +-- .../analysis/CUSPICE/cucktnonconupdate.c | 8 +- .../analysis/CUSPICE/cucktrhsoldupdate.c | 16 +-- src/spicelib/analysis/CUSPICE/cucktsetup.c | 112 +++++++++--------- .../analysis/CUSPICE/cucktstatesupdate.c | 52 ++++---- src/spicelib/analysis/CUSPICE/cucktsystem.c | 16 +-- src/spicelib/analysis/CUSPICE/cuckttrunc.cu | 90 +++++++------- src/spicelib/analysis/cktload.c | 76 ++++++------ src/spicelib/analysis/cktsetup.c | 42 +++---- 9 files changed, 213 insertions(+), 213 deletions(-) diff --git a/src/spicelib/analysis/CUSPICE/cucktflush.c b/src/spicelib/analysis/CUSPICE/cucktflush.c index fe6bf801f..7a5c21470 100644 --- a/src/spicelib/analysis/CUSPICE/cucktflush.c +++ b/src/spicelib/analysis/CUSPICE/cucktflush.c @@ -35,16 +35,16 @@ cuCKTflush CKTcircuit *ckt ) { - long unsigned int m, mRHS ; + 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)) ; return (OK) ; } diff --git a/src/spicelib/analysis/CUSPICE/cucktnonconupdate.c b/src/spicelib/analysis/CUSPICE/cucktnonconupdate.c index f56377153..7f7658fa2 100644 --- a/src/spicelib/analysis/CUSPICE/cucktnonconupdate.c +++ b/src/spicelib/analysis/CUSPICE/cucktnonconupdate.c @@ -47,8 +47,8 @@ CKTcircuit *ckt { cudaError_t status ; - status = cudaMemcpy (ckt->d_CKTnoncon, &(ckt->CKTnoncon), sizeof(int), cudaMemcpyHostToDevice) ; - CUDAMEMCPYCHECK (ckt->d_CKTnoncon, 1, int, status) + status = cudaMemcpy (ckt->d_CKTnoncon, &(ckt->CKTnoncon), sizeof(int), cudaMemcpyHostToDevice) ; + CUDAMEMCPYCHECK (ckt->d_CKTnoncon, 1, int, status) return (OK) ; } @@ -61,8 +61,8 @@ CKTcircuit *ckt { cudaError_t status ; - status = cudaMemcpy (&(ckt->CKTnoncon), ckt->d_CKTnoncon, sizeof(int), cudaMemcpyDeviceToHost) ; - CUDAMEMCPYCHECK (&(ckt->CKTnoncon), 1, int, status) + status = cudaMemcpy (&(ckt->CKTnoncon), ckt->d_CKTnoncon, sizeof(int), cudaMemcpyDeviceToHost) ; + CUDAMEMCPYCHECK (&(ckt->CKTnoncon), 1, int, status) return (OK) ; } diff --git a/src/spicelib/analysis/CUSPICE/cucktrhsoldupdate.c b/src/spicelib/analysis/CUSPICE/cucktrhsoldupdate.c index 1e07e0281..bdd7801c7 100644 --- a/src/spicelib/analysis/CUSPICE/cucktrhsoldupdate.c +++ b/src/spicelib/analysis/CUSPICE/cucktrhsoldupdate.c @@ -47,8 +47,8 @@ CKTcircuit *ckt { long unsigned int size ; - size = (long unsigned int)(ckt->d_MatrixSize + 1) ; - cudaMemset (ckt->d_CKTrhsOld, 0, size * sizeof(double)) ; + size = (long unsigned int)(ckt->d_MatrixSize + 1) ; + cudaMemset (ckt->d_CKTrhsOld, 0, size * sizeof(double)) ; return (OK) ; } @@ -62,9 +62,9 @@ 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) + 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 +78,9 @@ 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) + size = (long unsigned int)(ckt->d_MatrixSize + 1) ; + status = cudaMemcpy (ckt->CKTrhsOld, ckt->d_CKTrhsOld, size * sizeof(double), cudaMemcpyDeviceToHost) ; + CUDAMEMCPYCHECK (ckt->CKTrhsOld, size, double, status) return (OK) ; } diff --git a/src/spicelib/analysis/CUSPICE/cucktsetup.c b/src/spicelib/analysis/CUSPICE/cucktsetup.c index bd14fc15c..75d997465 100644 --- a/src/spicelib/analysis/CUSPICE/cucktsetup.c +++ b/src/spicelib/analysis/CUSPICE/cucktsetup.c @@ -76,88 +76,88 @@ 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) + /* 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) + /* 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 = 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) + 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) +// ckt->CKTtimeSteps = (double *) malloc (size3 * sizeof(double)) ; + status = cudaMalloc ((void **)&(ckt->d_CKTtimeSteps), size3 * sizeof(double)) ; + CUDAMALLOCCHECK (ckt->d_CKTtimeSteps, size3, double, status) + status = cudaMalloc ((void **)&(ckt->d_CKTtimeStepsOut), size3 * sizeof(double)) ; + CUDAMALLOCCHECK (ckt->d_CKTtimeStepsOut, size3, double, status) return (OK) ; } diff --git a/src/spicelib/analysis/CUSPICE/cucktstatesupdate.c b/src/spicelib/analysis/CUSPICE/cucktstatesupdate.c index 65ba5f6ec..2fb453807 100644 --- a/src/spicelib/analysis/CUSPICE/cucktstatesupdate.c +++ b/src/spicelib/analysis/CUSPICE/cucktstatesupdate.c @@ -47,8 +47,8 @@ CKTcircuit *ckt { long unsigned int size ; - size = (long unsigned int)ckt->CKTnumStates ; - cudaMemset (ckt->d_CKTstate0, 0, size * sizeof(double)) ; + size = (long unsigned int)ckt->CKTnumStates ; + cudaMemset (ckt->d_CKTstate0, 0, size * sizeof(double)) ; return (OK) ; } @@ -62,9 +62,9 @@ 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) + 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 +78,9 @@ 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) + 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 +91,12 @@ cuCKTstate01copy CKTcircuit *ckt ) { - long unsigned int size ; - cudaError_t status ; + 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) return (OK) ; } @@ -110,11 +110,11 @@ 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] ; + 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 ; return (OK) ; } @@ -125,16 +125,16 @@ cuCKTstate123copy CKTcircuit *ckt ) { - long unsigned int size ; - cudaError_t status ; + 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) return (OK) ; } @@ -147,8 +147,8 @@ CKTcircuit *ckt { cudaError_t status ; - status = cudaMemcpy (ckt->d_CKTdeltaOld, ckt->CKTdeltaOld, 7 * sizeof(double), cudaMemcpyHostToDevice) ; - CUDAMEMCPYCHECK (ckt->d_CKTdeltaOld, 7, double, status) + status = cudaMemcpy (ckt->d_CKTdeltaOld, ckt->CKTdeltaOld, 7 * sizeof(double), cudaMemcpyHostToDevice) ; + CUDAMEMCPYCHECK (ckt->d_CKTdeltaOld, 7, double, status) return (OK) ; } diff --git a/src/spicelib/analysis/CUSPICE/cucktsystem.c b/src/spicelib/analysis/CUSPICE/cucktsystem.c index 447a6d654..e512d4a4b 100644 --- a/src/spicelib/analysis/CUSPICE/cucktsystem.c +++ b/src/spicelib/analysis/CUSPICE/cucktsystem.c @@ -48,16 +48,16 @@ CKTcircuit *ckt long unsigned int nz, n ; cudaError_t status ; - nz = (long unsigned int)ckt->CKTmatrix->CKTklunz ; - n = (long unsigned int)ckt->CKTmatrix->CKTkluN ; + nz = (long unsigned int)ckt->CKTmatrix->CKTklunz ; + n = (long unsigned int)ckt->CKTmatrix->CKTkluN ; - /* Copy back the Matrix */ - status = cudaMemcpy (ckt->CKTmatrix->CKTkluAx, ckt->CKTmatrix->d_CKTkluAx, nz * sizeof(double), cudaMemcpyDeviceToHost) ; - CUDAMEMCPYCHECK (ckt->CKTmatrix->CKTkluAx, nz, double, status) + /* Copy back the Matrix */ + status = cudaMemcpy (ckt->CKTmatrix->CKTkluAx, ckt->CKTmatrix->d_CKTkluAx, nz * sizeof(double), cudaMemcpyDeviceToHost) ; + CUDAMEMCPYCHECK (ckt->CKTmatrix->CKTkluAx, nz, double, status) - /* Copy back the RHS */ - status = cudaMemcpy (ckt->CKTrhs, ckt->CKTmatrix->d_CKTrhs, (n + 1) * sizeof(double), cudaMemcpyDeviceToHost) ; - CUDAMEMCPYCHECK (ckt->CKTrhs, (n + 1), double, status) + /* Copy back the RHS */ + status = cudaMemcpy (ckt->CKTrhs, ckt->CKTmatrix->d_CKTrhs, (n + 1) * sizeof(double), cudaMemcpyDeviceToHost) ; + CUDAMEMCPYCHECK (ckt->CKTrhs, (n + 1), double, status) return (OK) ; } diff --git a/src/spicelib/analysis/CUSPICE/cuckttrunc.cu b/src/spicelib/analysis/CUSPICE/cuckttrunc.cu index b215b730f..f12acda2c 100644 --- a/src/spicelib/analysis/CUSPICE/cuckttrunc.cu +++ b/src/spicelib/analysis/CUSPICE/cuckttrunc.cu @@ -31,63 +31,63 @@ cuCKTtrunc CKTcircuit *ckt, double timetemp, double *timeStep ) { - long unsigned int size ; - double timetempGPU ; - int thread_x, thread_y, block_x ; + 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 ; - } else { - *timeStep = timetemp ; - } + /* Final Comparison */ + if (timetempGPU < timetemp) + { + timetemp = timetempGPU ; + } + if (2 * *timeStep < timetemp) + { + *timeStep = 2 * *timeStep ; + } else { + *timeStep = timetemp ; + } return 0 ; } diff --git a/src/spicelib/analysis/cktload.c b/src/spicelib/analysis/cktload.c index 8b4e2a697..7ffbe4016 100644 --- a/src/spicelib/analysis/cktload.c +++ b/src/spicelib/analysis/cktload.c @@ -127,50 +127,50 @@ 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) + /* 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) ; + /* 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) ; - } + 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) ; + /* 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) ; - } + if (cusparseStatus != CUSPARSE_STATUS_SUCCESS) + { + fprintf (stderr, "CUSPARSE RHS Call Error\n") ; + return (E_NOMEM) ; + } - cudaDeviceSynchronize () ; + cudaDeviceSynchronize () ; - status = cuCKTsystemDtoH (ckt) ; - if (status != 0) - return (E_NOMEM) ; + status = cuCKTsystemDtoH (ckt) ; + if (status != 0) + return (E_NOMEM) ; #endif #ifdef XSPICE diff --git a/src/spicelib/analysis/cktsetup.c b/src/spicelib/analysis/cktsetup.c index 8145f6e00..4d3d7366f 100644 --- a/src/spicelib/analysis/cktsetup.c +++ b/src/spicelib/analysis/cktsetup.c @@ -337,30 +337,30 @@ CKTsetup(CKTcircuit *ckt) } #ifdef USE_CUSPICE - ckt->d_MatrixSize = SMPmatSize (ckt->CKTmatrix) ; - status = cuCKTsetup (ckt) ; - if (status != 0) - return (E_NOMEM) ; + 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 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 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