Assert on out-of-range accesses

CMakeLists.txt

@@ -53,6 +53,7 @@ option(ENABLE_IMPLICIT_CONVERSIONS "Enable implicit conversions to-from raw poin
 option(DISABLE_RM "Make ManagedArray a thin wrapper" Off)
 mark_as_advanced(DISABLE_RM)
 option(ENABLE_UM "Use CUDA unified (managed) memory" Off)
+option(ENABLE_BOUNDS_CHECK "Enable bounds checking for chai::ManagedArray<T>::operator[]" Off)
 
 set(ENABLE_TESTS On CACHE BOOL "")
 set(ENABLE_EXAMPLES On CACHE BOOL "")

src/chai/CMakeLists.txt

@@ -47,6 +47,7 @@ set(CHAI_ENABLE_HIP ${ENABLE_HIP})
 set(CHAI_ENABLE_IMPLICIT_CONVERSIONS ${ENABLE_IMPLICIT_CONVERSIONS})
 set(CHAI_DISABLE_RM ${DISABLE_RM})
 set(CHAI_ENABLE_UM ${ENABLE_UM})
+set(CHAI_ENABLE_BOUNDS_CHECK ${ENABLE_BOUNDS_CHECK})
 
 configure_file(
   ${PROJECT_SOURCE_DIR}/src/chai/config.hpp.in

src/chai/ManagedArray.inl

@@ -46,6 +46,16 @@
 #include "ManagedArray.hpp"
 #include "ArrayManager.hpp"
 
+#if defined(CHAI_ENABLE_BOUNDS_CHECK)
+
+#if defined(NDEBUG)
+#undef NDEBUG // This makes sure assert is enabled even for release builds
+#endif // defined(NDEBUG)
+
+#include <cassert>
+
+#endif // defined(CHAI_ENABLE_BOUNDS_CHECK)
+
 namespace chai {
 
 template<typename T>
@@ -389,6 +399,10 @@ template<typename T>
 template<typename Idx>
 CHAI_INLINE
 CHAI_HOST_DEVICE T& ManagedArray<T>::operator[](const Idx i) const {
+#if defined(CHAI_ENABLE_BOUNDS_CHECK)
+  assert(i >= 0 && static_cast<size_t>(i) < m_elems);
+#endif
+
   return m_active_pointer[i];
 }
 

src/chai/config.hpp.in

@@ -49,5 +49,6 @@
 #cmakedefine CHAI_ENABLE_IMPLICIT_CONVERSIONS
 #cmakedefine CHAI_DISABLE_RM
 #cmakedefine CHAI_ENABLE_UM
+#cmakedefine CHAI_ENABLE_BOUNDS_CHECK
 
 #endif // CHAI_config_HPP

src/util/forall.hpp

@@ -62,7 +62,7 @@ template <typename LOOP_BODY>
 void forall_kernel_cpu(int begin, int end, LOOP_BODY body)
 {
   for (int i = 0; i < (end - begin); ++i) {
-    body(i);
+    body(begin+i);
   }
 }
 

tests/integration/managed_array_tests.cpp

@@ -1394,4 +1394,51 @@ GPU_TEST(ManagedArray, CopyZero)
 
   array.free();
 }
-#endif
+#endif
+
+#if defined(CHAI_ENABLE_BOUNDS_CHECK)
+
+TEST(ManagedArray, UpperOutOfRangeAccess)
+{
+   chai::ManagedArray<float> array(20);
+   array[19] = 0.0; // Should be fine
+   ASSERT_DEATH_IF_SUPPORTED( array[20] = 0.0, ".*" );
+}
+
+TEST(ManagedArray, LowerOutOfRangeAccess)
+{
+   chai::ManagedArray<float> array(20);
+   array[0] = 0.0; // Should be fine
+   ASSERT_DEATH_IF_SUPPORTED( array[-1] = 0.0, ".*" );
+}
+
+#if defined(CHAI_ENABLE_CUDA)
+
+GPU_TEST(ManagedArray, UpperOutOfRangeAccessGPU)
+{
+   chai::ManagedArray<float> array(20);
+
+   forall(gpu(), 0, 1, [=] __device__ (int) {
+      array[20] = 0.0;
+   });
+
+   cudaError_t errorCode = cudaGetLastError();
+   ASSERT_EQ(cudaErrorAssert, errorCode);
+   ASSERT_EQ(cudaDeviceReset(), cudaSuccess);
+}
+
+GPU_TEST(ManagedArray, LowerOutOfRangeAccessGPU)
+{
+   chai::ManagedArray<float> array(20);
+
+   forall(gpu(), 0, 1, [=] __device__ (int) {
+      array[-1] = 0.0;
+   });
+
+   cudaError_t errorCode = cudaGetLastError();
+   ASSERT_EQ(cudaErrorAssert, errorCode);
+   ASSERT_EQ(cudaDeviceReset(), cudaSuccess);
+}
+
+#endif // defined(CHAI_ENABLE_CUDA)
+#endif // defined(CHAI_ENABLE_BOUNDS_CHECK)