additional constrains for layout_stride_relaxed (#8354)

* layout_stride_relaxed additional constrains

* use cuda::std::in_range

Author: fbusato

libcudacxx/include/cuda/__mdspan/layout_stride_relaxed.h

@@ -25,7 +25,6 @@
 #include <cuda/__fwd/mdspan.h>
 #include <cuda/__numeric/add_overflow.h>
 #include <cuda/__numeric/mul_overflow.h>
-#include <cuda/__numeric/overflow_cast.h>
 #include <cuda/std/__concepts/concept_macros.h>
 #include <cuda/std/__cstddef/types.h>
 #include <cuda/std/__mdspan/concepts.h>
@@ -132,7 +131,7 @@ class _CCCL_DECLSPEC_EMPTY_BASES layout_stride_relaxed::mapping
     {
       for (rank_type __d = 0; __d < __rank_; ++__d)
       {
-        _CCCL_ASSERT(!::cuda::overflow_cast<offset_type>(__other.stride(__d)),
+        _CCCL_ASSERT(::cuda::std::in_range<offset_type>(__other.stride(__d)),
                      "layout_stride_relaxed::mapping: stride is out of range");
         __init_strides[__d] = static_cast<offset_type>(__other.stride(__d));
       }
@@ -261,25 +260,44 @@ class _CCCL_DECLSPEC_EMPTY_BASES layout_stride_relaxed::mapping
       // The dot product of indices and strides is linear.
       // Thus, over all valid indices, the max value of the dot product is achieved at the extrema: either the min
       // index (0) if the stride is negative, or the max index (extent(r) - 1) if the stride is non-negative.
-      // For non-negative stride: contribution is (extent - 1) * stride
-      // For negative stride: contribution is 0 (max achieved at index 0)
+      // For non-negative stride: max contribution is (extent - 1) * stride, min contribution is 0
+      // For negative stride: max contribution is 0 (max achieved at index 0), min is -(extent - 1) * |stride|
+      // __min_dot tracks the total positive magnitude of the negative contributions
       index_type __dot{1};
+      offset_type __min_dot{0};
       for (rank_type __r = 0; __r < __rank_; ++__r)
       {
         const auto __ext = extents().extent(__r);
         if (__ext == index_type{0})
         {
           return index_type{0};
         }
-        _CCCL_ASSERT(!::cuda::overflow_cast<index_type>(::cuda::uabs(strides().stride(__r))),
+        const auto __stride_val = strides().stride(__r);
+        _CCCL_ASSERT(::cuda::std::in_range<index_type>(::cuda::uabs(__stride_val)),
                      "layout_stride_relaxed::mapping: stride is out of range");
-        const auto __max_index = strides().stride(__r) < 0 ? index_type{0} : static_cast<index_type>(__ext - 1);
-        const auto __stride    = static_cast<index_type>(strides().stride(__r));
+        if (__stride_val < 0)
+        {
+          _CCCL_ASSERT(::cuda::std::in_range<offset_type>(__ext - 1),
+                       "layout_stride_relaxed::mapping: extent - 1 is not representable as offset_type");
+          const auto __min_extent   = static_cast<offset_type>(__ext - 1);
+          const auto __abs_stride_u = ::cuda::uabs(__stride_val);
+          _CCCL_ASSERT(::cuda::std::in_range<offset_type>(__abs_stride_u),
+                       "layout_stride_relaxed::mapping: absolute stride is not representable as offset_type");
+          const auto __abs_stride = static_cast<offset_type>(__abs_stride_u);
+          _CCCL_ASSERT(!::cuda::mul_overflow(__min_extent, __abs_stride)
+                         && !::cuda::add_overflow(__min_extent * __abs_stride, __min_dot),
+                       "layout_stride_relaxed::mapping: minimum mapped index is not representable");
+          __min_dot += __min_extent * __abs_stride;
+        }
+        const auto __max_index = __stride_val < 0 ? index_type{0} : static_cast<index_type>(__ext - 1);
+        const auto __stride    = static_cast<index_type>(__stride_val);
         _CCCL_ASSERT(!::cuda::mul_overflow<index_type>(__max_index, __stride)
                        && !::cuda::add_overflow(__max_index * __stride, __dot),
                      "layout_stride_relaxed::mapping: required_span_size is not representable as index_type");
         __dot += __max_index * __stride;
       }
+      _CCCL_ASSERT(::cuda::std::cmp_greater_equal(__offset_val, __min_dot),
+                   "layout_stride_relaxed::mapping: offset is insufficient for negative strides");
       _CCCL_ASSERT(!::cuda::add_overflow<index_type>(__offset_val, __dot),
                    "layout_stride_relaxed::mapping: required_span_size is not representable as index_type");
       return static_cast<index_type>(__offset_val + __dot);
@@ -291,7 +309,7 @@ class _CCCL_DECLSPEC_EMPTY_BASES layout_stride_relaxed::mapping
   {
     if constexpr (::cuda::std::__cccl_is_integer_v<_Index>)
     {
-      return ::cuda::std::cmp_greater_equal(__index, index_type{0}) && !::cuda::overflow_cast<index_type>(__index);
+      return ::cuda::std::cmp_greater_equal(__index, index_type{0}) && ::cuda::std::in_range<index_type>(__index);
     }
     else
     {

libcudacxx/include/cuda/__mdspan/strides.h

@@ -22,7 +22,6 @@
 #endif // no system header
 
 #include <cuda/__fwd/mdspan.h>
-#include <cuda/__numeric/overflow_cast.h>
 #include <cuda/std/__concepts/concept_macros.h>
 #include <cuda/std/__cstddef/types.h>
 #include <cuda/std/__mdspan/extents.h>
@@ -67,7 +66,7 @@ class strides
   [[nodiscard]] _CCCL_API static constexpr bool __is_representable_as(_From... __values) noexcept
   {
     return (
-      (!::cuda::overflow_cast<offset_type>(__values) || static_cast<::cuda::std::ptrdiff_t>(__values) == dynamic_stride)
+      (::cuda::std::in_range<offset_type>(__values) || static_cast<::cuda::std::ptrdiff_t>(__values) == dynamic_stride)
       && ...);
   }
 

libcudacxx/test/libcudacxx/cuda/containers/views/mdspan/layout_stride_relaxed/assertions.pass.cpp

@@ -63,10 +63,45 @@ void test_strides_narrowing_assertion()
   TEST_CCCL_ASSERT_FAILURE((small_strides(too_big)), "strides construction: stride is out of range");
 }
 
+void test_insufficient_offset_for_negative_strides()
+{
+  // 1D: extent=4, stride=-1 requires offset >= 3
+  {
+    using extents_t = cuda::std::extents<int, 4>;
+    using mapping_t = cuda::layout_stride_relaxed::mapping<extents_t>;
+    using strides_t = typename mapping_t::strides_type;
+    using offset_t  = typename mapping_t::offset_type;
+
+    TEST_CCCL_ASSERT_FAILURE(mapping_t(extents_t{}, strides_t(-1), static_cast<offset_t>(2)),
+                             "layout_stride_relaxed::mapping: offset is insufficient for negative strides");
+  }
+  // 2D: extents(3,4), strides(4,-1) requires offset >= (4-1)*1 = 3
+  {
+    using extents_t = cuda::std::extents<int, 3, 4>;
+    using mapping_t = cuda::layout_stride_relaxed::mapping<extents_t>;
+    using strides_t = typename mapping_t::strides_type;
+    using offset_t  = typename mapping_t::offset_type;
+
+    TEST_CCCL_ASSERT_FAILURE(mapping_t(extents_t{}, strides_t(4, -1), static_cast<offset_t>(2)),
+                             "layout_stride_relaxed::mapping: offset is insufficient for negative strides");
+  }
+  // 2D: extents(3,4), strides(-2,-1) requires offset >= (3-1)*2 + (4-1)*1 = 7
+  {
+    using extents_t = cuda::std::extents<int, 3, 4>;
+    using mapping_t = cuda::layout_stride_relaxed::mapping<extents_t>;
+    using strides_t = typename mapping_t::strides_type;
+    using offset_t  = typename mapping_t::offset_type;
+
+    TEST_CCCL_ASSERT_FAILURE(mapping_t(extents_t{}, strides_t(-2, -1), static_cast<offset_t>(6)),
+                             "layout_stride_relaxed::mapping: offset is insufficient for negative strides");
+  }
+}
+
 int main(int, char**)
 {
   test_negative_offset_assertion();
   test_static_stride_comparison();
   test_strides_narrowing_assertion();
+  test_insufficient_offset_for_negative_strides();
   return 0;
 }

libcudacxx/test/libcudacxx/cuda/containers/views/mdspan/layout_stride_relaxed/required_span_size.pass.cpp

@@ -172,6 +172,22 @@ __host__ __device__ constexpr bool test()
   test_required_span_size(cuda::std::extents<int, D, D, D>(5, 0, 3), cuda::std::array<intptr_t, 3>{-1, 0, 1}, 4, 0);
   test_required_span_size(cuda::std::extents<int, D, D, D>(5, 3, 0), cuda::std::array<intptr_t, 3>{-1, 0, 1}, 4, 0);
 
+  // ============================================================================
+  // Minimum valid offset for negative strides (boundary cases)
+  // ============================================================================
+
+  // 1D: extent=4, stride=-1 requires offset >= (4-1)*1 = 3
+  // mapped values: 3, 2, 1, 0 → required_span_size = 3 + 1 = 4
+  test_required_span_size(cuda::std::extents<int, 4>(), cuda::std::array<intptr_t, 1>{-1}, 3, 4);
+
+  // 2D: extents(3,4), strides(4,-1) requires offset >= (4-1)*1 = 3 (only dim 1 is negative)
+  // max mapped = 3 + (3-1)*4 + 0 = 11 → required_span_size = 12
+  test_required_span_size(cuda::std::extents<int, 3, 4>(), cuda::std::array<intptr_t, 2>{4, -1}, 3, 12);
+
+  // 2D: extents(3,4), strides(-2,-1), both negative, requires offset >= (3-1)*2 + (4-1)*1 = 7
+  // max mapped = 7 + 0 + 0 = 7 → required_span_size = 7 + 1 = 8
+  test_required_span_size(cuda::std::extents<int, 3, 4>(), cuda::std::array<intptr_t, 2>{-2, -1}, 7, 8);
+
   return true;
 }