3 ms·

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 divergenc

• 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.
    • raphlinus 6y ago
      Saying there is no execution mask from the ISA perspective isn't quite true either: `__activemask` in CUDA translates to a simple fetch of that mask, and the Vulkan subgroup operations also access it - I believe subgroupBallot(true) is equivalent. The way I would put this: the hardware is SIMD, and the number of operations executed per clock is the same as pure SIMD, but the independent instruction pointer per thread gives the scheduler considerably more flexibility, which is useful for avoiding deadlocks on synchronization, and also helpful for hiding memory access latency. On top of that, there's an abstraction of a large number of scalar threads running on the hardware. It's a leaky abstraction, though, as performance really requires sympathy with the underlying SIMD reality, and also the subgroup operations also expose a good deal of the implementation details of that abstraction.