4 ms·

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, instea

• 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.
      • my123 6y ago
        SIMT became an abstraction that was lowered to SIMD at compile time to instead being the abstraction used when targeting the underlying hardware. > Furthermore, the PTX is substantially similar PTX is just a (forwards-compatible) intermediate representation. > The per-thread instruction pointers is pretty transparent in most code Except control flow divergence, which is what changes there, yes. > 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...) That's a very good question. AVX-512 with its masking abilities was substantial progress over AVX2. For a CPU, having secondary instruction flows just for the vector units just isn't a (reasonable) option though. If there wasn't the 10nm issues, the next Xeon Phi would have been very interesting on that front. You might also want to look at the Fujitsu A64fx on the Arm side, used in Fugaku. (building a supercomputer with just CPUs, no GPUs) We'll see what will be there in the future... will certainly be very interesting.
• 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.