Changes

=The Implicitly Parallel SIMD SIMT Programming Model=

CUDA, OpenCL, HIP, and even other GPU languages like GLSL, HLSL, C++AMP and even non-GPU languages like Intel [https://ispc.github.io/ ISPC] all have the same model of implicitly parallel programming. Gangs of All threads called Warps are given an identifier: a threadIdx in CUDA or Wavefronts local_id in OpenCL . Aside from this index, all threads will execute on a SIMD unit concurrentlythe same code.  The GPU executes smallest group of threads is called a CUDA Warp, or OpenCL Wavefront. NVidia GPUs execute 32-threads per warp , while AMD GCN GPUs execute 64-threads per wavefront. All threads within a Warp or Wavefront share an instruction pointer. Consider the following CUDA code:  if(threadIdx.x == 0){ doA(NVidia) or wavefront ; } else { doB(AMD) at a time; } While there is only one thread in the warp that has threadIdx == 0, all 32 threads stepping with of the same program counter / warp will have their shared instruction pointerexecute doA() together. This causes issues with ifTo keep the code semantically correct, threads #1 through #31 will have their NVidia Predicate-statements register cleared (or whileAMD Execution Mask cleared), which means the thread will throw away the work after executing a specific statement. Once doA() is complete, the machine will continue and doB(). In this case, thread#0 will have its execution mask-loops: in the GPU hardwarecleared, while threads #1 through #31 will disable themselves if actually complete the results of doB(). This highlights the rest fundamental trade off of the gang needs GPU platform. GPUs have many threads of execution, but they are forced to execute an with their warps or wavefronts. In complicated loops or trees of if-statement. This is called statements, this thread divergence and is a common source of GPU inefficiencyproblem can cause your code to potentially leave many hardware threads idle. == Building up to larger thread groups ==