diff --git a/CMakeLists.txt b/CMakeLists.txt
index 231dc9bab..38201d244 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -741,6 +741,9 @@ if(CUDA_FOUND)
                        gpu_resources
                        gpu_smart_pointers
                        is_gpu_pointer
+                       throw_macros_on_gpu
+                       conditional_storage
+                       blackoilfluidstategpu
                        PROPERTIES LABELS ${gpu_label})
 endif()
 
diff --git a/CMakeLists_files.cmake b/CMakeLists_files.cmake
index f4aa0dec2..89f80f4b4 100644
--- a/CMakeLists_files.cmake
+++ b/CMakeLists_files.cmake
@@ -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_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_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
   if(NOT CONVERT_CUDA_TO_HIP)
@@ -490,6 +494,7 @@ if (HAVE_CUDA)
       tests/gpuistl/test_gpu_ad.cu
       tests/gpuistl/test_gpu_linear_two_phase_material.cu
       tests/gpuistl/test_gpuPvt.cu
+      tests/gpuistl/test_blackoilfluidstategpu
     )
 
     foreach(file ${CU_FILES_NEEDING_RELAXED_CONSTEXPR})
diff --git a/tests/gpuistl/test_blackoilfluidstategpu.cu b/tests/gpuistl/test_blackoilfluidstategpu.cu
new file mode 100644
index 000000000..99b59ee62
--- /dev/null
+++ b/tests/gpuistl/test_blackoilfluidstategpu.cu
@@ -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);
+
+}
\ No newline at end of file
diff --git a/tests/gpuistl/test_conditional_storage.cu b/tests/gpuistl/test_conditional_storage.cu
new file mode 100644
index 000000000..af0dead42
--- /dev/null
+++ b/tests/gpuistl/test_conditional_storage.cu
@@ -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);
+}
\ No newline at end of file
diff --git a/tests/gpuistl/test_throw_macros_on_gpu.cu b/tests/gpuistl/test_throw_macros_on_gpu.cu
new file mode 100644
index 000000000..800ed6149
--- /dev/null
+++ b/tests/gpuistl/test_throw_macros_on_gpu.cu
@@ -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);
+}