diff --git a/CMakeLists.txt b/CMakeLists.txt
index 89f0cca41..9b09042b2 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -653,6 +653,8 @@ if(CUDA_FOUND)
                        solver_adapter
                        GpuBuffer
                        GpuView
+                       gpu_ad
+                       gpu_linear_two_phase_material
                        PROPERTIES LABELS ${gpu_label})
 endif()
 
diff --git a/CMakeLists_files.cmake b/CMakeLists_files.cmake
index fe4e8ec74..4e5cd5e5c 100644
--- a/CMakeLists_files.cmake
+++ b/CMakeLists_files.cmake
@@ -448,6 +448,20 @@ if (HAVE_CUDA)
   ADD_CUDA_OR_HIP_FILE(TEST_SOURCE_FILES tests test_cuVector_operations.cpp)
   ADD_CUDA_OR_HIP_FILE(TEST_SOURCE_FILES tests test_safe_conversion.cpp)
   ADD_CUDA_OR_HIP_FILE(TEST_SOURCE_FILES tests test_solver_adapter.cpp)
+  ADD_CUDA_OR_HIP_FILE(TEST_SOURCE_FILES tests test_gpu_ad.cu)
+  ADD_CUDA_OR_HIP_FILE(TEST_SOURCE_FILES tests test_gpu_linear_two_phase_material.cu)
+
+  # for loop providing the flag --expt-relaxed-constexpr to fix some cuda issues with constexpr
+  if(NOT CONVERT_CUDA_TO_HIP)
+    set(CU_FILES_NEEDING_RELAXED_CONSTEXPR
+      tests/gpuistl/test_gpu_ad.cu
+      tests/gpuistl/gpu_linear_two_phase_material.cu
+    )
+
+    foreach(file ${LIST_OF_FILES})
+      set_source_file_properties(${file} "--expt-relaxed-constexpr")
+    endforeach()
+  endif()
 endif()
 
 if(USE_BDA_BRIDGE)
diff --git a/tests/gpuistl/test_gpu_ad.cu b/tests/gpuistl/test_gpu_ad.cu
new file mode 100644
index 000000000..2bb0075b0
--- /dev/null
+++ b/tests/gpuistl/test_gpu_ad.cu
@@ -0,0 +1,52 @@
+/*
+  Copyright 2024 SINTEF AS
+  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 TestGpuAD
+
+#include <boost/test/unit_test.hpp>
+#include <opm/material/densead/Evaluation.hpp>
+#include <opm/simulators/linalg/gpuistl/detail/gpu_safe_call.hpp>
+#include <cuda_runtime.h>
+
+namespace{
+__global__ void instansiate_ad_object(Opm::DenseAd::Evaluation<float, 3>* adObj, double value){
+    *adObj = Opm::DenseAd::Evaluation<float, 3>(value, 0);
+}
+
+//TODO add more comprenehsive AD tests
+
+} // END EMPTY NAMESPACE
+
+
+BOOST_AUTO_TEST_CASE(TestInstansiateADObject)
+{
+    using Evaluation = Opm::DenseAd::Evaluation<float, 3>;
+    double testValue = 123.456;
+    Evaluation cpuMadeAd = Evaluation(testValue, 0);
+
+    Evaluation gpuMadeAd[1]; // allocate space for one more AD object on the CPU
+    Evaluation *d_ad;
+
+    // allocate space on GPU, run kernel, and move results back to the CPU
+    OPM_GPU_SAFE_CALL(cudaMalloc(&d_ad, sizeof(Evaluation)));
+    instansiate_ad_object<<<1,1>>>(d_ad, testValue);
+    OPM_GPU_SAFE_CALL(cudaDeviceSynchronize());
+    OPM_GPU_SAFE_CALL(cudaMemcpy(&gpuMadeAd, d_ad, sizeof(Evaluation), cudaMemcpyDeviceToHost));
+    OPM_GPU_SAFE_CALL(cudaFree(d_ad));
+
+    // Check that the object made in a GPU kernel is equivalent to that made on the CPU
+    BOOST_CHECK(cpuMadeAd == gpuMadeAd[0]);
+}
diff --git a/tests/gpuistl/test_gpu_linear_two_phase_material.cu b/tests/gpuistl/test_gpu_linear_two_phase_material.cu
new file mode 100644
index 000000000..72c99e65a
--- /dev/null
+++ b/tests/gpuistl/test_gpu_linear_two_phase_material.cu
@@ -0,0 +1,102 @@
+/*
+  Copyright 2024 SINTEF AS
+  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 TestGpuLinearTwoPhaseMaterial
+
+#include <cuda_runtime.h>
+
+#include <boost/test/unit_test.hpp>
+
+#include <dune/common/fvector.hh>
+#include <dune/istl/bvector.hh>
+
+#include <opm/material/common/MathToolbox.hpp>
+#include <opm/material/densead/Evaluation.hpp>
+#include <opm/material/densead/Math.hpp>
+#include <opm/material/fluidmatrixinteractions/MaterialTraits.hpp>
+#include <opm/material/fluidmatrixinteractions/PiecewiseLinearTwoPhaseMaterial.hpp>
+#include <opm/material/fluidmatrixinteractions/PiecewiseLinearTwoPhaseMaterialParams.hpp>
+
+#include <opm/simulators/linalg/gpuistl/GpuView.hpp>
+#include <opm/simulators/linalg/gpuistl/GpuBuffer.hpp>
+#include <opm/simulators/linalg/gpuistl/detail/gpu_safe_call.hpp>
+
+#include <array>
+#include <algorithm>
+#include <type_traits>
+#include <vector>
+#include <iostream>
+
+  // these types are taken from Norne
+  using Scalar = float;
+  using ValueVector = std::vector<Scalar>;
+  using GPUBuffer = Opm::gpuistl::GpuBuffer<Scalar>;
+  using GPUView = Opm::gpuistl::GpuView<const Scalar>;
+
+  using TraitsT = Opm::TwoPhaseMaterialTraits<Scalar, 1, 2>;
+  using CPUParams = Opm::PiecewiseLinearTwoPhaseMaterialParams<TraitsT>;
+  using constGPUBufferParams = Opm::PiecewiseLinearTwoPhaseMaterialParams<TraitsT, const GPUBuffer>;
+  using GPUBufferParams = Opm::PiecewiseLinearTwoPhaseMaterialParams<TraitsT, GPUBuffer>;
+  using GPUViewParams = Opm::PiecewiseLinearTwoPhaseMaterialParams<TraitsT, GPUView>;
+
+  using CPUTwoPhaseMaterialLaw = Opm::PiecewiseLinearTwoPhaseMaterial<TraitsT, CPUParams>;
+  using GPUTwoPhaseViewMaterialLaw = Opm::PiecewiseLinearTwoPhaseMaterial<TraitsT, GPUViewParams>;
+  using NorneEvaluation = Opm::DenseAd::Evaluation<Scalar, 3, 0u>;
+
+__global__ void gpuTwoPhaseSatPcnwWrapper(GPUTwoPhaseViewMaterialLaw::Params params, NorneEvaluation Sw, NorneEvaluation* res){
+    *res = GPUTwoPhaseViewMaterialLaw::twoPhaseSatPcnw(params, Sw);
+}
+
+BOOST_AUTO_TEST_CASE(TestSimpleInterpolation)
+{
+    CPUParams cpuParams;
+    GPUViewParams gpuViewParams;
+
+    ValueVector cx = {0.0, 0.5, 1.0};
+    ValueVector cy = {0.0, 0.9, 1.0};
+    const GPUBuffer gx(cx);
+    const GPUBuffer gy(cy);
+
+    cpuParams.setPcnwSamples(cx, cy);
+    cpuParams.setKrwSamples(cx, cy);
+    cpuParams.setKrnSamples(cx, cy);
+    cpuParams.finalize();
+
+    constGPUBufferParams gpuBufferParams(gx, gy, gx, gy, gx, gy);
+
+    gpuViewParams = Opm::gpuistl::make_view<TraitsT, const GPUBuffer, GPUView>(gpuBufferParams);
+
+    ValueVector testXs = {-1.0, 0, 0.1, 0.3, 0.5, 0.7, 0.9, 0.99, 1.0, 1.1};
+
+    for (Scalar x_i : testXs){
+        auto cpuMadeAd = NorneEvaluation(x_i, 0);
+        NorneEvaluation cpuInterpolatedEval = CPUTwoPhaseMaterialLaw::twoPhaseSatPcnw<NorneEvaluation>(cpuParams, cpuMadeAd);
+
+        NorneEvaluation* gpuAdInput;
+        NorneEvaluation* gpuAdResOnGPU;
+        NorneEvaluation gpuAdResOnCPU[1];
+
+        OPM_GPU_SAFE_CALL(cudaMalloc(&gpuAdInput, sizeof(NorneEvaluation)));
+        OPM_GPU_SAFE_CALL(cudaMalloc(&gpuAdResOnGPU, sizeof(NorneEvaluation)));
+
+        OPM_GPU_SAFE_CALL(cudaMemcpy(gpuAdInput, &cpuMadeAd, sizeof(NorneEvaluation), cudaMemcpyHostToDevice));
+        gpuTwoPhaseSatPcnwWrapper<<<1,1>>>(gpuViewParams, *gpuAdInput, gpuAdResOnGPU);
+        OPM_GPU_SAFE_CALL(cudaDeviceSynchronize());
+        OPM_GPU_SAFE_CALL(cudaMemcpy(gpuAdResOnCPU, gpuAdResOnGPU, sizeof(NorneEvaluation), cudaMemcpyDeviceToHost));
+
+        BOOST_CHECK(gpuAdResOnCPU[0] == cpuInterpolatedEval);
+    }
+}