From b48479b7e45609dac1de6b6fc9f0ac4e30acff12 Mon Sep 17 00:00:00 2001 From: Tobias Meyer Andersen Date: Fri, 3 Jan 2025 15:27:31 +0100 Subject: [PATCH 1/5] 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 | 28 ++++++++++++++++--- opm/simulators/linalg/gpuistl/GpuView.hpp | 16 +++++++++-- tests/gpuistl/test_GpuView.cu | 10 ++----- .../test_gpu_linear_two_phase_material.cu | 27 +++++++++--------- 4 files changed, 55 insertions(+), 26 deletions(-) diff --git a/opm/simulators/linalg/gpuistl/GpuBuffer.hpp b/opm/simulators/linalg/gpuistl/GpuBuffer.hpp index 7272d3c28..d79bf3b19 100644 --- a/opm/simulators/linalg/gpuistl/GpuBuffer.hpp +++ b/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 } /** diff --git a/opm/simulators/linalg/gpuistl/GpuView.hpp b/opm/simulators/linalg/gpuistl/GpuView.hpp index acde482ee..530e2ae10 100644 --- a/opm/simulators/linalg/gpuistl/GpuView.hpp +++ b/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 } /** diff --git a/tests/gpuistl/test_GpuView.cu b/tests/gpuistl/test_GpuView.cu index 5a9a0c68d..9d21aab79 100644 --- a/tests/gpuistl/test_GpuView.cu +++ b/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 - bool correct_type_of_cpu_front = std::is_same_v; - bool correct_type_of_cpu_back = std::is_same_v; + // check that indexing a const view produces a value bool correct_type_of_const_cpu_front = std::is_same_v; bool correct_type_of_const_cpu_back = std::is_same_v; - 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(cpuview.back() == buf.back()); + BOOST_CHECK(const_cpuview.front() == buf.front()); + BOOST_CHECK(const_cpuview.back() == buf.back()); } BOOST_AUTO_TEST_CASE(TestGpuViewOnCPUWithSTLIteratorAlgorithm) diff --git a/tests/gpuistl/test_gpu_linear_two_phase_material.cu b/tests/gpuistl/test_gpu_linear_two_phase_material.cu index ebbd3ac99..6a6ad9c31 100644 --- a/tests/gpuistl/test_gpu_linear_two_phase_material.cu +++ b/tests/gpuistl/test_gpu_linear_two_phase_material.cu @@ -56,47 +56,48 @@ using GPUTwoPhaseViewMaterialLaw = Opm::PiecewiseLinearTwoPhaseMaterial; using NorneEvaluation = Opm::DenseAd::Evaluation; -__global__ void gpuTwoPhaseSatPcnwWrapper(GPUTwoPhaseViewMaterialLaw::Params params, NorneEvaluation Sw, NorneEvaluation* res){ - *res = GPUTwoPhaseViewMaterialLaw::twoPhaseSatPcnw(params, Sw); +__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 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(cpuParams); - gpuViewParams = Opm::gpuistl::make_view(gpuBufferParams); + GPUViewParams gpuViewParams = Opm::gpuistl::make_view(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(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)); } From 486f7cc819b9c8c7ad7108c00648facdbf3b319a Mon Sep 17 00:00:00 2001 From: Tobias Meyer Andersen Date: Wed, 8 Jan 2025 08:53:16 +0100 Subject: [PATCH 2/5] remove buffer and view cpu support --- opm/simulators/linalg/gpuistl/GpuBuffer.hpp | 20 -------------------- opm/simulators/linalg/gpuistl/GpuView.hpp | 12 ------------ 2 files changed, 32 deletions(-) diff --git a/opm/simulators/linalg/gpuistl/GpuBuffer.hpp b/opm/simulators/linalg/gpuistl/GpuBuffer.hpp index d79bf3b19..82cc152f9 100644 --- a/opm/simulators/linalg/gpuistl/GpuBuffer.hpp +++ b/opm/simulators/linalg/gpuistl/GpuBuffer.hpp @@ -127,11 +127,7 @@ public: #ifndef NDEBUG assertHasElements(); #endif -#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 } /** @@ -142,11 +138,7 @@ public: #ifndef NDEBUG assertHasElements(); #endif -#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 } /** @@ -157,13 +149,7 @@ public: #ifndef NDEBUG assertHasElements(); #endif -#if OPM_IS_INSIDE_DEVICE_FUNCTION return m_dataOnDevice[0]; -#else - T value; - cudaMemcpy(&value, &m_dataOnDevice[0], sizeof(T), cudaMemcpyDeviceToHost); - return value; -#endif } /** @@ -174,13 +160,7 @@ public: #ifndef NDEBUG assertHasElements(); #endif -#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 } /** diff --git a/opm/simulators/linalg/gpuistl/GpuView.hpp b/opm/simulators/linalg/gpuistl/GpuView.hpp index 530e2ae10..dd39288d9 100644 --- a/opm/simulators/linalg/gpuistl/GpuView.hpp +++ b/opm/simulators/linalg/gpuistl/GpuView.hpp @@ -151,13 +151,7 @@ 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 } /** @@ -168,13 +162,7 @@ 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 } /** From adc6afd749df12838e2dfe739fbe9664ec9d5181 Mon Sep 17 00:00:00 2001 From: Tobias Meyer Andersen Date: Wed, 8 Jan 2025 14:54:03 +0100 Subject: [PATCH 3/5] WIP --- opm/simulators/linalg/gpuistl/GpuBuffer.hpp | 44 --------------------- 1 file changed, 44 deletions(-) diff --git a/opm/simulators/linalg/gpuistl/GpuBuffer.hpp b/opm/simulators/linalg/gpuistl/GpuBuffer.hpp index 82cc152f9..369ea3ba2 100644 --- a/opm/simulators/linalg/gpuistl/GpuBuffer.hpp +++ b/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 From 6c29cf6962a2d4b03f805b22dde0b585a3663454 Mon Sep 17 00:00:00 2001 From: Tobias Meyer Andersen Date: Wed, 8 Jan 2025 15:08:50 +0100 Subject: [PATCH 4/5] use less template arguments --- tests/gpuistl/test_gpu_linear_two_phase_material.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tests/gpuistl/test_gpu_linear_two_phase_material.cu b/tests/gpuistl/test_gpu_linear_two_phase_material.cu index 6a6ad9c31..ee581d748 100644 --- a/tests/gpuistl/test_gpu_linear_two_phase_material.cu +++ b/tests/gpuistl/test_gpu_linear_two_phase_material.cu @@ -73,9 +73,9 @@ BOOST_AUTO_TEST_CASE(TestSimpleInterpolation) cpuParams.setKrnSamples(cx, cy); cpuParams.finalize(); - constGPUBufferParams gpuBufferParams = Opm::gpuistl::move_to_gpu(cpuParams); + constGPUBufferParams gpuBufferParams = Opm::gpuistl::move_to_gpu(cpuParams); - GPUViewParams gpuViewParams = Opm::gpuistl::make_view(gpuBufferParams); + GPUViewParams gpuViewParams = Opm::gpuistl::make_view(gpuBufferParams); ValueVector testXs = {-1.0, 0, 0.1, 0.3, 0.5, 0.7, 0.9, 0.99, 1.0, 1.1}; From cabc049e493547098f68f1a06a1a4e464055f539 Mon Sep 17 00:00:00 2001 From: Tobias Meyer Andersen Date: Thu, 9 Jan 2025 10:38:58 +0100 Subject: [PATCH 5/5] minor adjustments --- opm/simulators/linalg/gpuistl/GpuView.hpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/opm/simulators/linalg/gpuistl/GpuView.hpp b/opm/simulators/linalg/gpuistl/GpuView.hpp index dd39288d9..54cd855c3 100644 --- a/opm/simulators/linalg/gpuistl/GpuView.hpp +++ b/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) { } @@ -124,7 +124,7 @@ public: /** * @return fetch the first element in a GpuView */ - __device__ T& front() + __host__ __device__ T& front() { #ifndef NDEBUG assertHasElements(); @@ -135,7 +135,7 @@ public: /** * @return fetch the last element in a GpuView */ - __device__ T& back() + __host__ __device__ T& back() { #ifndef NDEBUG assertHasElements();