search

GEOS-DEV / LvArray

Portable HPC Containers (C++)
BSD 3-Clause "New" or "Revised" License
47 stars 10 forks source link

Added nicer errors when you do invalid operations on a ChaiBuffer from device #286

Closed corbett5 closed 1 year ago

corbett5 commented 1 year ago

Say you try and capture an Array by value in a device lambda, I think we've all made this mistake before

Array< int, 1, RAJA::PERM_I, std::ptrdiff_t, ChaiBuffer > phi( 10 );
  forall< RAJA::cuda_exec< 32 > >( 10, [=] LVARRAY_DEVICE ( std::ptrdiff_t const i )
  {
    phi( i ) = i;
  } );

Previous error message was the result of calling a host function on device

unknown file: Failure
C++ exception with description "campCudaErrchk(cudaStreamSynchronize(stream)) an illegal memory access was encountered /usr/gapps/GEOSX/thirdPartyLibs/2023-01-23/install-lassen-clang13-cuda11-release/raja/include/camp/resource/cuda.hpp:172" thrown in the test body.

Now you get something clearer, although albeit much longer

/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [0,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [1,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [2,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [3,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [4,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [5,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [6,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [7,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [8,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [9,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [10,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [11,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [12,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [13,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [14,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [15,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [16,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [17,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [18,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [19,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [20,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [21,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [22,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [23,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [24,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [25,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [26,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [27,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [28,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [29,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [30,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
/usr/WS2/corbett5/geosx/GEOSX/src/coreComponents/LvArray/unitTests/../src/ChaiBuffer.hpp:146: LvArray::ChaiBuffer<T>::ChaiBuffer(__nv_bool) [with T = int]: block: [0,0,0], thread: [31,0,0] Assertion `false && "EXP = " "true" "MSG = " "\"Creating a new ChaiBuffer on device is not supported.\""` failed.
unknown file: Failure
C++ exception with description "campCudaErrchk(cudaStreamSynchronize(stream)) device-side assert triggered /usr/gapps/GEOSX/thirdPartyLibs/2023-01-23/install-lassen-clang13-cuda11-release/raja/include/camp/resource/cuda.hpp:172" thrown in the test body.

The only downside I see is that I have to mark certain function in ChaiBuffer which should really be host only as host-device. But I don't really see this as a problem since clearly the compiler was fine calling these host functions from device.

corbett5 commented 1 year ago

If we really wanted to we could put error messages up higher in the call stack so that say the Array copy constructor would say something like "can't copy construct an Array< ..., ChaiBuffer> on device". But that seems a bit excessive to me.

corbett5 commented 1 year ago

I updated the error message you see when you capture an Array on device.

corbett5 commented 1 year ago

Now part of https://github.com/GEOS-DEV/LvArray/pull/288