Represents a string that owns its buffer and may be synchronised between host and GPU. More...

#include <StringBuffer.hpp>

Public Member Functions
int	size () const
	Get the length of the string, including any null characters. More...

	StringBuffer (int size)
	Construct an uninitialised string of the given length. More...

template<typename Traits >
	StringBuffer (const std::basic_string< CharT, Traits > &source)
	Construct a `StringBuffer` from an `std::string`. More...

template<typename Traits >
	StringBuffer (const std::basic_string_view< CharT, Traits > source)
	Construct a `StringBuffer` from an `std::string_view`. More...

	StringBuffer (sp::UniquePtr< CharT > &&buffer, int size)
	Construct a `StringBuffer` from a character allocation and size. More...

auto	deviceView (sp::Stream &s, bool move=false)
	Get an immutable `StringView` backed by a buffer that will, once all current work on stream `s` is complete, contain a copy of the string on the device associated with stream `s`. More...

auto	mutableDeviceView (sp::Stream &s, bool move=false)
	Get a mutable `StringView` backed by a buffer that will, once all current work on stream `s` is complete, contain a copy of the string on the device associated with stream `s`. More...

auto	hostView (sp::Stream &s, bool move=false)
	Get an immutable `StringView` backed by the host buffer, and if necessary enqueue a copy to the host on stream `s`. More...

auto	mutableHostView (sp::Stream &s, bool move=false)
	Get a mutable `StringView` backed by the host buffer, and if necessary enqueue a copy to the host on stream `s`. More...

Detailed Description

template<typename CharT = char, typename IntT = uint32_t>
class sp::StringBuffer< CharT, IntT >

Represents a string that owns its buffer and may be synchronised between host and GPU.

Use an sp::StringBuffer to represent ownership of the string, and extract sp::StringViews using mutableDeviceView() or deviceView() to process the string. This object will automatically take care of enqueuing CPU<->GPU copies as needed to satisfy your demands.

This is conceptually similar to sp::NomadicBuffer or sp::NomadicTensor, except it operates in terms of sp::StringViews.

Interoperation with STL strings

sp::StringBuffer interoperates smoothly with std::string, but some care is needed to do this as efficiently as possible.

Recall that:

Copies between host and GPU of pageable memory are considerably slower than copies involving page-locked memory.
Allocation of page-locked memory is more expensive than allocation of pageable memory.
std::string (unless you provide a custom allocator) allocates pageable memory.

Many std::string functions cause the allocation of a new string. Such operations would be made slower if the std::strings involved were made to use page-locked memory.

To avoid the large overhead of GPU data transfers involving pageable memory, sp::StringBuffer will, when constructed from an std::string, allocate a new page-locked host buffer and copy the std::string into it. The cost of this copy is less than the slowdown a GPU transfer from pageable memory would introduce, so it pays for itself provided you copy the string to the GPU at least once.

Note: sp::StringBuffer defaults to using an int32_t for length, but std:string uses an int64_t. This is done because GPUs don't have 64-bit integer hardware, so this is a significant performance benefit if strings are sufficiently small. Make sure to set the IntT template parameter to something longer if you have very large strings.

Best practices

Avoid using allocating std::string functions (like std::string::insert()) where possible. Allocation isn't cheap at the best of times. std::string_view and in-place modification can often help.
If you are using allocating std::string functions, complete your std::string processing and pass your finished string to sp::StringBuffer (where it will be copied to a page-locked buffer), and then do your GPU stuff.
If you are not using any allocating std::string functions, you can avoid the copy done by sp::StringBuffer by migrating even host-side string processing to use sp::StringBuffer and sp::StringView. Construct a suitably-sized sp::StringBuffer at the beginning do your thing.
Null terminators are unnecessary, ignored, and will just waste a byte.

Examples

Calling a kernel with a string of known length

// Create an uninitialised string of length 42. No allocation occurs until you request a view somewhere.
sp::StringBuffer<> theString{42};
 
// Get the mutable host-side view of the string. This causes the host buffer to be allocated.
sp::StringView<char> hostView = theString.mutableHostView();
 
// Do something with it. Let's be very silly and put in some ascendingly-numbered chars.
for (int i = 0; i < hostView.size(); i++) {
    hostView[i] = (char) i;
}
 
// Note that we could have used the custom literal syntax instead:
// sp::StringBuffer<> theString = "This is a string"_sp;
// Doing so causes `theString` to be initialised with that literal value in its host buffer.
 
// Asynchronously copy to the GPU and then run the kernel. Both the copy and the kernel launch are done on
// stream `s`, so this is safe and asynchronous with respect to the host.
myKernel<<<1, 1, 0, s>>>(theString.gpuView(s));

Using host-side reallocating-string APIs.

sp::Stream s = ...;
std::string myString = ...;
 
// Insert a substring into the middle of the string, (probably) causing a reallocation.
myString = myString.insert(1, "boink");
 
// This will allocate a page-locked host buffer and copy myString into it. `myString` can now be deallocated.
sp::StringBuffer<> buffer{myString};
 
// Asynchronously copy to the GPU and then run the kernel.
myKernel<<<1, 1, 0, s>>>(gpuString.gpuView(s));

Template Parameters

CharT	The type of characters the string is comprised of.
IntT	The integer type to use to represent the length.

Constructor & Destructor Documentation

◆ StringBuffer() [1/4]

template<typename CharT = char, typename IntT = uint32_t>

sp::StringBuffer< CharT, IntT >::StringBuffer ( int size )

explicit

Construct an uninitialised string of the given length.

You'll probably want to use the mutable*View() functions to get references to the underlying memory to put something into the string.

◆ StringBuffer() [2/4]

template<typename CharT = char, typename IntT = uint32_t>

template<typename Traits >

sp::StringBuffer< CharT, IntT >::StringBuffer ( const std::basic_string< CharT, Traits > & source )

explicit

Construct a StringBuffer from an std::string.

The contents of the source string is copied into a newly-allocated pinned host buffer owned by this object.

◆ StringBuffer() [3/4]

template<typename CharT = char, typename IntT = uint32_t>

template<typename Traits >

sp::StringBuffer< CharT, IntT >::StringBuffer ( const std::basic_string_view< CharT, Traits > source )

explicit

Construct a StringBuffer from an std::string_view.

The contents of the source view is copied into a newly-allocated pinned host buffer owned by this object.

◆ StringBuffer() [4/4]

template<typename CharT = char, typename IntT = uint32_t>

sp::StringBuffer< CharT, IntT >::StringBuffer	(	sp::UniquePtr< CharT > &&	buffer,
		int	size
	)

explicit

Construct a StringBuffer from a character allocation and size.

Ownership of the buffer is taken. The pointer should be to the start of an allocation. If appropriate, create desired StringView objects afterwards.

Member Function Documentation

◆ deviceView()

template<typename CharT = char, typename IntT = uint32_t>

auto sp::StringBuffer< CharT, IntT >::deviceView	(	sp::Stream &	s,
		bool	move = `false`
	)

Get an immutable StringView backed by a buffer that will, once all current work on stream s is complete, contain a copy of the string on the device associated with stream s.

If necessary, a copy to the device will be enqueued on the stream.

Parameters

s	The stream to use.
move	If true, deallocate all other copies of this string from other devices (and the host)

◆ hostView()

template<typename CharT = char, typename IntT = uint32_t>

auto sp::StringBuffer< CharT, IntT >::hostView	(	sp::Stream &	s,
		bool	move = `false`
	)

Get an immutable StringView backed by the host buffer, and if necessary enqueue a copy to the host on stream s.

Parameters

s	The stream to use.
move	If true, deallocate all GPU-resident copies of this string

◆ mutableDeviceView()

template<typename CharT = char, typename IntT = uint32_t>

auto sp::StringBuffer< CharT, IntT >::mutableDeviceView	(	sp::Stream &	s,
		bool	move = `false`
	)

Get a mutable StringView backed by a buffer that will, once all current work on stream s is complete, contain a copy of the string on the device associated with stream s.

If necessary, a copy to the device will be enqueued on the stream.

The target device will be considered to have made the buffer dirty, requiring future requests for a view on other devices to copy from this one to the target.

Parameters

s	The stream to use.
move	If true, deallocate all other copies of this string from other devices (and the host)

◆ mutableHostView()

template<typename CharT = char, typename IntT = uint32_t>

auto sp::StringBuffer< CharT, IntT >::mutableHostView	(	sp::Stream &	s,
		bool	move = `false`
	)

Get a mutable StringView backed by the host buffer, and if necessary enqueue a copy to the host on stream s.

The host will be considered to have made the buffer dirty, requiring future requests for a view on other devices to copy from the host to the target.

Parameters

s	The stream to use.
move	If true, deallocate all GPU-resident copies of this string

◆ size()

template<typename CharT = char, typename IntT = uint32_t>

int sp::StringBuffer< CharT, IntT >::size ( ) const

Get the length of the string, including any null characters.

Public Member Functions

Detailed Description

Interoperation with STL strings

Best practices

Examples

Calling a kernel with a string of known length

Using host-side reallocating-string APIs.

Constructor & Destructor Documentation

◆ StringBuffer() [1/4]

◆ StringBuffer() [2/4]

◆ StringBuffer() [3/4]

◆ StringBuffer() [4/4]

Member Function Documentation

◆ deviceView()

◆ hostView()

◆ mutableDeviceView()

◆ mutableHostView()

◆ size()