Changes

[[FILE:6600GT GPUNvidiaTesla.jpg|border|right|thumb| [https://en.wikipedia.org/wiki/GeForce_6_series GeForce 6600GT (NV43)Nvidia_Tesla Nvidia Tesla] GPU <ref>[https://commons.wikimedia.org/wiki/Graphics_processing_unit Graphics processing unit - File:NvidiaTesla.jpg Image] by Mahogny, February 09, 2008, [https://en.wikipedia.org/wiki/Wikimedia_Commons Wikimedia Commons]</ref> ]]

a specialized processor primarily initially intended to for 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 resp. texture 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 [https://en.wikipedia.org/wiki/IMPACT_(computer_graphics) SGI Impact] (1995) in 3D graphics-workstations or [https://en.wikipedia.org/wiki/3dfx#Voodoo_Graphics_PCI 3dfx Voodoo] (1996) for playing 3D games on PCs, emerged. Some game engines 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!]] for [https://en.wikipedia.org/wiki/Real-time_computer_graphics real-time rendering]. Sony's 3D capable chip [https://en.wikipedia.org/wiki/PlayStation_technical_specifications#Graphics_processing_unit_(GPU) GTE] used in the PlayStation (1994) and Nvidia's 2D/3D combi chips like [https://en.wikipedia.org/wiki/NV1 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), GPGPUframeworks like [https://en.wikipedia.org/wiki/CUDA CUDA] and [[OpenCL|OpenCL]] emerged and gained in popularity. =GPU in Computer Chess=  There are in main four ways how to use a GPU for chess: * As an accelerator in [[Leela_Chess_Zero|Lc0]]: run a neural network for position evaluation on GPU* Offload the search in [[Zeta|Zeta]]: run a parallel game tree search with move generation and position evaluation on GPU* As a hybrid in [http://www.talkchess.com/forum3/viewtopic.php?t=64983&start=4#p729152 perft_gpu]: expand the game tree to a certain degree on CPU and offload to GPU to compute the sub-tree* Neural network training such as [https://github.com/glinscott/nnue-pytorch Stockfish NNUE trainer in Pytorch]<ref>[http://www.talkchess.com/forum3/viewtopic.php?f=7&t= 75724 Pytorch NNUE training] by [[Gary Linscott]], [[CCC]], November 08, 2020</ref> or [https://github.com/LeelaChessZero/lczero-training Lc0 TensorFlow Training]

The traditional job of a =GPU is to take the Chess Engines=* [https://en.wikipedia.org/wiki/Three-dimensional_space x,y,z coordinates] of [https://en.wikipedia.org/wiki/Triangle_strip triangles], and [httpsCategory://en.wikipedia.org/wiki/3D_projection mapGPU] these triangles to [https://en.wikipedia.org/wiki/Glossary_of_computer_graphics#screen_space screen space] through a [https://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.

Eventually, the fixed-functionality of GPUs disappeared, and GPUs became primarily Early efforts to leverage a massively parallel GPU for general -purpose computerscomputing required reformulating computational problems in terms of graphics primitives via graphics APIs like [https://en.wikipedia.org/wiki/OpenGL OpenGL] or [https://en. Instead of using vertex shaders inside of wikipedia.org/wiki/DirectXDirextX], general compute languages are designed to make sense outside of a graphical settingfollowed by first GPGPU frameworks such as [https://en.wikipedia.org/wiki/Lib_Sh Sh/RapidMind] or [https://en.wikipedia.org/wiki/BrookGPU Brook] and finally [https://en.wikipedia.org/wiki/CUDA CUDA] and [https://www.chessprogramming.org/OpenCL OpenCL].

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|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[AMD]] supports language frontends like OpenCL, HIP, C++ AMP and with OpenMP offload directives.nvidia.com/cuda-zone NVidia CUDA Zone]* It offers with [https://docsrocmdocs.nvidiaamd.com/cudaen/parallel-thread-executionlatest/index.html NVidia PTX ISAROCm]* [https://docs.nvidia.com/cuda/indexits own parallel compute platform.html NVidia CUDA Toolkit Documentation]

== AMD Software Overview == [[AMD|AMD's]] original software stack, called * [https://encommunity.wikipediaamd.orgcom/wikit5/AMDGPU AMDGPUopencl/bd-pro], provides 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:p//en.wikipedia.org/wiki/OpenCL#Implementations ROCm]. ROCm is opencl-discussions 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), and OpenCL 2.0. ROCm only works on Linux machines supporting modern hardware, such as [https://en.wikipedia.org/wiki/PCI_Express#3.0 PCIe 3.0Developer Community] and relatively recent GPUs (such as the * [https://enrocmdocs.wikipediaamd.org/wiki/AMD_Radeon_500_series RX 580], and [https:/com/en.wikipedia.org/wikilatest/AMD_RX_Vega_series Vega] GPUs)index. html 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 ROCm™ 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 well.]* [https://rocm.github.io/ ROCm Homepage]* [http://developer.amdmanualzz.com/wordpressdoc/mediao/2013cggy6/07/AMD_Accelerated_Parallel_Processing_OpenCL_Programming_Guideamd-opencl-programming-user-revguide-2.7.pdf contents AMD OpenCL Programming Guide]

* [https://gpuopen.com/wpamd-content/uploads/2019/08/RDNA_Shader_ISA_5August2019.pdf RDNA Instruction Set]* [https://developer.amd.com/wpisa-contentdocumentation/resources/Vega_Shader_ISA_28July2017.pdf Vega Instruction Set] == Other 3rd party tools ==  * [https://en.wikipedia.org/wiki/DirectCompute DirectCompute] (GPGPU API by Microsoft)* OpenMP 4.5 Device Offload =The SIMT Programming Model= CUDA, OpenCL, ROCm HIP, all have the same model of implicitly parallel programming. All threads are given an identifier: 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. The executed code is always implicitly [[SIMD]]. Instead of thinking of SIMD-lanes, each lane is considered its own thread. The 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(); } 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() together. To keep the code semantically correct, threads #1 through #31 will have their NVidia Predicate-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 AVX. 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 instructions. Once doA() is complete, the machine will continue and doB(). In this case, thread#0 will have its execution mask-cleared, while threads #1 through #31 will actually complete the results of doB(). This highlights the fundamental trade off of the GPU platform. GPUs have many threads of execution, but they are forced to execute with their warps or wavefronts. In complicated loops or trees of if-statements, this thread divergence problem can cause your code to potentially leave many hardware threads idle. == Blocks and Workgroups ==  Programmers can group warps or wavefronts together into larger clusters, called CUDA Blocks or OpenCL Workgroups. 1024 threads can work together 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 and NDRange == CUDA Grids and OpenCL NDRange is the end of scaling for the programming model. Many blocks can be specified in a CUDA Grid, while many workgroups operate over an OpenCL NDRange. The underlying hardware supports running many workgroups in parallel, across different compute units. An AMD Vega64 has 64 compute units for example, while an NVidia RTX 2070 has 40 symmetric multiprocessors. The hardware scheduler can fit many blocks and workgroups per compute unit. The exact number is dependent on the amount of registers, memory, and wavefronts a particular workgroup uses. CUDA Grids and OpenCL NDRanges may operate in parallel, or may be traversed sequentially if the GPU doesn't have enough parallel resources. = Architectures and Physical Hardware = Each generation, the manufacturers create a series of cards, with set vRAM and SIMD Cores. The market is split into three categories: server, professional, and consumer. Consumer cards are cheapest and are primarily targeted for the video game market. Professional cards have better driver support for 3d programs like Autocad. Finally, server cards provide virtualization services, allowing cloud companies to virtually split their cards between customers. While server and professional cards have more vRAM, consumer cards are more than adequate starting points for GPU Programmers.ISA documentation]

GPUs use high-bandwidth RAM, such as GDDR6 or HBM2== Apple ==Since macOS 10. These specialized RAM are designed for the extremely parallel nature of GPUs, and can provide 200GBps 14 Mojave a transition from OpenCL to 1000GBps throughput[https://en. In comparison: a typical DDR4 channel can provide 20GBpswikipedia. A dual channel desktop will typically have under 50GBps bandwidth to DDR4 main memoryorg/wiki/Metal_(API) Metal] is recommended by [[Apple]].

NVidia's consumer line of cards is Geforce, branded == Intel ==Intel supports OpenCL with RTX or GTX labelsimplementations like BEIGNET and NEO for different GPU architectures and the [https://en. Nvidia's professional line of cards is Quadrowikipedia. Finally, Tesla cards constitute NVidia's server lineorg/wiki/OneAPI_(compute_acceleration) oneAPI] platform with [https://en.wikipedia.org/wiki/DPC++ DPC++] as frontend language.

NVidia's "Titan" line of Geforce cards use consumer drivers, but use professional or server class chips* [https://www. As such, the Titan line can cost anywhere from $1000 to $3000 per cardintel.com/content/www/us/en/developer/overview.html#gs.pu62bi Intel Developer Zone]* [https://www.intel.com/content/www/us/en/develop/documentation/oneapi-programming-guide/top.html Intel oneAPI Programming Guide]

=== Turing Architecture =Nvidia ==

[https://wwwen.nvidiawikipedia.comorg/wiki/contentCUDA CUDA] is the parallel computing platform by [[Nvidia]]. It supports language frontends like C, C++, Fortran, OpenCL and offload directives via [https:/dam/en-zz.wikipedia.org/Solutionswiki/design-visualizationOpenACC OpenACC] and [https:/technologies/turing-architectureen.wikipedia.org/NVIDIA-Turing-Architecture-Whitepaperwiki/OpenMP OpenMP].pdf Architectural Whitepaper]

Turing cards were first released in 2018* [https://developer. They are the first consumer cores to launch with RTX, or raytracing, featuresnvidia. RTX instructions will more quickly traverse an aabb tree to discover raycom/cuda-intersections with lists of objectszone Nvidia CUDA Zone]* [https://docs. These are also the first consumer cards to launch with Tensor cores, 4x4 matrix multiplication FP16 instructions to accelerate convolutional neural networksnvidia.com/cuda/parallel-thread-execution/index.html Nvidia PTX ISA]* [https://docs.nvidia.com/cuda/index.html Nvidia CUDA Toolkit Documentation]* [https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html Nvidia CUDA C++ Programming Guide]* [https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html Nvidia CUDA C++ Best Practices Guide]

* RTX 2080 Ti* RTX 2080* RTX 2070 Ti== Further == * RTX 2070 Super[https://en.wikipedia.org/wiki/Vulkan#Planned_features Vulkan] (OpenGL sucessor of Khronos Group)* RTX 2070 [https://en.wikipedia.org/wiki/DirectCompute DirectCompute] (Microsoft)* RTX 2060 Super[https://en.wikipedia.org/wiki/C%2B%2B_AMP C++ AMP] (Microsoft)* RTX 2060[https://en.wikipedia.org/wiki/OpenACC OpenACC] (offload directives)* GTX 1660[https://en.wikipedia.org/wiki/OpenMP OpenMP] (offload directives)

=== Volta Architecture ==Hardware Model=

A common scheme on GPUs with unified shader architecture is to run multiple threads in [https://en.wikipedia.org/wiki/Single_instruction,_multiple_threads SIMT] fashion and a multitude of SIMT waves on the same [https://imagesen.nvidiawikipedia.comorg/wiki/SIMD SIMD] unit to hide memory latencies. Multiple processing elements (GPU cores) are members of a SIMD unit, multiple SIMD units are coupled to a compute unit, with up to hundreds of compute units present on a discrete GPU. The actual SIMD units may have architecture dependent different numbers of cores (SIMD8, SIMD16, SIMD32), and different computation abilities - floating-point and/contentor integer with specific bit-width of the FPU/voltaALU and registers. There is a difference between a vector-processor with variable bit-width and SIMD units with fix bit-width cores. Different architecturewhite papers from different vendors leave room for speculation about the concrete underlying hardware implementation and the concrete classification as [https:/pdf/voltaen.wikipedia.org/wiki/Flynn%27s_taxonomy hardware architecture]. Scalar units present in the compute unit perform special functions the SIMD units are not capable of and MMAC units (matrix-architecturemultiply-whitepaperaccumulate units) are used to speed up neural networks further.pdf Architecture Whitepaper]

[https://images* 512 CUDA cores @1.nvidia.com/content/pdf/tesla/whitepaper/pascal544GHz* 16 SMs -architecture-whitepaper.pdf Architecture Whitepaper]Streaming Multiprocessors* organized in 2x16 CUDA cores per SM* Warp size of 32 threads

Pascal cards were first released in 2016AMD Radeon HD 7970 ([https://en.wikipedia.org/wiki/Graphics_Core_Next GCN)]<ref>[https://en.wikipedia.org/wiki/Graphics_Core_Next Graphics Core Next on Wikipedia]</ref><ref>[https://en.wikipedia.org/wiki/List_of_AMD_graphics_processing_units#Radeon_HD_7000_series Radeon HD 7000 series on Wikipedia]</ref>

* Tesla P1002048 Stream cores @0.925GHz* Titan Xp32 Compute Units* GTX 1080 Tiorganized in 4xSIMD16, each SIMT4, per Compute Unit* GTX 1080* GTX 1070 Ti* GTX 1060* GTX 1050* GTX 1030Wavefront size of 64 work-items

== AMD =Wavefront and Warp===Generalized the definition of the Wavefront and Warp size is the amount of threads executed in SIMT fashion on a GPU with unified shader architecture.

=== Navi RDNA 1.0 ==Programming Model=

A [https://gpuopenen.comwikipedia.org/wiki/Parallel_programming_model parallel programming model] for GPGPU can be [https://en.wikipedia.org/wpwiki/Data_parallelism data-contentparallel], [https://en.wikipedia.org/uploadswiki/2019Task_parallelism task-parallel], a mixture of both, or with libraries and offload-directives also [https:/08/RDNA_Architecture_publicen.pdf Architecture Slide Deckwikipedia.org/wiki/Implicit_parallelism implicitly-parallel]. Single GPU threads (work-items in OpenCL) contain the kernel to be computed and are coupled to a work-group, one or multiple work-groups form the NDRange to be executed on the GPU device. The members of a work-group execute the same kernel, can be usually synchronized and have access to the same scratch-pad memory, with an architecture limit of how many work-items a work-group can hold and how many threads can run in total concurrently on the device.

RDNA cards were first released in 2019. RDNA is a major change for AMD cards{| class="wikitable" style="margin: the underlying hardware supports both Wave32 and Wave64 gangs of threads. auto"|+ Terminology|-! OpenCL Terminology !! CUDA Terminology|-| Kernel || Kernel|-| Compute Unit || Streaming Multiprocessor|-| Processing Element || CUDA Core|-| Work-Item || Thread|-| Work-Group || Block|-| NDRange || Grid|-|}

[https* Warp size://www.techpowerup.com/gpu-specs/docs/amd-vega-architecture.pdf Architecture Whitepaper]32* Maximum number of threads per block: 1024* Maximum number of resident blocks per multiprocessor: 32* Maximum number of resident warps per multiprocessor: 64* Maximum number of resident threads per multiprocessor: 2048

* AMD Radeon VII* Vega64* Vega56HD 7970 (GCN) <ref>[https://www.olcf.ornl.gov/wp-content/uploads/2019/10/ORNL_Application_Readiness_Workshop-AMD_GPU_Basics.pdf AMD GPU Hardware Basics]</ref>

=Inside= Modern GPUs consist of up to hundreds of [[SIMD and SWAR Techniques|SIMD]] or [https://en.wikipedia* __private - usually registers, accessable only by a single work-item resp.org/wiki/Vector_processor Vector] units, coupled to compute unitsthread. Each compute unit processes multiple [https://en.wikipedia.org/wiki/Thread_block#Warps Warps] (Nvidia term) * __local - scratch-pad memory shared across work-items of a work-group resp. Wavefronts ([[AMD]] term) in [https://enthreads of block.wikipedia* __constant - read-only memory.org/wiki/Single_instruction* __global - usually VRAM,_multiple_threads SIMT] fashion. Each Warp accessable by all work-items resp. Wavefront runs n (32 or 64) [[Thread|threads]] simultaneously.

The Nvidia [https{| class="wikitable" style="margin://en.wikipedia.org/wiki/List_of_Nvidia_graphics_processing_units#GeForce_500_Series GeForce GTX 580], for example, is able to run 32 threads in one Warp, in total of 24576 threads, spread on 16 compute units with a total of 512 cores. <ref>auto"|+ Terminology|-! OpenCL Terminology !! CUDA C Programming Guide v7.0, Appendix G. COMPUTE CAPABILITIES, Table 12 Technical Specifications per Compute Capability</ref>Terminology|-| Private Memory || Registers|-| Local Memory || Shared Memory|-| Constant Memory || Constant MemoryThe AMD [https://en.wikipedia.org/wiki/Radeon_HD_7000_Series#Radeon_HD_7900 Radeon HD 7970] is able to run 64 threads in one Wavefront, in total of 81920 threads, spread on 32 compute units with a total of 2048 cores. <ref>AMD Accelerated Parallel Processing OpenCL Programming Guide rev2.7, Appendix D Device Parameters, Table D.1 Parameters for 7xxx Devices</ref>. In real life the register and shared memory size limits the amount of total threads.|-| Global Memory || Global Memory|}

===MemoryExamples=== The memory hierarchy of an GPU consists in main of private memory (registers accessed by an single thread resp. work-item), local memory (shared by threads of an block resp. work-items of an work-group ), constant memory, different types of cache and global memory. Size, latency and bandwidth vary between vendors and architectures.

Here the data for the Nvidia GeForce GTX 580 ([https://en.wikipedia.org/wiki/Fermi_%28microarchitecture%29 Fermi)] as an example: <ref>CUDA C Programming Guide v7.0, Appendix G.COMPUTE CAPABILITIES</ref>

* 768 KiB L2 cachein total

Here the data for the AMD Radeon HD 7970 ([https://en.wikipedia.org/wiki/Graphics_Core_Next GCN]) as an example: <ref>AMD Accelerated Parallel Processing OpenCL Programming Guide rev2.7, Appendix D Device Parameters, Table D.1 Parameters for 7xxx Devices</ref>

* 768 KiB L2 cachein total

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.

==Integer Instruction Throughput==* 32 bit Integer Performance INT32: The 32 -bit integer performance can be architecture and operation depended less than 32 -bit FLOP or 24 -bit integer performance.

* 64 bit Integer PerformanceINT64: Current GPU In general [https://en.wikipedia.org/wiki/Processor_register registers] and Vector-[https://en.wikipedia.org/wiki/Arithmetic_logic_unit ALUs] of consumer brand GPUs are 32 -bit wide and have to emulate 64 -bit integer operations.<ref>[https://en.wikichip.org/w/images/a/a1/vega-whitepaper.pdf AMD Vega White Paper]</ref> <ref>[https* INT8://wwwSome architectures offer higher throughput with lower precision.nvidiaThey quadruple the INT8 or octuple the INT4 throughput.com/content/dam/en-zz/Solutions/design-visualization/technologies/turing-architecture/NVIDIA-Turing-Architecture-Whitepaper.pdf Nvidia Turing White Paper]</ref>

* Mixed Precision Support: Newer architectures like Nvidia [https://en.wikipedia.org/wiki/Turing_(microarchitecture) Turing] and AMD [https://en.wikipedia.org/wiki/AMD_RX_Vega_series Vega] have mixed precision support. Vega doubles the [https://en.wikipedia.org/wiki/FP16 FP16] and quadruples the [https://en.wikipedia.org/wiki/Integer_(computer_science)#Common_integral_data_types INT8] throughput.<ref>[https://en.wikipedia.org/wiki/Graphics_Core_Next#fifth Vega (GCN 5th generation) from Wikipedia]</ref>Turing doubles the FP16 throughput of its [https://en.wikipedia.org/wiki/==Floating-point_unit FPUs].<ref>[https://www.anandtech.com/show/13282/nvidia-turing-architecture-deep-dive/4 AnandTech - Nvidia Turing Deep Dive page 4]</ref>Point Instruction Throughput==

* TensorCoresFP32: With Nvidia [https://en.wikipedia.org/wiki/Volta_Consumer GPU performance is measured usually in single-precision (microarchitecture32-bit) Volta] series TensorCores were introduced. They offer fp16*fp16+fp32, matrixfloating-multiplicationpoint FMA (fused-accumulatemultiply-units, used to accelerate neural networksadd) throughput.<ref>[https * FP64: Consumer GPUs have in general a lower ratio (FP32://onFP64) for double-demand.gputechconf.com/gtc/2017/presentation/s7798precision (64-lukebit) floating-durant-inside-voltapoint operations throughput than server brand GPUs.pdf INSIDE VOLTA]</ref> Turings 2nd gen TensorCores add  * FP16, INT8, INT4 optimized computation.<ref>[https://www.anandtech.com/show/13282/nvidiaSome GPGPU architectures offer half-turingprecision (16-architecturebit) floating-deep-dive/6 AnandTech - Nvidia Turing Deep Dive page 6]</ref>point operation throughput with an FP32:FP16 ratio of 1:2.

 Nvidia GeForce GTX 580 (Fermi, CC 2.0) - 32 -bit integer operations/clock cycle per compute unit <ref>CUDA C Programming Guide v7.0, Chapter 5.4.1. Arithmetic Instructions</ref>

Max theoretic ADD operation throughput: 32 Ops * x 16 CUs * x 1544 MHz = 790.528 GigaOps/sec

AMD Radeon HD 7970 (GCN 1.0) - 32 -bit integer operations/clock cycle per processing element <ref>AMD_OpenCL_Programming_Optimization_Guide.pdf 3.0beta, Chapter 2.7.1 Instruction Bandwidths</ref>

Max theoretic ADD operation throughput: 1 Op * x 2048 PEs * x 925 MHz = 1894.4 GigaOps/sec  =Tensors=MMAC (matrix-multiply-accumulate) units are used in consumer brand GPUs for neural network based upsampling of video game resolutions, in professional brands for upsampling of images and videos, and in server brand GPUs for accelerating convolutional neural networks in general. Convolutions can be implemented as a series of matrix-multiplications via Winograd-transformations <ref>[https://talkchess.com/forum3/viewtopic.php?f=7&t=66025&p=743355#p743355 Re: To TPU or not to TPU...] by [[Rémi Coulom]], [[CCC]], December 16, 2017</ref>. Mobile SoCs usually have an dedicated neural network engine as MMAC unit. ==Nvidia TensorCores==: With Nvidia [https://en.wikipedia.org/wiki/Volta_(microarchitecture) Volta] series TensorCores were introduced. They offer FP16xFP16+FP32, matrix-multiplication-accumulate-units, used to accelerate neural networks.<ref>[https://on-demand.gputechconf.com/gtc/2017/presentation/s7798-luke-durant-inside-volta.pdf INSIDE VOLTA]</ref> Turing's 2nd gen TensorCores add FP16, INT8, INT4 optimized computation.<ref>[https://www.anandtech.com/show/13282/nvidia-turing-architecture-deep-dive/6 AnandTech - Nvidia Turing Deep Dive page 6]</ref> Amperes's 3rd gen adds support for BF16, TF32, FP64 and sparsity acceleration.<ref>[https://en.wikipedia.org/wiki/Ampere_(microarchitecture)#Details Wikipedia - Ampere microarchitecture]</ref>Ada Lovelaces's 4th gen adds support for FP8.<ref>[https://en.wikipedia.org/wiki/Ada_Lovelace_(microarchitecture) - Ada Lovelace microarchitecture]</ref> ==AMD Matrix Cores==: AMD released 2020 its server-class [https://www.amd.com/system/files/documents/amd-cdna-whitepaper.pdf CDNA] architecture with Matrix Cores which support MFMA (matrix-fused-multiply-add) operations on various data types like INT8, FP16, BF16, FP32. AMD's CDNA 2 architecture adds FP64 optimized throughput for matrix operations. AMD's RDNA 3 architecture features dedicated AI tensor operation acceleration. AMD's CDNA 3 architecture adds support for FP8 and sparse matrix data (sparsity). ==Intel XMX Cores==: Intel added XMX, Xe Matrix eXtensions, cores to some of the [https://en.wikipedia.org/wiki/Intel_Xe Intel Xe] GPU series, like [https://en.wikipedia.org/wiki/Intel_Arc#Alchemist Arc Alchemist] and [https://www.intel.com/content/www/us/en/products/sku/232876/intel-data-center-gpu-max-1100/specifications.html Intel Data Center GPU Max Series].

One reason GPUs are not used as accelerators for chess engines is the host-device latency, aka. kernel-launch-overhead. Nvidia and AMD have not published official numbers, but in practice there is an a measurable latency for null-kernels of 5 microseconds <ref>[https://devtalk.nvidia.com/default/topic/1047965/cuda-programming-and-performance/host-device-latencies-/post/5318041/#5318041 host-device latencies?] by [[Srdja Matovic]], Nvidia CUDA ZONE, Feb 28, 2019</ref> up to 100s of microseconds <ref>[https://community.amd.com/thread/237337#comment-2902071 host-device latencies?] by [[Srdja Matovic]] AMD Developer Community, Feb 28, 2019</ref>. One solution to overcome this limitation is to couple tasks to batches to be executed in one run <ref>[http://www.talkchess.com/forum3/viewtopic.php?f=7&t=67347#p761239 Re: GPU ANN, how to deal with host-device latencies?] by [[Milos Stanisavljevic]], [[CCC]], May 06, 2018</ref>.

= Architectures =The market is split into two categories, integrated and discrete 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 multiplicationsThe first being the most important by quantity, the second by performance. Discrete GPUs are therefore a natural fit to parallelize divided as consumer brands for playing 3D games, professional brands for CAD/CGI programs and process CNNsserver brands for big-data and number-crunching workloads. Each brand offering different feature sets in driver, VRAM, or computation abilities.

== AMD ==AMD line of discrete GPUs traditionally operated is branded as Radeon for consumer, Radeon Pro for professional and Radeon Instinct for server. * [https://en.wikipedia.org/wiki/List_of_AMD_graphics_processing_units List of AMD graphics processing units on 32-bit floating point numbersWikipedia]  === CDNA3 === CDNA3 HPC architecture was unveiled in December, 2023. HoweverWith MI300A APU model (CPU+GPU+HBM) and MI300X GPU model, CNNs can make due both with 16multi-bit half floats chip modules design. Featuring Matrix Cores with support for a broad type of precision, as INT8, FP8, BF16, FP16, TF32, FP32, FP64, as well as sparse matrix data (FP16sparsity), or even 8. Supported by AMD's ROCm open software stack for AMD Instinct accelerators.  * [https://www.amd.com/content/dam/amd/en/documents/instinct-tech-docs/white-papers/amd-cdna-3-white-paper.pdf AMD CDNA3 Whitepaper]* [https://www.amd.com/content/dam/amd/en/documents/instinct-tech-docs/instruction-set-architectures/amd-instinct-mi300-cdna3-instruction-bit or 4set-bit numbersarchitecture.pdf AMD Instinct MI300/CDNA3 Instruction Set Architecture]* [https://www.amd. One thousand singlecom/en/developer/resources/rocm-precision floats will take up 4kB of spacehub.html AMD ROCm Developer Hub] === Navi 3x RDNA3 === RDNA3 architecture in Radeon RX 7000 series was announced on November 3, 2022, while onefeaturing dedicated AI tensor operation acceleration. * [https://en.wikipedia.org/wiki/Radeon_RX_7000_series AMD Radeon RX 7000 on Wikipedia]* [https://developer.amd.com/wp-thousand FP16 will take up 2kB of spacecontent/resources/RDNA3_Shader_ISA_December2022. A halfpdf RDNA3 Instruction Set Architecture] === CDNA2 === CDNA2 architecture in MI200 HPC-float uses half the memory, eats only half the memory bandwidthGPU with optimized FP64 throughput (matrix and vector), multi-chip-module design and only half the space Infinity Fabric was unveiled in cachesNovember, 2021. * [https://www.amd.com/system/files/documents/amd-cdna2-white-paper.pdf AMD CDNA2 Whitepaper]* [https://developer.amd.com/wp-content/resources/CDNA2_Shader_ISA_4February2022. As suchpdf CDNA2 Instruction Set Architecture] === CDNA === CDNA architecture in MI100 HPC-GPU with Matrix Cores was unveiled in November, GPUs such as 2020. * [https://www.amd.com/system/files/documents/amd-cdna-whitepaper.pdf AMD Vega or NVidia Volta added support for FP16 processingCDNA Whitepaper]* [https://developer.amd.com/wp-content/resources/CDNA1_Shader_ISA_14December2020.pdf CDNA Instruction Set Architecture]

Specialized units, such as NVidia Volta's "Tensor cores"=== Navi 2x RDNA2 === [https://en.wikipedia.org/wiki/RDNA_(microarchitecture)#RDNA_2 RDNA2] cards were unveiled on October 28, 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 accelerated2020.

GPUs are much more suited than CPUs to implement * [https://en.wikipedia.org/wiki/Radeon_RX_6000_series AMD Radeon RX 6000 on Wikipedia]* [https://developer.amd.com/wp-content/resources/RDNA2_Shader_ISA_November2020.pdf RDNA 2 Instruction Set Architecture] === Navi RDNA === [https://en.wikipedia.org/wiki/RDNA_(microarchitecture) RDNA] cards were unveiled on July 7, 2019. * [https://www.amd.com/system/files/documents/rdna-whitepaper.pdf RDNA Whitepaper]* [https://gpuopen.com/wp-content/uploads/2019/08/RDNA_Architecture_public.pdf Architecture Slide Deck]* [https://gpuopen.com/wp-content/uploads/2019/08/RDNA_Shader_ISA_5August2019.pdf RDNA Instruction Set Architecture] === Vega GCN 5th gen === [https://en.wikipedia.org/wiki/Radeon_RX_Vega_series Vega] cards were unveiled on August 14, 2017. * [https://www.techpowerup.com/gpu-specs/docs/amd-vega-architecture.pdf Architecture Whitepaper]* [https://developer.amd.com/wp-content/resources/Vega_Shader_ISA_28July2017.pdf Vega Instruction Set Architecture] === Polaris GCN 4th gen ===  [https://en.wikipedia.org/wiki/Graphics_Core_Next#Graphics_Core_Next_4 Polaris] cards were first released in 2016. * [https://www.amd.com/system/files/documents/polaris-whitepaper.pdf Architecture Whitepaper]* [https://developer.amd.com/wordpress/media/2013/12/AMD_GCN3_Instruction_Set_Architecture_rev1.1.pdf GCN3/4 Instruction Set Architecture] === Southern Islands GCN 1st gen === Southern Island cards introduced the [https://en.wikipedia.org/wiki/Graphics_Core_Next GCN] architecture in 2012. * [https://en.wikipedia.org/wiki/Radeon_HD_7000_series AMD Radeon HD 7000 on Wikipedia]* [https://www.amd.com/content/dam/amd/en/documents/radeon-tech-docs/programmer-references/si_programming_guide_v2.pdf Southern Islands Programming Guide]* [https://developer.amd.com/wordpress/media/2012/12/AMD_Southern_Islands_Instruction_Set_Architecture.pdf Southern Islands Instruction Set Architecture] == Apple == === M series === Apple released its M series SoC (system on a chip) with integrated GPU for desktops and train notebooks in 2020. * [https://en.wikipedia.org/wiki/Apple_silicon#M_series Apple M series on Wikipedia] == ARM ==The ARM Mali GPU variants can be found on various systems on chips (SoCs) from different vendors. Since Midgard (2012) with unified-shader-model OpenCL support is offered. * [Neural Networkshttps://en.wikipedia.org/wiki/Mali_(GPU)#Convolutional|Convolutional Neural NetworksVariants Mali variants on Wikipedia] === Valhall (2019) === * [https://developer.arm.com/documentation/101574/latest Bifrost and Valhall OpenCL Developer Guide] === Bifrost (2016) === * [https://developer.arm.com/documentation/101574/latest Bifrost and Valhall OpenCL Developer Guide] === Midgard (2012) ===* [https://developer.arm.com/documentation/100614/latest Midgard OpenCL Developer Guide] == Intel == === Xe === [https://en.wikipedia.org/wiki/Intel_Xe Intel Xe] line of GPUs (released since 2020) is divided as Xe-LP (low-power), Xe-HPG (high-performance-gaming), Xe-HP (CNNhigh-performace)and Xe-HPC (high-performance-computing). * [https://en.wikipedia.org/wiki/List_of_Intel_graphics_processing_units#Gen12 List of Intel Gen12 GPUs on Wikipedia]* [https://en.wikipedia.org/wiki/Intel_Arc#Alchemist Arc Alchemist series on Wikipedia] ==Nvidia==Nvidia line of discrete GPUs is branded as GeForce for consumer, Quadro for professional and were therefore also responsible Tesla for server. * [https://en.wikipedia.org/wiki/List_of_Nvidia_graphics_processing_units List of Nvidia graphics processing units on Wikipedia] === Grace Hopper Superchip ===The Nvidia GH200 Grace Hopper Superchip was unveiled August, 2023 and combines the Nvidia Grace CPU ([[Deep LearningARM|deep learningARM v9]] boom) and Nvidia Hopper GPU architectures via NVLink to deliver a CPU+GPU coherent memory model for accelerated AI and HPC applications. * [https://resources.nvidia.com/en-us-grace-cpu/grace-hopper-superchip NVIDIA Grace Hopper Superchip Data Sheet]* [https://resources.nvidia.com/en-us-grace-cpu/nvidia-grace-hopper NVIDIA Grace Hopper Superchip Architecture Whitepaper] === Ada Lovelace Architecture ===The [https://en.wikipedia.org/wiki/Ada_Lovelace_(microarchitecture) Ada Lovelace microarchitecture] was announced on September 20, 2022, featuring 4th-generation Tensor Cores with FP8, FP16, BF16, TF32 and sparsity acceleration. also affecting game playing programs combining CNN with * [https://images.nvidia.com/aem-dam/Solutions/geforce/ada/nvidia-ada-gpu-architecture.pdf Ada GPU Whitepaper]* [Montehttps://docs.nvidia.com/cuda/ada-Carlo Tree Search|MCTStuning-guide/index.html Ada Tuning Guide] === Hopper Architecture ===The [https://en.wikipedia.org/wiki/Hopper_(microarchitecture) Hopper GPU Datacenter microarchitecture]was announced on March 22, as pioneered by 2022, featuring Transfomer Engines for large language models. * [https://resources.nvidia.com/en-us-tensor-core Hopper H100 Whitepaper]* [Googlehttps://docs.nvidia.com/cuda/hopper-tuning-guide/index.html Hopper Tuning Guide] === Ampere Architecture ===The [https://en.wikipedia.org/wiki/Ampere_(microarchitecture) Ampere microarchitecture] was announced on May 14, 2020 <ref>[[DeepMindhttps://devblogs.nvidia.com/nvidia-ampere-architecture-in-depth/ NVIDIA Ampere Architecture In-Depth |DeepMind'sNVIDIA Developer Blog]by [https://people.csail.mit.edu/ronny/ Ronny Krashinsky] , [https://cppcast.com/guest/ogiroux/ Olivier Giroux], [AlphaGohttps://blogs.nvidia.com/blog/author/stephenjones/ Stephen Jones], [https://blogs.nvidia.com/blog/author/nick-stam/ Nick Stam] and [https://en.wikipedia.org/wiki/Sridhar_Ramaswamy Sridhar Ramaswamy], May 14, 2020</ref>. The Nvidia A100 GPU based on the Ampere architecture delivers a generational leap in accelerated computing in conjunction with CUDA 11 <ref>[AlphaZerohttps://devblogs.nvidia.com/cuda-11-features-revealed/ CUDA 11 Features Revealed | NVIDIA Developer Blog]by [https://devblogs.nvidia.com/author/pramarao/ Pramod Ramarao] entities in , May 14, 2020</ref>. * [https://www.nvidia.com/content/dam/en-zz/Solutions/Data-Center/nvidia-ampere-architecture-whitepaper.pdf Ampere GA100 Whitepaper]* [Gohttps://www.nvidia.com/content/PDF/nvidia-ampere-ga-102-gpu-architecture-whitepaper-v2.pdf Ampere GA102 Whitepaper]* [https://docs.nvidia.com/cuda/ampere-tuning-guide/index.html Ampere GPU Architecture Tuning Guide] === Turing Architecture ===[https://en.wikipedia.org/wiki/Turing_(microarchitecture) Turing] cards were first released in 2018. They are the first consumer cores to launch with RTX, for [https://en.wikipedia.org/wiki/Ray_tracing_(graphics) raytracing], features. These are also the first consumer cards to launch with TensorCores used for matrix multiplications to accelerate [[ShogiNeural Networks#Convolutional|convolutional neural networks]] and . The Turing GTX line of chips do not offer RTX or TensorCores. * [https://www.nvidia.com/content/dam/en-zz/Solutions/design-visualization/technologies/turing-architecture/NVIDIA-Turing-Architecture-Whitepaper.pdf Turing Architecture Whitepaper]* [Chesshttps://docs.nvidia.com/cuda/turing-tuning-guide/index.html Turing Tuning Guide]] using  === Volta Architecture === [https://en.wikipedia.org/wiki/Tensor_processing_unit TPUsVolta_(microarchitecture) Volta]cards were released in 2017. They were the first cards to launch with TensorCores, and the open source projects supporting matrix multiplications to accelerate [[Leela ZeroNeural Networks#Convolutional|convolutional neural networks]] headed by . * [https://images.nvidia.com/content/volta-architecture/pdf/volta-architecture-whitepaper.pdf Volta Architecture Whitepaper]* [https://docs.nvidia.com/cuda/volta-tuning-guide/index.html Volta Tuning Guide] === Pascal Architecture ===[https://en.wikipedia.org/wiki/Pascal_(microarchitecture) Pascal] cards were first released in 2016. * [https://images.nvidia.com/content/pdf/tesla/whitepaper/pascal-architecture-whitepaper.pdf Pascal Architecture Whitepaper]* [Gianhttps://docs.nvidia.com/cuda/pascal-tuning-Carlo Pascuttoguide/index.html Pascal Tuning Guide] === Maxwell Architecture ===[https://en.wikipedia.org/wiki/Maxwell(microarchitecture) Maxwell] for cards were first released in 2014. * [https://web.archive.org/web/20170721113746/http://international.download.nvidia.com/geforce-com/international/pdfs/GeForce_GTX_980_Whitepaper_FINAL.PDF Maxwell Architecture Whitepaper on archiv.org]* [Gohttps://docs.nvidia.com/cuda/maxwell-tuning-guide/index.html Maxwell Tuning Guide] == PowerVR ==PowerVR (Imagination Technologies) licenses IP to third parties (most notable Apple) used for system on a chip (SoC) designs. Since Series5 SGX OpenCL support via licensees is available. === PowerVR === * [https://en.wikipedia.org/wiki/PowerVR#PowerVR_Graphics PowerVR series on Wikipedia] and its  === IMG === * [https://en.wikipedia.org/wiki/PowerVR#IMG_A-Series_(Albiorix) IMG A series on Wikipedia]* [Leela Chess Zerohttps://en.wikipedia.org/wiki/PowerVR#IMG_B-Series IMG B series on Wikipedia] == Qualcomm ==Qualcomm offers Adreno GPUs in various types as a component of their Snapdragon SoCs. Since Adreno 300 series OpenCL support is offered. === Adreno ===* [https://en.wikipedia.org/wiki/Adreno#Variants Adreno variants on Wikipedia] adaption == Vivante Corporation ==Vivante licenses IP to third parties for embedded systems, the GC series offers optional OpenCL support. === GC-Series === * [https://en.wikipedia.org/wiki/Vivante_Corporation#Products GC series on Wikipedia]

 ==20091986==* [[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==1990==* [[Mathematician#GEBlelloch|Guy E. Blelloch]] ('''1990'''). ''[https://dl.acm.org/citation.cfm?id=91254 Vector Models for Data-Parallel Computing]''. [https://en.wikipedia.org/wiki/MIT_Press MIT Press], [https://www.cs.cmu.edu/~guyb/papers/Ble90.pdf pdf]==2008 ...==* [[Vlad Stamate]] ('''2008'''). ''Real Time Photon Mapping Approximation on the GPU''. in [http://shaderx6.com/TOC.html ShaderX6 - Advanced Rendering Techniques] <ref>[https://en.wikipedia.org/wiki/Photon_mapping Photon mapping from Wikipedia]</ref>

* [https://groups.google.com/d/msg/lczero/I0lTgR-fFFU/NGC3kJDzAwAJ Michael Larabel benches lc0 on various GPUs] by [[Warren D. Smith]], [[Computer Chess Forums|LCZero Forum]], January 14, 2019 » [[Leela Chess Zero#Lc0|Lc0]] <ref>[https://en.wikipedia.org/wiki/Phoronix_Test_Suite Phoronix Test Suite from Wikipedia]</ref>* [http://www.talkchess.com/forum3/viewtopic.php?f=2&t=70362 Using LC0 with one or two GPUs - a guide] by [[Srdja Matovic]], [[CCC]], March 30, 2019 » [[Leela Chess Zero#Lc0|Lc0]]* [http://www.talkchess.com/forum3/viewtopic.php?f=7&t=70584 Wouldn't it be nice if C++ GPU] by [[Chris Whittington]], [[CCC]], April 25, 2019» [[Cpp|C++]]