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.

The examples use HIP syntax but the patterns work for CUDA, SYCL, Vulkan, or anything else where you push a C++ struct across an API boundary into device-side code.

Note: All code in this post requires C++26 reflection. As of writing, GCC 16 ships P2996 as a partial implementation in mainline (both the core language and the <meta> library), which is what every godbolt link in this post uses. Bloomberg’s clang-p2996 fork was the original prototype and still covers a few corners GCC has not caught up to yet. Mainline Clang and MSVC have nothing yet. You can compile any example on Compiler Explorer with -std=c++26 -freflection -fexpansion-statements on a GCC 16 trunk build.

A two-minute reflection primer

If you’ve already read a reflection tutorial, skip this. If not, the entire API fits on a postcard.

You reflect a type, value, or template with the reflection operator ^^, which produces a std::meta::info. You can splice an info back into ordinary C++ with the splicer [: r :]. Also, you iterate over reflections at compile time with template for. Everything in std::meta is consteval.

struct Particle { float x, y, z; float mass; };

template <typename T>
consteval auto field_names()
{
    constexpr auto ctx = std::meta::access_context::current();
    static constexpr auto members = std::define_static_array(
                                        std::meta::nonstatic_data_members_of(^^T, ctx)
                                    );
    std::array<std::string_view, members.size()> out{};
    std::size_t i = 0;
    template for (constexpr auto m : members)
    {
        out[i++] = std::meta::identifier_of(m);
    }
    return out;
}

static_assert(field_names<Particle>()[3] == "mass");

nonstatic_data_members_of(^^T) returns a vector<info> of the data members in declaration order. template for is the new compile-time loop that instantiates its body once per element, so m is a constexpr reflection inside each iteration. identifier_of(m) gives you the spelling as a string_view. The whole thing is evaluated in a consteval context, so the result is baked into
the binary.

Note: std::meta::nonstatic_data_members_of returns std::vector<std::meta::info>, and because std::vector can allocate, it cannot be constexpr as a non-temporary. To work around this you wrap the result in std::define_static_array (P3491), which produces a compile-time-friendly static array.
This example can be found here.

That’s enough vocabulary for everything that follows. The metafunctions worth remembering are members_of, nonstatic_data_members_of, enumerators_of, parameters_of, bases_of, identifier_of, type_of, offset_of, size_of, alignment_of, and annotations_of. The synthesis side has define_aggregate for building new types and data_member_spec for describing the members you want to inject.

The crucial property is that splicing turns a reflection back into a real language entity. obj.[:m:] accesses a member. typename [: t :] names a type. That round-trip between “the meta universe” and “the regular universe”, to borrow Lemire’s phrasing in his equality-check post, is the whole engine.

Automatic struct-to-device layout validation

The classic GPU bug. You have a host struct, you memcpy it into a constant buffer, and your kernel reads it back through a struct with the same name but slightly different layout. Maybe device code aligned a double3 differently. Maybe somebody removed a float pad that was load-bearing. The kernel reads the wrong offsets and produces plausible-looking garbage.

The pattern is straightforward: write a consteval function that compares two types field by field and fails the build if they don’t agree.

template <typename Host, typename Device>
consteval bool layouts_match()
{
    constexpr auto ctx = std::meta::access_context::current();
    static constexpr auto h = std::define_static_array(
                                  std::meta::nonstatic_data_members_of(^^Host, ctx)
                              );
    static constexpr auto d = std::define_static_array(
                                  std::meta::nonstatic_data_members_of(^^Device, ctx)
                              );

    if (h.size() != d.size()) return false;
    if (sizeof(Host) != sizeof(Device)) return false;
    if (alignof(Host) != alignof(Device)) return false;

    for (std::size_t i = 0; i < h.size(); ++i)
    {
        if (std::meta::identifier_of(h[i]) !=
            std::meta::identifier_of(d[i])) return false;
        if (std::meta::offset_of(h[i]) !=
            std::meta::offset_of(d[i])) return false;
        if (std::meta::size_of(h[i]) !=
            std::meta::size_of(d[i])) return false;
    }
    return true;
}

// In the header that defines both:
struct LaunchParams        { int n; float dt; float3 origin; uint32_t flags; };
struct LaunchParams_device { int n; float dt; float3 origin; uint32_t flags; };

static_assert(layouts_match<LaunchParams, LaunchParams_device>(),
              "Host and device layouts diverged. Check field order, "
              "padding, and alignment.");

Note: The above code example can be found here.

A few things to notice. The check is consteval, so it runs at compile time with zero runtime cost. The static_assert triggers in the translation unit that includes both definitions, which means the build breaks at the point of mismatch, not at the kernel launch site three hundred files away. The error message is yours, written in English, and it points at the kind of problem rather than at a template instantiation stack.

In practice the device struct is usually the same type as the host struct, in which case this check is redundant. The interesting case is when the kernel is written in a vendor IR (SPIR-V, AMDGCN assembly) or in a separate compilation unit with different target-specific types. There the two structs really are different, and the static_assert is the cheapest insurance you’ll ever buy.

For the more useful case where the host and device struct are nominally the same, the check is still valuable as a padding guard. The same function works across compiler versions and target flags:

template <typename T>
consteval bool is_packed()
{
    constexpr auto ctx = std::meta::access_context::current();
    static constexpr auto m = std::define_static_array(
                                  std::meta::nonstatic_data_members_of(^^T, ctx)
                              );
    std::size_t expected = 0;
    template for (constexpr auto f : m)
    {
        if (std::meta::offset_of(f).bytes != expected) return false;
        expected += std::meta::size_of(f);
    }
    return expected == sizeof(T);
}

struct PushConstants { uint32_t a; uint64_t b; uint32_t c; };
static_assert(is_packed<PushConstants>(),
              "PushConstants has internal padding. Reorder members or "
              "add explicit pad fields.");

Note: The above code example can be found here.

This catches the case where somebody adds a uint64_t after a uint32_t and the compiler silently inserts four bytes. On a discrete GPU those four bytes are a real DMA transfer.

Generating SoA structs from AoS definitions

Almost every GPU kernel that takes a list of “things” wants the data in structure of arrays form rather than array of structures. You declared std::vector<Particle> because that’s what made sense for the rest of your code, but the kernel wants four pointers: float* x, float* y, float* z,
float* mass. The traditional workaround is a hand-written shadow struct that gets out of sync the moment somebody adds a member to Particle.

Reflection lets you generate the SoA type from the AoS type. The standard synthesis function is define_aggregate, which takes a forward declaration of a class and a list of data_member_spec descriptions and finishes the definition in place.

template <typename Aos>
struct soa_of_helper {
    struct type;
    consteval {
        constexpr auto ctx = std::meta::access_context::current();
        std::vector<std::meta::info> specs;
        template for (constexpr auto m : std::define_static_array(
                                             std::meta::nonstatic_data_members_of(^^Aos, ctx)
                                         ))
        {
            // Build a pointer-to-T member with the same name as the AoS field.
            auto member_type = std::meta::add_pointer(std::meta::type_of(m));
            specs.push_back(std::meta::data_member_spec(member_type, {
                .name = std::meta::identifier_of(m)
            }));
        }
        std::meta::define_aggregate(^^type, specs);
    }
};

template <typename Aos>
using soa_of = typename soa_of_helper<Aos>::type;

struct Particle { float x, y, z; float mass; };

// Generated: struct { float* x; float* y; float* z; float* mass; };
using ParticleSoA = soa_of<Particle>;

__global__ void update(ParticleSoA p, float dt, int n)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= n) return;
    p.x[i] += p.vx_or_whatever_you_added_to_Particle[i] * dt;
    // ... etc.
}

Note: The above code example can be found here.

The consteval { ... } block is a consteval block from the same proposal family. It’s a piece of compile-time imperative code that runs at the point where it appears in the program, which is what lets you call define_aggregate to retroactively finish the forward-declared type. The body builds a vector of data_member_spec, each one consisting of “a pointer to the original field type, named the same as the original field.” Once the consteval block returns, type is a complete aggregate and can be used like any other struct.

Add a member to Particle, recompile, and ParticleSoA grows a matching pointer. The kernel that takes ParticleSoA by value sees the new field immediately. There is no generator script, no shadow header, no #define FOREACH_PARTICLE_FIELD(X) list to keep in sync.

The same machinery handles the inverse direction. Given an SoA struct of pointers, you can synthesize an indexing helper that takes an int i and returns a populated AoS struct, with each member loaded by index. The point is not the specific gadget. The point is that the relationship “field X of the AoS type corresponds to a pointer to X in the SoA type” stops being a comment in a README and becomes a property the compiler enforces.

Compile-time kernel argument binding and push-constant packing

A HIP or CUDA kernel launch is just a function call with extra ceremony. The arguments are copied into a small parameter buffer that the driver hands to the GPU. For Vulkan compute or D3D12 the analogous thing is a push constant: a small chunk of memory the application packs at the host and the shader unpacks on the device. In both cases the contract is “the bytes in this buffer, at these offsets, with these types.”

Hand-written push-constant packing looks like this, and it is exactly as awful as it sounds:

std::byte buf[128];
std::memcpy(buf + 0, &transform, sizeof(transform));
std::memcpy(buf + 64, &lighting, sizeof(lighting));
std::memcpy(buf + 96, &debug_flag, sizeof(debug_flag));
// hope you got the offsets right

With reflection you describe the layout with a regular struct and generate the packer from it:

template <typename Layout>
void pack_push_constants(std::byte* dst, const Layout& src)
{
    constexpr auto ctx = std::meta::access_context::current();
    template for (constexpr auto m : std::define_static_array(
                                         std::meta::nonstatic_data_members_of(^^Layout, ctx)
                                     ))
    {
        constexpr std::size_t off = std::meta::offset_of(m).bytes;
        constexpr std::size_t sz  = std::meta::size_of(m);
        std::memcpy(dst + off, &src.[:m:], sz);
    }
}

struct DrawPC {
    glm::mat4 transform;
    glm::vec4 tint;
    uint32_t  material_id;
};

DrawPC pc{...};
pack_push_constants(scratch, pc);
vkCmdPushConstants(cmd, layout, stage, 0, sizeof(DrawPC), scratch);

Note: The above code example can be found here.

Three things to notice. First, the offsets and sizes are constants. They are known at compile time because they come from std::meta::offset_of and std::meta::size_of, both of which return consteval values. A good optimizer turns the whole loop into a single memcpy of sizeof(Layout) bytes, which is exactly what you’d write by hand. Second, the access src.[:m:] is the splicer doing real work: it expands to src.transform, src.tint, src.material_id across the three iterations of the template for. Third, the function is generic over Layout. One function template covers every push constant struct in your codebase.

You can take this further. If your backend cares about ordering within the buffer (some APIs do), you can sort the members at compile time before emitting, and assert that the user-facing struct matches the chosen order:

template <typename Layout>
consteval auto layout_descriptor()
{
    constexpr auto ctx = std::meta::access_context::current();
    static constexpr auto m = std::define_static_array(
                                  std::meta::nonstatic_data_members_of(^^Layout, ctx)
                              );
    std::array<std::pair<std::string_view, std::size_t>, m.size()> out{};
    for (std::size_t i = 0; i < m.size(); ++i)
    {
        out[i] = { std::meta::identifier_of(m[i]),
                   std::meta::offset_of(m[i]).bytes };
    }

    // Sort by offset ascending so the descriptor reflects buffer order,
    // regardless of declaration order in the struct.
    std::ranges::sort(out,
              [](const auto& a, const auto& b) { return a.second < b.second; });
    return out;
}

The returned array is a constexpr value. You can use it to drive runtime validation (compare against the layout the driver reports), or you can dump it into a static_assert that diffs your struct against a reference layout the shader author committed.

Automatic enum-to-dispatch-table for kernel variants

GPU kernels live and die by specialization. A reduction kernel for float is different from a reduction kernel for __bf16. A GEMM with M=128, N=128, K=32 is a different binary from M=64, N=64, K=128. You write the kernel as a template and instantiate it thousands of times. Then you need a dispatch table that maps “the runtime tile size the user asked for” to “the specific instantiation.”

The boilerplate version is a switch:

switch (cfg.tile) {
    case TileSize::T64:  return launch<TileSize::T64>(args);
    case TileSize::T128: return launch<TileSize::T128>(args);
    case TileSize::T256: return launch<TileSize::T256>(args);
    default: std::unreachable();
}

This works. It also breaks every time you add a new tile size and forget to add the case. With reflection, the switch writes itself from the enum definition:

enum class TileSize { T64 = 64, T128 = 128, T256 = 256, T512 = 512 };

template <typename Enum, typename Args>
void dispatch(Enum value, const Args& args)
{
    template for (constexpr auto e : std::define_static_array(
                                         std::meta::enumerators_of(^^Enum)
                                     ))
    {
        if (value == [:e:])
        {
            launch<([:e:])>(args);
            return;
        }
    }
    std::unreachable();
}

// Call site:
dispatch(cfg.tile, args);

Note: The above code example can be found here.

enumerators_of(^^TileSize) returns reflections for the four enumerators. [:e:] splices each one as a constexpr value, which is exactly what launch<...> needs as a non-type template argument. Each iteration of template for produces a separate if-branch with its own template instantiation, so the resulting code is identical to the hand-written switch.

Add T1024 to the enum, recompile, and the dispatch picks it up automatically. You’ll get a hard error in the same translation unit if launch<T1024> fails to instantiate, which is the right time to learn about it.

For two-dimensional dispatch (tile size and data type, say) you nest the template for and the combinatorial explosion is now the compiler’s problem instead of yours:

template <typename Tile, typename Dtype, typename Args>
void dispatch2d(Tile t, Dtype d, const Args& args)
{
    template for (constexpr auto te : std::define_static_array(
                                          std::meta::enumerators_of(^^Tile)
                                      ))
    {
        if (t != [:te:]) continue;
        template for (constexpr auto de : std::define_static_array(
                                              std::meta::enumerators_of(^^Dtype)
                                          ))
        {
            if (d == [:de:])
            {
                launch<([:te:]), ([:de:])>(args);
                return;
            }
        }
    }
    std::unreachable();
}

Note: The above code example can be found here.

Note: This pattern can balloon binary size quickly. A 4-tile x 6-dtype x 3-layout dispatch is 72 kernel instantiations. Use it deliberately, and consider partial-specialisation policies that exclude combinations the user can prove they don’t need.

GPU memory layout inspector and padding detector

The host and device disagree on alignment more often than they agree on it. A bool is one byte on the host and four bytes on most shading languages. A double is sometimes 8-byte aligned, sometimes 16. The std140 and std430 layouts have their own opinions. Even within HIP, __align__(16) decorations on device structs can move offsets without changing the source.

A reflection-driven layout report is a small, useful tool. Stick it in a unit test and dump the layout of every kernel-facing struct. Compare two builds of the same struct, or compare against a checked-in golden file.

struct field_layout {
    std::string_view name;
    std::string_view type;
    std::size_t offset;
    std::size_t size;
    std::size_t align;
};

template <typename T>
consteval auto inspect()
{
    constexpr auto ctx = std::meta::access_context::current();
    static constexpr auto m = std::define_static_array(
                                  std::meta::nonstatic_data_members_of(^^T, ctx)
                              );
    std::array<field_layout, m.size()> out{};
    for (std::size_t i = 0; i < m.size(); ++i)
    {
        out[i] = {
            .name   = std::meta::identifier_of(m[i]),
            .type   = std::meta::display_string_of(std::meta::type_of(m[i])),
            .offset = std::meta::offset_of(m[i]).bytes,
            .size   = std::meta::size_of(m[i]),
            .align  = std::meta::alignment_of(std::meta::type_of(m[i]))
        };
    }
    return out;
}

template <typename T>
void print_layout(const char* label)
{
    constexpr auto layout = inspect<T>();
    std::println("=== {} (size={}, align={}) ===",
                 label, sizeof(T), alignof(T));
    std::size_t expected = 0;
    for (auto f : layout)
    {
        if (f.offset != expected)
        {
            std::println("  -- {} bytes of padding --", f.offset - expected);
        }
        std::println("  +{:>4}  {:<20} {} (size={}, align={})",
                     f.offset, f.name, f.type, f.size, f.align);
        expected = f.offset + f.size;
    }
    if (expected != sizeof(T))
    {
        std::println("  -- {} bytes of trailing padding --",
                     sizeof(T) - expected);
    }
}

Note: The above code example can be found here.

=== MaterialParams (size=96, align=16) ===
  +   0  base_color           float [4] (size=16, align=4)
  +  16  emissive             float [3] (size=12, align=4)
  +  28  emissive_strength    float (size=4, align=4)
  +  32  roughness            float (size=4, align=4)
  +  36  metallic             float (size=4, align=4)
  +  40  ao                   float (size=4, align=4)
  +  44  normal_scale         float (size=4, align=4)
  +  48  albedo_tex_idx       unsigned int (size=4, align=4)
  +  52  normal_tex_idx       unsigned int (size=4, align=4)
  +  56  roughness_tex_idx    unsigned int (size=4, align=4)
  +  60  ao_tex_idx           unsigned int (size=4, align=4)
  +  64  emissive_tex_idx     unsigned int (size=4, align=4)
  +  68  flags                unsigned int (size=4, align=4)
  +  72  alpha_cutoff         float (size=4, align=4)
  +  76  clearcoat            float (size=4, align=4)
  +  80  dummy_test_pad       char [9] (size=9, align=1)
  -- 7 bytes of trailing padding --

Running print_layout<MaterialParams>("MaterialParams") produces a compact report you can paste into a PR description, diff against last week’s build, or hand to a shader author who is convinced the bug is on the C++ side. The padding detection makes the cost of a poorly-ordered struct visible. Visible is the first step to fixable.

inspect returns its result by value, but the result is a consteval expression of std::array, so the compiler folds the entire layout description into static initializer data. The runtime cost of print_layout is the cost of the prints, nothing more.

Shader and kernel metadata via annotations

The most underrated feature in C++26 is P3394, annotations for reflection. The syntax is [[=value]] on a declaration, and the value is any structural type. Reflection queries can read the attached annotations back. This is the missing piece that turns reflection from a “compute things about types” tool into a “carry intent from the source into the compile-time program” tool.

The GPU use case is screaming for it. A kernel argument struct has all sorts of metadata that today lives in comments or in a parallel YAML file: which register set, which binding slot, which memory space, whether a buffer is read-only, whether a sampler is anisotropic. Annotations let you put it on the declaration and pull it back out at codegen time.

// Annotation vocabulary
namespace gpu::annotations {
    struct binding   { uint32_t set, slot; };
    struct readonly  {};
    struct storage   {};
    struct uniform   {};
}

// Domain structs
struct Particle { float x, y, z; float mass; };
struct DebugCounter { uint32_t draws, culled, overdraw_samples; uint32_t _pad; };

struct DrawResources {
    [[=gpu::annotations::binding(0, 0), =gpu::annotations::uniform{}]]
    glm::mat4*    camera;

    [[=gpu::annotations::binding(0, 1), =gpu::annotations::storage{}, =gpu::annotations::readonly{}]]
    Particle*     particles;

    [[=gpu::annotations::binding(0, 2), =gpu::annotations::storage{}]]
    DebugCounter* counters;
};

The reflection side reads them back. std::meta::annotations_of(m) returns the list of annotations on a member, and value_of (or extract<T>) recovers the typed value:

// Annotation helper
template <typename Annotation, std::meta::info m>
consteval auto find_annotation() -> std::optional<Annotation>
{
    static constexpr auto annotations = std::define_static_array(
                                            std::meta::annotations_of(m)
                                        );
    template for (constexpr auto a : annotations)
    {
        if (std::meta::remove_const(std::meta::type_of(a)) == ^^Annotation)
        {
            return std::meta::extract<Annotation>(a);
        }
    }
    return std::nullopt;
}

// Descriptor builder
template <typename Resources>
consteval auto build_descriptor_set()
{
    constexpr auto ctx = std::meta::access_context::current();
    static constexpr auto m = std::define_static_array(
                                  std::meta::nonstatic_data_members_of(^^Resources, ctx)
                              );
    std::array<VkDescriptorSetLayoutBinding, m.size()> out{};
    std::size_t i = 0;
    template for (constexpr auto member : m)
    {
        constexpr auto bind = find_annotation<gpu::annotations::binding, member>();
        static_assert(bind.has_value(),
                      "Every resource must declare a gpu::annotations::binding.");
        out[i].binding = bind->slot;
        out[i].descriptorType =
            find_annotation<gpu::annotations::storage, member>().has_value()
                ? VK_DESCRIPTOR_TYPE_STORAGE_BUFFER
                : VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER;
        ++i;
    }
    return out;
}

constexpr auto draw_bindings = build_descriptor_set<DrawResources>();

Note: The above code example can be found here.

The whole descriptor-set layout is computed at compile time, from declarations that read like documentation. A new resource means adding a member with annotations. The descriptor set updates itself. The static_assert ensures that nobody adds a binding without saying where it goes.

The same trick works for any “out of band” piece of information you’d otherwise encode in naming conventions or stringly-typed registration calls:

• [[=workgroup_size{256, 1, 1}]] on a kernel function template, read by the launch wrapper to derive grid dimensions.
• [[=numa{0}]] or [[=on_device{1}]] on a buffer member, read by the allocator to pick the right pool.
• [[=docstring{"Energy of the system in joules."}]] on a field, read by the layout-report generator above so the dumped output explains itself.

The Qt people have been particularly vocal about this. Their QRangeModel post shows the same pattern in a totally different domain: meta-properties for UI binding, declared on the data type, harvested by the framework at compile time. The GPU case is structurally identical.

Conclusion

Six examples, one shape. You take metadata that used to live in a code generator, a macro list, a comment, or a wiki page, and you put it on the declaration where it belongs. The compiler reads it back, generates the code that used to be hand-written boilerplate, and breaks the build when the
declaration and the kernel disagree.

Current support

The tooling story has moved fast. GCC 16 is the first mainline compiler to ship P2996, and as of writing it covers the core language pieces (reflection operator, splicers, template for, annotations from P3394, expansion statements from P1306) plus the <meta> library, all marked partial because a handful of corners are still being filled in. Bloomberg’s clang fork is still useful as a cross-check and was the implementation everything was prototyped against. Mainline Clang and MSVC have nothing yet, which means cross-compiler portability is, for the moment, theoretical. Error messages when a consteval block fails halfway through still test your patience. Linker diagnostics for define_aggregate are, in 2026, an acquired taste.

The bigger picture

Having said that, the underlying model is correct and the GPU use case is one of the most compelling targets for it. Layouts, bindings, kernel variants, and metadata are exactly the things reflection was designed to manipulate, and they happen to be exactly the things GPU programmers spend half their time getting wrong by hand. Whether you adopt reflection on the day your toolchain ships it or wait until C++29 fills in the gaps (code injection from P2237, richer parameter reflection from P3096), it is worth learning now.

The shift this represents is bigger than any one feature. GPU library authors have spent the last decade papering over the lack of language-level metadata with code generators, embedded DSLs, and Python build steps. Reflection does not retire any of those tools overnight, but it shrinks the surface area they need to cover. The parts of those systems that really do amount to “describe this struct to the runtime” can now be plain C++, sitting next to the kernel that uses them. That is a much smaller, much friendlier build, and it is the first time in a long while that the host side of a GPU codebase gets to feel like a modern language instead of a preprocessor exhibit.

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.

Continue reading “Lazy Ranges in C++23 with std::generator” →

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:

Continue reading “GPU Day 2026” →

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()));

Continue reading “Thinking with iterators in CUDA and HIP” →

Via our Twitter channel we have various polls. Not always have we shared the full background of these polls, so we’ve taken the polls of the past half year and put them here. The first half of the year there were no polls, in case you wanted to know.

As inclusive polls are not focused (and thus difficult to answer), most polls are incomplete by design. Still insights can be given. Or comments given.

Below’s polls have given us insight and we hope they give you insights too how our industry is developing. It’s sorted on date from oldest first.

It was very interesting that the percentage of votes per choice did not change much after 30 votes. Even when it was retweeted by a large account, opinions had the same distribution.

Is HIP (a clone of CUDA) an option?

Continue reading “The 12 latest Twitter Poll Results of 2018” →

This series “Basic concepts” is based on GPGPU-questions we get via email more than once, or when the question is not clearly explained in the books. For one it is obvious, for the other just what they’re missing.

They say that learning a new technique is best done by playing around with working code and then try to combine it. The idea is that when you have Stackoverflowed and Githubed code together, you’ve created so many bugs by design that you’ll learn a lot if you make it work. When applying this to OpenCL, you quickly get to a situation that you want to run one.cl file and then another.cl file. Almost all beginner’s material discuss a single OpenCL-file, so how to do this elegantly?

Continue reading “OpenCL Basics: Running multiple kernels in OpenCL” →

After the last post on Excel (“Accelerating an Excel Sheet with OpenCL“), there have been various request and discussions how we do “the miracle”. Short story: we only apply proper engineering tactics. Below I’ll explain how you can also speed up Excel and when you actually have to call us (last step).

A computer can handle 10s of gigabytes per second. Now look how big your Excel-sheet is and how much time it takes. Now you understand that the problem is probably not your computer.

Excel is a special piece of software from a developer’s perspective. An important rule of software engineering is to keep functionality (code) and data separate. Excel mixes these two as no other, which actually goes pretty well in many cases unless the data gets too big or the computations too heavy. In that case you’ve reached Excel’s limits and need to properly solve it.

An Excel-file often does things one-by-one, with a new command in every cell. This prevents any kind of automatic optimizations – besides that, Excel-sheets are very prone to errors.

Below are the steps to go through, of which most you can do yourself!

Continue reading “How to speed up Excel in 6 steps” →

The goal of this writing is to explain which applications are suitable to be ported to OpenCL and run on GPU (or multiple GPUs). It is done by showing the main differences between GPU and CPU, and by listing features and characteristics of problems and algorithms, which can make use of highly parallel architecture of GPU and simply run faster on graphic cards. Additionally, there is a list of issues that can decrease potential speed-up.

It does not try to be complete, but tries to focus on the most essential parts of assessing if code is a good candidate for porting to the GPU.

GPU vs CPU

The biggest difference between a GPU and a CPU is how they process tasks, due to different purposes. A CPU has a few (usually 4 or 8, but up to 32) ”fat” cores optimized for sequential serial processing like running an operating system, Microsoft Word, a web browser etc, while a GPU has a thousands of ”thin” cores designed to be very efficient when running hundreds of thousands of alike tasks simultaneously.

A CPU is very good at multi-tasking, whereas a GPU is very good at repetitive tasks. GPUs offer much more raw computational power compared to CPUs, but they would completely fail to run an operating system. Compare this to 4 motor cycles (CPU) of 1 truck (GPU) delivering goods – when the goods have to be delivered to customers throughout the city the motor cycles win, when all goods have to be delivered to a few supermarkets the truck wins.

Most problems need both processors to deliver the best value of system performance, price, and power. The GPU does the heavy lifting (truck brings goods to distribution centers) and the CPU does the flexible part of the job (motor cycles distributing doing deliveries).

Assessing software for GPU-porting fitness

Software that does not meet the performance requirement (time taken / time available), is always a potential candidate for being ported to a GPU. Continue reading “Selecting Applications Suitable for Porting to the GPU” →

When CUDA kept having a dominance over OpenCL, AMD introduced HIP – a programming language that closely resembles CUDA. Now it doesn’t take months to port code to AMD hardware, but more and more CUDA-software converts to HIP without problems. The real large and complex code-bases only take a few weeks max, where we found that solved problems also made the CUDA-code run faster.

The only problem is that CUDA-libraries need to have their HIP-equivalent to be able to port all CUDA-software.

Here is where we come in. We helped AMD make a high-performance Pseudo Random Generator (PRNG) Library, called rocRAND. Random number generation is important in many fields, from finance (Monte Carlo simulations) to Cryptographics, and from procedural generation in games to providing white noise. For some applications it’s enough to have some data, but for large simulations the PRNG is the limiting factor. Continue reading “Learn about AMD’s PRNG library we developed: rocRAND – includes benchmarks” →

While going through my email, I found out about the third “HiPEAC Student Heterogeneous Programming Challenge”. Unfortunately the deadline was last week, but just got an email: if you register by this weekend (17 September), you can still join.

EDIT: if you joined, be sure to comment in early November how it was. This would hopefully motivate others to join in next year. Continue reading “GPU and FPGA challenge for MSc and PhD students” →