2. Buffer Classes

The buffer classes are the backbone of every LvArray class. A buffer class is responsible for allocating, reallocating and de-allocating a chunk of memory as well as moving it between memory spaces. A buffer is not responsible for managing the lifetime of the objects in their allocation. In general buffer classes have shallow copy semantics and do not de-allocate their allocations upon destruction. Buffer classes implement the copy and move constructors as well as the copy and move assignment operators. They also have a default constructor that leaves them in an uninitialized state. In general it is only safe to assign to an uninitialized buffer although different buffer implementations may allow other operations. To construct an initialized buffer pass a dummy boolean argument, this value of the parameter is not important and it only exists to differentiate it from the default constructor. Once created an initialized buffer must be free’d, either directly or though one of its copies. There are currently three buffer implementations: LvArray::MallocBuffer, LvArray::ChaiBuffer and LvArray::StackBuffer.

2.1. LvArray::MallocBuffer

As you might have guessed LvArray::MallocBuffer uses malloc and free to handle its allocation. Copying a LvArray::MallocBuffer does not copy the allocation. The allocation of a LvArray::MallocBuffer lives exclusively on the host and as such it will abort the program if you try to move it to or touch it in any space other than MemorySpace::host.

TEST( MallocBuffer, copy )
{
  constexpr std::ptrdiff_t size = 55;
  LvArray::MallocBuffer< int > buffer( true );
  buffer.reallocate( 0, LvArray::MemorySpace::host, size );

  for( int i = 0; i < size; ++i )
  {
    buffer[ i ] = i;
  }

  for( int i = 0; i < size; ++i )
  {
    EXPECT_EQ( buffer[ i ], i );
  }

  // MallocBuffer has shallow copy semantics.
  LvArray::MallocBuffer< int > copy = buffer;
  EXPECT_EQ( copy.data(), buffer.data() );

  // Must be manually free'd.
  buffer.free();
}

TEST( MallocBuffer, nonPOD )
{
  constexpr std::ptrdiff_t size = 4;
  LvArray::MallocBuffer< std::string > buffer( true );
  buffer.reallocate( 0, LvArray::MemorySpace::host, size );

  // Buffers don't initialize data so placement new must be used.
  for( int i = 0; i < size; ++i )
  {
    new ( buffer.data() + i ) std::string( std::to_string( i ) );
  }

  for( int i = 0; i < size; ++i )
  {
    EXPECT_EQ( buffer[ i ], std::to_string( i ) );
  }

  // Buffers don't destroy the objects in free.
  // The using statement is needed to explicitly call the destructor
  using std::string;
  for( int i = 0; i < size; ++i )
  {
    buffer[ i ].~string();
  }

  buffer.free();
}

[Source: examples/exampleBuffers.cpp]

2.2. LvArray::ChaiBuffer

LvArray::ChaiBuffer uses CHAI to manage an allocation which can exist on both the host and device, it functions similarly to the chai::ManagedArray. Like the LvArray::MallocBuffer copying a LvArray::ChaiBuffer via the assignment operators or the move constructor do not copy the allocation. The unique feature of the LvArray::ChaBuffer is that when it is copy constructed if the CHAI execution space is set it will move its allocation to the appropriate space creating an allocation there if it did not already exist.

CUDA_TEST( ChaiBuffer, captureOnDevice )
{
  constexpr std::ptrdiff_t size = 55;
  LvArray::ChaiBuffer< int > buffer( true );
  buffer.reallocate( 0, LvArray::MemorySpace::host, size );

  for( int i = 0; i < size; ++i )
  {
    buffer[ i ] = i;
  }

  // Capture buffer in a device kernel which creates an allocation on device
  // and copies the data there.
  RAJA::forall< RAJA::cuda_exec< 32 > >(
    RAJA::TypedRangeSegment< std::ptrdiff_t >( 0, size ),
    [buffer] __device__ ( std::ptrdiff_t const i )
  {
    buffer[ i ] += i;
  } );

  // Capture buffer in a host kernel moving the data back to the host allocation.
  RAJA::forall< RAJA::loop_exec >(
    RAJA::TypedRangeSegment< std::ptrdiff_t >( 0, size ),
    [buffer] ( std::ptrdiff_t const i )
  {
    EXPECT_EQ( buffer[ i ], 2 * i );
  } );

  buffer.free();
}

[Source: examples/exampleBuffers.cpp]

In order to prevent unnecessary memory motion if the type contained in the LvArray::ChaiBuffer is const then the data is not touched in any space it is moved to.

CUDA_TEST( ChaiBuffer, captureOnDeviceConst )
{
  constexpr std::ptrdiff_t size = 55;
  LvArray::ChaiBuffer< int > buffer( true );
  buffer.reallocate( 0, LvArray::MemorySpace::host, size );

  for( int i = 0; i < size; ++i )
  {
    buffer[ i ] = i;
  }

  // Create a const buffer and capture it in a device kernel which
  // creates an allocation on device and copies the data there.
  LvArray::ChaiBuffer< int const > const constBuffer( buffer );
  RAJA::forall< RAJA::cuda_exec< 32 > >(
    RAJA::TypedRangeSegment< std::ptrdiff_t >( 0, size ),
    [constBuffer] __device__ ( std::ptrdiff_t const i )
  {
    const_cast< int & >( constBuffer[ i ] ) += i;
  } );

  // Capture buffer in a host kernel moving the data back to the host allocation.
  // If constBuffer didn't contain "int const" then this check would fail because
  // the data would be copied back from device.
  RAJA::forall< RAJA::loop_exec >(
    RAJA::TypedRangeSegment< std::ptrdiff_t >( 0, size ),
    [buffer] ( std::ptrdiff_t const i )
  {
    EXPECT_EQ( buffer[ i ], i );
  } );

  buffer.free();
}

[Source: examples/exampleBuffers.cpp]

LvArray::ChaiBuffer supports explicit movement and touching as well via the methods move and registerTouch.

Whenever a LvArray::ChaiBuffer is moved between memory spaces it will print the size of the allocation, the type of the buffer and the name. Both the name and the type can be set with the setName method. If this behavior is not desired it can be disabled with chai::ArrayManager::getInstance()->disableCallbacks().

TEST( ChaiBuffer, setName )
{
  LvArray::ChaiBuffer< int > buffer( true );
  buffer.reallocate( 0, LvArray::MemorySpace::host, 1024 );

  // Move to the device.
  buffer.move( LvArray::MemorySpace::cuda, true );

  // Give buffer a name and move back to the host.
  buffer.setName( "my_buffer" );
  buffer.move( LvArray::MemorySpace::host, true );

  // Rename buffer and override the default type.
  buffer.setName< double >( "my_buffer_with_a_nonsensical_type" );
  buffer.move( LvArray::MemorySpace::cuda, true );
}

[Source: examples/exampleBuffers.cpp]

Output

Moved    4.0 KB to the DEVICE: LvArray::ChaiBuffer<int>
Moved    4.0 KB to the HOST  : LvArray::ChaiBuffer<int> my_buffer
Moved    4.0 KB to the DEVICE: double my_buffer_with_a_nonsensical_type

2.3. LvArray::StackBuffer

The LvArray::StackBuffer is unique among the buffer classes because it wraps a c-array of objects whose size is fixed at compile time. It is so named because if you declare a LvArray::StackBuffer on the stack its allocation will also live on the stack. Unlike the other buffer classes by nature copying a LvArray::StackBuffer is a deep copy, furthermore a LvArray::StackBuffer can only contain trivially destructible types, so no putting a std::string in one. If you try to grow the allocation beyond the fixed size it will abort the program.

TEST( StackBuffer, example )
{
  constexpr std::ptrdiff_t size = 55;
  LvArray::StackBuffer< int, 55 > buffer( true );

  static_assert( buffer.capacity() == size, "Capacity is fixed at compile time." );

  for( std::ptrdiff_t i = 0; i < size; ++i )
  {
    buffer[ i ] = i;
  }

  for( std::ptrdiff_t i = 0; i < size; ++i )
  {
    EXPECT_EQ( buffer[ i ], i );
  }

  EXPECT_DEATH_IF_SUPPORTED( buffer.reallocate( size, LvArray::MemorySpace::host, 2 * size ), "" );

  // Not necessary with the StackBuffer but it's good practice.
  buffer.free();
}

[Source: examples/exampleBuffers.cpp]