Remove ldg_ptr and Replace Functionality by const_ptr_deref (#1810)

No performance change in the GridTools performance tests visible, however this might change performance of external code that does not rely on GridTools storages but uses `sid::composite` or `sid::dimension_to_tuple_like` which now use `__ldg` but previously did not.
GridTools · Oct 30, 2024 · 049d6fe · 049d6fe
1 parent 6a0bf2f
commit 049d6fe
Show file tree

Hide file tree

Showing 14 changed files with 57 additions and 274 deletions.
diff --git a/include/gridtools/common/const_ptr_deref.hpp b/include/gridtools/common/const_ptr_deref.hpp
@@ -0,0 +1,38 @@
+/*
+ * GridTools
+ *
+ * Copyright (c) 2014-2023, ETH Zurich
+ * All rights reserved.
+ *
+ * Please, refer to the LICENSE file in the root directory.
+ * SPDX-License-Identifier: BSD-3-Clause
+ */
+#pragma once
+
+#include "host_device.hpp"
+
+#ifdef GT_CUDACC
+#include "cuda_type_traits.hpp"
+#endif
+
+namespace gridtools {
+
+#ifdef GT_CUDACC
+
+    template <class T>
+    GT_FUNCTION constexpr std::enable_if_t<is_texture_type<T>::value, T> const_ptr_deref(T const *ptr) {
+#ifdef GT_CUDA_ARCH
+        return __ldg(ptr);
+#else
+        return *ptr;
+#endif
+    }
+
+#endif
+
+    template <class T>
+    GT_FUNCTION constexpr decltype(auto) const_ptr_deref(T &&ptr) {
+        return *ptr;
+    }
+
+} // namespace gridtools
diff --git a/include/gridtools/common/ldg_ptr.hpp b/include/gridtools/common/ldg_ptr.hpp
diff --git a/include/gridtools/fn/cartesian.hpp b/include/gridtools/fn/cartesian.hpp
@@ -11,7 +11,7 @@
 
 #include <functional>
 
-#include "../common/ldg_ptr.hpp"
+#include "../common/const_ptr_deref.hpp"
 #include "../common/tuple_util.hpp"
 #include "../sid/concept.hpp"
 #include "./common_interface.hpp"
@@ -45,7 +45,7 @@ namespace gridtools::fn {
 
         template <class Tag, class Ptr, class Strides>
         GT_FUNCTION auto deref(iterator<Tag, Ptr, Strides> const &it) {
-            return *as_ldg_ptr(it.m_ptr);
+            return const_ptr_deref(it.m_ptr);
         }
 
         template <class Tag, class Ptr, class Strides, class Dim, class Offset, class... Offsets>

diff --git a/include/gridtools/fn/neighbor_table.hpp b/include/gridtools/fn/neighbor_table.hpp
@@ -11,7 +11,7 @@
 
 #include <type_traits>
 
-#include "../common/ldg_ptr.hpp"
+#include "../common/const_ptr_deref.hpp"
 #include "../common/tuple_util.hpp"
 #include "../meta/logical.hpp"
 
@@ -57,12 +57,12 @@ namespace gridtools::fn::neighbor_table {
 
         template <class T, std::enable_if_t<is_neighbor_list<T>::value, int> = 0>
         GT_FUNCTION T const &neighbor_table_neighbors(T const *table, int index) {
-            return *as_ldg_ptr(&table[index]);
+            return const_ptr_deref(&table[index]);
         }
 
         template <class NeighborTable>
-        GT_FUNCTION constexpr auto neighbors(NeighborTable const &nt, int index)
-            -> decltype(neighbor_table_neighbors(nt, index)) {
+        GT_FUNCTION constexpr auto neighbors(
+            NeighborTable const &nt, int index) -> decltype(neighbor_table_neighbors(nt, index)) {
             return neighbor_table_neighbors(nt, index);
         }
 

diff --git a/include/gridtools/fn/sid_neighbor_table.hpp b/include/gridtools/fn/sid_neighbor_table.hpp
@@ -13,7 +13,7 @@
 #include <type_traits>
 
 #include "../common/array.hpp"
-#include "../common/ldg_ptr.hpp"
+#include "../common/const_ptr_deref.hpp"
 #include "../fn/unstructured.hpp"
 #include "../sid/concept.hpp"
 
@@ -47,7 +47,7 @@ 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] = *as_ldg_ptr(ptr);
+                neighbors[element_idx] = const_ptr_deref(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
@@ -11,9 +11,9 @@
 
 #include <functional>
 
+#include "../common/const_ptr_deref.hpp"
 #include "../common/defs.hpp"
 #include "../common/hymap.hpp"
-#include "../common/ldg_ptr.hpp"
 #include "../meta/logical.hpp"
 #include "../sid/concept.hpp"
 #include "./common_interface.hpp"
@@ -79,7 +79,7 @@ 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 *as_ldg_ptr(sid::shifted(it.m_ptr, stride, it.m_index));
+            return const_ptr_deref(sid::shifted(it.m_ptr, stride, it.m_index));
         }
 
         template <class Tag, class Ptr, class Strides, class Domain, class Conn, class Offset>

diff --git a/include/gridtools/sid/composite.hpp b/include/gridtools/sid/composite.hpp
@@ -14,6 +14,7 @@
 #include <type_traits>
 #include <utility>
 
+#include "../common/const_ptr_deref.hpp"
 #include "../common/defs.hpp"
 #include "../common/for_each.hpp"
 #include "../common/host_device.hpp"
@@ -211,7 +212,7 @@ namespace gridtools {
 #if defined(__clang__) || !defined(__GNUC__) || (__GNUC__ != 9 && __GNUC__ != 10)
                                                                GT_FORCE_INLINE_LAMBDA
 #endif
-                            -> decltype(auto) { return *ptr; },
+                            -> decltype(auto) { return const_ptr_deref(ptr); },
                             m_vals));
                 }
 
@@ -476,5 +477,5 @@ namespace gridtools {
                 friend values tuple_getter(values) { return {}; }
             };
         } // namespace composite
-    }     // namespace sid
+    } // namespace sid
 } // namespace gridtools
diff --git a/include/gridtools/sid/simple_ptr_holder.hpp b/include/gridtools/sid/simple_ptr_holder.hpp
@@ -16,7 +16,6 @@
 
 #include "../common/defs.hpp"
 #include "../common/host_device.hpp"
-#include "../common/ldg_ptr.hpp"
 
 #define GT_FILENAME <gridtools/sid/simple_ptr_holder.hpp>
 #include GT_ITERATE_ON_TARGETS()
@@ -39,7 +38,7 @@ namespace gridtools {
                 simple_ptr_holder() = default;
                 GT_TARGET GT_FORCE_INLINE constexpr simple_ptr_holder(T const &ptr) : m_val{ptr} {}
 #endif
-                GT_TARGET GT_FORCE_INLINE constexpr decltype(auto) operator()() const { return as_ldg_ptr(m_val); }
+                GT_TARGET GT_FORCE_INLINE constexpr T const &operator()() const { return m_val; }
             };
 
             template <class T>

diff --git a/include/gridtools/stencil/gpu/entry_point.hpp b/include/gridtools/stencil/gpu/entry_point.hpp
@@ -13,12 +13,12 @@
 #include <type_traits>
 #include <utility>
 
+#include "../../common/const_ptr_deref.hpp"
 #include "../../common/cuda_type_traits.hpp"
 #include "../../common/cuda_util.hpp"
 #include "../../common/defs.hpp"
 #include "../../common/hymap.hpp"
 #include "../../common/integral_constant.hpp"
-#include "../../common/ldg_ptr.hpp"
 #include "../../common/tuple_util.hpp"
 #include "../../meta.hpp"
 #include "../../sid/allocator.hpp"
@@ -136,7 +136,7 @@ namespace gridtools {
                 template <class Key, class T>
                 GT_FUNCTION std::enable_if_t<meta::st_contains<Keys, Key>::value, T> operator()(
                     Key, T const *ptr) const {
-                    return *as_ldg_ptr(ptr);
+                    return const_ptr_deref(ptr);
                 }
                 template <class Key, class Ptr>
                 GT_FUNCTION decltype(auto) operator()(Key, Ptr ptr) const {

diff --git a/include/gridtools/stencil/gpu_horizontal/entry_point.hpp b/include/gridtools/stencil/gpu_horizontal/entry_point.hpp
@@ -12,6 +12,7 @@
 #include <type_traits>
 #include <utility>
 
+#include "../../common/const_ptr_deref.hpp"
 #include "../../common/cuda_type_traits.hpp"
 #include "../../common/cuda_util.hpp"
 #include "../../common/defs.hpp"
@@ -20,7 +21,6 @@
 #include "../../common/host_device.hpp"
 #include "../../common/hymap.hpp"
 #include "../../common/integral_constant.hpp"
-#include "../../common/ldg_ptr.hpp"
 #include "../../common/tuple_util.hpp"
 #include "../../meta.hpp"
 #include "../../sid/as_const.hpp"
@@ -45,7 +45,7 @@ namespace gridtools {
                 template <class Key, class T>
                 GT_FUNCTION std::enable_if_t<meta::st_contains<Keys, Key>::value, T> operator()(
                     Key, T const *ptr) const {
-                    return *as_ldg_ptr(ptr);
+                    return const_ptr_deref(ptr);
                 }
                 template <class Key, class Ptr>
                 GT_FUNCTION decltype(auto) operator()(Key, Ptr ptr) const {

diff --git a/include/gridtools/storage/sid.hpp b/include/gridtools/storage/sid.hpp
@@ -18,7 +18,6 @@
 #include "../common/hymap.hpp"
 #include "../common/integral_constant.hpp"
 #include "../common/layout_map.hpp"
-#include "../common/ldg_ptr.hpp"
 #include "../common/tuple.hpp"
 #include "../common/tuple_util.hpp"
 #include "../meta.hpp"
@@ -37,7 +36,7 @@ namespace gridtools {
             template <class T>
             struct ptr_holder {
                 T *m_val;
-                GT_FUNCTION constexpr auto operator()() const { return as_ldg_ptr(m_val); }
+                GT_FUNCTION constexpr T *operator()() const { return m_val; }
 
                 friend GT_FORCE_INLINE constexpr ptr_holder operator+(ptr_holder obj, int_t arg) {
                     return {obj.m_val + arg};

diff --git a/tests/unit_tests/common/CMakeLists.txt b/tests/unit_tests/common/CMakeLists.txt
@@ -17,7 +17,6 @@ gridtools_add_unit_test(test_gt_math SOURCES test_gt_math.cpp NO_NVCC)
 gridtools_add_unit_test(test_hypercube_iterator SOURCES test_hypercube_iterator.cpp NO_NVCC)
 gridtools_add_unit_test(test_tuple SOURCES test_tuple.cpp NO_NVCC)
 gridtools_add_unit_test(test_int_vector SOURCES test_int_vector.cpp NO_NVCC)
-gridtools_add_unit_test(test_ldg_ptr SOURCES test_ldg_ptr.cpp NO_NVCC)
 
 if(TARGET _gridtools_cuda)
     gridtools_check_compilation(test_cuda_type_traits test_cuda_type_traits.cu)
@@ -50,8 +49,4 @@ if(TARGET _gridtools_cuda)
             SOURCES test_tuple.cu
             LIBRARIES _gridtools_cuda
             LABELS cuda)
-    gridtools_add_unit_test(test_ldg_ptr_cuda
-            SOURCES test_ldg_ptr.cu
-            LIBRARIES _gridtools_cuda
-            LABELS cuda)
 endif()