> I guess we aren’t solving the same problems: memory allocation is trivial is my domain, mapping complex nested control flow in SIMD is the hard part.
Those statements are just confusing to me. In most GPU-code, you have crazy amounts of parallelism and therefore don't really care what order those statements execute in.
As such, if you can allocate memory, and map those if/else statements into a collection of queues and/or stacks (depending on if you want a breadth-first search, or depth-first search pattern).
In effect: you use memory allocation to solve the complex control flow issue. Maybe its more obvious with code:
Instead of doing:
if(baz()){
foo();
} else {
bar(); // Thread divergence!!
}
Do:
if(baz()){
pushIntoFooQueue();
} else {
pushIntoBarQueue();
// Thread divergence, but not much of a penalty
}
while(fooIsNotEmpty()){ // No thread divergence at all
task = SIMDPopFoo();
task.execute();
}
while(barIsNotEmpty()){
task = SIMDPopBar();
task.execute();
}
This is heavier on the memory units, because you now have to manage the data-structures. But this style practically negates the thread-divergence penalty completely. If you're lucky, your fooQueue and barQueue fit in __shared__ memory.
Bonus points: Not only is thread-divergence negated, but you also achieve effective load-balancing across your workgroup. If Thread#0 spawns 20 items for FooQueue, after pushing/popping from the queue, those 20 Foo-tasks will be assigned to 20 different threads.
-----
From there on out, you're just pushing / popping different parts of your code to various queues and/or stacks. But this is only really a valid solution if you have a decent memory allocator that knows where and how to clean up these queues / stacks (especially if you have nodes starting to point to each other for dependency management)
I haven't really solved this problem "in general", but is clear that the queues should be sorted into topological order, and that any tasks that depend onto each other need to be run in different iterations. It really depends on how much you're willing to spend on organizing this execution information.
In any case: the memory allocation issue is one-and-the-same with the thread-divergence / complex instruction flow issue to me. You need to create a data-structure that organizes the instruction flow, and any complex data-structure will need memory management as soon as you start linking things together. The above uses a queue or stack, but things can get more complicated.
--------
Not that Julia, Python, ISPC or anything really... solves this problem. But memory management is very useful for this "style" (I'm mostly doing ref-counted C++ with a custom allocator myself. But such trees or graphs of links can grow into the CPU-side and end up using the default memory allocator)