2 ms·

> you use memory allocation to solve the complex control flow issue This is an interesting remark I have to sleep on. I usually don't see this as possible, si

> you use memory allocation to solve the complex control flow issue This is an interesting remark I have to sleep on. I usually don't see this as possible, since I have a bunch of arrays which are allocated once, then a bunch of nested but static control flow. The only way I've found to go beyond single thread performance is with "whole program" SIMD which only worked in the ISPC/OpenCL/CUDA programming models.

dragontamer 6y ago
It doesn't work all the time. Sometimes its just faster to have thread-divergence than to hit VRAM constantly. But its a different paradigm to do things: something to try if your standard if/else stuff isn't working out. EDIT: If you can make due with just stack-allocation (or queue-allocation), if you have singular-sized tasks with predictable sizes, if all the information fits inside of __shared__ memory, and if thread-divergence is normally a problem... this methodology will probably work. Oh, and for a hint: myIdx = prefixSum(__activemask()); myTask = stack[stackTail + myIdx]; if(myIdx == 0){ stackTail -= __popc(__activemask()); // Lulz, I made a bug. Whatever, I think you understand... } __syncThreads(); // Barrier is important Stack pops and pushes are pretty easy. activemask() provides your execution mask, and the logic needed to synchronously push / pop items to a stack (and probably to more complicated memory allocation functions that I haven't figured out yet)