Assert on out-of-range accesses

CMakeLists.txt

@@ -71,6 +71,9 @@ endif()
 set(ENABLE_COPY_HEADERS On CACHE BOOL "")
 set(BLT_CXX_STD c++11 CACHE STRING "")
 
+if( CMAKE_BUILD_TYPE STREQUAL "Debug" )
+  set( CHAI_ENABLE_BOUNDS_CHECK ON CACHE BOOL "Enable Bounds Checking for chai::ManagedArray<>::operator[]" FORCE )
+endif()
 
 ################################
 # BLT

cmake/ChaiBasics.cmake

@@ -43,3 +43,4 @@
 
 set (CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --expt-extended-lambda")
 include(${CMAKE_SOURCE_DIR}/cmake/thirdparty/SetupChaiThirdparty.cmake)
+

src/ManagedArray.inl

@@ -46,6 +46,8 @@
 #include "ManagedArray.hpp"
 #include "ArrayManager.hpp"
 
+#include <cassert>
+
 namespace chai {
 
 template<typename T>
@@ -206,6 +208,9 @@ 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/config.hpp.in

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

src/tests/managed_array_tests.cpp

@@ -393,20 +393,44 @@ CUDA_TEST(ManagedArray, UserCallback)
 #endif
 #endif
 
+#if defined(CHAI_ENABLE_BOUNDS_CHECK)
+TEST(ManagedArray, UpperOutOfRangeAccess)
+{
+  chai::ManagedArray<float> array(20);
 
-#if defined(CHAI_ENABLE_CUDA)
-CUDA_TEST(ManagedArray, Move)
+  array[19] = 0.0;
+  EXPECT_DEATH_IF_SUPPORTED( array[20] = 0.0, ".*" );
+}
+
+TEST(ManagedArray, LowerOutOfRangeAccess)
 {
-  chai::ManagedArray<float> array(10, chai::GPU);
+  chai::ManagedArray<float> array(20);
 
-  forall(cuda(), 0, 10, [=] __device__ (int i) {
-      array[i] = i;
-  });
+  array[0] = 0.0;
+  EXPECT_DEATH_IF_SUPPORTED( array[-1] = 0.0, ".*" );
+}
 
-  array.move(chai::CPU);
+#if defined(CHAI_ENABLE_CUDA)
+CUDA_TEST(ManagedArray, UpperOutOfRangeAccessGPU)
+{
+  chai::ManagedArray<float> array(20);
 
-  ASSERT_EQ(array[5], 5);
+  EXPECT_DEATH(
+               forall(cuda(), 19, 20, [=] __device__ (int i) {
+                   array[i] = 0.0f;
+                 });,
+               "i > m_elems");
+}
 
-  array.free();
+CUDA_TEST(ManagedArray, LowerOutOfRangeAccessGPU)
+{
+  chai::ManagedArray<float> array(20);
+
+  EXPECT_DEATH(
+               forall(cuda(), -1, 0, [=] __device__ (int i) {
+                   array[i] = 0.0f;
+                 });,      
+               "i < 0");
 }
 #endif // defined(CHAI_ENABLE_CUDA)
+#endif // !defined(NDEBUG)

src/util/forall.hpp

@@ -60,7 +60,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);
   }
 }