Add tets for make_view of BOFS

This commit is contained in:
Tobias Meyer Andersen authored and Kjetil Olsen Lye committed 2025-07-31 13:22:39 +02:00
1 parent 7161ef7412
commit c6380bad39
7 files changed
+214 -121

No files matched your search

+8 -3
View File
@@ -53,7 +53,6 @@ template <class T>
GpuBuffer<T>::GpuBuffer(const GpuBuffer<T>& other)
: GpuBuffer(other.m_numberOfElements)
{
// assertHasElements();
assertSameSize(other);
if (m_numberOfElements == 0) {
return;
@@ -201,10 +200,8 @@ template class GpuBuffer<float>;
template class GpuBuffer<int>;
template class GpuBuffer<std::array<double, 3>>;
template class GpuBuffer<std::array<float, 3>>;
template class GpuBuffer<std::array<int, 3>>;
template class GpuBuffer<std::array<double, 9>>;
template class GpuBuffer<std::array<float, 9>>;
template class GpuBuffer<std::array<int, 9>>;
template <class T>
GpuView<T> make_view(GpuBuffer<T>& buf) {
@@ -214,6 +211,10 @@ GpuView<T> make_view(GpuBuffer<T>& buf) {
template GpuView<double> make_view(GpuBuffer<double>&);
template GpuView<float> make_view(GpuBuffer<float>&);
template GpuView<int> make_view(GpuBuffer<int>&);
template GpuView<std::array<double, 3>> make_view(GpuBuffer<std::array<double, 3>>&);
template GpuView<std::array<float, 3>> make_view(GpuBuffer<std::array<float, 3>>&);
template GpuView<std::array<double, 9>> make_view(GpuBuffer<std::array<double, 9>>&);
template GpuView<std::array<float, 9>> make_view(GpuBuffer<std::array<float, 9>>&);
template <class T>
GpuView<const T> make_view(const GpuBuffer<T>& buf) {
@@ -223,5 +224,9 @@ GpuView<const T> make_view(const GpuBuffer<T>& buf) {
template GpuView<const double> make_view(const GpuBuffer<double>&);
template GpuView<const float> make_view(const GpuBuffer<float>&);
template GpuView<const int> make_view(const GpuBuffer<int>&);
template GpuView<const std::array<double, 3>> make_view(const GpuBuffer<std::array<double, 3>>&);
template GpuView<const std::array<float, 3>> make_view(const GpuBuffer<std::array<float, 3>>&);
template GpuView<const std::array<double, 9>> make_view(const GpuBuffer<std::array<double, 9>>&);
template GpuView<const std::array<float, 9>> make_view(const GpuBuffer<std::array<float, 9>>&);
} // namespace Opm::gpuistl
@@ -78,5 +78,9 @@ GpuView<T>::copyToHost(std::vector<T>& data) const
template class GpuView<double>;
template class GpuView<float>;
template class GpuView<int>;
template class GpuView<std::array<double, 3>>;
template class GpuView<std::array<float, 3>>;
template class GpuView<std::array<double, 9>>;
template class GpuView<std::array<float, 9>>;
} // namespace Opm::gpuistl
-32
View File
@@ -34,33 +34,6 @@
namespace Opm::gpuistl
{
template<class ViewType>
class ViewPointer {
public:
__host__ __device__ ViewPointer (ViewType view) : view_(view) {}
__host__ __device__ ViewPointer() : view_() {}
__host__ __device__ ViewType& operator*(){
return view_;
}
__host__ __device__ const ViewType& operator*() const {
return view_;
}
__host__ __device__ ViewType* operator->(){
return &view_;
}
__host__ __device__ const ViewType* operator->() const {
return &view_;
}
private:
ViewType view_;
};
/**
* @brief The GpuView class is provides a view of some data allocated on the GPU
* Essenstially is only stores a pointer and a size.
@@ -134,11 +107,6 @@ public:
*/
~GpuView() = default;
// make a copy of this view that behaves like a pointer
ViewPointer<GpuView<T>> pointer(){
return ViewPointer<GpuView<T>>(*this);
}
/**
* @return the raw pointer to the GPU data
*/
@@ -225,6 +225,7 @@ copyToGPU(const T& value, const std::unique_ptr<T, Deleter>& ptr)
* being compatible with the requirements of the GPU kernels. This is useful when we want to pass
* a smart pointer to a GPU kernel, but we do not want to transfer the ownership of the memory.
*/
template <class T>
class PointerView
{
@@ -252,7 +253,12 @@ public:
return ptr_;
}
OPM_HOST_DEVICE T& operator*() const
OPM_HOST_DEVICE const T& operator*() const
{
return *ptr_;
}
OPM_HOST_DEVICE T& operator*()
{
return *ptr_;
}
@@ -266,6 +272,47 @@ private:
T* ptr_;
};
/**
* @brief Specialization of PointerView for void type
* This is needed beause we cannot have a PointerView<void> specialization
* due to dereferincing a void ptr
*/
template <>
class PointerView<void>
{
public:
PointerView(const PointerView& other) = default;
PointerView(const std::shared_ptr<void>& ptr)
: ptr_(ptr.get())
{
}
template <class Deleter>
PointerView(const std::unique_ptr<void, Deleter>& ptr)
: ptr_(ptr.get())
{
}
PointerView(void* ptr)
: ptr_(ptr)
{
}
OPM_HOST_DEVICE void* get() const
{
return ptr_;
}
OPM_HOST_DEVICE void* operator->() const
{
return ptr_;
}
private:
void* ptr_;
};
template <class T>
PointerView<T>
make_view(const std::shared_ptr<T>& ptr)
@@ -279,5 +326,43 @@ make_view(const std::unique_ptr<T, Deleter>& ptr)
{
return PointerView<T>(ptr);
}
/**
* @brief A value stored with a pointer interface.
* Can be used to wrap objects in GPU kernels that were otherwise stored as pointers
*/
template<class T>
class ValueAsPointer {
public:
using element_type = T;
ValueAsPointer(const T& t) : value(t) {}
OPM_HOST_DEVICE T* operator->() {
return &value;
}
OPM_HOST_DEVICE T* get() {
return &value;
}
OPM_HOST_DEVICE const T* operator->() const {
return &value;
}
OPM_HOST_DEVICE const T* get() const {
return &value;
}
OPM_HOST_DEVICE T& operator*() {
return value;
}
OPM_HOST_DEVICE const T& operator*() const {
return value;
}
private:
T value;
};
} // namespace Opm::gpuistl
#endif
-25
View File
@@ -34,7 +34,6 @@
using GpuViewDouble = ::Opm::gpuistl::GpuView<double>;
using GpuBufferDouble = ::Opm::gpuistl::GpuBuffer<double>;
using ViewPointerDouble = ::Opm::gpuistl::ViewPointer<GpuViewDouble>;
__global__ void useGpuViewOnGPU(GpuViewDouble a, GpuViewDouble b){
b[0] = a.front();
@@ -110,27 +109,3 @@ BOOST_AUTO_TEST_CASE(TestGpuViewOnGPU)
// checks that view[0] = view[2] works
BOOST_CHECK(buf[2] == vecA[0]);
}
__global__ void useViewPointer(ViewPointerDouble ptr, double* out){
out[0] = ptr->front();
out[1] = (*ptr).front();
}
BOOST_AUTO_TEST_CASE(TestGpuViewPointer)
{
auto buf = std::vector<double>({1.0, 2.0, 42.0, 59.9451743, 10.7132692});
auto gpubuf = GpuBufferDouble(buf);
auto gpuview = GpuViewDouble(gpubuf.data(), gpubuf.size());
auto gpuViewPtr = ::Opm::gpuistl::ViewPointer<decltype(gpuview)>(gpuview);
std::vector<double> out(2, -1.0);
auto gpuBufOut = GpuBufferDouble(out);
auto gpuViewOut = GpuViewDouble(gpuBufOut.data(), gpuBufOut.size());
useViewPointer<<<1,1>>>(gpuViewPtr, gpuViewOut.data());
auto outValues = gpuViewOut.asStdVector();
BOOST_CHECK(outValues[0] == buf.front());
BOOST_CHECK(outValues[1] == buf.front());
}
+103 -48
View File
@@ -31,6 +31,8 @@
* This test requires the presence of opm-parser.
*/
#include "config.h"
#include <fmt/format.h>
#include <iostream>
#if !HAVE_ECL_INPUT
#error "The test for the black oil fluid system classes requires ecl input support in opm-common"
@@ -44,6 +46,8 @@
#include <opm/material/fluidsystems/BlackOilFluidSystem.hpp>
#include <opm/material/fluidsystems/BlackOilFluidSystemNonStatic.hpp>
#include <opm/material/fluidstates/BlackOilFluidState.hpp>
#include <opm/material/fluidsystems/blackoilpvt/GasPvtMultiplexer.hpp>
#include <opm/material/fluidsystems/blackoilpvt/WaterPvtMultiplexer.hpp>
#include <opm/material/components/CO2Tables.hpp>
#include <opm/material/densead/Evaluation.hpp>
@@ -60,6 +64,7 @@
#include <opm/simulators/linalg/gpuistl/GpuView.hpp>
#include <opm/simulators/linalg/gpuistl/GpuBuffer.hpp>
#include <opm/simulators/linalg/gpuistl/gpu_smart_pointer.hpp>
#include <opm/simulators/linalg/gpuistl/detail/gpu_safe_call.hpp>
static constexpr const char* deckString1 =
@@ -113,32 +118,37 @@ static constexpr const char* deckString1 =
"TSTEP\n"
"1 /";
using Types = boost::mpl::list<double,Opm::DenseAd::Evaluation<double,2>>;
// using GpuB = ::Opm::gpuistl::GpuBuffer;
// using GpuPointer = ::Opm::gpuistl::ViewPointer;
using GpuB = Opm::gpuistl::GpuBuffer<double>;
using GpuV = Opm::gpuistl::GpuView<double>;
using GpuBufCo2Tables = Opm::CO2Tables<double, GpuB>;
using GpuBufBrineCo2Pvt = Opm::BrineCo2Pvt<double, GpuBufCo2Tables, GpuB>;
using FluidSystem = Opm::BlackOilFluidSystem<double>;
using Evaluation = Opm::DenseAd::Evaluation<double,2>;
using Scalar = typename Opm::MathToolbox<Evaluation>::Scalar;
BOOST_AUTO_TEST_CASE_TEMPLATE(BlackOil, Evaluation, Types)
// checks that we can access value stored as scalar
template <class IndexTraits>
__global__ void getReservoirTemperature(Opm::BlackOilFluidSystemNonStatic<double, IndexTraits, Opm::gpuistl::GpuView, Opm::gpuistl::ValueAsPointer> fs, double* res)
{
// test the black-oil specific methods of BlackOilFluidSystem. The generic methods
// for fluid systems are already tested by the generic test for all fluidsystems.
*res = fs.reservoirTemperature();
}
using Scalar = typename Opm::MathToolbox<Evaluation>::Scalar;
using FluidSystem = Opm::BlackOilFluidSystem<double>;
// checks that we can access value stored in vectors/buffer/views
template <class IndexTraits>
__global__ void getReferenceDensity(Opm::BlackOilFluidSystemNonStatic<double, IndexTraits, Opm::gpuistl::GpuView, Opm::gpuistl::ValueAsPointer> fs, double* res)
{
*res = fs.referenceDensity(0, 0);
}
using GpuB = Opm::gpuistl::GpuBuffer<double>;
using GpuBufCo2Tables = Opm::CO2Tables<double, GpuB>;
using GpuBufBrineCo2Pvt = Opm::BrineCo2Pvt<double, GpuBufCo2Tables, GpuB>;
static constexpr int numPhases = FluidSystem::numPhases;
static constexpr int gasPhaseIdx = FluidSystem::gasPhaseIdx;
static constexpr int oilPhaseIdx = FluidSystem::oilPhaseIdx;
static constexpr int waterPhaseIdx = FluidSystem::waterPhaseIdx;
static constexpr int gasCompIdx = FluidSystem::gasCompIdx;
static constexpr int oilCompIdx = FluidSystem::oilCompIdx;
static constexpr int waterCompIdx = FluidSystem::waterCompIdx;
// check that we can correctly compute values that require using pvt multiplexers
template <class IndexTraits>
__global__ void getReferenceDensityFromGasPvt(Opm::BlackOilFluidSystemNonStatic<double, IndexTraits, Opm::gpuistl::GpuView, Opm::gpuistl::ValueAsPointer> fs, double* res)
{
*res = fs.gasPvt().gasReferenceDensity(0);
}
BOOST_AUTO_TEST_CASE(BlackOilFluidSystemOnGpu)
{
Opm::Parser parser;
auto deck = parser.parseString(deckString1);
@@ -148,38 +158,83 @@ BOOST_AUTO_TEST_CASE_TEMPLATE(BlackOil, Evaluation, Types)
FluidSystem::initFromState(eclState, schedule);
// TODO: get the type of the fluid system gasvpt multiplexer approach
auto& dynamicFluidSystem = FluidSystem::getNonStaticInstance();
using GpuCo2Tables = Opm::CO2Tables<double, ::Opm::gpuistl::GpuBuffer<double>>;
using CpuCo2Tables = Opm::CO2Tables<double>;
using GpuGasPvt = Opm::Co2GasPvt<double, GpuCo2Tables, ::Opm::gpuistl::GpuBuffer<double>>;
using GpuWaterPvt = Opm::BrineCo2Pvt<double, GpuCo2Tables, ::Opm::gpuistl::GpuBuffer<double>>;
// auto cpuGasPvt = dynamicFluidSystem.gasPvt().realGasPvt();
auto gpuCo2GasPvt = ::Opm::gpuistl::copy_to_gpu<::Opm::gpuistl::GpuBuffer, double>(dynamicFluidSystem);
// auto gpuGasPvt = ::Opm::gpuistl::copy_to_gpu<::Opm::gpuistl::GpuBuffer, CO2Tables<double, ::Opm::gpuistl::GpuBuffer<double>>(dynamicFluidSystem->getPvt());
// auto dynamicGpuFluidSystem = ::Opm::gpuistl::copy_to_gpu<::Opm::gpuistl::GpuBuffer>(dynamicFluidSystem, gpuGasPvt);
// auto dynamicGpuFluidSystemView = ::Opm::gpuistl::make_view<::Opm::gpuistl::GpuView, ::Opm::gpuistl::ViewPointer>(dynamicGpuFluidSystem);
auto dynamicGpuFluidSystemBuffer = ::Opm::gpuistl::copy_to_gpu<::Opm::gpuistl::GpuBuffer, double>(dynamicFluidSystem);
auto dynamicGpuFluidSystemView = ::Opm::gpuistl::make_view<::Opm::gpuistl::GpuView, ::Opm::gpuistl::ValueAsPointer>(dynamicGpuFluidSystemBuffer);
// create a parameter cache
// using ParamCache = typename FluidSystem::template ParameterCache<Scalar>;
// ParamCache paramCache(/*maxOilSat=*/0.5, /*regionIdx=*/1);
// BOOST_CHECK_EQUAL(paramCache.regionIndex(), 1);
using ParamCache = typename FluidSystem::template ParameterCache<Scalar>;
ParamCache paramCache(/*maxOilSat=*/0.5, /*regionIdx=*/1);
BOOST_CHECK_EQUAL(paramCache.regionIndex(), 1);
BOOST_CHECK_EQUAL(FluidSystem::numRegions(), 1);
BOOST_CHECK_EQUAL(FluidSystem::numActivePhases(), 2);
// BOOST_CHECK_SMALL(Opm::abs(FluidSystem::reservoirTemperature() - (273.15 + 15.555)), 1e-10);
// BOOST_CHECK_EQUAL(FluidSystem::numRegions(), 1);
// BOOST_CHECK_EQUAL(FluidSystem::numActivePhases(), 2);
double GpuComputedVal = 0.0;
double* gpuComputedValPtr = nullptr;
OPM_GPU_SAFE_CALL(cudaMalloc(&gpuComputedValPtr, sizeof(double)));
getReservoirTemperature<<<1, 1>>>(dynamicGpuFluidSystemView, gpuComputedValPtr);
OPM_GPU_SAFE_CALL(cudaMemcpy(&GpuComputedVal, gpuComputedValPtr, sizeof(double), cudaMemcpyDeviceToHost));
BOOST_CHECK_CLOSE(FluidSystem::reservoirTemperature(), GpuComputedVal, 1e-10);
// BOOST_CHECK(FluidSystem::phaseIsActive(0));
// BOOST_CHECK(FluidSystem::phaseIsActive(2));
getReferenceDensityFromGasPvt<<<1, 1>>>(dynamicGpuFluidSystemView, gpuComputedValPtr);
OPM_GPU_SAFE_CALL(cudaMemcpy(&GpuComputedVal, gpuComputedValPtr, sizeof(double), cudaMemcpyDeviceToHost));
BOOST_CHECK_CLOSE(FluidSystem::gasPvt().gasReferenceDensity(0), GpuComputedVal, 1e-10);
// // make sure that the {oil,gas,water}Pvt() methods are available
// [[maybe_unused]] const auto& gPvt = FluidSystem::gasPvt();
// [[maybe_unused]] const auto& oPvt = FluidSystem::oilPvt();
// [[maybe_unused]] const auto& wPvt = FluidSystem::waterPvt();
getReferenceDensity<<<1, 1>>>(dynamicGpuFluidSystemView, gpuComputedValPtr);
OPM_GPU_SAFE_CALL(cudaMemcpy(&GpuComputedVal, gpuComputedValPtr, sizeof(double), cudaMemcpyDeviceToHost));
BOOST_CHECK_CLOSE(FluidSystem::referenceDensity(0, 0), GpuComputedVal, 1e-10);
OPM_GPU_SAFE_CALL(cudaFree(gpuComputedValPtr));
}
__global__ void useGasPvtMultiplexer(Opm::GasPvtMultiplexer<double, true, GpuV, GpuV, false> gasMultiplexer, double* refTemp)
{
*refTemp = gasMultiplexer.gasReferenceDensity(0);
}
__global__ void useWaterPvtMultiplexer(Opm::WaterPvtMultiplexer<double, true, true, GpuV, GpuV, false> waterMultiplexer, double* refTemp)
{
*refTemp = waterMultiplexer.waterReferenceDensity(0);
}
BOOST_AUTO_TEST_CASE(GasPvtMultiplexer)
{
Opm::Parser parser;
auto deck = parser.parseString(deckString1);
auto python = std::make_shared<Opm::Python>();
Opm::EclipseState eclState(deck);
Opm::Schedule schedule(deck, eclState, python);
FluidSystem::initFromState(eclState, schedule);
// get and compute CPU pvts
auto gaspvt = FluidSystem::gasPvt();
auto cpuGasRefDensity = gaspvt.gasReferenceDensity(0);
auto waterpvt = FluidSystem::waterPvt();
auto cpuWaterRefDensity = waterpvt.waterReferenceDensity(0);
// move pvts to gpu
auto gpuGasPvtBuf = ::Opm::gpuistl::copy_to_gpu<GpuB, GpuB>(gaspvt);
auto gpuGasPvtView = ::Opm::gpuistl::make_view<GpuV, GpuV>(gpuGasPvtBuf);
auto gpuWaterPvtBuf = ::Opm::gpuistl::copy_to_gpu<GpuB, GpuB>(waterpvt);
auto gpuWaterPvtView = ::Opm::gpuistl::make_view<GpuV, GpuV>(gpuWaterPvtBuf);
double gpuRefDensity = 0.0;
double* gpuRefDensityPtr = nullptr;
OPM_GPU_SAFE_CALL(cudaMalloc(&gpuRefDensityPtr, sizeof(double)));
// check that the GPU computed GAS reference density is correct
useGasPvtMultiplexer<<<1, 1>>>(gpuGasPvtView, gpuRefDensityPtr);
OPM_GPU_SAFE_CALL(cudaMemcpy(&gpuRefDensity, gpuRefDensityPtr, sizeof(double), cudaMemcpyDeviceToHost));
BOOST_CHECK_CLOSE(cpuGasRefDensity, gpuRefDensity, 1e-10);
// check that the GPU computed WATER reference density is correct
useWaterPvtMultiplexer<<<1, 1>>>(gpuWaterPvtView, gpuRefDensityPtr);
OPM_GPU_SAFE_CALL(cudaMemcpy(&gpuRefDensity, gpuRefDensityPtr, sizeof(double), cudaMemcpyDeviceToHost));
BOOST_CHECK_CLOSE(cpuWaterRefDensity, gpuRefDensity, 1e-10);
OPM_GPU_SAFE_CALL(cudaFree(gpuRefDensityPtr));
}
+13 -12
View File
@@ -26,8 +26,8 @@
#include <opm/simulators/linalg/gpuistl/detail/gpu_safe_call.hpp>
namespace {
// NOTE: We have to split this into a separate function due
// to some weirdness of hipcc. Note however that this is
// NOTE: We have to split this into a separate function due
// to some weirdness of hipcc. Note however that this is
// the realistic use case of the macro.
__device__ __host__ void functionThatContainsMacros(bool call) {
if (call) {
@@ -44,19 +44,20 @@ __global__ void codeThatContainsMacros(bool call) {
BOOST_AUTO_TEST_CASE(TestKernel)
{
// OPM_GPU_SAFE_CALL(cudaDeviceSynchronize());
// OPM_GPU_SAFE_CALL(cudaGetLastError());
// codeThatContainsMacros<<<1, 1>>>(false);
// OPM_GPU_SAFE_CALL(cudaDeviceSynchronize());
// OPM_GPU_SAFE_CALL(cudaGetLastError());
// TODO: Figure out why this test halts when run in debug mode on AMD GPUs
OPM_GPU_SAFE_CALL(cudaDeviceSynchronize());
OPM_GPU_SAFE_CALL(cudaGetLastError());
codeThatContainsMacros<<<1, 1>>>(false);
OPM_GPU_SAFE_CALL(cudaDeviceSynchronize());
OPM_GPU_SAFE_CALL(cudaGetLastError());
}
BOOST_AUTO_TEST_CASE(TestOutsideKernel)
BOOST_AUTO_TEST_CASE(TestOutsideKernel)
{
// This is to make sure that the macros work outside of kernels but inside a .cu file
// ie. inside a file compiled by nvcc/hipcc.
// BOOST_CHECK_THROW(OPM_THROW(std::runtime_error, "THROW"), std::runtime_error);
// BOOST_CHECK_THROW(OPM_THROW_NOLOG(std::runtime_error, "THROW_NOLOG"), std::runtime_error);
// BOOST_CHECK_THROW(OPM_THROW_PROBLEM(std::runtime_error, "THROW_PROBLEM"), std::runtime_error);
// BOOST_CHECK_THROW(OPM_ERROR_IF(true, "ERROR_IF"), std::logic_error);
BOOST_CHECK_THROW(OPM_THROW(std::runtime_error, "THROW"), std::runtime_error);
BOOST_CHECK_THROW(OPM_THROW_NOLOG(std::runtime_error, "THROW_NOLOG"), std::runtime_error);
BOOST_CHECK_THROW(OPM_THROW_PROBLEM(std::runtime_error, "THROW_PROBLEM"), std::runtime_error);
BOOST_CHECK_THROW(OPM_ERROR_IF(true, "ERROR_IF"), std::logic_error);
}