4 ms·

This looks really interesting - but I have to admit it's gone completely over my head. Can someone break this down into simpler terms?

Const-me 8y ago
TLDR: registers are faster to access than memory. Even warp-local on-chip memory. While CUDA looks like scalar code, GPUs are vector based, that scalar code is executed by 32-wide vector units. The threads running on the same units are called warp. There're little-known intrinsics in CUDA, __shfl_sync, __shfl_up_sync, __shfl_down_sync, that allow to permute data across lanes, thus exchanging data across threads of the warp. Replacing local memory access with registers is what gave the profit.
dragontamer 8y ago
In CUDA PTX, there is an instruction called "shfl.sync.idx", which allows you to shuffle data between threads arbitrarily. However, this shuffling is only efficient if you follow a variety of rules: avoiding bank conflicts and so forth. Lets say you have 32-threads of a Warp running. Thread0, Thread1, Thread2, etc. etc. Now lets say they all have a variable named "x". Thread0.x, Thread1.x, Thread2.x, etc. etc. What "shfl.sync" does, is it allows these "x" variables to be shuffled between all threads in one step. // Parallel and synchronous: All threads swap // values arbitrarily. This is a simple "rotate", // but all combinations are legal for a shuffle Thread0.x = Thread13.x Thread1.x = Thread14.x Thread2.x = Thread15.x Thread3.x = Thread16.x ... Thread30.x = Thread11.x Thread32.x = Thread12.x That's a swizzle. Now the descriptions on the right? The 15, 16, 17, etc. etc. can be any combination. You can broadcast (have Thread1.x copied to all 32 other threads). You can rotate, you can shuffle. All 32^32 specifications of CUDA-thread transfers is legal. Now, it may be legal, but broadcasts are slower due to bank conflicts. So following the special rules to make fast swizzles is an annoying task for sure. This paper helps the programmer create optimal swizzles for their program. For more information: * https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-shfl-sync https://docs.nvidia.com/cuda/parallel-thread-execution/index... * https://devblogs.nvidia.com/using-cuda-warp-level-primitives/ https://devblogs.nvidia.com/using-cuda-warp-level-primitives... TL;DR: GPUs have a rich set of highly-efficient data movement operations. They are easy to use, but hard to optimize. Paper presents concepts that help optimize them automatically.
- avinium 8y ago
  Great explanation, appreciate the effort. Will dig deeper into the provided links.

Const-me 8y ago

TLDR: registers are faster to access than memory. Even warp-local on-chip memory. While CUDA looks like scalar code, GPUs are vector based, that scalar code is executed by 32-wide vector units. The threads running on the same units are called warp. There're little-known intrinsics in CUDA, __shfl_sync, __shfl_up_sync, __shfl_down_sync, that allow to permute data across lanes, thus exchanging data across threads of the warp. Replacing local memory access with registers is what gave the profit.

dragontamer 8y ago

In CUDA PTX, there is an instruction called "shfl.sync.idx", which allows you to shuffle data between threads arbitrarily. However, this shuffling is only efficient if you follow a variety of rules: avoiding bank conflicts and so forth. Lets say you have 32-threads of a Warp running. Thread0, Thread1, Thread2, etc. etc. Now lets say they all have a variable named "x". Thread0.x, Thread1.x, Thread2.x, etc. etc. What "shfl.sync" does, is it allows these "x" variables to be shuffled between all threads in one step. // Parallel and synchronous: All threads swap // values arbitrarily. This is a simple "rotate", // but all combinations are legal for a shuffle Thread0.x = Thread13.x Thread1.x = Thread14.x Thread2.x = Thread15.x Thread3.x = Thread16.x ... Thread30.x = Thread11.x Thread32.x = Thread12.x That's a swizzle. Now the descriptions on the right? The 15, 16, 17, etc. etc. can be any combination. You can broadcast (have Thread1.x copied to all 32 other threads). You can rotate, you can shuffle. All 32^32 specifications of CUDA-thread transfers is legal. Now, it may be legal, but broadcasts are slower due to bank conflicts. So following the special rules to make fast swizzles is an annoying task for sure. This paper helps the programmer create optimal swizzles for their program. For more information: * https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-shfl-sync https://docs.nvidia.com/cuda/parallel-thread-execution/index... * https://devblogs.nvidia.com/using-cuda-warp-level-primitives/ https://devblogs.nvidia.com/using-cuda-warp-level-primitives... TL;DR: GPUs have a rich set of highly-efficient data movement operations. They are easy to use, but hard to optimize. Paper presents concepts that help optimize them automatically.

avinium 8y ago
Great explanation, appreciate the effort. Will dig deeper into the provided links.