4 ms·

Apart from tensor cores and some compressed instructions on small data types, they are not SIMD instructions. Each individual SIMT lane can actually diverge fr

by my123 6y ago

Apart from tensor cores and some compressed instructions on small data types, they are not SIMD instructions. Each individual SIMT lane can actually diverge from the other on the underlying hardware. Active threads are dynamically mapped to SIMT units. You can use __syncwarp() to force reconvergence, to stall until all the threads in a warp are at the same location. The disassembled code targeting the underlying ISA for a given CUDA kernel is also easily accessible through nvdisasm.

dragontamer 6y ago
A 32-wide SIMD core executing 16 or 8 or 1-thread at a time is just SIMD execution that's running at 50%, 25%, or 3% utilization. See figure 22 and figure 23 in: https://images.nvidia.com/content/volta-architecture/pdf/volta-architecture-whitepaper.pdf https://images.nvidia.com/content/volta-architecture/pdf/vol... Thread divergence is bad: very very bad. Running 2x, 4x, or 32x slower (or less parallel technically). You want to avoid those situations as much as possible. What NVidia noticed is that thread divergence is necessary for many classical locking algorithms: where serialized code must execute one-at-a-time to run an algorithm correctly. Under these conditions, tracking the instruction pointer, and turning off 97% of your cores to execute 1-at-a-time (instead of 32-at-a-time SIMD) is done. That's WHY __syncwarp() exists, so that you can return to 32-at-a-time execution as soon as possible. Its not always possible for the compiler to figure it out, so the programmer can put a __syncwarp() as a compiler hint that 32-at-a-time is safe again.
- my123 6y ago
  You claimed that SASS was a SIMD instruction set, which is false. (through "NVidia's SASS on Volta is pretty clearly SIMD-instructions") Yes, thread divergence on GPUs can come with performance downsides and is a tool to use carefully.
  - dragontamer 6y ago
    > You claimed that SASS was a SIMD instruction set, which is false. Well, they execute 32-at-a-time SIMD, do they not? There are tricks to split up the execution mask and execute 1-at-a-time, 2-at-a-time, 4, 8, 16, or 31 at a time based on if-statements or thread-locks or whatever. But a "SASS" instruction of "R1 = R2 + R3" is implicitly executed across a 32-wide warp if its execution mask is set as such. ---------- EDIT: Lets put it this way: If you see a "SASS" assembly for R1 = R2 + R3 in isolation, how many adds take place on that clock tick? Somewhere between 1 add, and 32-adds. No more than 32-at-a-time. Seems pretty SIMD to me.
    - my123 6y ago
      There is no execution mask from the ISA perspective, each thread has its own instruction pointer, and the hardware deals with which threads get executed, instead of a mask being set/unset. I don't know why you continue to argue this...
      - dragontamer 6y ago
        Look, lets say we have: if( threadIdx.x % 2 == 0){ // True for 16-threads doA(); // 100 clock ticks } else { // The other 16 threads doB(); // 150 clock ticks } Lets say doA() takes 100 clock ticks, and doB() takes 150 clock ticks on NVidia Volta. How much time does the above code take to run? Answer: 250 clock ticks: doA() is run with 16-threads, then doB() is run on the 16-other threads afterwards. That's how SIMD systems work. If this were non-SIMD 32-core turing machine, it'd take 150 clock ticks (And the threads doing doA() would have spend 50-clock ticks doing something else). The most important thing about SIMD from a performance perspective is that you "add" both halves of the branch from a performance point of view.
        my123 6y ago
        Unlike a vanilla SIMD machine, you can diverge more cheaply when you need to. Scheduler flow control is different from a more conventional masks approach. The underlying implementation is still just lane masking and walking down separate flow control branches sequentially of course. The advantage of such a model however is that the complexity of this is abstracted away from even the compiler. This makes it a change in the programming model. That's why it's not _just_ named SIMD. (I fear that this discussion went too far away arguing semantics... sigh)
        dragontamer 6y ago
        > The advantage of such a model however is that the complexity of this is abstracted away from even the compiler. This makes it a change in the programming model. That's why it's not _just_ named SIMD. That changes the history of the term SIMT. SIMT was first used for NVidia Tesla in 2006: over a decade older than the incremental changes made between Pascal -> Volta/Turing. There has been no name change from Pascal -> Volta/Turing, as far as I'm aware. Furthermore, the PTX is substantially similar. The per-thread instruction pointers is pretty transparent in most code. > The advantage of such a model however is that the complexity of this is abstracted away from even the compiler. This makes it a change in the programming model. That's why it's not _just_ named SIMD. Have you looked at AMD's GPU ISA? Its just jump instructions, extremely similar to NVidia's PTX / SASS instruction set. AMD SIMD doesn't juggle execution masks explicitly either: its handled at a lower level (probably the decoder or something). You still pull the execution mask for things like ballot instructions, but... both AMD and NVidia SIMD are pretty similar. (Similarly, NVidia PTX can still access the execution mask for ballot instructions as well) ------------ Although, I guess both of those GPUs can laugh at AVX512, where the execution masks are explicitly handled by the assembly programmer. But I don't know if explicit execution masks is necessarily a bad thing (its the job of the compiler instead of the decoder or whatever...) > Unlike a vanilla SIMD machine, you can diverge more cheaply when you need to If we take CM2 from 1985 as a "vanilla SIMD machine", it had execution masks and diverged extremely cheaply, just like modern machines. C-Star and star-Lisp even had a programming model very similar to modern CUDA. http://bitsavers.informatik.uni-stuttgart.de/pdf/thinkingMachines/CM2/starLisp_Referenece_Version_5.0_Sep1988.pdf http://bitsavers.informatik.uni-stuttgart.de/pdf/thinkingMac... Back then, you'd use "when" statements to do a parallel divergent branch, while "if" was only for uniform branches. But it wasn't like "when" statements were expensive, they just diverged.

by my123 6y ago

dragontamer 6y ago

A 32-wide SIMD core executing 16 or 8 or 1-thread at a time is just SIMD execution that's running at 50%, 25%, or 3% utilization. See figure 22 and figure 23 in: https://images.nvidia.com/content/volta-architecture/pdf/volta-architecture-whitepaper.pdf https://images.nvidia.com/content/volta-architecture/pdf/vol... Thread divergence is bad: very very bad. Running 2x, 4x, or 32x slower (or less parallel technically). You want to avoid those situations as much as possible. What NVidia noticed is that thread divergence is necessary for many classical locking algorithms: where serialized code must execute one-at-a-time to run an algorithm correctly. Under these conditions, tracking the instruction pointer, and turning off 97% of your cores to execute 1-at-a-time (instead of 32-at-a-time SIMD) is done. That's WHY __syncwarp() exists, so that you can return to 32-at-a-time execution as soon as possible. Its not always possible for the compiler to figure it out, so the programmer can put a __syncwarp() as a compiler hint that 32-at-a-time is safe again.

my123 6y ago
You claimed that SASS was a SIMD instruction set, which is false. (through "NVidia's SASS on Volta is pretty clearly SIMD-instructions") Yes, thread divergence on GPUs can come with performance downsides and is a tool to use carefully.
- dragontamer 6y ago
  > You claimed that SASS was a SIMD instruction set, which is false. Well, they execute 32-at-a-time SIMD, do they not? There are tricks to split up the execution mask and execute 1-at-a-time, 2-at-a-time, 4, 8, 16, or 31 at a time based on if-statements or thread-locks or whatever. But a "SASS" instruction of "R1 = R2 + R3" is implicitly executed across a 32-wide warp if its execution mask is set as such. ---------- EDIT: Lets put it this way: If you see a "SASS" assembly for R1 = R2 + R3 in isolation, how many adds take place on that clock tick? Somewhere between 1 add, and 32-adds. No more than 32-at-a-time. Seems pretty SIMD to me.
  - my123 6y ago
    There is no execution mask from the ISA perspective, each thread has its own instruction pointer, and the hardware deals with which threads get executed, instead of a mask being set/unset. I don't know why you continue to argue this...
    - dragontamer 6y ago
      Look, lets say we have: if( threadIdx.x % 2 == 0){ // True for 16-threads doA(); // 100 clock ticks } else { // The other 16 threads doB(); // 150 clock ticks } Lets say doA() takes 100 clock ticks, and doB() takes 150 clock ticks on NVidia Volta. How much time does the above code take to run? Answer: 250 clock ticks: doA() is run with 16-threads, then doB() is run on the 16-other threads afterwards. That's how SIMD systems work. If this were non-SIMD 32-core turing machine, it'd take 150 clock ticks (And the threads doing doA() would have spend 50-clock ticks doing something else). The most important thing about SIMD from a performance perspective is that you "add" both halves of the branch from a performance point of view.
      - my123 6y ago
        Unlike a vanilla SIMD machine, you can diverge more cheaply when you need to. Scheduler flow control is different from a more conventional masks approach. The underlying implementation is still just lane masking and walking down separate flow control branches sequentially of course. The advantage of such a model however is that the complexity of this is abstracted away from even the compiler. This makes it a change in the programming model. That's why it's not _just_ named SIMD. (I fear that this discussion went too far away arguing semantics... sigh)
        dragontamer 6y ago
        > The advantage of such a model however is that the complexity of this is abstracted away from even the compiler. This makes it a change in the programming model. That's why it's not _just_ named SIMD. That changes the history of the term SIMT. SIMT was first used for NVidia Tesla in 2006: over a decade older than the incremental changes made between Pascal -> Volta/Turing. There has been no name change from Pascal -> Volta/Turing, as far as I'm aware. Furthermore, the PTX is substantially similar. The per-thread instruction pointers is pretty transparent in most code. > The advantage of such a model however is that the complexity of this is abstracted away from even the compiler. This makes it a change in the programming model. That's why it's not _just_ named SIMD. Have you looked at AMD's GPU ISA? Its just jump instructions, extremely similar to NVidia's PTX / SASS instruction set. AMD SIMD doesn't juggle execution masks explicitly either: its handled at a lower level (probably the decoder or something). You still pull the execution mask for things like ballot instructions, but... both AMD and NVidia SIMD are pretty similar. (Similarly, NVidia PTX can still access the execution mask for ballot instructions as well) ------------ Although, I guess both of those GPUs can laugh at AVX512, where the execution masks are explicitly handled by the assembly programmer. But I don't know if explicit execution masks is necessarily a bad thing (its the job of the compiler instead of the decoder or whatever...) > Unlike a vanilla SIMD machine, you can diverge more cheaply when you need to If we take CM2 from 1985 as a "vanilla SIMD machine", it had execution masks and diverged extremely cheaply, just like modern machines. C-Star and star-Lisp even had a programming model very similar to modern CUDA. http://bitsavers.informatik.uni-stuttgart.de/pdf/thinkingMachines/CM2/starLisp_Referenece_Version_5.0_Sep1988.pdf http://bitsavers.informatik.uni-stuttgart.de/pdf/thinkingMac... Back then, you'd use "when" statements to do a parallel divergent branch, while "if" was only for uniform branches. But it wasn't like "when" statements were expensive, they just diverged.