From 30889b5aa910abdc43b41bbb917f2d51f6b8db64 Mon Sep 17 00:00:00 2001
From: Jacques Xing <k2258928@erc-hpc-login1.create.kcl.ac.uk>
Date: Thu, 18 Jan 2024 09:27:07 +0000
Subject: [PATCH] Some tidy-up for the Mass and Matrix operators
---
Operators/Mass/MassCUDA.hpp | 6 +-----
Operators/Matrix/MatrixCUDA.hpp | 21 ++++-----------------
Operators/Matrix/MatrixStdMat.hpp | 2 --
3 files changed, 5 insertions(+), 24 deletions(-)
diff --git a/Operators/Mass/MassCUDA.hpp b/Operators/Mass/MassCUDA.hpp
index d48b50c..8575a96 100644
--- a/Operators/Mass/MassCUDA.hpp
+++ b/Operators/Mass/MassCUDA.hpp
@@ -18,8 +18,6 @@ public:
Field<TData, FieldState::Phys>::template create<MemoryRegionCUDA>(
GetBlockAttributes(FieldState::Phys, expansionList)))
{
- auto nCoord = this->m_expansionList->GetCoordim(0);
-
m_BwdTransOp = BwdTrans<>::create(this->m_expansionList, "CUDA");
m_IProductWRTBaseOp =
IProductWRTBase<>::create(this->m_expansionList, "CUDA");
@@ -44,12 +42,10 @@ public:
static std::string className;
-private:
+protected:
std::shared_ptr<OperatorBwdTrans<TData>> m_BwdTransOp;
std::shared_ptr<OperatorIProductWRTBase<TData>> m_IProductWRTBaseOp;
Field<TData, FieldState::Phys> m_field;
- size_t m_blockSize = 32;
- size_t m_gridSize;
};
} // namespace Nektar::Operators::detail
diff --git a/Operators/Matrix/MatrixCUDA.hpp b/Operators/Matrix/MatrixCUDA.hpp
index 40dfd5f..40910be 100644
--- a/Operators/Matrix/MatrixCUDA.hpp
+++ b/Operators/Matrix/MatrixCUDA.hpp
@@ -1,3 +1,5 @@
+#pragma once
+
#include "MemoryRegionCUDA.hpp"
#include "Operators/Matrix/MatrixCUDAKernels.cuh"
#include "Operators/OperatorHelper.cuh"
@@ -6,11 +8,6 @@
namespace Nektar::Operators::detail
{
-template <typename TData>
-void MatrixKernel(const size_t gridSize, const size_t blockSize,
- const size_t numPts, const size_t nElmts, const size_t size,
- const TData *mat, const TData *in, TData *out);
-
// Matrix implementation
template <typename TData, FieldState TFieldState>
class OperatorMatrixImpl<TData, TFieldState, ImplCUDA>
@@ -34,7 +31,6 @@ public:
~OperatorMatrixImpl()
{
cudaFree(m_matrix);
- m_matrix = nullptr;
}
size_t size()
@@ -99,8 +95,8 @@ public:
m_gridSize = nElmts / m_blockSize;
m_gridSize += (nElmts % m_blockSize == 0) ? 0 : 1;
- MatrixKernel(m_gridSize, m_blockSize, numPts, nElmts, m_size,
- m_matrix, inptr, outptr);
+ MatrixKernel<<<m_gridSize, m_blockSize>>>(numPts, nElmts, m_size,
+ m_matrix, inptr, outptr);
// Increment pointer and index for next element type.
inptr += numPts * nElmts;
@@ -126,13 +122,4 @@ private:
size_t m_blockSize = 32;
};
-template <typename TData>
-void MatrixKernel(const size_t gridSize, const size_t blockSize,
- const size_t numPts, const size_t nElmts, const size_t size,
- const TData *mat, const TData *in, TData *out)
-{
- MatrixKernel<TData>
- <<<gridSize, blockSize>>>(numPts, nElmts, size, mat, in, out);
-}
-
} // namespace Nektar::Operators::detail
diff --git a/Operators/Matrix/MatrixStdMat.hpp b/Operators/Matrix/MatrixStdMat.hpp
index 10c10f7..fd045a1 100644
--- a/Operators/Matrix/MatrixStdMat.hpp
+++ b/Operators/Matrix/MatrixStdMat.hpp
@@ -1,8 +1,6 @@
#pragma once
#include <numeric>
-
-#include "Field.hpp"
#include "Operators/OperatorMatrix.hpp"
namespace Nektar::Operators::detail
--
GitLab