fix piecewiselineartwophasematerial tests

Fixing the test required avoid accesses to a
mutable gpu buffer on the CPU. Removing function
from the GPUBuffer and GPUView class also required
a minor adjustment in the GPUView test.

opm/simulators/linalg/gpuistl/GpuBuffer.hpp

@@ -127,7 +127,11 @@ public:
 #ifndef NDEBUG
         assertHasElements();
 #endif
-    return m_dataOnDevice[0];
+#if OPM_IS_INSIDE_DEVICE_FUNCTION
+        return m_dataOnDevice[0];
+#else
+        assert(false && "Getting a reference to a value in a GPUBuffer is not allowed in CPU code");
+#endif
     }
 
     /**
@@ -138,7 +142,11 @@ public:
 #ifndef NDEBUG
         assertHasElements();
 #endif
-    return m_dataOnDevice[m_numberOfElements-1];
+#if OPM_IS_INSIDE_DEVICE_FUNCTION
+        return m_dataOnDevice[m_numberOfElements-1];
+#else
+        assert(false && "Getting a reference to a value in a GPUBuffer is not allowed in CPU code");
+#endif
     }
 
     /**
@@ -149,7 +157,13 @@ public:
 #ifndef NDEBUG
         assertHasElements();
 #endif
-    return m_dataOnDevice[0];
+#if OPM_IS_INSIDE_DEVICE_FUNCTION
+        return m_dataOnDevice[0];
+#else
+        T value;
+        cudaMemcpy(&value, &m_dataOnDevice[0], sizeof(T), cudaMemcpyDeviceToHost);
+        return value;
+#endif
     }
 
     /**
@@ -160,7 +174,13 @@ public:
 #ifndef NDEBUG
         assertHasElements();
 #endif
-    return m_dataOnDevice[m_numberOfElements-1];
+#if OPM_IS_INSIDE_DEVICE_FUNCTION
+        return m_dataOnDevice[m_numberOfElements-1];
+#else
+        T value;
+        cudaMemcpy(&value, &m_dataOnDevice[m_numberOfElements - 1], sizeof(T), cudaMemcpyDeviceToHost);
+        return value;
+#endif
     }
 
     /**

opm/simulators/linalg/gpuistl/GpuView.hpp

@@ -124,7 +124,7 @@ public:
     /**
      * @return fetch the first element in a GpuView
      */
-    __host__ __device__ T& front()
+    __device__ T& front()
     {
 #ifndef NDEBUG
         assertHasElements();
@@ -135,7 +135,7 @@ public:
     /**
      * @return fetch the last element in a GpuView
      */
-    __host__ __device__ T& back()
+    __device__ T& back()
     {
 #ifndef NDEBUG
         assertHasElements();
@@ -151,7 +151,13 @@ public:
 #ifndef NDEBUG
         assertHasElements();
 #endif
+#if OPM_IS_INSIDE_DEVICE_FUNCTION
         return m_dataPtr[0];
+#else
+        T value;
+        cudaMemcpy(&value, &m_dataPtr[0], sizeof(T), cudaMemcpyDeviceToHost);
+        return value;
+#endif
     }
 
     /**
@@ -162,7 +168,13 @@ public:
 #ifndef NDEBUG
         assertHasElements();
 #endif
+#if OPM_IS_INSIDE_DEVICE_FUNCTION
         return m_dataPtr[m_numberOfElements-1];
+#else
+        T value;
+        cudaMemcpy(&value, &m_dataPtr[m_numberOfElements - 1], sizeof(T), cudaMemcpyDeviceToHost);
+        return value;
+#endif
     }
 
     /**

tests/gpuistl/test_GpuView.cu

@@ -67,20 +67,16 @@ BOOST_AUTO_TEST_CASE(TestGpuViewOnCPUTypes)
     auto cpuview = GpuViewDouble(buf.data(), buf.size());
     const auto const_cpuview = GpuViewDouble(buf.data(), buf.size());
 
-    // check that indexing a mutable view gives references when indexing it
+    // check that indexing a const view produces a value
-    bool correct_type_of_cpu_front = std::is_same_v<double&, decltype(cpuview.front())>;
-    bool correct_type_of_cpu_back = std::is_same_v<double&, decltype(cpuview.back())>;
     bool correct_type_of_const_cpu_front = std::is_same_v<double, decltype(const_cpuview.front())>;
     bool correct_type_of_const_cpu_back = std::is_same_v<double, decltype(const_cpuview.back())>;
 
-    BOOST_CHECK(correct_type_of_cpu_front);
-    BOOST_CHECK(correct_type_of_cpu_back);
     BOOST_CHECK(correct_type_of_const_cpu_front);
     BOOST_CHECK(correct_type_of_const_cpu_back);
 
     // check that the values are correct
-    BOOST_CHECK(cpuview.front() == buf.front());
+    BOOST_CHECK(const_cpuview.front() == buf.front());
-    BOOST_CHECK(cpuview.back() == buf.back());
+    BOOST_CHECK(const_cpuview.back() == buf.back());
 }
 
 BOOST_AUTO_TEST_CASE(TestGpuViewOnCPUWithSTLIteratorAlgorithm)

tests/gpuistl/test_gpu_linear_two_phase_material.cu

@@ -56,47 +56,48 @@
   using GPUTwoPhaseViewMaterialLaw = Opm::PiecewiseLinearTwoPhaseMaterial<TraitsT, GPUViewParams>;
   using NorneEvaluation = Opm::DenseAd::Evaluation<Scalar, 3, 0u>;
 
-__global__ void gpuTwoPhaseSatPcnwWrapper(GPUTwoPhaseViewMaterialLaw::Params params, NorneEvaluation Sw, NorneEvaluation* res){
+__global__ void gpuTwoPhaseSatPcnwWrapper(GPUTwoPhaseViewMaterialLaw::Params params, NorneEvaluation* Sw, NorneEvaluation* res){
-    *res = GPUTwoPhaseViewMaterialLaw::twoPhaseSatPcnw(params, Sw);
+    *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 cpuParams;
     cpuParams.setPcnwSamples(cx, cy);
     cpuParams.setKrwSamples(cx, cy);
     cpuParams.setKrnSamples(cx, cy);
     cpuParams.finalize();
 
-    constGPUBufferParams gpuBufferParams(gx, gy, gx, gy, gx, gy);
+    constGPUBufferParams gpuBufferParams = Opm::gpuistl::move_to_gpu<TraitsT, const GPUBuffer>(cpuParams);
 
-    gpuViewParams = Opm::gpuistl::make_view<TraitsT, const GPUBuffer, GPUView>(gpuBufferParams);
+    GPUViewParams 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};
 
+    NorneEvaluation* gpuAdInput;
+    NorneEvaluation* gpuAdResOnGPU;
+    OPM_GPU_SAFE_CALL(cudaMalloc(&gpuAdInput, sizeof(NorneEvaluation)));
+    OPM_GPU_SAFE_CALL(cudaMalloc(&gpuAdResOnGPU, sizeof(NorneEvaluation)));
+
     for (Scalar x_i : testXs){
         auto cpuMadeAd = NorneEvaluation(x_i, 0);
         NorneEvaluation cpuInterpolatedEval = CPUTwoPhaseMaterialLaw::twoPhaseSatPcnw<NorneEvaluation>(cpuParams, cpuMadeAd);
 
-        NorneEvaluation* gpuAdInput;
-        NorneEvaluation* gpuAdResOnGPU;
         NorneEvaluation gpuAdResOnCPU;
 
-        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);
+        gpuTwoPhaseSatPcnwWrapper<<<1,1>>>(gpuViewParams, gpuAdInput, gpuAdResOnGPU);
         OPM_GPU_SAFE_CALL(cudaDeviceSynchronize());
         OPM_GPU_SAFE_CALL(cudaMemcpy(&gpuAdResOnCPU, gpuAdResOnGPU, sizeof(NorneEvaluation), cudaMemcpyDeviceToHost));
 
         BOOST_CHECK(gpuAdResOnCPU == cpuInterpolatedEval);
     }
+
+    OPM_GPU_SAFE_CALL(cudaFree(gpuAdInput));
+    OPM_GPU_SAFE_CALL(cudaFree(gpuAdResOnGPU));
 }