OilfieldToolsNet/opm-simulators
4
0
Fork 0
mirror of https://github.com/OPM/opm-simulators.git synced 2025-02-25 18:55:30 -06:00
Code Issues Packages Projects Releases Wiki Activity

Merge c39a4750a6 into c015df56f9

Browse Source
This commit is contained in:
Kjetil Olsen Lye 2025-02-13 20:15:15 +01:00 committed by GitHub
commit fd0afa00ac
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
5 changed files with 426 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
CMakeLists.txt
View File

@ -741,6 +741,9 @@ if(CUDA_FOUND)
gpu_resources gpu_resources
gpu_smart_pointers gpu_smart_pointers
is_gpu_pointer is_gpu_pointer
throw_macros_on_gpu
conditional_storage
blackoilfluidstategpu
PROPERTIES LABELS ${gpu_label}) PROPERTIES LABELS ${gpu_label})
endif() endif()

5
CMakeLists_files.cmake
View File

@ -483,6 +483,10 @@ if (HAVE_CUDA)
ADD_CUDA_OR_HIP_FILE(TEST_SOURCE_FILES tests test_gpu_smart_pointers.cu) ADD_CUDA_OR_HIP_FILE(TEST_SOURCE_FILES tests test_gpu_smart_pointers.cu)
ADD_CUDA_OR_HIP_FILE(TEST_SOURCE_FILES tests test_gpu_resources.cu) ADD_CUDA_OR_HIP_FILE(TEST_SOURCE_FILES tests test_gpu_resources.cu)
ADD_CUDA_OR_HIP_FILE(TEST_SOURCE_FILES tests test_is_gpu_pointer.cpp) ADD_CUDA_OR_HIP_FILE(TEST_SOURCE_FILES tests test_is_gpu_pointer.cpp)
ADD_CUDA_OR_HIP_FILE(TEST_SOURCE_FILES tests test_throw_macros_on_gpu.cu)
ADD_CUDA_OR_HIP_FILE(TEST_SOURCE_FILES tests test_blackoilfluidstategpu.cu)
ADD_CUDA_OR_HIP_FILE(TEST_SOURCE_FILES tests test_conditional_storage.cu)
# for loop providing the flag --expt-relaxed-constexpr to fix some cuda issues with constexpr # for loop providing the flag --expt-relaxed-constexpr to fix some cuda issues with constexpr
if(NOT CONVERT_CUDA_TO_HIP) if(NOT CONVERT_CUDA_TO_HIP)
@ -490,6 +494,7 @@ if (HAVE_CUDA)
tests/gpuistl/test_gpu_ad.cu tests/gpuistl/test_gpu_ad.cu
tests/gpuistl/test_gpu_linear_two_phase_material.cu tests/gpuistl/test_gpu_linear_two_phase_material.cu
tests/gpuistl/test_gpuPvt.cu tests/gpuistl/test_gpuPvt.cu
tests/gpuistl/test_blackoilfluidstategpu
) )
foreach(file ${CU_FILES_NEEDING_RELAXED_CONSTEXPR}) foreach(file ${CU_FILES_NEEDING_RELAXED_CONSTEXPR})

186
tests/gpuistl/test_blackoilfluidstategpu.cu Normal file
View File

@ -0,0 +1,186 @@
/*
Copyright 2025 Equinor ASA
This file is part of the Open Porous Media project (OPM).
OPM is free software: you can redistribute it and/or modify
it under the terms of the GNU General Public License as published by
the Free Software Foundation, either version 3 of the License, or
(at your option) any later version.
OPM is distributed in the hope that it will be useful,
but WITHOUT ANY WARRANTY; without even the implied warranty of
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
GNU General Public License for more details.
You should have received a copy of the GNU General Public License
along with OPM. If not, see <http://www.gnu.org/licenses/>.
*/
#include <config.h>
#define BOOST_TEST_MODULE TestBlackOilFluidStateGPU
#include <cuda.h>
#include <cuda_runtime.h>
#include <boost/test/unit_test.hpp>
#include <opm/material/fluidstates/BlackOilFluidState.hpp>
#include <opm/simulators/linalg/gpuistl/detail/gpu_safe_call.hpp>
#include <opm/material/common/HasMemberGeneratorMacros.hpp>
#include <opm/simulators/linalg/gpuistl/gpu_smart_pointer.hpp>
namespace
{
template<class ScalarT>
struct DummyFluidSystem {
static constexpr auto numPhases = 3u;
static constexpr auto numComponents = 1;
static constexpr auto waterPhaseIdx = 0;
static constexpr auto oilPhaseIdx = 1;
static constexpr auto gasPhaseIdx = 2;
static constexpr auto waterCompIdx = 0;
static constexpr auto oilCompIdx = 0;
static constexpr auto gasCompIdx = 0;
static auto reservoirTemperature(int) { return ScalarT{ 0.0 }; }
static auto enthalpyEqualEnergy() { return true; }
static auto molarMass(int, int) { return ScalarT{ 0.0 }; }
template<class T>
static auto viscosity(const T&, int, int) { return ScalarT{ 0.0 }; }
static auto convertRsToXoG(ScalarT, int) { return ScalarT{ 0.0 }; }
static auto convertRvToXg0(ScalarT, int) { return ScalarT{ 0.0 }; }
static auto convertXoGToRs(ScalarT, int) { return ScalarT{ 0.0 }; }
template<class T>
static auto fugacityCoefficient(const T&, int, int, int) { return ScalarT{ 0.0 }; }
static auto activeToCanonicalPhaseIdx(int) { return 0u; }
static auto canonicalToActivePhaseIdx(int) { return 0u; }
};
template<class ScalarT>
struct DummyFluidSystemDynamic {
static constexpr auto numPhases = 3u;
static constexpr auto numComponents = 1;
static constexpr auto waterPhaseIdx = 0;
static constexpr auto oilPhaseIdx = 1;
static constexpr auto gasPhaseIdx = 2;
static constexpr auto waterCompIdx = 0;
static constexpr auto oilCompIdx = 0;
static constexpr auto gasCompIdx = 0;
OPM_HOST_DEVICE auto reservoirTemperature(int) const { return ScalarT{ 0.0 }; }
OPM_HOST_DEVICE auto enthalpyEqualEnergy() const { return true; }
OPM_HOST_DEVICE auto molarMass(int, int) const { return ScalarT{ 0.0 }; }
template<class T>
OPM_HOST_DEVICE auto viscosity(const T&, int, int) const { return ScalarT{ someVariable }; }
OPM_HOST_DEVICE auto convertRsToXoG(ScalarT, int) const { return ScalarT{ 0.0 }; }
OPM_HOST_DEVICE auto convertRvToXg0(ScalarT, int) const { return ScalarT{ 0.0 }; }
OPM_HOST_DEVICE auto convertXoGToRs(ScalarT, int) const { return ScalarT{ 0.0 }; }
template<class T>
OPM_HOST_DEVICE auto fugacityCoefficient(const T&, int, int, int) const { return ScalarT{ 0.0 }; }
OPM_HOST_DEVICE auto activeToCanonicalPhaseIdx(int) const { return 0u; }
OPM_HOST_DEVICE auto canonicalToActivePhaseIdx(int) const { return 0u; }
double someVariable = 43.2;
};
template <class FluidState>
__global__ void kernelCreatingBlackoilFluidState() {
FluidState state;
}
template<class FluidState, class FluidSystem>
__global__ void kernelCreatingBlackoilFluidStateDynamic() {
FluidSystem system;
FluidState state(system);
}
template<class FluidState>
__global__ void kernelSetAndGetTotalSaturation(double saturation, double* readSaturation) {
FluidState state;
state.setTotalSaturation(saturation);
*readSaturation = state.totalSaturation();
}
template<class FluidState>
__global__ void getPressure(Opm::gpuistl::PointerView<FluidState> input, std::array<double, 3>* output) {
for (int i = 0; i < 3; ++i) {
(*output)[i] = input->pressure(i);
}
}
template<class FluidState, class FluidSystem>
__global__ void getViscosity(FluidSystem input, double* output) {
FluidState state(input);
*output = state.viscosity(0);
}
} // namespace
using ScalarT = double;
using FluidState = Opm::BlackOilFluidState<ScalarT, DummyFluidSystem<ScalarT>>;
using FluidStateDynamic = Opm::BlackOilFluidState<ScalarT, DummyFluidSystemDynamic<ScalarT>>;
BOOST_AUTO_TEST_CASE(TestCreation)
{
kernelCreatingBlackoilFluidState<FluidState><<<1, 1>>>();
OPM_GPU_SAFE_CALL(cudaDeviceSynchronize());
OPM_GPU_SAFE_CALL(cudaGetLastError());
}
BOOST_AUTO_TEST_CASE(TestSaturation)
{
const double saturation = 0.5;
auto saturationRead = Opm::gpuistl::make_gpu_unique_ptr<double>(0.0);
kernelSetAndGetTotalSaturation<FluidState><<<1, 1>>>(saturation, saturationRead.get());
auto saturationFromGPU = Opm::gpuistl::copyFromGPU(saturationRead);
BOOST_CHECK_EQUAL(saturationFromGPU, saturation);
OPM_GPU_SAFE_CALL(cudaDeviceSynchronize());
OPM_GPU_SAFE_CALL(cudaGetLastError());
}
BOOST_AUTO_TEST_CASE(TestPressure)
{
FluidState state;
state.setPressure(0, 1.0);
state.setPressure(1, 2.0);
state.setPressure(2, 3.0);
auto stateGPU = Opm::gpuistl::make_gpu_unique_ptr<FluidState>(state);
auto output = Opm::gpuistl::make_gpu_unique_ptr<std::array<double, 3>>();
getPressure<<<1, 1>>>(Opm::gpuistl::make_view(stateGPU), output.get());
auto outputCPU = Opm::gpuistl::copyFromGPU(output);
BOOST_CHECK_EQUAL(1.0, outputCPU[0]);
BOOST_CHECK_EQUAL(2.0, outputCPU[1]);
BOOST_CHECK_EQUAL(3.0, outputCPU[2]);
}
BOOST_AUTO_TEST_CASE(TestDynamicCreation)
{
kernelCreatingBlackoilFluidStateDynamic<FluidStateDynamic, DummyFluidSystemDynamic<ScalarT>><<<1, 1>>>();
OPM_GPU_SAFE_CALL(cudaDeviceSynchronize());
OPM_GPU_SAFE_CALL(cudaGetLastError());
}
BOOST_AUTO_TEST_CASE(TestPassByValueToGPUDynamic)
{
DummyFluidSystemDynamic<ScalarT> system;
system.someVariable = 1234;
auto output = Opm::gpuistl::make_gpu_unique_ptr<double>();
getViscosity<FluidStateDynamic><<<1, 1>>>(system, output.get());
auto outputCPU = Opm::gpuistl::copyFromGPU(output);
BOOST_CHECK_EQUAL(1234, outputCPU);
}

123
tests/gpuistl/test_conditional_storage.cu Normal file
View File

@ -0,0 +1,123 @@
/*
Copyright 2025 Equinor ASA
This file is part of the Open Porous Media project (OPM).
OPM is free software: you can redistribute it and/or modify
it under the terms of the GNU General Public License as published by
the Free Software Foundation, either version 3 of the License, or
(at your option) any later version.
OPM is distributed in the hope that it will be useful,
but WITHOUT ANY WARRANTY; without even the implied warranty of
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
GNU General Public License for more details.
You should have received a copy of the GNU General Public License
along with OPM. If not, see <http://www.gnu.org/licenses/>.
*/
#include <config.h>
#define BOOST_TEST_MODULE TestConditionalStorageGPU
#include <cuda.h>
#include <cuda_runtime.h>
#include <boost/test/unit_test.hpp>
#include <opm/material/common/ConditionalStorage.hpp>
#include <opm/simulators/linalg/gpuistl/detail/gpu_safe_call.hpp>
#include <opm/simulators/linalg/gpuistl/gpu_smart_pointer.hpp>
namespace {
template<bool enabled, class T>
__global__ void createCondtionalStorage() {
// just make sure we can run the constructor
Opm::ConditionalStorage<enabled, T> someStorage;
Opm::ConditionalStorage<enabled, T> other = someStorage;
other = someStorage;
}
template<class T>
__global__ void testEnabledStorage(Opm::ConditionalStorage<true, T> storage, T* output) {
output[0] = *storage;
}
template<class T, class S>
__global__ void testEnabledStorageArrow(Opm::ConditionalStorage<true, T> storage, S* output) {
output[0] = storage->someFunc();
}
template<class T>
__global__ void testEnabledStorage(Opm::ConditionalStorage<true, T>* storage, T* output) {
output[0] = **storage;
}
template<class T, class S>
__global__ void testEnabledStorageArrow(Opm::ConditionalStorage<true, T>* storage, S* output) {
output[0] = (*storage)->someFunc();
}
struct SomeStruct {
OPM_HOST_DEVICE int someFunc() {
return 123;
}
};
} // namespace
BOOST_AUTO_TEST_CASE(TestRunConstructor)
{
OPM_GPU_SAFE_CALL(cudaDeviceSynchronize());
OPM_GPU_SAFE_CALL(cudaGetLastError());
createCondtionalStorage<true, double><<<1, 1>>>();
OPM_GPU_SAFE_CALL(cudaDeviceSynchronize());
OPM_GPU_SAFE_CALL(cudaGetLastError());
createCondtionalStorage<false, double><<<1, 1>>>();
OPM_GPU_SAFE_CALL(cudaDeviceSynchronize());
OPM_GPU_SAFE_CALL(cudaGetLastError());
}
BOOST_AUTO_TEST_CASE(TestEnabledStoragePointer)
{
using namespace Opm;
using CS = ConditionalStorage<true, double>;
auto storage = Opm::gpuistl::make_gpu_unique_ptr<CS>(CS(32.2));
auto numberFromGPU = Opm::gpuistl::make_gpu_unique_ptr<double>(0.0);
OPM_GPU_SAFE_CALL(cudaGetLastError());
testEnabledStorage<<<1, 1>>>(storage.get(), numberFromGPU.get());
OPM_GPU_SAFE_CALL(cudaDeviceSynchronize());
OPM_GPU_SAFE_CALL(cudaGetLastError());
auto number = Opm::gpuistl::copyFromGPU(numberFromGPU);
BOOST_CHECK_EQUAL(32.2, number);
auto numberFromGPUFromCall = Opm::gpuistl::make_gpu_unique_ptr<int>(0);
auto storageSomeStruct = Opm::gpuistl::make_gpu_unique_ptr<ConditionalStorage<true, SomeStruct>>(ConditionalStorage<true, SomeStruct>());
testEnabledStorageArrow<<<1, 1>>>(storageSomeStruct.get(), numberFromGPUFromCall.get());
OPM_GPU_SAFE_CALL(cudaDeviceSynchronize());
OPM_GPU_SAFE_CALL(cudaGetLastError());
auto numberFromCall = Opm::gpuistl::copyFromGPU(numberFromGPUFromCall);
BOOST_CHECK_EQUAL(123, numberFromCall);
}
BOOST_AUTO_TEST_CASE(TestEnabledStorageCopy)
{
using namespace Opm;
using CS = ConditionalStorage<true, double>;
auto storage = CS(32.2);
auto numberFromGPU = Opm::gpuistl::make_gpu_unique_ptr<double>(0.0);
testEnabledStorage<<<1, 1>>>(storage, numberFromGPU.get());
OPM_GPU_SAFE_CALL(cudaDeviceSynchronize());
OPM_GPU_SAFE_CALL(cudaGetLastError());
auto number = Opm::gpuistl::copyFromGPU(numberFromGPU);
BOOST_CHECK_EQUAL(32.2, number);
auto numberFromGPUFromCall = Opm::gpuistl::make_gpu_unique_ptr<int>(0);
auto storageSomeStruct = ConditionalStorage<true, SomeStruct>();
testEnabledStorageArrow<<<1, 1>>>(storageSomeStruct, numberFromGPUFromCall.get());
OPM_GPU_SAFE_CALL(cudaDeviceSynchronize());
OPM_GPU_SAFE_CALL(cudaGetLastError());
auto numberFromCall = Opm::gpuistl::copyFromGPU(numberFromGPUFromCall);
BOOST_CHECK_EQUAL(123, numberFromCall);
}

109
tests/gpuistl/test_throw_macros_on_gpu.cu Normal file
View File

@ -0,0 +1,109 @@
/*
Copyright 2025 Equinor ASA
This file is part of the Open Porous Media project (OPM).
OPM is free software: you can redistribute it and/or modify
it under the terms of the GNU General Public License as published by
the Free Software Foundation, either version 3 of the License, or
(at your option) any later version.
OPM is distributed in the hope that it will be useful,
but WITHOUT ANY WARRANTY; without even the implied warranty of
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
GNU General Public License for more details.
You should have received a copy of the GNU General Public License
along with OPM. If not, see <http://www.gnu.org/licenses/>.
*/
#include <boost/test/tools/old/interface.hpp>
#include <config.h>
#include <stdexcept>
#define BOOST_TEST_MODULE TestThrowMacrosOnGPU
#include <cuda.h>
#include <cuda_runtime.h>
#include <boost/test/unit_test.hpp>
#include <opm/common/ErrorMacros.hpp>
#include <opm/simulators/linalg/gpuistl/detail/gpu_safe_call.hpp>
namespace {
__global__ void codeThatContainsMacros(bool call) {
if (call) {
OPM_THROW(std::logic_error, "Something went wrong");
OPM_THROW_NOLOG(std::logic_error, "Something went wrong");
OPM_THROW_PROBLEM(std::logic_error, "Something went wrong");
}
OPM_ERROR_IF(!call, "Something went horribly wrong");
}
// TODO: Check if this is better on HIP
#if 0 // I am leaving this here to show that this is not possible due to limitations in CUDA
// the assert will indeed cause an error, but the CUDA context will be broken for
// the rest of the lifetime of the process, see
// https://forums.developer.nvidia.com/t/how-to-clear-cuda-errors/296393/5
__global__ void checkThrow() {
OPM_THROW(std::logic_error, "Something went wrong");
}
__global__ void checkThrowNoLog() {
OPM_THROW_NOLOG(std::logic_error, "Something went wrong");
}
__global__ void checkThrowProblem() {
OPM_THROW_PROBLEM(std::logic_error, "Something went wrong");
}
__global__ void checkErrorIf() {
OPM_ERROR_IF(true, "Something went horribly wrong");
}
#endif
}
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());
#if 0 // I am leaving this here to show that this is not possible due to limitations in CUDA
// the assert will indeed cause an error, but the CUDA context will be broken for
// the rest of the lifetime of the process, see
// https://forums.developer.nvidia.com/t/how-to-clear-cuda-errors/296393/5
codeThatContainsMacros<<<1, 1>>>(true);
// Make sure this actually throws
BOOST_CHECK_THROW(OPM_GPU_SAFE_CALL(cudaDeviceSynchronize()), std::runtime_error);
OPM_GPU_SAFE_CALL(cudaDeviceReset());
OPM_GPU_SAFE_CALL(cudaGetLastError());
checkThrow<<<1, 1>>>();
BOOST_CHECK_THROW(OPM_GPU_SAFE_CALL(cudaDeviceSynchronize()), std::runtime_error);
OPM_GPU_SAFE_CALL(cudaDeviceReset());
OPM_GPU_SAFE_CALL(cudaGetLastError());
checkThrowNoLog<<<1, 1>>>();
BOOST_CHECK_THROW(OPM_GPU_SAFE_CALL(cudaDeviceSynchronize()), std::runtime_error);
OPM_GPU_SAFE_CALL(cudaDeviceReset());
OPM_GPU_SAFE_CALL(cudaGetLastError());
checkThrowProblem<<<1, 1>>>();
BOOST_CHECK_THROW(OPM_GPU_SAFE_CALL(cudaDeviceSynchronize()), std::runtime_error);
OPM_GPU_SAFE_CALL(cudaDeviceReset());
OPM_GPU_SAFE_CALL(cudaGetLastError());
checkErrorIf<<<1, 1>>>();
BOOST_CHECK_THROW(OPM_GPU_SAFE_CALL(cudaDeviceSynchronize()), std::runtime_error);
OPM_GPU_SAFE_CALL(cudaDeviceReset());
OPM_GPU_SAFE_CALL(cudaGetLastError());
#endif
}
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);
}