Anyone who has written a non-trivial GPU kernel has hit the same wall. You define a struct in C++ and pass it to a kernel. Six months later, somebody adds a member, forgets to update the matching layout on the device side, and a 200-line shader reads garbage. The compiler doesn’t help. The runtime doesn’t help. You find out from a customer.

The traditional answer is more macros. BOOST_PP_SEQ_FOR_EACH, X-macros, FOREACH_FIELD. They work, but they make the source unreadable and they cost nothing to misuse. The other answer is a code generator, which means a build-time dependency, a generated header nobody reads, and a debugging experience that involves three files at once.

C++26 ships P2996, static reflection, which was voted into the working draft at St. Louis in June 2024 and finalized at the Hagenberg meeting in 2025. For most of the C++ world this is an enabler for serialization, ORMs, and CLI parsers. For people writing GPU code, it is also one of the most practical things to happen to the toolchain in years. The reason is simple. GPU programming is about layouts, bindings, kernel variants, and metadata. All four of those are exactly what reflection is good at.

This post is a tour through six concrete things you can do once your compiler speaks reflection:

• validating that host and device struct layouts agree at compile time,
• generating SoA types from AoS definitions,
• packing push constants without hand-computed offsets,
• building kernel-variant dispatch tables from enum definitions,
• inspecting memory layout and detecting padding, and
• attaching binding/memory-space metadata directly to declarations via annotations.

Disclaimer: LLMs were used for proof-reading and grammar check.

C++20 gave us coroutines. The machinery needed was included: co_yield, co_return, co_await but the standard library had no concrete coroutine types. You had to write your own promise type, your own iterator, your own bookkeeping. The complexity of this boilerplate was intimidating for most people.

C++23 fixes the most obvious gap: std::generator<T>. It’s a synchronous, pull-based, lazy sequence. You write a function that co_yields values, and the caller iterates over them with a range-for. No allocator gymnastics, no hand-rolled promise types. See appendix for an example of creating your own generator.

Let’s see this feature in detail, with some examples and illustrate some caveats.

What std::generator actually is

An std::generator<T> is a view. It satisfies std::ranges::input_range. That means it plugs into the entire ranges pipeline. You can use it with useful functions like views::take, views::filter, views::transform and without ever materializing the full sequence in memory.

The function suspends at each co_yield and resumes when the caller asks for the next element. Values are produced one at a time, on demand. If the caller stops iterating, the remaining values are never computed.

This makes it trivial to express infinite sequences, stateful streams, or anything where computing the entire result set upfront is wasteful or impossible.

At Stream HPC, we enjoy opportunities to connect with the HPC and accelerator community, exchange ideas, and learn from engineers and researchers working across the GPU ecosystem. Later this month, several members of our team will be attending GPU Day 2026 in Budapest, Hungary.

Now in its 16th edition, GPU Day has become an established annual conference focused on massively parallel computing in science and industrial applications. Organized by the Wigner Scientific Computation Laboratory, the event brings together researchers, developers, students, and industry experts to discuss technologies spanning GPUs, compilers, machine learning, visualization, and emerging accelerator platforms.

For Stream HPC, events like GPU Day are a natural fit. We work with clients across a wide range of hardware platforms and software ecosystems, helping optimize and accelerate applications where performance matters. Conferences like this provide an opportunity to share practical experiences from real-world projects and contribute to the broader HPC community.

This year two Stream HPC engineers will present their work at the conference:

Parallel primitives are the ubiquitous building blocks of GPU programming with CUDA and HIP, to make your life as a programmer easier. Primitives like scans, reductions, and sorts operate in parallel over large data inputs. The basic use case has input and output residing in device memory as an array of values. However, the libraries provided by NVIDIA and AMD allow the use of iterators, which abstract the concept of input and output. An iterator is a type that behaves as a pointer, but overrides part of the dereferencing and arithmetic logic. With creative use of iterators, using the parallel primitives can become simpler and more performant. Assuming basic familiarity with the primitives, in this article we will show two examples of how iterators can be used to create better GPU programs: finding the arguments of the maximum of a function using a zip iterator in HIP and fast reduction with equally-sized segments in CUDA.

Leveraging the zip iterator to find the maximum argument

A reduction with a binary maximum operator finds the maximum element in the input. But what if we are also interested in where this element is located in the input? For the first example, let us find the argument of the maximum of a function with HIP. To make it a bit more interesting, suppose we want to find an index of an integer value that has most “set” bits.

A naive approach might be to first perform a reduction over the entire input, and then find the index that produced that input. However, by making clever use of iterators we can use just a single reduction. This is where the zip iterator comes in, which “zips” two iterators together: the dereferenced type is a tuple of both iterator’s value types. By zipping a counting iterator and an input array, we can enumerate the input elements.

thrust::device_vector<unsigned int> d_in(num_elements);
auto iter = rocprim::make_zip_iterator(
    rocprim::make_tuple(rocprim::make_counting_iterator(0), d_in.data()));