removed unnecessary cmake option file, fixed bug in forall.hpp, added upper and lower bounds tests.

Author: rrsettgast

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

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

cmake/ChaiOptions.cmake

@@ -208,7 +208,7 @@ template<typename T>
 template<typename Idx>
 CHAI_INLINE
 CHAI_HOST_DEVICE T& ManagedArray<T>::operator[](const Idx i) const {
-#if defined(CHAI_ARRAY_BOUNDS_CHECK)
+#if defined(CHAI_ENABLE_BOUNDS_CHECK)
   assert( i>=0 && static_cast<size_t>(i) < m_elems );
 #endif
   return m_active_pointer[i];

src/ManagedArray.inl

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

src/config.hpp.in

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

src/tests/managed_array_tests.cpp

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