Skip to content

GitLab

Commit e58afb5a authored by Markus Mirz's avatar Markus Mirz
Browse files

MNASolver: fix tasks and adapt GPU version

parent fdb7d12b
Hide whitespace changes
Inline Side-by-side
Include/dpsim/MNASolver.h
View file @ e58afb5a
......@@ -181,7 +181,7 @@ namespace DPsim {
///
class SolveTask : public CPS::Task {
public:
SolveTask(MnaSolver<VarType>& solver, Bool steadyStateInit) :
SolveTask(MnaSolver<VarType>& solver) :
Task(solver.mName + ".Solve"), mSolver(solver) {
for (auto it : solver.mMNAComponents) {
......@@ -203,7 +203,7 @@ namespace DPsim {
///
class SolveTaskHarm : public CPS::Task {
public:
SolveTaskHarm(MnaSolver<VarType>& solver, Bool steadyStateInit, UInt freqIdx) :
SolveTaskHarm(MnaSolver<VarType>& solver, UInt freqIdx) :
Task(solver.mName + ".Solve"), mSolver(solver), mFreqIdx(freqIdx) {
for (auto it : solver.mMNAComponents) {
......
Include/dpsim/MNASolverGpu.h
View file @ e58afb5a
......@@ -48,6 +48,8 @@ namespace DPsim {
void copySystemMatrixToDevice();
/// LU factorization
void LUfactorization();
///
void solve(Real time, Int timeStepCount);
public:
MnaSolverGpu(String name,
......@@ -60,12 +62,12 @@ namespace DPsim {
class SolveTask : public CPS::Task {
public:
SolveTask(MnaSolverGpu<VarType>& solver, Bool steadyStateInit) :
Task(solver.mName + ".Solve"), mSolver(solver), mSteadyStateInit(steadyStateInit) {
SolveTask(MnaSolverGpu<VarType>& solver) :
Task(solver.mName + ".Solve"), mSolver(solver) {
for (auto it : solver.mMNAComponents) {
if (it->template attribute<Matrix>("right_vector")->get().size() != 0) {
if (it->template attribute<Matrix>("right_vector")->get().size() != 0)
mAttributeDependencies.push_back(it->attribute("right_vector"));
}
}
for (auto node : solver.mNodes) {
mModifiedAttributes.push_back(node->attribute("v"));
......@@ -73,11 +75,10 @@ namespace DPsim {
mModifiedAttributes.push_back(solver.attribute("left_vector"));
}
void execute(Real time, Int timeStepCount);
void execute(Real time, Int timeStepCount) { mSolver.solve(time, timeStepCount); }
private:
MnaSolverGpu<VarType>& mSolver;
Bool mSteadyStateInit;
};
class LogTask : public CPS::Task {
......@@ -88,10 +89,10 @@ namespace DPsim {
mModifiedAttributes.push_back(Scheduler::external);
}
void execute(Real time, Int timeStepCount);
void execute(Real time, Int timeStepCount) { mSolver.log(time, timeStepCount); }
private:
MnaSolverGpu<VarType>& mSolver;
};
};
}
\ No newline at end of file
}
Source/MNASolver.cpp
View file @ e58afb5a
......@@ -544,7 +544,7 @@ void MnaSolver<VarType>::steadyStateInitialization() {
tasks.push_back(task);
}
}
tasks.push_back(std::make_shared<MnaSolver<VarType>::SolveTask>(*this, true));
tasks.push_back(std::make_shared<MnaSolver<VarType>::SolveTask>(*this));
sched.resolveDeps(tasks, inEdges, outEdges);
sched.createSchedule(tasks, inEdges, outEdges);
......@@ -612,9 +612,9 @@ Task::List MnaSolver<VarType>::getTasks() {
}
if (mFrequencyParallel) {
for (UInt i = 0; i < mSystem.mFrequencies.size(); i++)
l.push_back(std::make_shared<MnaSolver<VarType>::SolveTaskHarm>(*this, false, i));
l.push_back(std::make_shared<MnaSolver<VarType>::SolveTaskHarm>(*this, i));
} else {
l.push_back(std::make_shared<MnaSolver<VarType>::SolveTask>(*this, false));
l.push_back(std::make_shared<MnaSolver<VarType>::SolveTask>(*this));
l.push_back(std::make_shared<MnaSolver<VarType>::LogTask>(*this));
}
return l;
......
Source/MNASolverGpu.cpp
View file @ e58afb5a
......@@ -52,7 +52,7 @@ void MnaSolverGpu<VarType>::initialize() {
allocateDeviceMemory();
//Copy Systemmatrix to device
copySystemMatrixToDevice();
// Debug logging, whether LU-factorization and copying was successfull
/*DPsim::Matrix mat;
mat.resize(mDeviceCopy.size, mDeviceCopy.size);
......@@ -81,7 +81,7 @@ void MnaSolverGpu<VarType>::allocateDeviceMemory() {
//Workspace
int workSpaceSize = 0;
cusolverStatus_t status = CUSOLVER_STATUS_SUCCESS;
if((status =
if((status =
cusolverDnDgetrf_bufferSize(
mCusolverHandle,
mDeviceCopy.size,
......@@ -118,7 +118,7 @@ void MnaSolverGpu<VarType>::LUfactorization() {
mDeviceCopy.errInfo);
CUDA_ERROR_HANDLER(cudaDeviceSynchronize())
if(status != CUSOLVER_STATUS_SUCCESS) {
std::cerr << "cusolverDnDgetrf() failed (calculating LU-factorization)" << std::endl;
}
......@@ -147,68 +147,63 @@ Task::List MnaSolverGpu<VarType>::getTasks() {
l.push_back(task);
}
}
l.push_back(std::make_shared<MnaSolverGpu<VarType>::SolveTask>(*this, false));
l.push_back(std::make_shared<MnaSolverGpu<VarType>::SolveTask>(*this));
l.push_back(std::make_shared<MnaSolverGpu<VarType>::LogTask>(*this));
return l;
}
template <typename VarType>
void MnaSolverGpu<VarType>::SolveTask::execute(Real time, Int timeStepCount) {
void MnaSolverGpu<VarType>::solve(Real time, Int timeStepCount) {
// Reset source vector
mSolver.mRightSideVector.setZero();
this->mRightSideVector.setZero();
// Add together the right side vector (computed by the components'
// pre-step tasks)
for (const auto &stamp : mSolver.mRightVectorStamps)
mSolver.mRightSideVector += *stamp;
for (const auto &stamp : this->mRightVectorStamps)
this->mRightSideVector += *stamp;
//Copy right vector to device
CUDA_ERROR_HANDLER(cudaMemcpy(mSolver.mDeviceCopy.vector, &mSolver.mRightSideVector(0), mSolver.mDeviceCopy.size * sizeof(Real), cudaMemcpyHostToDevice))
CUDA_ERROR_HANDLER(cudaMemcpy(mDeviceCopy.vector, &this->mRightSideVector(0), mDeviceCopy.size * sizeof(Real), cudaMemcpyHostToDevice))
// Solve
if (mSolver.mSwitchedMatrices.size() > 0) {
if (this->mSwitchedMatrices.size() > 0) {
cusolverStatus_t status = cusolverDnDgetrs(
mSolver.mCusolverHandle,
mCusolverHandle,
CUBLAS_OP_N,
mSolver.mDeviceCopy.size,
mDeviceCopy.size,
1, /* nrhs */
mSolver.mDeviceCopy.matrix,
mSolver.mDeviceCopy.size,
mSolver.mDeviceCopy.pivSeq,
mSolver.mDeviceCopy.vector,
mSolver.mDeviceCopy.size,
mSolver.mDeviceCopy.errInfo);
mDeviceCopy.matrix,
mDeviceCopy.size,
mDeviceCopy.pivSeq,
mDeviceCopy.vector,
mDeviceCopy.size,
mDeviceCopy.errInfo);
CUDA_ERROR_HANDLER(cudaDeviceSynchronize())
if(status != CUSOLVER_STATUS_SUCCESS)
std::cerr << "cusolverDnDgetrs() failed (Solving A*x = b)" << std::endl;
int info;
CUDA_ERROR_HANDLER(cudaMemcpy(&info, mSolver.mDeviceCopy.errInfo, sizeof(int), cudaMemcpyDeviceToHost))
CUDA_ERROR_HANDLER(cudaMemcpy(&info, mDeviceCopy.errInfo, sizeof(int), cudaMemcpyDeviceToHost))
if(0 > info) {
std::cerr << -info << "-th parameter is wrong" << std::endl;
}
}
}
//Copy Solution back
CUDA_ERROR_HANDLER(cudaMemcpy(&mSolver.mLeftSideVector(0), mSolver.mDeviceCopy.vector, mSolver.mDeviceCopy.size * sizeof(Real), cudaMemcpyDeviceToHost))
CUDA_ERROR_HANDLER(cudaMemcpy(&this->mLeftSideVector(0), mDeviceCopy.vector, mDeviceCopy.size * sizeof(Real), cudaMemcpyDeviceToHost))
// TODO split into separate task? (dependent on x, updating all v attributes)
for (UInt nodeIdx = 0; nodeIdx < mSolver.mNumNetNodes; nodeIdx++)
mSolver.mNodes[nodeIdx]->mnaUpdateVoltage(mSolver.mLeftSideVector);
for (UInt nodeIdx = 0; nodeIdx < this->mNumNetNodes; nodeIdx++)
this->mNodes[nodeIdx]->mnaUpdateVoltage(this->mLeftSideVector);
if (!mSteadyStateInit)
mSolver.updateSwitchStatus();
if (!this->mIsInInitialization)
this->updateSwitchStatus();
// Components' states will be updated by the post-step tasks
}
template <typename VarType>
void MnaSolverGpu<VarType>::LogTask::execute(Real time, Int timeStepCount) {
mSolver.log(time);
}
}
template class DPsim::MnaSolverGpu<Real>;
template class DPsim::MnaSolverGpu<Complex>;
\ No newline at end of file
template class DPsim::MnaSolverGpu<Complex>;
Supports Markdown
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment