From 272f320b1ad814b061c82701b998be0dde966693 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Cezary=20Skrzy=C5=84ski?=
 <cezary.skrzynski@ng-analytics.com>
Date: Wed, 18 Sep 2024 17:46:38 +0200
Subject: [PATCH] #11: fix "calling a constexpr __host__ function from a
 __host__ __device__ function" warnings

---
 src/bvh/traits.hpp    | 40 ++++++++++++++++++++++------------------
 src/bvh/util/span.hpp |  3 ++-
 2 files changed, 24 insertions(+), 19 deletions(-)

diff --git a/src/bvh/traits.hpp b/src/bvh/traits.hpp
index 5149bc8..67554d2 100644
--- a/src/bvh/traits.hpp
+++ b/src/bvh/traits.hpp
@@ -33,21 +33,25 @@
 #ifndef INC_BVH_TRAITS_HPP
 #define INC_BVH_TRAITS_HPP
 
+#include <Kokkos_Macros.hpp>
+
+#include <type_traits>
+
 namespace bvh
 {
   template< typename... >
   using void_t = void;
-  
+
   template< unsigned N >
   struct overload_priority
     : overload_priority< N - 1 >
   {
   };
-  
+
   template<>
   struct overload_priority< 0 >
   {};
-  
+
   namespace detail
   {
     template< typename Element >
@@ -56,79 +60,79 @@ namespace bvh
     {
       return get_entity_kdop( _element );
     }
-  
+
     template< typename Element >
     constexpr auto get_kdop_impl( const Element &_element, overload_priority< 0 > )
       -> decltype( _element.kdop() )
     {
       return _element.kdop();
     }
-    
+
     template< typename Element >
     constexpr auto get_kdop( const Element &_element )
     {
       return get_kdop_impl( _element, overload_priority< 1 >{} );
     }
-  
+
     template< typename Element >
     constexpr auto get_centroid_impl( const Element &_element, overload_priority< 1 > )
       -> decltype( get_entity_centroid( _element ) )
     {
       return get_entity_centroid( _element );
     }
-  
+
     template< typename Element >
     constexpr auto get_centroid_impl( const Element &_element, overload_priority< 0 >  )
       -> decltype( _element.centroid() )
     {
       return _element.centroid();
     }
-  
+
     template< typename Element >
     constexpr auto get_centroid( const Element &_element )
     {
       return get_centroid_impl( _element, overload_priority< 1 >{} );
     }
-  
+
     template< typename Element >
     constexpr auto get_global_id_impl( const Element &_element, overload_priority< 1 > )
       -> decltype( get_entity_global_id( _element ) )
     {
       return get_entity_global_id( _element );
     }
-  
+
     template< typename Element >
     constexpr auto get_global_id_impl( const Element &_element, overload_priority< 0 >  )
       -> decltype( _element.global_id() )
     {
       return _element.global_id();
     }
-  
+
     template< typename Element >
     constexpr auto get_global_id( const Element &_element )
     {
       return get_global_id_impl( _element, overload_priority< 1 >{} );
     }
   }
-  
+
   template< typename Element >
   struct element_traits
   {
     using kdop_type = std::remove_cv_t< std::remove_reference_t< decltype( detail::get_kdop( std::declval< Element >() ) ) > >;
     using centroid_type = std::remove_cv_t< std::remove_reference_t< decltype( detail::get_centroid( std::declval< Element >() ) ) > >;
     using global_id_type = std::remove_cv_t< std::remove_reference_t< decltype( detail::get_global_id( std::declval< Element >() ) ) > >;
-    
-    static constexpr kdop_type get_kdop( const Element &_element )
+
+    static constexpr KOKKOS_INLINE_FUNCTION kdop_type get_kdop( const Element &_element )
     {
       return detail::get_kdop( _element );
     }
-    
-    static constexpr centroid_type get_centroid( const Element &_element )
+
+    static constexpr KOKKOS_INLINE_FUNCTION centroid_type get_centroid( const Element &_element )
     {
       return detail::get_centroid( _element );
     }
-    
-    static constexpr global_id_type get_global_id( const Element &_element )
+
+    static constexpr KOKKOS_INLINE_FUNCTION global_id_type get_global_id( const Element &_element )
     {
       return detail::get_global_id( _element );
     }
diff --git a/src/bvh/util/span.hpp b/src/bvh/util/span.hpp
index 568aa7c..7156133 100644
--- a/src/bvh/util/span.hpp
+++ b/src/bvh/util/span.hpp
@@ -37,6 +37,7 @@
 #include <type_traits>
 #include "assert.hpp"
 #include <array>
+#include <Kokkos_Macros.hpp>
 #include "../range.hpp"
 
 namespace bvh
@@ -174,7 +175,7 @@ namespace bvh
     constexpr index_type size_bytes() const noexcept { return m_count * sizeof( element_type ); }
     constexpr bool empty() const noexcept { return m_count == 0; }
 
-    constexpr reference operator[]( index_type _idx ) const
+    constexpr KOKKOS_INLINE_FUNCTION reference operator[]( index_type _idx ) const
     {
       BVH_ASSERT( _idx < m_count );
       return m_data[_idx];