test gpu instantiation of BlackOilFluidSystem

CMakeLists.txt

@@ -749,6 +749,7 @@ if(CUDA_FOUND)
                        gpu_smart_pointers
                        is_gpu_pointer
                        throw_macros_on_gpu
+                       gpuBlackOilFluidSystem
                        PROPERTIES LABELS ${gpu_label})
 endif()
 

CMakeLists_files.cmake

@@ -488,6 +488,7 @@ if (HAVE_CUDA)
   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_gpuBlackOilFluidSystem.cu)
 
 
   # for loop providing the flag --expt-relaxed-constexpr to fix some cuda issues with constexpr

opm/simulators/linalg/gpuistl/GpuBuffer.cpp

@@ -37,9 +37,6 @@ template <class T>
 GpuBuffer<T>::GpuBuffer(const size_t numberOfElements)
     : m_numberOfElements(numberOfElements)
 {
-    if (numberOfElements < 1) {
-        OPM_THROW(std::invalid_argument, "Setting a GpuBuffer size to a non-positive number is not allowed");
-    }
     OPM_GPU_SAFE_CALL(cudaMalloc(&m_dataOnDevice, sizeof(T) * m_numberOfElements));
 }
 
@@ -56,8 +53,11 @@ template <class T>
 GpuBuffer<T>::GpuBuffer(const GpuBuffer<T>& other)
     : GpuBuffer(other.m_numberOfElements)
 {
-    assertHasElements();
+    // assertHasElements();
     assertSameSize(other);
+    if (m_numberOfElements == 0) {
+        return;
+    }
     OPM_GPU_SAFE_CALL(cudaMemcpy(m_dataOnDevice,
                                   other.m_dataOnDevice,
                                   m_numberOfElements * sizeof(T),
@@ -199,6 +199,12 @@ GpuBuffer<T>::copyToHost(std::vector<T>& data) const
 template class GpuBuffer<double>;
 template class GpuBuffer<float>;
 template class GpuBuffer<int>;
+template class GpuBuffer<std::array<double, 3>>;
+template class GpuBuffer<std::array<float, 3>>;
+template class GpuBuffer<std::array<int, 3>>;
+template class GpuBuffer<std::array<double, 9>>;
+template class GpuBuffer<std::array<float, 9>>;
+template class GpuBuffer<std::array<int, 9>>;
 
 template <class T>
 GpuView<T> make_view(GpuBuffer<T>& buf) {
@@ -208,6 +214,12 @@ GpuView<T> make_view(GpuBuffer<T>& buf) {
 template GpuView<double> make_view(GpuBuffer<double>&);
 template GpuView<float> make_view(GpuBuffer<float>&);
 template GpuView<int> make_view(GpuBuffer<int>&);
+template GpuView<std::array<double, 3>> make_view(GpuBuffer<std::array<double, 3>>&);
+template GpuView<std::array<float, 3>> make_view(GpuBuffer<std::array<float, 3>>&);
+template GpuView<std::array<int, 3>> make_view(GpuBuffer<std::array<int, 3>>&);
+template GpuView<std::array<double, 9>> make_view(GpuBuffer<std::array<double, 9>>&);
+template GpuView<std::array<float, 9>> make_view(GpuBuffer<std::array<float, 9>>&);
+template GpuView<std::array<int, 9>> make_view(GpuBuffer<std::array<int, 9>>&);
 
 template <class T>
 GpuView<const T> make_view(const GpuBuffer<T>& buf) {
@@ -217,5 +229,11 @@ GpuView<const T> make_view(const GpuBuffer<T>& buf) {
 template GpuView<const double> make_view(const GpuBuffer<double>&);
 template GpuView<const float> make_view(const GpuBuffer<float>&);
 template GpuView<const int> make_view(const GpuBuffer<int>&);
+template GpuView<const std::array<double, 3>> make_view(const GpuBuffer<std::array<double, 3>>&);
+template GpuView<const std::array<float, 3>> make_view(const GpuBuffer<std::array<float, 3>>&);
+template GpuView<const std::array<int, 3>> make_view(const GpuBuffer<std::array<int, 3>>&);
+template GpuView<const std::array<double, 9>> make_view(const GpuBuffer<std::array<double, 9>>&);
+template GpuView<const std::array<float, 9>> make_view(const GpuBuffer<std::array<float, 9>>&);
+template GpuView<const std::array<int, 9>> make_view(const GpuBuffer<std::array<int, 9>>&);
 
 } // namespace Opm::gpuistl

opm/simulators/linalg/gpuistl/GpuView.cpp

@@ -78,5 +78,11 @@ GpuView<T>::copyToHost(std::vector<T>& data) const
 template class GpuView<double>;
 template class GpuView<float>;
 template class GpuView<int>;
+template class GpuView<std::array<double, 3>>;
+template class GpuView<std::array<float, 3>>;
+template class GpuView<std::array<int, 3>>;
+template class GpuView<std::array<double, 9>>;
+template class GpuView<std::array<float, 9>>;
+template class GpuView<std::array<int, 9>>;
 
 } // namespace Opm::gpuistl

opm/simulators/linalg/gpuistl/GpuView.hpp

@@ -34,6 +34,48 @@
 namespace Opm::gpuistl
 {
 
+/// @brief The ViewPointer is a class that is an alternative to smart pointers on the GPU
+/// A value wrapped in a shared_ptr cannot be used on the GPU, we then create a ViewPointer
+/// around a view to the avlue instead. This means that we have an object here
+/// that just accesses a view through the -> and * operators of regular pointers
+/// Example ViewPointer<GpuView<double>*> ptr(view);
+/// @tparam ViewType 
+template<class ViewType>
+class ViewPointer {
+public:
+    using element_type = ViewType;
+    __host__ __device__ ViewPointer (ViewType view) : view_(view) {}
+
+    __host__ __device__ ViewPointer() : view_() {}
+
+    __host__ __device__ ViewType& operator*(){
+        return view_;
+    }
+
+    __host__ __device__ const ViewType& operator*() const {
+        return view_;
+    }
+
+    __host__ __device__ ViewType* operator->(){
+        return &view_;
+    }
+
+    __host__ __device__ const ViewType* operator->() const {
+        return &view_;
+    }
+
+    __host__ __device__ ViewType& get(){
+        return &view_;
+    }
+
+    __host__ __device__ const ViewType& get() const {
+        return &view_;
+    }
+
+private:
+    ViewType view_;
+};
+
 /**
  * @brief The GpuView class is provides a view of some data allocated on the GPU
  * Essenstially is only stores a pointer and a size.
@@ -107,6 +149,11 @@ class GpuView
      */
     ~GpuView() = default;
 
+    // make a copy of this view that behaves like a pointer
+    ViewPointer<GpuView<T>> pointer(){
+        return ViewPointer<GpuView<T>>(*this);
+    }
+
     /**
      * @return the raw pointer to the GPU data
      */

opm/simulators/linalg/gpuistl/gpu_smart_pointer.hpp

@@ -225,6 +225,47 @@ copyToGPU(const T& value, const std::unique_ptr<T, Deleter>& ptr)
  * being compatible with the requirements of the GPU kernels. This is useful when we want to pass
  * a smart pointer to a GPU kernel, but we do not want to transfer the ownership of the memory.
  */
+// template <class T>
+// class PointerView
+// {
+// public:
+//     PointerView(const PointerView& other) = default;
+
+//     PointerView(const std::shared_ptr<T>& ptr)
+//         : ptr_(ptr.get())
+//     {
+//     }
+
+//     template <class Deleter>
+//     PointerView(const std::unique_ptr<T, Deleter>& ptr)
+//         : ptr_(ptr.get())
+//     {
+//     }
+
+//     PointerView(T* ptr)
+//         : ptr_(ptr)
+//     {
+//     }
+
+//     OPM_HOST_DEVICE T* get() const
+//     {
+//         return ptr_;
+//     }
+
+//     OPM_HOST_DEVICE T& operator*() const
+//     {
+//         return *ptr_;
+//     }
+
+//     OPM_HOST_DEVICE T* operator->() const
+//     {
+//         return ptr_;
+//     }
+
+// private:
+//     T* ptr_;
+// };
+
 template <class T>
 class PointerView
 {
@@ -252,7 +293,12 @@ class PointerView
         return ptr_;
     }
 
-    OPM_HOST_DEVICE T& operator*() const
+    OPM_HOST_DEVICE const T& operator*() const
+    {
+        return *ptr_;
+    }
+
+    OPM_HOST_DEVICE T& operator*()
     {
         return *ptr_;
     }
@@ -266,6 +312,44 @@ class PointerView
     T* ptr_;
 };
 
+
+template <>
+class PointerView<void>
+{
+public:
+    PointerView(const PointerView& other) = default;
+
+    PointerView(const std::shared_ptr<void>& ptr)
+        : ptr_(ptr.get())
+    {
+    }
+
+    template <class Deleter>
+    PointerView(const std::unique_ptr<void, Deleter>& ptr)
+        : ptr_(ptr.get())
+    {
+    }
+
+    PointerView(void* ptr)
+        : ptr_(ptr)
+    {
+    }
+
+    OPM_HOST_DEVICE void* get() const
+    {
+        return ptr_;
+    }
+
+
+    OPM_HOST_DEVICE void* operator->() const
+    {
+        return ptr_;
+    }
+
+private:
+    void* ptr_;
+};
+
 template <class T>
 PointerView<T>
 make_view(const std::shared_ptr<T>& ptr)

tests/gpuistl/test_GpuView.cu

@@ -34,6 +34,7 @@
 
 using GpuViewDouble = ::Opm::gpuistl::GpuView<double>;
 using GpuBufferDouble = ::Opm::gpuistl::GpuBuffer<double>;
+using ViewPointerDouble = ::Opm::gpuistl::ViewPointer<GpuViewDouble>;
 
 __global__ void useGpuViewOnGPU(GpuViewDouble a, GpuViewDouble b){
     b[0] = a.front();
@@ -109,3 +110,27 @@ BOOST_AUTO_TEST_CASE(TestGpuViewOnGPU)
     // checks that view[0] = view[2] works
     BOOST_CHECK(buf[2] == vecA[0]);
 }
+
+__global__ void useViewPointer(ViewPointerDouble ptr, double* out){
+    out[0] = ptr->front();
+    out[1] = (*ptr).front();
+}
+
+BOOST_AUTO_TEST_CASE(TestGpuViewPointer)
+{
+    auto buf = std::vector<double>({1.0, 2.0, 42.0, 59.9451743, 10.7132692});
+    auto gpubuf = GpuBufferDouble(buf);
+    auto gpuview = GpuViewDouble(gpubuf.data(), gpubuf.size());
+    auto gpuViewPtr = ::Opm::gpuistl::ViewPointer<decltype(gpuview)>(gpuview);
+
+    std::vector<double> out(2, -1.0);
+    auto gpuBufOut = GpuBufferDouble(out);
+    auto gpuViewOut = GpuViewDouble(gpuBufOut.data(), gpuBufOut.size());
+
+    useViewPointer<<<1,1>>>(gpuViewPtr, gpuViewOut.data());
+
+    auto outValues = gpuViewOut.asStdVector();
+
+    BOOST_CHECK(outValues[0] == buf.front());
+    BOOST_CHECK(outValues[1] == buf.front());
+}