What is the right way to implement a stack push on OpenCL 1.2?

Question

Question

Suppose multiple work-items want to append to a global stack:

void kernel(__global int* stack) {
    ... do stuff ...
    push(stack, value);
    ... do stuff ...
    return y;
}

It is desirable that, after the kernel runs, stack contains every value pushed to it. Order does not matter. What is the proper way to do it in OpenCL 1.2?

What I've tried

An obvious idea would be to use atomic_inc to get the length and just write to it:

void push(__global int* stack, int val) {
    int idx = atomic_inc(stack) + 1; // first element is the stack length
    stack[idx] = val;
}

But I speculate having all work-items call atomic_inc separately on the same memory position ruins the parallelism. A separate idea would be to just write to a temporary array larger than the number of work items:

void push(__global int* stack, int val) {
    stack[get_global_id(0)] = val;
}

That'd leave us with a sparse array of values:

[0, 0, 0, 7, 0, 0, 0, 2, 0, 0, 3, 0, 0, 0, 9, 0, 0, ...]

Which could then be compacted using "stream compaction". I, thus, wonder what of those ideas is the most efficient, and if perhaps there is a third option I'm not aware of.

Answer 1

I can't give you a definite answer here, but I can make a few suggestions of things to try - if you have the resources, try to implement more than one of them and profile their performance on all the different types of OpenCL implementation you're planning to deploy on. You might find that different solutions perform differently on different hardware/software.

1. Create a stack per work-group in local memory (either explicitly or by compacting after all values have been generated) and only increment the global stack by the per-work-group count and copy the whole local stack into the global one. This means you only have one global atomic add per work-group. Works better for large groups of course.
2. Your biggest source of atomic contention in the naive approach will be from items on the same work-group. So you could create as many stacks as items per work group, and have each item in the group submit to its "own" stack. You'll still need a compaction step after this to combine it all into one list. Vary group size if you try this. I'm not sure to what extent current GPUs suffer from false sharing (atomics locking a whole cache line, not just that word) so you'll want to check that and/or experiment with different gaps between stack counters in memory.
3. Write all results to fixed offsets (based on global id) an array large enough to catch the worst case, and queue a separate compaction kernel that post-processes the result into a contiguous array.
4. Don't bother with a compact representation of the result. Instead, use the sparse array as the input for the next stage of computation. This next stage's work group can compact a fixed subset of the sparse array into local memory. When that's done, each work item then works on one item of the compacted array. Iterate inside the kernel until all have been processed. How well this works will depend on how predictable the statistical distribution of the sparse items in the array is, and your choice of work group size and how much of the sparse array each work group processes. This version also avoids the round trip to the host processor.
5. On Intel IGPs specifically, I have heard that DirectX/OpenGL/Vulkan geometry shaders with variable number of outputs perform exceptionally well. If you can write your algorithm in the format of a geometry shader, this might be worth a try if you're targeting those devices. For nvidia/AMD, don't bother with this.

There are probably other options, but those should give you some ideas.