Explicitly use __ldg on pointer derefs

GridTools · Jul 24, 2024 · c591f79 · c591f79
1 parent 0b0a306
commit c591f79
Show file tree

Hide file tree

Showing 3 changed files with 33 additions and 2 deletions.
diff --git a/include/gridtools/fn/cartesian.hpp b/include/gridtools/fn/cartesian.hpp
@@ -16,6 +16,10 @@
 #include "./common_interface.hpp"
 #include "./executor.hpp"
 
+#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350
+#include "../common/cuda_type_traits.hpp"
+#endif
+
 namespace gridtools::fn {
     namespace cartesian::dim {
         using i = integral_constant<int, 0>;
@@ -44,6 +48,12 @@ namespace gridtools::fn {
 
         template <class Tag, class Ptr, class Strides>
         GT_FUNCTION auto deref(iterator<Tag, Ptr, Strides> const &it) {
+#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350
+            if constexpr (std::is_pointer_v<decltype(it.m_ptr)> &&
+                          is_texture_type<std::decay_t<std::remove_pointer_t<decltype(it.m_ptr)>>>::value) {
+                return __ldg(it.m_ptr);
+            }
+#endif
             return *it.m_ptr;
         }
 

diff --git a/include/gridtools/fn/sid_neighbor_table.hpp b/include/gridtools/fn/sid_neighbor_table.hpp
@@ -16,6 +16,10 @@
 #include "../fn/unstructured.hpp"
 #include "../sid/concept.hpp"
 
+#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350
+#include "../common/cuda_type_traits.hpp"
+#endif
+
 namespace gridtools::fn::sid_neighbor_table {
     namespace sid_neighbor_table_impl_ {
         template <class IndexDimension,
@@ -46,7 +50,13 @@ namespace gridtools::fn::sid_neighbor_table {
 
             sid::shift(ptr, sid::get_stride<IndexDimension>(table.strides), index);
             for (std::size_t element_idx = 0; element_idx < MaxNumNeighbors; ++element_idx) {
-                neighbors[element_idx] = *ptr;
+#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350
+                if constexpr (std::is_pointer_v<decltype(ptr)> &&
+                              is_texture_type<std::decay_t<std::remove_pointer_t<decltype(ptr)>>>::value)
+                    neighbors[element_idx] = __ldg(ptr);
+                else
+#endif
+                    neighbors[element_idx] = *ptr;
                 sid::shift(ptr, sid::get_stride<NeighborDimension>(table.strides), 1_c);
             }
             return neighbors;

diff --git a/include/gridtools/fn/unstructured.hpp b/include/gridtools/fn/unstructured.hpp
@@ -20,6 +20,10 @@
 #include "./executor.hpp"
 #include "./neighbor_table.hpp"
 
+#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350
+#include "../common/cuda_type_traits.hpp"
+#endif
+
 namespace gridtools::fn {
     namespace unstructured::dim {
         using horizontal = integral_constant<int, 0>;
@@ -80,7 +84,14 @@ namespace gridtools::fn {
         GT_FUNCTION constexpr auto deref(iterator<Tag, Ptr, Strides, Domain> const &it) {
             GT_PROMISE(can_deref(it));
             decltype(auto) stride = host_device::at_key<Tag>(sid::get_stride<dim::horizontal>(it.m_strides));
-            return *sid::shifted(it.m_ptr, stride, it.m_index);
+            auto ptr = sid::shifted(it.m_ptr, stride, it.m_index);
+#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350
+            if constexpr (std::is_pointer_v<decltype(ptr)> &&
+                          is_texture_type<std::decay_t<std::remove_pointer_t<decltype(ptr)>>>::value) {
+                return __ldg(ptr);
+            }
+#endif
+            return *ptr;
         }
 
         template <class Tag, class Ptr, class Strides, class Domain, class Conn, class Offset>