Recently, I had an exchange with a colleague of mine whose background is primarily in conventional graphics APIs with limited compute capabilities, he was complaining about the lack of an intrinsic append buffer in MSL. Certainly not the first time this was brought up to me or somewhere else online. I thought it would make for a good article to show both why an intrinsic one is not necessary and how you can roll your own if you're working with Apple platforms. If you're not familiar with an append buffer in the context of GPGPU, it's a structure geared for queuing entities processed in parallel without in-fighting between threads for a slot, which is to say, it supports temporally unordered access. When it comes down to it, it's really just a simple frontend consisting of a single method which manipulates the underlying buffer and its associated atomic counter.
In HLSL, which is considerably less expressive than MSL, the structure in question is provided intrinsically (AppendStructuredBuffer<T>) as it is impossible to link a counter and a buffer together in a custom generic structure. The buffer which is to be augmented with append functionality needs to have its unordered access view marked as such through the DirectX API on the CPU side. This enables the runtime to procure the necessary atomic counters which are hidden from the user. For us, everything will be a bit more explicit. In the next few paragraphs, you will find out how to hide the atomic counter so the user doesn't have to lug anything around as long as they size the underlying buffer appropriately.
Understanding the limitations of MSL
As you're probably familiar, the Metal Shading Language is based on the C++14 specification and a subset of the functionality specified therein, which sounds very impressive and is bound to give people the wrong impression. Before we get to implementing our append_buffer<T>, we need to reiterate on the limitations of MSL so that certain implementation decisions don't feel jarring to those versed in C++, but not necessarily MSL.
Forget about inheritance
It is not possible to extend the functionality of your classes or those provided by the Metal standard library. Naturally, from this follows that virtual specifier (dynamic dispatch) and dynamic_cast are not allowed. Mind you, even if we had some modicum of inheritance, the latter wouldn't work due to the complete lack of runtime type information.
new/delete are not available
As with any other shading language, all persistent allocations are performed through the appropriate APIs on the CPU side and fed into the appropriate shaders running on the GPU, and as usual, both sides need to agree on layout, offsets and bind points. Pointers and references are also limited only to more persistent data locations and require an appropriate address space attribute prefix, one of device/constant/threadgroup/threadgroup_imageblock. Temporaries are passed around solely by value, just like in HLSL. So, if you thought you could pass a struct by constant reference or maybe even move it, nope.
But wait, there is more
Obviously, exceptions are not supported and thus noexcept is meaningless. You can't perform recursive function calls as GPUs are not quite there yet. goto is not available, now if only we could get rid of it in C++20y as well. Storage class specifiers register and thread_local are not available in MSL, former of which was also deprecated in C++11.
Shopping at the atomics store
The part of the Metal standard library (metal_stdlib) dedicated to dealing with atomics is contained within metal_atomic. Don't worry, you don't have to explicitly include it. There is a myriad of operations available, just like in the C++ standard library itself and intrinsics that come with HLSL. All devices which are certified for Metal 1.0 support relaxed memory ordering and atomic_(u)int. Relaxed memory order makes no guarantees on the ordering of the memory operations as observed from any one thread, but it does guarantee atomicity which is all we need to enable an atomic counter. We'll be using atomic_uint for the counter type and avail ourselves of a read-modify-write operation to increment the counter in one whole swoop, returning the previous value of the atomic data (to index into the underlying buffer). The specific overload is effectively the same as the one in the C++ standard library:
template<typename T> // For atomic integers, atomic<T>::OperandType == T
T atomic_fetch_key_explicit(device atomic<T>* object, atomic<T>::OperandType operand, memory_order order);
// Or more to the point, for low-tier atomics and T is uint:
uint atomic_fetch_key_explicit(device atomic_uint* object, uint operand, memory_order order);
This function is all we need to enable an append buffer on any Metal 1.0 device, going all the way back to an iPhone 5S should that matter to your project. I might cover atomics in general in a future article as I've seen a lot of confusion online regarding why they're needed and how they're used (especially when). For now, let's proceed to write the append buffer implementation.
Putting it all together
We've got to address the elephant in the room... Sourcing the counter as it needs to have persistent storage to be of any use. In a lesser system, we'd have to lug around at least \(n + 1\) buffers around where the extra one is storing \(n\) counters, one for each append buffer. Passing that around would be a chore and just plain ugly. Remember that allocation is requested by the CPU and resizing on the GPU is not possible, thus we always provide some considerable upper bound which we never expect to reach. So, if we were to tear out sizeof(atomic_uint) from the start of each buffer and use it as the counter, this lets us avoid a lot of boilerplate. Actually, we push up the start of the actual buffer to 16, which is the safest alignment:
template<typename T>
struct append_buffer {
append_buffer(device T* underlying_buffer) {
auto untyped_underlying_buffer = reinterpret_cast<device char*>(underlying_buffer);
_counter = reinterpret_cast<device atomic_uint*>(untyped_underlying_buffer);
_underlying_buffer = reinterpret_cast<device T*>(untyped_underlying_buffer + 16);
}
auto append(T t) {
auto index_before_increment = atomic_fetch_add_explicit(_counter, 1, memory_order_relaxed);
_underlying_buffer[index_before_increment] = t;
}
private:
device atomic_uint* _counter;
device T* _underlying_buffer;
};
From here, we can use it in a compute kernel. Every buffer that we bind on a compute command encoder is passed to the compute kernel entrypoint as an argument. I prefer to provide an overarching struct which I consider the compute kernel, which is really "the way it's meant to be played" and refreshing after years of global constant buffers. This helps me consolidate shared data without passing it around a bunch of free functions.
struct compute_kernel {
compute_kernel(uint2 thread_id, device element* elements_buffer)
: _thread_id(thread_id), _elements_append_buffer(elements_buffer) {}
auto execute() -> void;
private:
uint2 _thread_id;
append_structured_buffer<element> _elements_append_buffer;
};
auto compute_kernel::execute() -> void {
// ...
_elements_append_buffer.append(...);
// ...
}
kernel auto compute_kernel_entrypoint(
uint2 thread_id [[thread_position_in_grid]],
device element* elements_buffer [[buffer(0)]],
) -> void {
compute_kernel compute_kernel_instance { thread_id, elements_buffer };
compute_kernel_instance.execute();
}
Regarding the allocation request for a buffer on the CPU side via Metal API, it's fairly straightforward:
id<MTLDevice> _device;
NSUInteger bufferLength = sizeof(element) * elementCountWithAReasonableMargin;
[_device newBufferWithLength:bufferLength options:MTLResourceStorageModePrivate];
A simple "consume" buffer
For the sake of completeness, the following snippet also describes a very simple "consume" buffer, quotes required as it isn't quite what the HLSL counterpart has in mind, which I find quite useless. It's simply a layer above the underlying buffer which can decode the utilized count of the buffer which can then be used to limit threads from accessing out of bounds:
template<typename T>
struct consume_buffer {
consume_buffer(constant T* underlying_buffer) {
auto untyped_underlying_buffer = reinterpret_cast<constant char*>(underlying_buffer);
_counter = reinterpret_cast<constant uint*>(untyped_underlying_buffer);
_underlying_buffer = reinterpret_cast<constant T*>(untyped_underlying_buffer + 16);
}
auto get_count() {
return *_counter;
}
auto operator[](uint i) {
return _underlying_buffer[i];
}
private:
constant uint* _counter;
constant T* _underlying_buffer;
};
I used a similar setup when I was working on the port of my adaptive terrain solution for Apple platforms for numerous stages which were computed in the geometry shader. You can see the region subdivision stage below.
Until next time.
