From e1f32c43fa01d1a5e94419a18e6a7640f7a45399 Mon Sep 17 00:00:00 2001 From: David Beckingsale Date: Thu, 24 May 2018 10:15:20 -0700 Subject: [PATCH 01/16] Add assert for out of range --- src/ManagedArray.inl | 5 +++++ src/tests/managed_array_tests.cpp | 23 +++++++++++++++++++++++ 2 files changed, 28 insertions(+) diff --git a/src/ManagedArray.inl b/src/ManagedArray.inl index 52afae55..43d132c1 100644 --- a/src/ManagedArray.inl +++ b/src/ManagedArray.inl @@ -46,6 +46,8 @@ #include "ManagedArray.hpp" #include "ArrayManager.hpp" +#include + namespace chai { template @@ -175,6 +177,9 @@ template template CHAI_INLINE CHAI_HOST_DEVICE T& ManagedArray::operator[](const Idx i) const { +#if !defined(NDEBUG) + HOST_DEVICE_ASSERT(i < m_elems); +#endif return m_active_pointer[i]; } diff --git a/src/tests/managed_array_tests.cpp b/src/tests/managed_array_tests.cpp index b801b753..32a2759c 100644 --- a/src/tests/managed_array_tests.cpp +++ b/src/tests/managed_array_tests.cpp @@ -393,3 +393,26 @@ CUDA_TEST(ManagedArray, UserCallback) #endif #endif +#if !defined(NDEBUG) +TEST(ManagedArray, OutOfRangeAccess) +{ + chai::ManagedArray array(20); + + EXPECT_DEATH( + forall(sequential(), 10, 50, [=] (int i) { + array[i] = 0.0f; + });, + "i < m_elems"); +} + +#if defined(CHAI_ENABLE_CUDA) +CUDA_TEST(ManagedArray, OutOfRangeAccessGPU) +{ + chai::ManagedArray array(20); + + forall(cuda(), 10, 50, [=] __device__ (int i) { + array[i] = 0.0f; + }); +} +#endif // defined(CHAI_ENABLE_CUDA) +#endif // !defined(NDEBUG) From a7bde928063ca0c1f3ba79ed32b614dd9b89ff49 Mon Sep 17 00:00:00 2001 From: "Randolph R. Settgast" Date: Thu, 7 Jun 2018 11:15:43 -0700 Subject: [PATCH 02/16] added compile time option to perform range checking on ManagedArray --- cmake/ChaiBasics.cmake | 2 ++ cmake/ChaiOptions.cmake | 2 ++ src/ManagedArray.inl | 5 ++++- src/config.hpp.in | 1 + src/tests/managed_array_tests.cpp | 12 ++++++++++++ 5 files changed, 21 insertions(+), 1 deletion(-) create mode 100644 cmake/ChaiOptions.cmake diff --git a/cmake/ChaiBasics.cmake b/cmake/ChaiBasics.cmake index de46b224..98a24975 100644 --- a/cmake/ChaiBasics.cmake +++ b/cmake/ChaiBasics.cmake @@ -43,3 +43,5 @@ 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") diff --git a/cmake/ChaiOptions.cmake b/cmake/ChaiOptions.cmake new file mode 100644 index 00000000..510fd1c1 --- /dev/null +++ b/cmake/ChaiOptions.cmake @@ -0,0 +1,2 @@ + +option( CHAI_ARRAY_BOUNDS_CHECK "Enables Bounds Checking for chai::ManagedArray<>::operator[]" OFF ) diff --git a/src/ManagedArray.inl b/src/ManagedArray.inl index 4c9b98c1..93552f9a 100644 --- a/src/ManagedArray.inl +++ b/src/ManagedArray.inl @@ -117,7 +117,7 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray(ManagedArray const& other): */ if (!std::is_const::value) { CHAI_LOG("ManagedArray", "T is non-const, registering touch of pointer" << m_active_pointer); - T_non_const* non_const_pointer = const_cast(m_active_pointer); +// T_non_const* non_const_pointer = const_cast(m_active_pointer); m_resource_manager->registerTouch(m_pointer_record); } #endif @@ -207,6 +207,9 @@ template template CHAI_INLINE CHAI_HOST_DEVICE T& ManagedArray::operator[](const Idx i) const { +#if defined(CHAI_ARRAY_BOUNDS_CHECK) + assert( i>=0 && static_cast(i) < m_elems ); +#endif return m_active_pointer[i]; } diff --git a/src/config.hpp.in b/src/config.hpp.in index 75610721..810e48af 100644 --- a/src/config.hpp.in +++ b/src/config.hpp.in @@ -47,5 +47,6 @@ #cmakedefine CHAI_ENABLE_IMPLICIT_CONVERSIONS #cmakedefine CHAI_DISABLE_RM #cmakedefine CHAI_ENABLE_UM +#cmakedefine CHAI_ARRAY_BOUNDS_CHECK #endif // CHAI_config_HPP diff --git a/src/tests/managed_array_tests.cpp b/src/tests/managed_array_tests.cpp index c47d58cc..113f5b8a 100644 --- a/src/tests/managed_array_tests.cpp +++ b/src/tests/managed_array_tests.cpp @@ -238,6 +238,9 @@ TEST(ManagedArray, ImplicitConversions) { chai::ManagedArray a2 = a; + forall(sequential(), 0, 10, [=] (int i) { + ASSERT_EQ( a[i], a2[i]); + }); SUCCEED(); } #endif @@ -410,3 +413,12 @@ CUDA_TEST(ManagedArray, Move) array.free(); } #endif // defined(CHAI_ENABLE_CUDA) + +#if defined(CHAI_ARRAY_BOUNDS_CHECK) +TEST(ManagedArray_DeathTest, RangeCheck) { + chai::ManagedArray array(10); + + EXPECT_DEATH_IF_SUPPORTED( array[-1], ".*" ); + EXPECT_DEATH_IF_SUPPORTED( array[10], ".*" ); +} +#endif From 6a7b9299c05ec47195dceefa23255f732fa6104c Mon Sep 17 00:00:00 2001 From: Randolph Settgast Date: Fri, 8 Jun 2018 00:08:53 -0700 Subject: [PATCH 03/16] undo previous change to implicit conversion test --- src/tests/managed_array_tests.cpp | 3 --- 1 file changed, 3 deletions(-) diff --git a/src/tests/managed_array_tests.cpp b/src/tests/managed_array_tests.cpp index 113f5b8a..2eaa9019 100644 --- a/src/tests/managed_array_tests.cpp +++ b/src/tests/managed_array_tests.cpp @@ -238,9 +238,6 @@ TEST(ManagedArray, ImplicitConversions) { chai::ManagedArray a2 = a; - forall(sequential(), 0, 10, [=] (int i) { - ASSERT_EQ( a[i], a2[i]); - }); SUCCEED(); } #endif From a989c60f90f7127f542502c280b003b21572b5e7 Mon Sep 17 00:00:00 2001 From: David Beckingsale Date: Tue, 12 Jun 2018 13:38:05 -0700 Subject: [PATCH 04/16] Updates for test --- src/ManagedArray.inl | 5 +++-- src/tests/managed_array_tests.cpp | 9 ++++++--- 2 files changed, 9 insertions(+), 5 deletions(-) diff --git a/src/ManagedArray.inl b/src/ManagedArray.inl index 43d132c1..1fe08814 100644 --- a/src/ManagedArray.inl +++ b/src/ManagedArray.inl @@ -46,7 +46,8 @@ #include "ManagedArray.hpp" #include "ArrayManager.hpp" -#include +//#include +#include namespace chai { @@ -178,7 +179,7 @@ template CHAI_INLINE CHAI_HOST_DEVICE T& ManagedArray::operator[](const Idx i) const { #if !defined(NDEBUG) - HOST_DEVICE_ASSERT(i < m_elems); + assert((i >= 0) && (i < m_elems)); #endif return m_active_pointer[i]; } diff --git a/src/tests/managed_array_tests.cpp b/src/tests/managed_array_tests.cpp index 32a2759c..779cc849 100644 --- a/src/tests/managed_array_tests.cpp +++ b/src/tests/managed_array_tests.cpp @@ -410,9 +410,12 @@ CUDA_TEST(ManagedArray, OutOfRangeAccessGPU) { chai::ManagedArray array(20); - forall(cuda(), 10, 50, [=] __device__ (int i) { - array[i] = 0.0f; - }); + EXPECT_DEATH( + forall(cuda(), 10, 50, [=] __device__ (int i) { + array[i] = 0.0f; + }); + , + "i < m_elems"); } #endif // defined(CHAI_ENABLE_CUDA) #endif // !defined(NDEBUG) From 5bbb458af3fddcecff8c6cd9f54d9cf206c61baf Mon Sep 17 00:00:00 2001 From: Arturo Vargas Date: Tue, 12 Jun 2018 14:08:12 -0700 Subject: [PATCH 05/16] added cassert header --- src/ManagedArray.inl | 1 + 1 file changed, 1 insertion(+) diff --git a/src/ManagedArray.inl b/src/ManagedArray.inl index 93552f9a..b71152f6 100644 --- a/src/ManagedArray.inl +++ b/src/ManagedArray.inl @@ -45,6 +45,7 @@ #include "ManagedArray.hpp" #include "ArrayManager.hpp" +#include namespace chai { From 908df3ff939f3054c10692ac10316aaf1361b194 Mon Sep 17 00:00:00 2001 From: Arturo Vargas Date: Wed, 13 Jun 2018 09:04:49 -0700 Subject: [PATCH 06/16] reverted managed_array_test.cpp and added --- src/ManagedArray.inl | 2 +- src/tests/managed_array_tests.cpp | 38 +++++++++++++++---------------- 2 files changed, 20 insertions(+), 20 deletions(-) diff --git a/src/ManagedArray.inl b/src/ManagedArray.inl index de06f565..db543b48 100644 --- a/src/ManagedArray.inl +++ b/src/ManagedArray.inl @@ -46,7 +46,7 @@ #include "ManagedArray.hpp" #include "ArrayManager.hpp" -#include +#include namespace chai { diff --git a/src/tests/managed_array_tests.cpp b/src/tests/managed_array_tests.cpp index e5d3c05d..8015648e 100644 --- a/src/tests/managed_array_tests.cpp +++ b/src/tests/managed_array_tests.cpp @@ -393,28 +393,28 @@ CUDA_TEST(ManagedArray, UserCallback) #endif #endif -#if defined(CHAI_ENABLE_CUDA) -CUDA_TEST(ManagedArray, Move) +#if !defined(NDEBUG) +TEST(ManagedArray, OutOfRangeAccess) { - chai::ManagedArray array(10, chai::GPU); - - forall(cuda(), 0, 10, [=] __device__ (int i) { - array[i] = i; - }); - - array.move(chai::CPU); - - ASSERT_EQ(array[5], 5); + chai::ManagedArray array(20); - array.free(); + EXPECT_DEATH( + forall(sequential(), 10, 50, [=] (int i) { + array[i] = 0.0f; + });, + "i < m_elems"); } -#endif // defined(CHAI_ENABLE_CUDA) -#if defined(CHAI_ARRAY_BOUNDS_CHECK) -TEST(ManagedArray_DeathTest, RangeCheck) { - chai::ManagedArray array(10); +#if defined(CHAI_ENABLE_CUDA) +CUDA_TEST(ManagedArray, OutOfRangeAccessGPU) +{ + chai::ManagedArray array(20); - EXPECT_DEATH_IF_SUPPORTED( array[-1], ".*" ); - EXPECT_DEATH_IF_SUPPORTED( array[10], ".*" ); + EXPECT_DEATH( + forall(cuda(), 10, 50, [=] __device__ (int i) { + array[i] = 0.0f; + });, + "i < m_elems"); } -#endif +#endif // defined(CHAI_ENABLE_CUDA) +#endif // !defined(NDEBUG) From f38337d2f4d348b0bf02fd7b38c5af7342a7f96d Mon Sep 17 00:00:00 2001 From: "Randolph R. Settgast" Date: Thu, 14 Jun 2018 16:27:17 -0700 Subject: [PATCH 07/16] removed unnecessary cmake option file, fixed bug in forall.hpp, added upper and lower bounds tests. --- CMakeLists.txt | 3 +++ cmake/ChaiBasics.cmake | 1 - cmake/ChaiOptions.cmake | 2 -- src/ManagedArray.inl | 2 +- src/config.hpp.in | 2 +- src/tests/managed_array_tests.cpp | 32 +++++++++++++++++++++++-------- src/util/forall.hpp | 2 +- 7 files changed, 30 insertions(+), 14 deletions(-) delete mode 100644 cmake/ChaiOptions.cmake diff --git a/CMakeLists.txt b/CMakeLists.txt index 4a7a29a8..72494e7a 100644 --- a/CMakeLists.txt +++ b/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 diff --git a/cmake/ChaiBasics.cmake b/cmake/ChaiBasics.cmake index 98a24975..951f59b6 100644 --- a/cmake/ChaiBasics.cmake +++ b/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") diff --git a/cmake/ChaiOptions.cmake b/cmake/ChaiOptions.cmake deleted file mode 100644 index 510fd1c1..00000000 --- a/cmake/ChaiOptions.cmake +++ /dev/null @@ -1,2 +0,0 @@ - -option( CHAI_ARRAY_BOUNDS_CHECK "Enables Bounds Checking for chai::ManagedArray<>::operator[]" OFF ) diff --git a/src/ManagedArray.inl b/src/ManagedArray.inl index 7bb8409a..71871735 100644 --- a/src/ManagedArray.inl +++ b/src/ManagedArray.inl @@ -208,7 +208,7 @@ template template CHAI_INLINE CHAI_HOST_DEVICE T& ManagedArray::operator[](const Idx i) const { -#if defined(CHAI_ARRAY_BOUNDS_CHECK) +#if defined(CHAI_ENABLE_BOUNDS_CHECK) assert( i>=0 && static_cast(i) < m_elems ); #endif return m_active_pointer[i]; diff --git a/src/config.hpp.in b/src/config.hpp.in index 810e48af..ece92b86 100644 --- a/src/config.hpp.in +++ b/src/config.hpp.in @@ -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 diff --git a/src/tests/managed_array_tests.cpp b/src/tests/managed_array_tests.cpp index 8015648e..56d5f9ac 100644 --- a/src/tests/managed_array_tests.cpp +++ b/src/tests/managed_array_tests.cpp @@ -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 array(20); + + array[19] = 0.0; + EXPECT_DEATH_IF_SUPPORTED( array[20] = 0.0, ".*" ); +} + +TEST(ManagedArray, LowerOutOfRangeAccess) +{ + chai::ManagedArray 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 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 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) diff --git a/src/util/forall.hpp b/src/util/forall.hpp index 7a6c043d..868db2bd 100644 --- a/src/util/forall.hpp +++ b/src/util/forall.hpp @@ -60,7 +60,7 @@ template void forall_kernel_cpu(int begin, int end, LOOP_BODY body) { for (int i = 0; i < (end - begin); ++i) { - body(i); + body(begin+i); } } From 6c55bf98e41b1540881cc8eeaf1ff047ff28bb28 Mon Sep 17 00:00:00 2001 From: Alan Dayton Date: Mon, 16 Sep 2019 12:28:49 -0700 Subject: [PATCH 08/16] Fix bad merge --- src/ManagedArray.inl | 297 -------------------------------------- src/chai/ManagedArray.inl | 6 + 2 files changed, 6 insertions(+), 297 deletions(-) delete mode 100644 src/ManagedArray.inl diff --git a/src/ManagedArray.inl b/src/ManagedArray.inl deleted file mode 100644 index 71871735..00000000 --- a/src/ManagedArray.inl +++ /dev/null @@ -1,297 +0,0 @@ -// --------------------------------------------------------------------- -// Copyright (c) 2016, Lawrence Livermore National Security, LLC. All -// rights reserved. -// -// Produced at the Lawrence Livermore National Laboratory. -// -// This file is part of CHAI. -// -// LLNL-CODE-705877 -// -// For details, see https:://github.com/LLNL/CHAI -// Please also see the NOTICE and LICENSE files. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions -// are met: -// -// - Redistributions of source code must retain the above copyright -// notice, this list of conditions and the following disclaimer. -// -// - Redistributions in binary form must reproduce the above copyright -// notice, this list of conditions and the following disclaimer in the -// documentation and/or other materials provided with the -// distribution. -// -// - Neither the name of the LLNS/LLNL nor the names of its contributors -// may be used to endorse or promote products derived from this -// software without specific prior written permission. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS -// "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT -// LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR -// A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT -// HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, -// INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, -// BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS -// OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED -// AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT -// LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY -// WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE -// POSSIBILITY OF SUCH DAMAGE. -// --------------------------------------------------------------------- -#ifndef CHAI_ManagedArray_INL -#define CHAI_ManagedArray_INL - -#include "ManagedArray.hpp" -#include "ArrayManager.hpp" - -#include - -namespace chai { - -template -CHAI_INLINE -CHAI_HOST_DEVICE ManagedArray::ManagedArray(): - m_active_pointer(nullptr), - m_resource_manager(nullptr), - m_elems(0), - m_pointer_record(nullptr) -{ -#if !defined(__CUDA_ARCH__) - m_resource_manager = ArrayManager::getInstance(); -#endif -} - -template -CHAI_INLINE -CHAI_HOST_DEVICE ManagedArray::ManagedArray( - size_t elems, ExecutionSpace space): - m_active_pointer(nullptr), - m_resource_manager(nullptr), - m_elems(elems), - m_pointer_record(nullptr) -{ -#if !defined(__CUDA_ARCH__) - m_resource_manager = ArrayManager::getInstance(); - this->allocate(elems, space); -#endif -} - -template -CHAI_INLINE -CHAI_HOST_DEVICE ManagedArray::ManagedArray(std::nullptr_t) : - m_active_pointer(nullptr), - m_resource_manager(nullptr), - m_elems(0), - m_pointer_record(nullptr) -{ -} - -template -CHAI_INLINE -CHAI_HOST_DEVICE ManagedArray::ManagedArray(PointerRecord* record, ExecutionSpace space): - m_active_pointer(static_cast(record->m_pointers[space])), - m_resource_manager(ArrayManager::getInstance()), - m_elems(record->m_size/sizeof(T)), - m_pointer_record(record) -{ -} - - -template -CHAI_INLINE -CHAI_HOST_DEVICE ManagedArray::ManagedArray(ManagedArray const& other): - m_active_pointer(other.m_active_pointer), - m_resource_manager(other.m_resource_manager), - m_elems(other.m_elems), - m_pointer_record(other.m_pointer_record) -{ -#if !defined(__CUDA_ARCH__) - CHAI_LOG("ManagedArray", "Moving " << m_active_pointer); - - m_active_pointer = static_cast(m_resource_manager->move(const_cast(m_active_pointer), m_pointer_record)); - - CHAI_LOG("ManagedArray", "Moved to " << m_active_pointer); - - /* - * Register touch - */ - if (!std::is_const::value) { - CHAI_LOG("ManagedArray", "T is non-const, registering touch of pointer" << m_active_pointer); - m_resource_manager->registerTouch(m_pointer_record); - } -#endif -} - -template -CHAI_INLINE -CHAI_HOST_DEVICE ManagedArray::ManagedArray(T* data, ArrayManager* array_manager, size_t elems, PointerRecord* pointer_record) : - m_active_pointer(data), - m_resource_manager(array_manager), - m_elems(elems), - m_pointer_record(pointer_record) -{ -} - -template -CHAI_INLINE -CHAI_HOST void ManagedArray::allocate(size_t elems, ExecutionSpace space, UserCallback const &cback) { - CHAI_LOG("ManagedArray", "Allocating array of size " << elems << " in space " << space); - - if (space == NONE) { - space = m_resource_manager->getDefaultAllocationSpace(); - } - - m_elems = elems; - m_pointer_record = m_resource_manager->allocate(elems, space, cback); - m_active_pointer = static_cast(m_pointer_record->m_pointers[space]); - - CHAI_LOG("ManagedArray", "m_active_ptr allocated at address: " << m_active_pointer); -} - -template -CHAI_INLINE -CHAI_HOST void ManagedArray::reallocate(size_t elems) -{ - CHAI_LOG("ManagedArray", "Reallocating array of size " << m_elems << " with new size" << elems); - - m_elems = elems; - m_active_pointer = - static_cast(m_resource_manager->reallocate(m_active_pointer, elems, - m_pointer_record)); - - CHAI_LOG("ManagedArray", "m_active_ptr reallocated at address: " << m_active_pointer); -} - -template -CHAI_INLINE -CHAI_HOST void ManagedArray::free() -{ - m_resource_manager->free(m_pointer_record); -} - -template -CHAI_INLINE -CHAI_HOST void ManagedArray::reset() -{ - m_resource_manager->resetTouch(m_pointer_record); -} - -template -CHAI_INLINE -CHAI_HOST size_t ManagedArray::size() const { - return m_elems; -} - -template -CHAI_INLINE -CHAI_HOST void ManagedArray::registerTouch(ExecutionSpace space) { - m_resource_manager->registerTouch(m_pointer_record, space); -} - -template -CHAI_INLINE -CHAI_HOST -void ManagedArray::move(ExecutionSpace space) -{ - m_active_pointer = static_cast(m_resource_manager->move(m_active_pointer, m_pointer_record, space)); - - if (!std::is_const::value) { - CHAI_LOG("ManagedArray", "T is non-const, registering touch of pointer" << m_active_pointer); - T_non_const* non_const_pointer = const_cast(m_active_pointer); - m_resource_manager->registerTouch(m_pointer_record); - } -} - -template -template -CHAI_INLINE -CHAI_HOST_DEVICE T& ManagedArray::operator[](const Idx i) const { -#if defined(CHAI_ENABLE_BOUNDS_CHECK) - assert( i>=0 && static_cast(i) < m_elems ); -#endif - return m_active_pointer[i]; -} - -#if defined(CHAI_ENABLE_IMPLICIT_CONVERSIONS) -template -CHAI_INLINE -CHAI_HOST_DEVICE ManagedArray::operator T*() const { -#if !defined(__CUDA_ARCH__) - ExecutionSpace prev_space = m_resource_manager->getExecutionSpace(); - m_resource_manager->setExecutionSpace(CPU); - auto non_const_active_pointer = const_cast(static_cast(m_active_pointer)); - m_active_pointer = static_cast(m_resource_manager->move(non_const_active_pointer, m_pointer_record)); - - m_resource_manager->registerTouch(m_pointer_record); - - - // Reset to whatever space we rode in on - m_resource_manager->setExecutionSpace(prev_space); - - return m_active_pointer; -#else - return m_active_pointer; -#endif -} - - -template -template -CHAI_INLINE -CHAI_HOST_DEVICE ManagedArray::ManagedArray(T* data, bool ) : - m_active_pointer(data), -#if !defined(__CUDA_ARCH__) - m_resource_manager(ArrayManager::getInstance()), - m_elems(m_resource_manager->getSize(m_active_pointer)), - m_pointer_record(m_resource_manager->getPointerRecord(data)) -#else - m_resource_manager(nullptr), - m_elems(0), - m_pointer_record(nullptr) -#endif -{ -} -#endif - -template -T* -ManagedArray::getActivePointer() const -{ - return m_active_pointer; -} - - -template -ManagedArray::operator ManagedArray< - typename std::conditional::value, - const T, - InvalidConstCast>::type> ()const -{ - return ManagedArray(const_cast(m_active_pointer), - m_resource_manager, m_elems, m_pointer_record); -} - -template -CHAI_INLINE -CHAI_HOST_DEVICE -ManagedArray& -ManagedArray::operator= (std::nullptr_t from) { - m_active_pointer = from; - m_elems = 0; - return *this; -} - -template -CHAI_INLINE -CHAI_HOST_DEVICE -bool -ManagedArray::operator== (ManagedArray& rhs) -{ - return (m_active_pointer == rhs.m_active_pointer); -} - -} // end of namespace chai - -#endif // CHAI_ManagedArray_INL diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index 18f4db7a..94939d56 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -46,6 +46,8 @@ #include "ManagedArray.hpp" #include "ArrayManager.hpp" +#include + namespace chai { template @@ -389,6 +391,10 @@ template template CHAI_INLINE CHAI_HOST_DEVICE T& ManagedArray::operator[](const Idx i) const { +#if defined(CHAI_ENABLE_BOUNDS_CHECK) + assert(i >= 0 && static_cast(i) < m_elems); +#endif + return m_active_pointer[i]; } From dfa4b556dfff5c0ceccc0f9c19f70e91d578df27 Mon Sep 17 00:00:00 2001 From: Alan Dayton Date: Mon, 16 Sep 2019 12:53:47 -0700 Subject: [PATCH 09/16] Add cmake option for bounds checking --- CMakeLists.txt | 1 + src/chai/CMakeLists.txt | 1 + 2 files changed, 2 insertions(+) diff --git a/CMakeLists.txt b/CMakeLists.txt index b10d9fb8..eff16368 100644 --- a/CMakeLists.txt +++ b/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::operator[]" Off) set(ENABLE_TESTS On CACHE BOOL "") set(ENABLE_EXAMPLES On CACHE BOOL "") diff --git a/src/chai/CMakeLists.txt b/src/chai/CMakeLists.txt index 7821fa3f..64577fd9 100644 --- a/src/chai/CMakeLists.txt +++ b/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 From deba461d6283130281348b2c9858cd9e26d8ae49 Mon Sep 17 00:00:00 2001 From: Alan Dayton Date: Mon, 16 Sep 2019 12:54:44 -0700 Subject: [PATCH 10/16] Remove unnecessary variable --- CMakeLists.txt | 4 ---- 1 file changed, 4 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index eff16368..9208f28d 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -77,10 +77,6 @@ else () cmake_minimum_required(VERSION 3.8) endif() -if( CMAKE_BUILD_TYPE STREQUAL "Debug" ) - set( CHAI_ENABLE_BOUNDS_CHECK ON CACHE BOOL "Enable Bounds Checking for chai::ManagedArray<>::operator[]" FORCE ) -endif() - ################################ # BLT ################################ From 183358a42d5bddeba2529ff7a0ac08215afea525 Mon Sep 17 00:00:00 2001 From: Alan Dayton Date: Mon, 16 Sep 2019 14:42:52 -0700 Subject: [PATCH 11/16] Another fix after bad merge --- tests/integration/managed_array_tests.cpp | 49 ++++++++++++++++++++++- 1 file changed, 48 insertions(+), 1 deletion(-) diff --git a/tests/integration/managed_array_tests.cpp b/tests/integration/managed_array_tests.cpp index 842293e7..c8fc74d0 100644 --- a/tests/integration/managed_array_tests.cpp +++ b/tests/integration/managed_array_tests.cpp @@ -1394,4 +1394,51 @@ GPU_TEST(ManagedArray, CopyZero) array.free(); } -#endif \ No newline at end of file +#endif + +#if defined(CHAI_ENABLE_BOUNDS_CHECK) + +TEST(ManagedArray, UpperOutOfRangeAccess) +{ + chai::ManagedArray array(20); + array[19] = 0.0; // Should be fine + EXPECT_DEATH_IF_SUPPORTED( array[20] = 0.0, ".*" ); +} + +TEST(ManagedArray, LowerOutOfRangeAccess) +{ + chai::ManagedArray array(20); + array[0] = 0.0; // Should be fine + EXPECT_DEATH_IF_SUPPORTED( array[-1] = 0.0, ".*" ); +} + +#if defined(CHAI_ENABLE_CUDA) + +GPU_TEST(ManagedArray, UpperOutOfRangeAccessGPU) +{ + chai::ManagedArray array(20); + + forall(gpu(), 0, 1, [=] __device__ (int) { + array[20] = 0.0; + }); + + cudaError_t errorCode = cudaGetLastError(); + EXPECT_EQ(errorCode, 59, "No device side assert was triggered!"); + cudaDeviceReset(); +} + +GPU_TEST(ManagedArray, LowerOutOfRangeAccessGPU) +{ + chai::ManagedArray array(20); + + forall(gpu(), 0, 1, [=] __device__ (int) { + array[-1] = 0.0; + }); + + cudaError_t errorCode = cudaGetLastError(); + EXPECT_EQ(errorCode, 59, "No device side assert was triggered!"); + cudaDeviceReset(); +} + +#endif // defined(CHAI_ENABLE_CUDA) +#endif // defined(CHAI_ENABLE_BOUNDS_CHECK) From 4eb66361349aed5ba5c6de404ef50c1e7bdd4029 Mon Sep 17 00:00:00 2001 From: Alan Dayton Date: Mon, 16 Sep 2019 14:44:43 -0700 Subject: [PATCH 12/16] Remove old tests in incorrect location --- src/tests/managed_array_tests.cpp | 436 ------------------------------ 1 file changed, 436 deletions(-) delete mode 100644 src/tests/managed_array_tests.cpp diff --git a/src/tests/managed_array_tests.cpp b/src/tests/managed_array_tests.cpp deleted file mode 100644 index 56d5f9ac..00000000 --- a/src/tests/managed_array_tests.cpp +++ /dev/null @@ -1,436 +0,0 @@ -// --------------------------------------------------------------------- -// Copyright (c) 2016, Lawrence Livermore National Security, LLC. All -// rights reserved. -// -// Produced at the Lawrence Livermore National Laboratory. -// -// This file is part of CHAI. -// -// LLNL-CODE-705877 -// -// For details, see https:://github.com/LLNL/CHAI -// Please also see the NOTICE and LICENSE files. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions -// are met: -// -// - Redistributions of source code must retain the above copyright -// notice, this list of conditions and the following disclaimer. -// -// - Redistributions in binary form must reproduce the above copyright -// notice, this list of conditions and the following disclaimer in the -// documentation and/or other materials provided with the -// distribution. -// -// - Neither the name of the LLNS/LLNL nor the names of its contributors -// may be used to endorse or promote products derived from this -// software without specific prior written permission. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS -// "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT -// LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR -// A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT -// HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, -// INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, -// BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS -// OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED -// AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT -// LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY -// WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE -// POSSIBILITY OF SUCH DAMAGE. -// --------------------------------------------------------------------- -#include "gtest/gtest.h" - -#define CUDA_TEST(X, Y) \ -static void cuda_test_ ## X ## Y();\ -TEST(X,Y) { cuda_test_ ## X ## Y();}\ -static void cuda_test_ ## X ## Y() - -#include "../util/forall.hpp" - -#include "chai/ManagedArray.hpp" - -#include "chai/config.hpp" - -struct my_point { - double x; - double y; -}; - -TEST(ManagedArray, DefaultConstructor) { - chai::ManagedArray array; - ASSERT_EQ(array.size(), 0u); -} - -TEST(ManagedArray, SizeConstructor) { - chai::ManagedArray array(10); - ASSERT_EQ(array.size(), 10u); - array.free(); -} - -TEST(ManagedArray, SpaceConstructorCPU) { - chai::ManagedArray array(10, chai::CPU); - ASSERT_EQ(array.size(), 10u); - array.free(); -} - -#if defined(CHAI_ENABLE_CUDA) -TEST(ManagedArray, SpaceConstructorGPU) { - chai::ManagedArray array(10, chai::GPU); - ASSERT_EQ(array.size(), 10u); - array.free(); -} - -#if defined(CHAI_ENABLE_UM) -TEST(ManagedArray, SpaceConstructorUM) { - chai::ManagedArray array(10, chai::UM); - ASSERT_EQ(array.size(), 10u); - array.free(); -} -#endif -#endif - -TEST(ManagedArray, SetOnHost) { - chai::ManagedArray array(10); - - forall(sequential(), 0, 10, [=] (int i) { - array[i] = i; - }); - - forall(sequential(), 0, 10, [=] (int i) { - ASSERT_EQ(array[i], i); - }); - - array.free(); -} - -TEST(ManagedArray, Const) { - chai::ManagedArray array(10); - - forall(sequential(), 0, 10, [=] (int i) { - array[i] = i; - }); - - chai::ManagedArray array_const(array); - - forall(sequential(), 0, 10, [=] (int i) { - ASSERT_EQ(array_const[i], i); - }); -} - -#if defined(CHAI_ENABLE_CUDA) -CUDA_TEST(ManagedArray, SetOnDevice) { - chai::ManagedArray array(10); - - forall(cuda(), 0, 10, [=] __device__ (int i) { - array[i] = i; - }); - - forall(sequential(), 0, 10, [=] (int i) { - ASSERT_EQ(array[i], i); - }); - - array.free(); -} - -CUDA_TEST(ManagedArray, GetGpuOnHost) { - chai::ManagedArray array(10, chai::GPU); - - forall(cuda(), 0, 10, [=] __device__ (int i) { - array[i] = i; - }); - - forall(sequential(), 0, 10, [=] (int i) { - ASSERT_EQ(array[i], i); - }); - - array.free(); -} - -#if defined(CHAI_ENABLE_UM) -CUDA_TEST(ManagedArray, SetOnDeviceUM) { - chai::ManagedArray array(10, chai::UM); - - forall(cuda(), 0, 10, [=] __device__ (int i) { - array[i] = i; - }); - - forall(sequential(), 0, 10, [=] (int i) { - ASSERT_EQ(array[i], i); - }); - - array.free(); -} -#endif -#endif - -TEST(ManagedArray, Allocate) { - chai::ManagedArray array; - ASSERT_EQ(array.size(), 0u); - - array.allocate(10); - ASSERT_EQ(array.size(), 10u); -} - -TEST(ManagedArray, ReallocateCPU) { - chai::ManagedArray array(10); - ASSERT_EQ(array.size(), 10u); - - forall(sequential(), 0, 10, [=](int i) { - array[i] = i; - }); - - array.reallocate(20); - ASSERT_EQ(array.size(), 20u); - - forall(sequential(), 0, 20, [=](int i) { - if (i < 10) { - ASSERT_EQ(array[i], i); - } else { - array[i] = i; - ASSERT_EQ(array[i], i); - } - }); -} - -#if defined(CHAI_ENABLE_CUDA) -CUDA_TEST(ManagedArray, ReallocateGPU) { - chai::ManagedArray array(10); - ASSERT_EQ(array.size(), 10u); - - forall(cuda(), 0, 10, [=] __device__ (int i) { - array[i] = i; - }); - - array.reallocate(20); - ASSERT_EQ(array.size(), 20u); - - forall(sequential(), 0, 20, [=] (int i) { - if ( i < 10) { - ASSERT_EQ(array[i], i); - } else { - array[i] = i; - ASSERT_EQ(array[i], i); - } - }); -} -#endif - -TEST(ManagedArray, NullpointerConversions) { - chai::ManagedArray a; - a = nullptr; - - chai::ManagedArray b; - b = nullptr; - - ASSERT_EQ(a.size(), 0u); - ASSERT_EQ(b.size(), 0u); - - chai::ManagedArray c(nullptr); - - ASSERT_EQ(c.size(), 0u); -} - -#if defined(CHAI_ENABLE_IMPLICIT_CONVERSIONS) -TEST(ManagedArray, ImplicitConversions) { - chai::ManagedArray a(10); - - chai::ManagedArray a2 = a; - - SUCCEED(); -} -#endif - -TEST(ManagedArray, PodTest) { - chai::ManagedArray array(1); - - forall(sequential(), 0, 1, [=] (int i) { - array[i].x = (double) i; - array[i].y = (double) i*2.0; - }); - - forall(sequential(), 0, 1, [=] (int i) { - ASSERT_EQ(array[i].x, i); - ASSERT_EQ(array[i].y, i*2.0); - }); -} - -#if defined(CHAI_ENABLE_CUDA) -CUDA_TEST(ManagedArray, PodTestGPU) { - chai::ManagedArray array(1); - - forall(cuda(), 0, 1, [=] __device__ (int i) { - array[i].x = (double) i; - array[i].y = (double) i*2.0; - }); - - forall(sequential(), 0, 1, [=] (int i) { - ASSERT_EQ(array[i].x, i); - ASSERT_EQ(array[i].y, i*2.0); - }); -} -#endif - -TEST(ManagedArray, ExternalConstructorUnowned) { - float* data = new float[100]; - - for (int i = 0; i < 100; i++) { - data[i] = 1.0f*i; - } - - chai::ManagedArray array = chai::makeManagedArray(data, 100, chai::CPU, false); - - forall(sequential(), 0, 20, [=] (int i) { - ASSERT_EQ(data[i], array[i]); - }); - - array.free(); - - ASSERT_NE(nullptr, data); -} - -TEST(ManagedArray, ExternalConstructorOwned) { - float* data = new float[20]; - - for (int i = 0; i < 20; i++) { - data[i] = 1.0f*i; - } - - chai::ManagedArray array = chai::makeManagedArray(data, 20, chai::CPU, true); - - forall(sequential(), 0, 20, [=] (int i) { - ASSERT_EQ(data[i], array[i]); - }); - - array.free(); -} - -TEST(ManagedArray, Reset) -{ - chai::ManagedArray array(20); - - forall(sequential(), 0, 20, [=] (int i) { - array[i] = 1.0f*i; - }); - - array.reset(); -} - -#if defined(CHAI_ENABLE_CUDA) -#ifndef CHAI_DISABLE_RM -CUDA_TEST(ManagedArray, ResetDevice) -{ - chai::ManagedArray array(20); - - forall(sequential(), 0, 20, [=] (int i) { - array[i] = 0.0f; - }); - - forall(cuda(), 0, 20, [=] __device__ (int i) { - array[i] = 1.0f*i; - }); - - array.reset(); - - forall(sequential(), 0, 20, [=] (int i) { - ASSERT_EQ(array[i], 0.0f); - }); -} -#endif -#endif - - - -#if defined(CHAI_ENABLE_CUDA) -#ifndef CHAI_DISABLE_RM -CUDA_TEST(ManagedArray, UserCallback) -{ - int num_h2d = 0; - int num_d2h = 0; - size_t bytes_h2d = 0; - size_t bytes_d2h = 0; - size_t bytes_alloc = 0; - size_t bytes_free = 0; - - chai::ManagedArray array; - array.allocate(20, chai::CPU, [&](chai::Action act, chai::ExecutionSpace s, size_t bytes){ - //printf("cback: act=%d, space=%d, bytes=%ld\n", - // (int)act, (int)s, (long)bytes); - if(act == chai::ACTION_MOVE){ - if(s == chai::CPU){ ++num_d2h; bytes_d2h += bytes;} - if(s == chai::GPU){ ++num_h2d; bytes_h2d += bytes;} - } - if(act == chai::ACTION_ALLOC){ - bytes_alloc += bytes; - } - if(act == chai::ACTION_FREE){ - bytes_free += bytes; - } - }); - - for(int iter = 0;iter < 10;++ iter){ - forall(sequential(), 0, 20, [=] (int i) { - array[i] = 0.0f; - }); - - forall(cuda(), 0, 20, [=] __device__ (int i) { - array[i] = 1.0f*i; - }); - } - - - ASSERT_EQ(num_d2h, 9); - ASSERT_EQ(bytes_d2h, 9*sizeof(float)*20); - ASSERT_EQ(num_h2d, 10); - ASSERT_EQ(bytes_h2d, 10*sizeof(float)*20); - - array.free(); - - ASSERT_EQ(bytes_alloc, 2*20*sizeof(float)); - ASSERT_EQ(bytes_free, 2*20*sizeof(float)); -} -#endif -#endif - -#if defined(CHAI_ENABLE_BOUNDS_CHECK) -TEST(ManagedArray, UpperOutOfRangeAccess) -{ - chai::ManagedArray array(20); - - array[19] = 0.0; - EXPECT_DEATH_IF_SUPPORTED( array[20] = 0.0, ".*" ); -} - -TEST(ManagedArray, LowerOutOfRangeAccess) -{ - chai::ManagedArray 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 array(20); - - EXPECT_DEATH( - forall(cuda(), 19, 20, [=] __device__ (int i) { - array[i] = 0.0f; - });, - "i > m_elems"); -} - -CUDA_TEST(ManagedArray, LowerOutOfRangeAccessGPU) -{ - chai::ManagedArray 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) From 495e9b48aa58d03f2392e3fd8a9e529213de366f Mon Sep 17 00:00:00 2001 From: Alan Dayton Date: Mon, 16 Sep 2019 15:31:35 -0700 Subject: [PATCH 13/16] Make sure asserts are turned on, even for release builds --- src/chai/ManagedArray.inl | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index 94939d56..70af6c2b 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -46,8 +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 +#endif // defined(CHAI_ENABLE_BOUNDS_CHECK) + namespace chai { template From 56044d4d25f7f0e8f0147893172e01633d96cb65 Mon Sep 17 00:00:00 2001 From: Alan Dayton Date: Mon, 16 Sep 2019 15:32:51 -0700 Subject: [PATCH 14/16] Use ASSERT instead of EXPECT to match the rest of the tests --- tests/integration/managed_array_tests.cpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/tests/integration/managed_array_tests.cpp b/tests/integration/managed_array_tests.cpp index c8fc74d0..d0853bca 100644 --- a/tests/integration/managed_array_tests.cpp +++ b/tests/integration/managed_array_tests.cpp @@ -1402,14 +1402,14 @@ TEST(ManagedArray, UpperOutOfRangeAccess) { chai::ManagedArray array(20); array[19] = 0.0; // Should be fine - EXPECT_DEATH_IF_SUPPORTED( array[20] = 0.0, ".*" ); + ASSERT_DEATH_IF_SUPPORTED( array[20] = 0.0, ".*" ); } TEST(ManagedArray, LowerOutOfRangeAccess) { chai::ManagedArray array(20); array[0] = 0.0; // Should be fine - EXPECT_DEATH_IF_SUPPORTED( array[-1] = 0.0, ".*" ); + ASSERT_DEATH_IF_SUPPORTED( array[-1] = 0.0, ".*" ); } #if defined(CHAI_ENABLE_CUDA) @@ -1423,8 +1423,8 @@ GPU_TEST(ManagedArray, UpperOutOfRangeAccessGPU) }); cudaError_t errorCode = cudaGetLastError(); - EXPECT_EQ(errorCode, 59, "No device side assert was triggered!"); - cudaDeviceReset(); + ASSERT_EQ(cudaErrorAssert, errorCode, "No device side assert was triggered!"); + ASSERT_EQ(cudaDeviceReset(), cudaSuccess); } GPU_TEST(ManagedArray, LowerOutOfRangeAccessGPU) @@ -1436,8 +1436,8 @@ GPU_TEST(ManagedArray, LowerOutOfRangeAccessGPU) }); cudaError_t errorCode = cudaGetLastError(); - EXPECT_EQ(errorCode, 59, "No device side assert was triggered!"); - cudaDeviceReset(); + ASSERT_EQ(cudaErrorAssert, errorCode, "No device side assert was triggered!"); + ASSERT_EQ(cudaDeviceReset(), cudaSuccess); } #endif // defined(CHAI_ENABLE_CUDA) From a2526d8b05ff36e766d970df11f91cdfd18c2bb7 Mon Sep 17 00:00:00 2001 From: Alan Dayton Date: Mon, 16 Sep 2019 17:17:43 -0700 Subject: [PATCH 15/16] Fix cuda build --- tests/integration/managed_array_tests.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tests/integration/managed_array_tests.cpp b/tests/integration/managed_array_tests.cpp index d0853bca..65cf0b1e 100644 --- a/tests/integration/managed_array_tests.cpp +++ b/tests/integration/managed_array_tests.cpp @@ -1423,7 +1423,7 @@ GPU_TEST(ManagedArray, UpperOutOfRangeAccessGPU) }); cudaError_t errorCode = cudaGetLastError(); - ASSERT_EQ(cudaErrorAssert, errorCode, "No device side assert was triggered!"); + ASSERT_EQ(cudaErrorAssert, errorCode); ASSERT_EQ(cudaDeviceReset(), cudaSuccess); } @@ -1436,7 +1436,7 @@ GPU_TEST(ManagedArray, LowerOutOfRangeAccessGPU) }); cudaError_t errorCode = cudaGetLastError(); - ASSERT_EQ(cudaErrorAssert, errorCode, "No device side assert was triggered!"); + ASSERT_EQ(cudaErrorAssert, errorCode); ASSERT_EQ(cudaDeviceReset(), cudaSuccess); } From 6f9e0383dea8fb6e0bcdab333b448f2a7e18cddc Mon Sep 17 00:00:00 2001 From: Alan Dayton Date: Wed, 25 Sep 2019 15:20:18 -0700 Subject: [PATCH 16/16] Add always assert function --- src/chai/ManagedArray.inl | 30 ++++++++++++++++++++--- tests/integration/managed_array_tests.cpp | 24 ++++++++++++------ 2 files changed, 42 insertions(+), 12 deletions(-) diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index 70af6c2b..ecfe52b5 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -47,13 +47,25 @@ #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) + +#define HOST_ALWAYS_ASSERT(EXP) if (!(EXP)) { \ + printf("Assert triggered at %s:%d\n", __FILE__, __LINE__); \ + abort(); \ +} + +#define DEVICE_ALWAYS_ASSERT(EXP) if (!(EXP)) { \ + asm("trap;") \ +} + +#else // defined(NDEBUG) #include +#define HOST_ALWAYS_ASSERT(EXP) assert(EXP) +#define DEVICE_ALWAYS_ASSERT(EXP) assert(EXP) + +#endif // defined(NDEBUG) #endif // defined(CHAI_ENABLE_BOUNDS_CHECK) namespace chai { @@ -400,8 +412,18 @@ template CHAI_INLINE CHAI_HOST_DEVICE T& ManagedArray::operator[](const Idx i) const { #if defined(CHAI_ENABLE_BOUNDS_CHECK) - assert(i >= 0 && static_cast(i) < m_elems); +#if defined(__CUDA_ARCH__) + DEVICE_ALWAYS_ASSERT(i >= 0 && static_cast(i) < m_elems); +#else + //std::cout << "m_elems = " << m_elems << std::endl; + //std::cout << "i = " << i << std::endl; + //std::cout << "(size_t) i = " << static_cast(i) << std::endl; + + //printf("m_elems = %d\n", m_elems); + //printf("i = %d\n", i); + HOST_ALWAYS_ASSERT(i >= 0 && static_cast(i) < m_elems); #endif +#endif // defined(CHAI_ENABLE_BOUNDS_CHECK) return m_active_pointer[i]; } diff --git a/tests/integration/managed_array_tests.cpp b/tests/integration/managed_array_tests.cpp index d0853bca..e271668c 100644 --- a/tests/integration/managed_array_tests.cpp +++ b/tests/integration/managed_array_tests.cpp @@ -1402,42 +1402,50 @@ TEST(ManagedArray, UpperOutOfRangeAccess) { chai::ManagedArray array(20); array[19] = 0.0; // Should be fine - ASSERT_DEATH_IF_SUPPORTED( array[20] = 0.0, ".*" ); + ASSERT_DEATH_IF_SUPPORTED(array[20] = 0.0, ".*"); } TEST(ManagedArray, LowerOutOfRangeAccess) { chai::ManagedArray array(20); array[0] = 0.0; // Should be fine - ASSERT_DEATH_IF_SUPPORTED( array[-1] = 0.0, ".*" ); + ASSERT_DEATH_IF_SUPPORTED(array[-1] = 0.0, ".*"); } #if defined(CHAI_ENABLE_CUDA) GPU_TEST(ManagedArray, UpperOutOfRangeAccessGPU) { + ASSERT_EQ(cudaSuccess, cudaDeviceReset()); + ASSERT_EQ(cudaSuccess, cudaDeviceSynchronize()); + chai::ManagedArray array(20); - forall(gpu(), 0, 1, [=] __device__ (int) { + ASSERT_DEATH_IF_SUPPORTED(forall(gpu(), 0, 1, [=] __device__ (int) { array[20] = 0.0; - }); + }), ".*"); - cudaError_t errorCode = cudaGetLastError(); - ASSERT_EQ(cudaErrorAssert, errorCode, "No device side assert was triggered!"); + //cudaError_t errorCode = cudaGetLastError(); + cudaError_t errorCode = cudaDeviceSynchronize(); + ASSERT_EQ(cudaErrorAssert, errorCode); ASSERT_EQ(cudaDeviceReset(), cudaSuccess); + ASSERT_EQ(cudaSuccess, cudaDeviceSynchronize()); } GPU_TEST(ManagedArray, LowerOutOfRangeAccessGPU) { chai::ManagedArray array(20); +#if 0 forall(gpu(), 0, 1, [=] __device__ (int) { array[-1] = 0.0; }); - cudaError_t errorCode = cudaGetLastError(); - ASSERT_EQ(cudaErrorAssert, errorCode, "No device side assert was triggered!"); + //cudaError_t errorCode = cudaGetLastError(); + cudaError_t errorCode = cudaDeviceSynchronize(); + ASSERT_EQ(cudaErrorAssert, errorCode); ASSERT_EQ(cudaDeviceReset(), cudaSuccess); +#endif } #endif // defined(CHAI_ENABLE_CUDA)