From 5a1c7b542f11568b186f24773f832497e55e7b23 Mon Sep 17 00:00:00 2001
From: Felix Thaler <thaler@cscs.ch>
Date: Wed, 24 Jul 2024 07:37:14 +0200
Subject: [PATCH] Use __ldg on scalar read-only data stores

---
 include/gridtools/storage/sid.hpp | 38 ++++++++++++++++++++++++++++++-
 1 file changed, 37 insertions(+), 1 deletion(-)

diff --git a/include/gridtools/storage/sid.hpp b/include/gridtools/storage/sid.hpp
index 7ee0807c2..a3819f27e 100644
--- a/include/gridtools/storage/sid.hpp
+++ b/include/gridtools/storage/sid.hpp
@@ -23,6 +23,10 @@
 #include "../meta.hpp"
 #include "data_store.hpp"
 
+#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350
+#include "../common/cuda_type_traits.hpp"
+#endif
+
 namespace gridtools {
     namespace storage {
         namespace storage_sid_impl_ {
@@ -32,11 +36,43 @@ namespace gridtools {
                     return lhs;
                 }
             };
+            template <class T>
+            struct const_ptr_wrapper {
+                T const *m_val;
+
+                GT_FUNCTION constexpr const_ptr_wrapper(T const *val) : m_val(val) {}
+
+                GT_FUNCTION constexpr T operator*() const {
+#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350
+                    if constexpr (is_texture_type<T>::value)
+                        return __ldg(m_val);
+#endif
+                    return *m_val;
+                }
+                GT_FUNCTION constexpr const_ptr_wrapper &operator+=(int_t arg) {
+                    m_val += arg;
+                    return *this;
+                }
+                GT_FUNCTION constexpr const_ptr_wrapper &operator-=(int_t arg) {
+                    m_val -= arg;
+                    return *this;
+                }
+                friend GT_FUNCTION constexpr const_ptr_wrapper operator+(const_ptr_wrapper obj, int_t arg) {
+                    return {obj.m_val + arg};
+                }
+                friend GT_FUNCTION constexpr const_ptr_wrapper operator-(const_ptr_wrapper obj, int_t arg) {
+                    return {obj.m_val - arg};
+                }
+            };
 
             template <class T>
             struct ptr_holder {
                 T *m_val;
-                GT_FUNCTION constexpr T *operator()() const { return m_val; }
+                GT_FUNCTION constexpr std::
+                    conditional_t<std::is_const_v<T>, const_ptr_wrapper<std::remove_const_t<T>>, 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};