Changes

a specialized processor primarily intended to fast [https://en.wikipedia.org/wiki/Image_processing image processing]. GPUs may have more raw computing power than general purpose [https://en.wikipedia.org/wiki/Central_processing_unit CPUs] but need a specialized and massive parallelized way of programming. [[Leela Chess Zero]] has proven that a [[Best-First|Best-first]] [[Monte-Carlo Tree Search|Monte-Carlo Tree Search]] (MCTS) with [[Deep Learning|deep learning]] methodology will work with GPU architectures. =History=In the 1970s and 1980s RAM was expensive and Home Computers used custom graphics chips to operate directly on registers/memory without a dedicated frame buffer, like [https://en.wikipedia.org/wiki/Television_Interface_Adaptor TIA]in the [[Atari 8-bit|Atari VCS]] gaming system, [https://en.wikipedia.org/wiki/CTIA_and_GTIA GTIA]+[https://en.wikipedia.org/wiki/ANTIC ANTIC] in the [[Atari 8-bit|Atari 400/800]] series, or [https://en.wikipedia.org/wiki/Original_Chip_Set#Denise Denise]+[https://en.wikipedia.org/wiki/Original_Chip_Set#Agnus Agnus] in the [[Amiga|Commodore Amiga]] series. The 1990s would make 3D graphics and 3D modeling more popular, especially for video games. Cards specifically designed to accelerate 3D math, such as the [https://en.wikipedia.org/wiki/Voodoo2 3dfx Voodoo2], were used by the video game community to play 3D graphics. Some game engines, such as [https://en.wikipedia.org/wiki/Quake_(video_game) Quake], could use instead the [[SIMD and SWAR Techniques|SIMD-capabilities]] of CPUs such as the [[Intel]] [[MMX]] instruction set or [[AMD|AMD's]] [[X86#3DNow!|3DNow!]]. Sony's 3D capable chip used in the PlayStation (1994) and Nvidia's 2D/3D combi chips like NV1 (1995) coined the term GPU for 3D graphics hardware acceleration. With the advent of the [https://en.wikipedia.org/wiki/Unified_shader_model unified shader architecture], like in Nvidia [https://en.wikipedia.org/wiki/Tesla_(microarchitecture) Tesla] (2006), ATI/AMD [https://en.wikipedia.org/wiki/TeraScale_(microarchitecture) TeraScale] (2007) or Intel [https://en.wikipedia.org/wiki/Intel_GMA#GMA_X3000 GMA X3000] (2006), GPGPU frameworks like CUDA and OpenCL emerged and gained in popularity.

The traditional job of Early efforts to leverage a GPU is to take the [https://en.wikipedia.org/wiki/Threefor general-dimensional_space x,y,z coordinates] purpose computing required reformulating computational problems in terms of graphics primitives via graphics APIs like [https://en.wikipedia.org/wiki/Triangle_strip trianglesOpenGL OpenGL], and or [https://en.wikipedia.org/wiki/3D_projection mapDirectX DirextX] these triangles to , followed by first GPGPU frameworks such as [https://en.wikipedia.org/wiki/Glossary_of_computer_graphics#screen_space screen space] through a [https:/Lib_Sh Sh/en.wikipedia.org/wiki/Matrix_multiplication matrix multiplication]. As video game graphics grew more sophisticated, the number of triangles per scene grew larger. GPUs similarly grew in size to massively parallel behemoths capable of performing billions of transformations hundreds of times per second. These lists of triangles were specified in Graphics APIs like [https://en.wikipedia.org/wiki/OpenGL OpenGLRapidMind] or [https://en.wikipedia.org/wiki/DirectX DirectXBrookGPU Brook]. But video game programmers demanded more flexibility from their hardware: such as lighting, transparency, and reflections. This flexibility was granted with specialized programming languages, called finally [https://en.wikipedia.org/wiki/Shader#Vertex_shaders vertex shadersCUDA CUDA] or and [https://enwww.wikipediachessprogramming.org/wiki/Shader#Pixel_shaders pixel shadersOpenCL OpenCL]. GPUs evolved to accelerate general purpose compute from pixel shader and vertex shader programmers, and even merged the functionality into "universal" shaders (which can perform either vertex shading or pixel shading). Today, these universal shaders are flexible enough to provide General Purpose compute for GPUs (GPGPU). GPGPU languages, such as OpenCL or CUDA, is how the programmer can access this capability.

The * [https://enwww.wikipediakhronos.org/wikiconformance/Khronos_Group Khronos group] is a committee formed to oversee the [https://en.wikipedia.org/wiki/OpenGL OpenGL], [[OpenCL]], and [https://en.wikipedia.orgadopters/wikiconformant-products/Vulkan_(API) Vulkan] standards. Although compute shaders exist in all languages, OpenCL is the designated general purpose compute language. opencl List of OpenCL 1.2 is widely supported by [[AMDConformant Products]], [[Nvidia]], and [[Intel]]. OpenCL 2.0, although specified in 2013, has had a slow rollout, and the specific features aren't necessarily widespread in modern GPUs yet. AMD continues to target OpenCL 2.0 support in their ROCm environment, while Nvidia has implemented some OpenCL 2.0 features.

* [https://developer.nvidia.com/cuda-zone Nvidia CUDA Zone]* [https://docs.nvidia.com/cuda/parallel-thread-execution/index.html Nvidia PTX ISA]* [https://docs.nvidia.com/cuda/index.html Nvidia CUDA Toolkit Documentation] == AMD Software Overview == [[AMD|AMD's]] original software stack, called [https://en.wikipedia.org/wiki/AMDGPU AMDGPU-pro], provides supports language frontends like OpenCL 1.2 and 2.0 capabilities on [[Linux]] and [[Windows]]. However, most of AMD's efforts today is on an experimental framework called [https://en.wikipedia.org/wiki/OpenCL#Implementations ROCm]. ROCm is AMD's open source compiler and device driver stack intended for general purpose compute. ROCm supports two languages: [https://en.wikipedia.org/wiki/GPUOpen#AMD_Boltzmann_Initiative HIP] (a CUDA-like single-source , C++ compiler also based on LLVM/clang), AMP and OpenCL 2with OpenMP offload directives.0. ROCm only works on Linux machines supporting modern hardware, such as It offers with [https://enrocmdocs.wikipediaamd.org/wiki/PCI_Express#3.0 PCIe 3.0] and relatively recent GPUs (such as the [https:/com/en.wikipedia.org/wikilatest/AMD_Radeon_500_series RX 580ROCm], and [https://en.wikipedia.org/wiki/AMD_RX_Vega_series Vega] GPUs). AMD regularly publishes the assembly language details of their architectures. Their "GCN Assembly" changes slightly from generation to generation, but the fundamental principles have remained the same. AMD's OpenCL documentation, especially the "OpenCL Programming Guide" and the "Optimization Guide" are good places to start for beginners looking to program their GPUs. For Linux developers, the ROCm environment is under active development and has enough features to get code working wellits own parallel compute platform.

== Other 3rd party tools Nvidia ==

* [https://en.wikipedia.org/wiki/DirectCompute DirectComputeCUDA CUDA] (GPGPU API is the parallel computing platform by Microsoft)* [https://en[Nvidia]].wikipedia.org/wiki/OpenMP OpenMP] Device Offload* It supports languages frontends like C, C++, Fortran, OpenCL and offload directives via [https://en.wikipedia.org/wiki/OpenACC OpenACC] Device Offload* and [https://en.wikipedia.org/wiki/Metal_(API) MetalOpenMP OpenMP] (GPU and GPGPU API by Apple)* [https://en.wikipedia.org/wiki/OneAPI_(programming_model) oneAPI] (Data Parallel C++ by Intel)

=The SIMT Programming Model= CUDA, OpenCL, ROCm HIP, all have the same model of implicitly parallel programming. All threads are given an identifier* [https: a threadIdx in CUDA or local_id in OpenCL. Aside from this index, all threads of a kernel will execute the same code. The only way to alter the behavior of code is to use this threadIdx to access different data//developer. The executed code is always implicitly [[SIMD]]nvidia. Instead of thinking of SIMDcom/cuda-lanes, each lane is considered its own thread. The smallest group of threads is called a CUDA Warp, or OpenCL Wavefront. zone 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 codeZone]* [https:  if(threadIdx//docs.x == 0){ doA(); } else { doB(); } While there is only one thread in the warp that has threadIdx == 0, all 32 threads of the warp will have their shared instruction pointer execute doA() togethernvidia. To keep the code semantically correct, threads #1 through #31 will have their Nvidia Predicatecom/cuda/parallel-register cleared (or AMD Execution Mask cleared), which means the thread will throw away the work after executing a specific statement. For those familiar with x64 AVX code, a GPU thread is comparable to a SIMD-lane in AVXexecution/index. All lanes of an AVX instruction will execute any particular instruction, but you may throw away the results of some registers using mask or comparison instructionshtml Nvidia PTX ISA]* [https://docs. Once doA() is complete, the machine will continue and doB()nvidia. In this case, thread#0 will have its execution mask-cleared, while threads #1 through #31 will actually complete the results of doB()com/cuda/index.html Nvidia CUDA Toolkit Documentation]

Programmers can group warps or wavefronts together into larger clusters, called CUDA Blocks or OpenCL Workgroups. 1024 threads can work together =SIMT and SIMD on a modern GPU Compute Unit (AMD) or Symmetric Multiprocessor (Nvidia), sharing L1 cache, shared memory and other resources. Because of the tight coupling of L1 cache and Shared Memory, these 1024 threads can communicate extremely efficiently. Case in point: both Nvidia PTX and AMD GCN implement thread barriers as a singular assembly language instruction, as long as those threads are within the same workgroup. Atomic operations, memory fences, and other synchronization primitives are extremely fast and well optimized in these cases.=

== Grids GPUs run multiple threads in [https://en.wikipedia.org/wiki/Single_instruction,_multiple_threads SIMT] fashion and NDRange ==are capable to hide memory latencies by running a multitude of SIMT waves on the same [https://en.wikipedia.org/wiki/SIMD SIMD] unit.

The programmer may choose to cut up OpenCL offers the 1920x1080 screen into blocks of size 32x32 pixels. Or maybe an algorithm is horizontal in nature, and it may be more convenient to work with blocks of 1x1024 pixels instead. Or maybe the block-sizes have been set to some video standards, and maybe 8x8 blocks (64-threads) are the biggest you can practically work with (say MPEG-2 decoder 8x8 macroblocks). Regardless, following memory model for the programmer chooses a block size which is most convenient and optimized for their purposes. To complete this hypothetical example, a 1920x1080 screen could be split up into 60x34 CUDA Blocks (or OpenCL Workgroups), each covering 32x32 pixels with 1024 CUDA Threads (or OpenCL Workitems) each. :

These blocks and workgroups will execute with as much parallel processing as the underlying hardware can support. Roughly 150 CUDA Blocks or OpenCL Workgroups at a time on a typical midrange GPU circa from 2019 (such as a Nvidia 2060 Super or AMD 5700). The most important note is that blocks within a grid (or workgroups within an NDRange) may not execute concurrently with each other. Some degree of sequential processing may happen. If thread #0 creates a Spinlock waiting for thread #1000000 to communicate with it, modern hardware will probably never have the two threads executing concurrently with each other, and the code would likely timeout. In practice, the easiest mechanism for Grid or NDRange sized synchronization is to wait for the kernel to finish executing: to have the CPU wait and process the results in between Grid or NDRanges. For example: LeelaZero will schedule an NDRange for each [https://github.com/leela* __private -zero/leela-zero/blob/next/src/kernels/convolve1.opencl Convolve operation]usually registers, as well as merge and other primitives. The convolve operation is over accessable only by a 3single work-dimensional NDRange for <channel, output, row_batch>item resp. To build up a full CNN operation, the CPU will schedule different operations for the GPU: convolve, merge, transform and morethread. ==Memory Model== OpenCL, CUDA, ROCM, and other GPU-languages all have a similar memory model. * __device__ (CUDA) or __global (OpenCL) memory -- OpenCL __global and CUDA __device__ memory exists on the GPU's VRAM. Any threads can access any part of __device__ or __global memory, although memory__local -ordering and caching details can get quite complicated if multiple threads simultaneously read and write to a particular memory location. Proper memory ordering with __threadfence() (CUDA) or mem_fence() (OpenCL) is essential to preventing memoryscratch-consistency issues. * __constant__ (CUDA) or __constant (OpenCL) pad memory shared across work-- Constants are not allowed to change during the execution items of a particular kernel. Historically, this was used by Pixel Shaders as they read texture data. The texturework-data could be computed and loaded onto the GPU, but the data was not allowed to change during the Pixel Shader's execution. Both Nvidia and AMD GPUs have special caches (and in AMD's case: special registers called sGPRs) which accelerate constant-data. The caches associated with this memory space is sometimes called K$ (Konstant-cache), and has to be independently flushed if its data ever changes. The main benefit in both AMD and Nvidia systems is that K$ values are broadcast extremely efficiently to all threads in a wavefront, but only if all threads in a wavefront are reading from the same memory locationgroup resp. Instead of haing 32-memory reads (Nvidia) or 64-memory reads (AMD GCN), a read from K$ can be optimized into a single-read, broadcast to all 32 or 64-threads of a Warp or Wavefrontblock. * __shared__ (CUDA) or __local (OpenCL) memory __constant -read- This is highly-accelerated memory regions designed for threads to exchange data within a CUDA Block or OpenCL Workgroup. On AMD Systems, there is more Local "LDS" memory than even L1 Cache (GCN) or L0 Cache (RDNA)only variable. * Default (CUDA) or __private (OpenCL) Memory __global -- Private memory typically maps to a GPU-registerusually VRAM, and is inaccessible to other threads. If a kernel requires more memory than what can exist in GPUaccessable by all work-registers, the data will automatically spill over into global VRAM (with an associated performance penalty)items resp. In practice, this spillover is well interleaved, well-optimized, and reduced to as small a subset as possible through compiler optimizationsthreads.

The market is split into two categories, integrated and discrete GPUs. The first being the most important by quantity, the second by performance. Discrete GPUs are divided as consumer brands for playing 3D games, professional brands for 3D CAD/CGI programs and server brands for big-data and number-crunching workloads. Each brand offering different feature sets in driversdriver, VRAM, or computation abilities.

* [https://en.wikipedia.org/wiki/RDNA_(microarchitecture)#RDNA_2 RDNA 2 from Wikipedua.0]cards were unveiled on October 28, 2020. 

[https://wwwen.techpowerupwikipedia.comorg/gpu-specswiki/docs/amd-vega-architecture.pdf Architecture WhitepaperRadeon_RX_Vega_series Vega] Vega cards were first released in unveiled on August 14, 2017. Vega is the last in the line of the GCN Architecture: 64 threads per wavefront. Each compute unit contains 4x SIMD units, supporting a total of 40 wavefronts per compute unit (a queue of 10-wavefronts per SIMD Unit). Each SIMD unit contains 16 vALUs for general compute + 1 sALU for branching and constant logic. Each SIMD unit executes the same instruction over four clock ticks (16 vALUs x 4 clock ticks == 64 threads per Wavefront). Vega specifically added Packed FP16 instructions, such as dot-product and packed add and packed multiply. From a programming level, these packed FP16 instructions are SIMD-within-SIMD, each SIMD thread could operate its own SIMD FP16 instruction akin to AVX or SSE from the x64 architecture.

* Radeon VII* Vega64* Vega56[https://www.techpowerup.com/gpu-specs/docs/amd-vega-architecture.pdf Architecture Whitepaper]

[https://en.wikipedia.org/wiki/Graphics_Core_Next#Graphics_Core_Next_4 Polaris ] cards were first released in 2016 under the AMD Radeon 400 series name.

* [https://www.amd.com/system/files/documents/polaris-whitepaper.pdf Architecture Whitepaper] * RX 580* RX 570* RX 560

Apple released its M1 SoC (system on a chip) with integrated GPU for desktops and notebooks in 2020.

* [https://en.wikipedia.org/wiki/List_of_Intel_graphics_processing_units#Gen12 List of Intel Xe 'Gen12' GPUs on Wikipedia]

GPUs are used in [https://en.wikipedia.org/wiki/High-performance_computing HPC] environments because of their good [https://en.wikipedia.org/wiki/FLOP FLOP]/Watt ratio. The instruction throughput in general depends on the architecture (like Nvidia's [https://en.wikipedia.org/wiki/Tesla_%28microarchitecture%29 Tesla], [https://en.wikipedia.org/wiki/Fermi_%28microarchitecture%29 Fermi], [https://en.wikipedia.org/wiki/Kepler_%28microarchitecture%29 Kepler], [https://en.wikipedia.org/wiki/Maxwell_%28microarchitecture%29 Maxwell] or AMD's [https://en.wikipedia.org/wiki/TeraScale_%28microarchitecture%29 TerascaleTeraScale], [https://en.wikipedia.org/wiki/Graphics_Core_Next GCN], [https://en.wikipedia.org/wiki/AMD_RDNA_Architecture RDNA]), the brand (like Nvidia [https://en.wikipedia.org/wiki/GeForce GeForce], [https://en.wikipedia.org/wiki/Nvidia_Quadro Quadro], [https://en.wikipedia.org/wiki/Nvidia_Tesla Tesla] or AMD [https://en.wikipedia.org/wiki/Radeon Radeon], [https://en.wikipedia.org/wiki/Radeon_Pro Radeon Pro], [https://en.wikipedia.org/wiki/Radeon_Instinct Radeon Instinct]) and the specific model.

GPUs were originally intended to process matrix multiplications for graphical transformations and rendering. [[Neural Networks#Convolutional|Convolutional Neural Networks]] can have their operations interpreted as a series of matrix multiplications. GPUs are therefore a natural fit to parallelize and process CNNs. GPUs traditionally operated on 32-bit floating point numbers. However, CNNs can make due with 16-bit half floats (FP16), or even 8-bit or 4-bit numbers. One thousand single-precision floats will take up 4kB of space, while one-thousand FP16 will take up 2kB of space. A half-float uses half the memory, eats only half the memory bandwidth, and only half the space in caches. As such, GPUs such as AMD Vega or Nvidia Volta added support for FP16 processing. Specialized units, such as Nvidia Volta's "Tensor cores", can perform an entire 4x4 block of FP16 matrix multiplications in just one PTX assembly language statement. It is with these instructions that CNN operations are accelerated.  GPUs are much more suited than CPUs to implement and train [[Neural Networks#Convolutional|Convolutional Neural Networks]] (CNN), and were therefore also responsible for the [[Deep Learning|deep learning]] boom, also affecting game playing programs combining CNN with [[Monte-Carlo Tree Search|MCTS]], as pioneered by [[Google]] [[DeepMind|DeepMind's]] [[AlphaGo]] and [[AlphaZero]] entities in [[Go]], [[Shogi]] and [[Chess]] using [https://en.wikipedia.org/wiki/Tensor_processing_unit TPUs], and the open source projects [[Leela Zero]] headed by [[Gian-Carlo Pascutto]] for [[Go]] and its [[Leela Chess Zero]] adaption. =History=In the 1970s and 1980s RAM was expensive and Home Computers used custom graphics chips to operate directly on registers/memory without a dedicated frame buffer, like [https://en.wikipedia.org/wiki/Television_Interface_Adaptor TIA]in the [[Atari 8-bit|Atari VCS]] gaming system, [https://en.wikipedia.org/wiki/CTIA_and_GTIA GTIA]+[https://en.wikipedia.org/wiki/ANTIC ANTIC] in the [[Atari 8-bit|Atari 400/800]] series, or [https://en.wikipedia.org/wiki/Original_Chip_Set#Denise Denise]+[https://en.wikipedia.org/wiki/Original_Chip_Set#Agnus Agnus] in the [[Amiga|Commodore Amiga]] series. The 1990s would make 3D graphics and 3D modeling more popular, especially for video games. Cards specifically designed to accelerate 3D math, such as the [https://en.wikipedia.org/wiki/Voodoo2 3dfx Voodoo2], were used by the video game community to play 3D graphics. Some game engines, such as [https://en.wikipedia.org/wiki/Quake_(video_game) Quake], could use instead the [[SIMD and SWAR Techniques|SIMD-capabilities]] of CPUs such as the [[Intel]] [[MMX]] instruction set or [[AMD|AMD's]] [[X86#3DNow!|3DNow!]]. Sony's 3D capable chip used in the PlayStation (1994) and Nvidia's 2D/3D combi chips like NV1 (1995) coined the term GPU for 3D graphics hardware acceleration. With the advent of the [https://en.wikipedia.org/wiki/Unified_shader_model unified shader architecture], like in Nvidia [https://en.wikipedia.org/wiki/Tesla_(microarchitecture) Tesla] (2006), ATI/AMD [https://en.wikipedia.org/wiki/TeraScale_(microarchitecture) TeraScale] (2007) or Intel [https://en.wikipedia.org/wiki/Intel_GMA#GMA_X3000 GMA X3000] (2006), GPGPU frameworks like CUDA and OpenCL emerged and gained in popularity. The large number of regular [https://en.wikipedia.org/wiki/Matrix_multiplication matrix multiplications] led to natural SIMD-style algorithms. The 3D graphics community drew upon the rich history of vector-compute and SIMD-compute from 1980s and 1970s supercomputers. As such, many publications relating to [[Cray X-MP|Cray-vector supercomputers]] or the [[Connection Machine]] supercomputer easily apply to modern GPUs. For example, all the algorithms described in the 1986 publication ''Data Parallel Algorithms'' <ref>[[Mathematician#Hillis|W. Daniel Hillis]], [[Mathematician#GSteele|Guy L. Steele, Jr.]] ('''1986'''). ''[https://dl.acm.org/citation.cfm?id=7903 Data parallel algorithms]''. [[ACM#Communications|Communications of the ACM]], Vol. 29, No. 12, Special Issue on Parallelism</ref> can be efficiently executed on a modern GPU workgroup (roughly ~256x GPU threads). The ''Data Parallel Algorithms'' paper is a beginner-level algorithms paper, demonstrating simple and efficient [[Parallel Prefix Algorithms|parallel-prefix sum]], parallel-linked list traversal, parallel RegEx matching on the 4096x parallel Connection Machine-2 supercomputer. Modern papers on GPUs, such as Nvidia's excellent ''Parallel Prefix Sum (Scan) with CUDA (GPU Gems 3)'' <ref>[https://developer.nvidia.com/gpugems/GPUGems3/gpugems3_ch39.html Chapter 39. Parallel Prefix Sum (Scan) with CUDA (GPU Gems 3)]</ref>, are built on top of these papers from the 1980s or 1990s. As such, the beginner will find it far easier to read the papers from the 1980s or 90s before attempting to read a modern piece like GPU Gems 3.