Merge pull request #5838 from multitalentloes/fix_gpu_linear_two_phase_material

fix gpu_linear... test

opm/simulators/linalg/gpuistl/GpuBuffer.hpp

@@ -119,50 +119,6 @@ public:
      */
     const T* data() const;
 
-    /**
-     * @return fetch the first element in a GpuBuffer
-     */
-    __host__ __device__ T& front()
-    {
-#ifndef NDEBUG
-        assertHasElements();
-#endif
-    return m_dataOnDevice[0];
-    }
-
-    /**
-     * @return fetch the last element in a GpuBuffer
-     */
-    __host__ __device__ T& back()
-    {
-#ifndef NDEBUG
-        assertHasElements();
-#endif
-    return m_dataOnDevice[m_numberOfElements-1];
-    }
-
-    /**
-     * @return fetch the first element in a GpuBuffer
-     */
-    __host__ __device__ T front() const
-    {
-#ifndef NDEBUG
-        assertHasElements();
-#endif
-    return m_dataOnDevice[0];
-    }
-
-    /**
-     * @return fetch the last element in a GpuBuffer
-     */
-    __host__ __device__ T back() const
-    {
-#ifndef NDEBUG
-        assertHasElements();
-#endif
-    return m_dataOnDevice[m_numberOfElements-1];
-    }
-
     /**
      * @brief copyFromHost copies data from a Dune::BlockVector
      * @param bvector the vector to copy from

opm/simulators/linalg/gpuistl/GpuView.hpp

@@ -97,7 +97,7 @@ public:
      * @param numberOfElements number of T elements to allocate
      * @param dataOnHost data on host/CPU
      */
-    __host__ __device__ GpuView(T* dataOnHost, size_t numberOfElements)
+    GpuView(T* dataOnHost, size_t numberOfElements)
         : m_dataPtr(dataOnHost), m_numberOfElements(numberOfElements)
     {
     }

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<const GPUBuffer>(cpuParams);
 
-    gpuViewParams = Opm::gpuistl::make_view<TraitsT, const GPUBuffer, GPUView>(gpuBufferParams);
+    GPUViewParams gpuViewParams = Opm::gpuistl::make_view<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));
 }