Changes

Jump to: navigation, search

GPU

17,166 bytes added, 24 January
m
2020 ...: link to Chinese gpus
'''[[Main Page|Home]] * [[Hardware]] * GPU'''
[[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> ]]
'''GPU''' (Graphics Processing Unit),<br/>
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.
These lists of triangles were specified in Graphics APIs like [https://en.wikipedia.org/wiki/DirectX DirectX]. 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 [https://en.wikipedia.org/wiki/Shader#Vertex_shaders vertex shaders] or [https://en.wikipedia.org/wiki/Shader#Pixel_shaders pixel shaders].=GPGPU=
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].
== Khronos OpenCL ==
[[OpenCL|OpenCL]] specified by the [https://en.wikipedia.org/wiki/Khronos_Group Khronos Group] is widely adopted across all kind of hardware accelerators from different vendors.
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://www.khronos.org/registry/OpenCL/specs/opencl-1.2.pdf OpenCL 1.2 Specification]
* [https://www.khronos.org/registry/OpenCL//sdk/2.0/docs/man/xhtml/ OpenCL 2.0 Reference]
* [https://www.khronos.org/registry/OpenCL/specs/3.0-unified/pdf/ OpenCL 3.0 Specifications] == NVidia Software overview AMD == [[AMD]] supports language frontends like OpenCL, HIP, C++ AMP and with OpenMP offload directives. It offers with [https://rocmdocs.amd.com/en/latest/ ROCm] its own parallel compute platform. * [https://community.amd.com/t5/opencl/bd-p/opencl-discussions AMD OpenCL Developer Community]* [https://rocmdocs.amd.com/en/latest/index.html AMD ROCm™ documentation]* [https://manualzz.com/doc/o/cggy6/amd-opencl-programming-user-guide-contents AMD OpenCL Programming Guide]* [http://developer.amd.com/wordpress/media/2013/12/AMD_OpenCL_Programming_Optimization_Guide2.pdf AMD OpenCL Optimization Guide]* [https://gpuopen.com/amd-isa-documentation/ AMD GPU ISA documentation] == Apple ==Since macOS 10.14 Mojave a transition from OpenCL to [https://en.wikipedia.org/wiki/Metal_(API) Metal] is recommended by [[Apple]]. * [https://developer.apple.com/opencl/ Apple OpenCL Developer] * [https://developer.apple.com/metal/ Apple Metal Developer]* [https://developer.apple.com/library/archive/documentation/Miscellaneous/Conceptual/MetalProgrammingGuide/Introduction/Introduction.html Apple Metal Programming Guide]* [https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf Metal Shading Language Specification] == Intel ==Intel supports OpenCL with implementations like BEIGNET and NEO for different GPU architectures and the [https://en.wikipedia.org/wiki/OneAPI_(compute_acceleration) oneAPI] platform with [https://en.wikipedia.org/wiki/DPC++ DPC++] as frontend language. * [https://www.intel.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]
[[== Nvidia|NVidia]] [https://en.wikipedia.org/wiki/CUDA CUDA] is their general purpose compute framework. CUDA has a [[Cpp|C++]] compiler based on [https://en.wikipedia.org/wiki/LLVM LLVM] / [https://en.wikipedia.org/wiki/Clang clang], which compiles into an assembly-like language called [https://en.wikipedia.org/wiki/Parallel_Thread_Execution PTX]. NVidia device drivers take PTX and compile that down to the final machine code (called NVidia SASS). NVidia keeps PTX portable between its GPUs, while its SASS assembly language may change from year-to-year as NVidia releases new GPUs. A defining feature of CUDA was the "single source" C++ compiler, the same compiler would work with both CPU host-code and GPU device-code. This meant that the data-structures and even pointers from the CPU can be shared directly with the GPU code.==
* [https://developeren.nvidiawikipedia.comorg/cuda-zone NVidia wiki/CUDA CUDA Zone]* is the parallel computing platform by [[Nvidia]]. It supports language frontends like C, C++, Fortran, OpenCL and offload directives via [https://docsen.nvidiawikipedia.com/cudaorg/parallel-thread-executionwiki/index.html NVidia PTX ISAOpenACC OpenACC]* and [https://docsen.nvidiawikipedia.comorg/cudawiki/indexOpenMP OpenMP].html NVidia CUDA Toolkit Documentation]
== AMD Software Overview ==* [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]* [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]
[[AMD|AMD's]] original software stack, called == Further == * [https://en.wikipedia.org/wiki/AMDGPU AMDGPU-proVulkan#Planned_features Vulkan], provides OpenCL 1.2 and 2.0 capabilities on [[Linux]] and [[Windows]]. However, most (OpenGL sucessor of AMD's efforts today is on an experimental framework called Khronos Group)* [https://en.wikipedia.org/wiki/OpenCL#Implementations ROCmDirectCompute DirectCompute]. ROCm is AMD's open source compiler and device driver stack intended for general purpose compute. ROCm supports two languages: (Microsoft)* [https://en.wikipedia.org/wiki/GPUOpen#AMD_Boltzmann_Initiative HIP] (a CUDA-like single-source C%2B%2B_AMP C++ compiler also based on LLVM/clangAMP] (Microsoft), 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.0OpenACC OpenACC] and relatively recent GPUs (such as the offload directives)* [https://en.wikipedia.org/wiki/AMD_Radeon_500_series RX 580OpenMP OpenMP], and [https://en.wikipedia.org/wiki/AMD_RX_Vega_series Vega] GPUs(offload directives).
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.=Hardware Model=
AMD's OpenCL documentationA common scheme on GPUs with unified shader architecture is to run multiple threads in [https://en.wikipedia.org/wiki/Single_instruction, especially the "OpenCL Programming Guide" _multiple_threads SIMT] fashion and a multitude of SIMT waves on the "Optimization Guide" same [https://en.wikipedia.org/wiki/SIMD SIMD] unit to hide memory latencies. Multiple processing elements (GPU cores) are members of a SIMD unit, multiple SIMD units are good places coupled to start for beginners looking a compute unit, with up to program their GPUshundreds of compute units present on a discrete GPU. For Linux developersThe actual SIMD units may have architecture dependent different numbers of cores (SIMD8, SIMD16, SIMD32), and different computation abilities - floating-point and/or integer with specific bit-width of the ROCm environment FPU/ALU and registers. There is under active development a difference between a vector-processor with variable bit-width and has enough features SIMD units with fix bit-width cores. Different architecture white papers from different vendors leave room for speculation about the concrete underlying hardware implementation and the concrete classification as [https://en.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-multiply-accumulate units) are used to get code working wellspeed up neural networks further.
* [https{| class="wikitable" style="margin://rocm.github.io/ ROCm Homepage]auto"* [http://developer.amd.com/wordpress/media/2013/07/AMD_Accelerated_Parallel_Processing_OpenCL_Programming_Guide|+ Vendor Terminology|-rev! AMD Terminology !! Nvidia Terminology|-2.7.pdf AMD OpenCL Programming Guide]* [http://developer.amd.com/wordpress/media/2013/12/AMD_OpenCL_Programming_Optimization_Guide2.pdf AMD OpenCL Optimization Guide]| Compute Unit || Streaming Multiprocessor* [https://gpuopen.com/wp|-content/uploads/2019/08/RDNA_Shader_ISA_5August2019.pdf RDNA Instruction Set]* [https://developer.amd.com/wp| Stream Core || CUDA Core|-content/resources/Vega_Shader_ISA_28July2017.pdf Vega Instruction Set]| Wavefront || Warp|}
== Other 3rd party tools =Hardware Examples===
* Nvidia GeForce GTX 580 ([https://en.wikipedia.org/wiki/DirectCompute DirectComputeFermi_%28microarchitecture%29 Fermi] (GPGPU API by Microsoft)* OpenMP 4<ref>[https://www.nvidia.com/content/PDF/fermi_white_papers/NVIDIA_Fermi_Compute_Architecture_Whitepaper.pdf Fermi white paper from Nvidia]</ref><ref>[https://en.5 Device Offloadwikipedia.org/wiki/List_of_Nvidia_graphics_processing_units#GeForce_500_series GeForce 500 series on Wikipedia]</ref>
=The SIMT Programming Model=* 512 CUDA cores @1.544GHz* 16 SMs - Streaming Multiprocessors* organized in 2x16 CUDA cores per SM* Warp size of 32 threads
CUDA, OpenCL, ROCm HIP, all have the same model of implicitly parallel programmingAMD Radeon HD 7970 ([https://en.wikipedia. All threads are given an identifierorg/wiki/Graphics_Core_Next GCN)]<ref>[https: a threadIdx in CUDA or local_id in OpenCL//en.wikipedia. Aside from this index, all threads of a kernel will execute the same codeorg/wiki/Graphics_Core_Next Graphics Core Next on Wikipedia]</ref><ref>[https://en. The only way to alter the behavior of code is to use this threadIdx to access different datawikipedia.org/wiki/List_of_AMD_graphics_processing_units#Radeon_HD_7000_series Radeon HD 7000 series on Wikipedia]</ref>
The executed code is always implicitly [[SIMD]]* 2048 Stream cores @0. Instead of thinking of SIMD-lanes925GHz* 32 Compute Units* organized in 4xSIMD16, each lane is considered its own thread. The smallest group of threads is called a CUDA WarpSIMT4, or OpenCL per Compute Unit* Wavefront. NVidia GPUs execute 32-threads per warp, while AMD GCN GPUs execute size of 64work-threads per wavefront. All threads within a Warp or Wavefront share an instruction pointer. Consider the following CUDA code:items
if(threadIdx.x == 0){=Wavefront and Warp=== doA(); } else { doB(); }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.
While there is only one thread in the warp that has threadIdx =Programming Model= 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 doAA [https://en.wikipedia.org/wiki/Parallel_programming_model parallel programming model] for GPGPU can be [https://en.wikipedia.org/wiki/Data_parallelism data-parallel], [https://en.wikipedia.org/wiki/Task_parallelism task-parallel], a mixture of both, or with libraries and offload-directives also [https://en.wikipedia.org/wiki/Implicit_parallelism implicitly-parallel]. Single GPU threads (work-items in OpenCL) is completecontain 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 machine will continue and doB()GPU device. In this caseThe members of a work-group execute the same kernel, thread#0 will can be usually synchronized and have its execution maskaccess to the same scratch-clearedpad memory, while with an architecture limit of how many work-items a work-group can hold and how many threads #1 through #31 will actually complete can run in total concurrently on the results of doB()device.
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{| class="wikitable" style="margin:auto"|+ Terminology|-statements, this thread divergence problem can cause your code to potentially leave many hardware threads idle.! OpenCL Terminology !! CUDA Terminology|-| Kernel || Kernel|-| Compute Unit || Streaming Multiprocessor|-| Processing Element || CUDA Core|-| Work-Item || Thread|-| Work-Group || Block|-| NDRange || Grid|-|}
== Blocks and Workgroups Thread Examples==
Programmers can group warps or wavefronts together into larger clustersNvidia GeForce GTX 580 (Fermi, called CUDA Blocks or OpenCL Workgroups. 1024 threads can work together on a modern GPU Compute Unit (AMDCC2) 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<ref>[https: 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//en. Atomic operations, memory fences, and other synchronization primitives are extremely fast and well optimized in these caseswikipedia.org/wiki/CUDA#Technical_Specification CUDA Technical_Specification on Wikipedia]</ref>
== Grids and NDRange ==* Warp size: 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
While warps, blocks, wavefronts and workgroups are concepts that the machine executes... Grids and NDRanges are the scope of the problem specified by a programmer. For example, the 1920x1080 screen could be defined as a Grid with 2073600 threads to execute (likely organized as a 2-dimensional grid for convenience). Specifying these 2,073,600 work items is the purpose of a CUDA Grid or OpenCL NDRange.
A typical midrange GPU will "only" be able to process tens-of-thousands of threads at a timeAMD Radeon HD 7970 (GCN) <ref>[https://www.olcf. In practice, the device driver will cut up a Grid or NDRange into Blocks or Workgroupsornl. These blocks and workgroups will execute with as much parallel processing as the underlying hardware can support (maybe gov/wp-content/uploads/2019/10,000 at a time on a midrange GPU). The device driver will implicitly iterate these blocks over the entire Grid or NDRange to complete the task the programmer has specified, similar to a for/ORNL_Application_Readiness_Workshop-loopAMD_GPU_Basics.pdf AMD GPU Hardware Basics]</ref>
The most important note is that blocks within a Grid (or workgroups within an NDRange) may not execute concurrently with each other. Some degree * Wavefront size: 64* Maximum number of work-items per work-group: 1024* Maximum number of work-groups per compute unit: 40* Maximum number of Wavefronts per compute unit: 40* Maximum number of sequential processing may happen. As such, communication across a Grid or NDRange is difficult to achieve. 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 executingwork-items per compute unit: to have the CPU wait and process the results in between Grid or NDRanges.2560
For example: LeelaZero will schedule an NDRange for each [https://github.com/leela-zero/leela-zero/blob/next/src/kernels/convolve1.opencl Convolve operation], as well as merge and other primitives. The convolve operation is over a 3-dimensional NDRange for <channel, output, row_batch>. To build up a full CNN operation, the CPU will schedule different operations for the GPU: convolve, merge, transform and more.=Memory Model=
= Architectures and Physical Hardware =OpenCL offers the following memory model for the programmer:
The market is split into three categories: server* __private - usually registers, professional, and consumeraccessable only by a single work-item resp. thread.* __local - scratch-pad memory shared across work-items of a work-group resp. Consumer cards are cheapest and are primarily targeted for the video game marketthreads of block. Professional cards have better driver support for 3d programs like Autocad* __constant - read-only memory. Finally* __global - usually VRAM, server cards provide virtualization services, allowing cloud companies to virtually split their cards between customersaccessable by all work-items resp. threads.
Consumer {| class GPUs cost anywhere from $100 to $1000. Professional cards can run to $2000, while server class cards can cost as much as $10,000.="wikitable" style="margin:auto"|+ Terminology|-! OpenCL Terminology !! CUDA Terminology|-| Private Memory || Registers|-| Local Memory || Shared Memory|-| Constant Memory || Constant Memory|-| Global Memory || Global Memory|}
GPUs use high-bandwidth RAM, such as GDDR6 or HBM2. GDDR6 and HBM2 are designed for the extremely parallel nature of GPUs, and can provide 200GBps to 1000GBps throughput. In comparison: a typical DDR4 channel can provide 20GBps. A dual channel desktop will typically have under 50GBps bandwidth to DDR4 main memory.===Memory Examples===
== NVidia ==Nvidia GeForce GTX 580 ([https://en.wikipedia.org/wiki/Fermi_%28microarchitecture%29 Fermi)] <ref>CUDA C Programming Guide v7.0, Appendix G.COMPUTE CAPABILITIES</ref>* 128 KiB private memory per compute unit* 48 KiB (16 KiB) local memory per compute unit (configurable)* 64 KiB constant memory* 8 KiB constant cache per compute unit* 16 KiB (48 KiB) L1 cache per compute unit (configurable)* 768 KiB L2 cache in total* 1.5 GiB to 3 GiB global memoryAMD Radeon HD 7970 ([https://en.wikipedia.org/wiki/Graphics_Core_Next GCN]) <ref>AMD Accelerated Parallel Processing OpenCL Programming Guide rev2.7, Appendix D Device Parameters, Table D.1 Parameters for 7xxx Devices</ref>* 256 KiB private memory per compute unit* 64 KiB local memory per compute unit* 64 KiB constant memory* 16 KiB constant cache per four compute units* 16 KiB L1 cache per compute unit* 768 KiB L2 cache in total* 3 GiB to 6 GiB global memory
NVidia's consumer line of cards is Geforce, branded with RTX or GTX labels. Nvidia's professional line of cards is "Quadro". Finally, Nvidia's server line of cards is "Tesla".===Unified Memory===
NVidia's "Titan" line of Geforce cards use consumer driversUsually data has to be copied between a CPU host and a discrete GPU device, but use professional or server class chips. As such, the Titan line can cost anywhere different architectures from $1000 to $3000 per carddifferent vendors with different frameworks on different operating systems may offer a unified and accessible address space between CPU and GPU.
=Instruction Throughput=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 TeraScale], [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. = Turing Architecture =Integer Instruction Throughput==* INT32: The 32-bit integer performance can be architecture and operation depended less than 32-bit FLOP or 24-bit integer performance. * INT64: 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.* INT8: Some architectures offer higher throughput with lower precision. They quadruple the INT8 or octuple the INT4 throughput. ==Floating-Point Instruction Throughput== * FP32: Consumer GPU performance is measured usually in single-precision (32-bit) floating-point FMA (fused-multiply-add) throughput.
[https* FP64://www.nvidia.com/content/dam/en-zz/Solutions/design-visualization/technologies/turing-architecture/NVIDIAConsumer GPUs have in general a lower ratio (FP32:FP64) for double-Turingprecision (64-Architecturebit) floating-Whitepaperpoint operations throughput than server brand GPUs.pdf Architectural Whitepaper]
Turing cards were first released in 2018. They are the first consumer cores to launch with RTX, or raytracing, features. RTX instructions will more quickly traverse an aabb tree to discover ray* FP16: Some GPGPU architectures offer half-precision (16-intersections with lists of boundingbit) floating-boxes, accelerating raytracing performance. These are also the first consumer cards to launch point operation throughput with Tensor cores, 4x4 matrix multiplication an FP32:FP16 instructions to accelerate convolutional neural networksratio of 1:2.
* RTX 2080 Ti==Throughput Examples==* RTX 2080* RTX 2070 Ti* RTX 2070 Super* RTX 2070 * RTX 2060 Super* RTX 2060* Nvidia GeForce GTX 1660 580 (Fermi, CC 2.0) -32- Low-end GPU without Tensor cores or RTX Coresbit integer operations/clock cycle per compute unit <ref>CUDA C Programming Guide v7.0, Chapter 5.4.1.Arithmetic Instructions</ref>
MAD 16 MUL 16 ADD 32 Bit-shift 16 Bitwise XOR 32 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>  MAD 1/4 MUL 1/4 ADD 1 Bit-shift 1 Bitwise XOR 1 Max theoretic ADD operation throughput: 1 Op x 2048 PEs x 925 MHz == Volta Architecture === 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://images.nvidiatalkchess.com/contentforum3/volta-architectureviewtopic.php?f=7&t=66025&p=743355#p743355 Re: To TPU or not to TPU...] by [[Rémi Coulom]], [[CCC]], December 16, 2017</pdf/volta-architecture-whitepaperref>. Mobile SoCs usually have an dedicated neural network engine as MMAC unit.pdf Architecture Whitepaper]
==Nvidia TensorCores==: With Nvidia [https://en.wikipedia.org/wiki/Volta_(microarchitecture) Volta cards ] series TensorCores were released in 2017. Only Tesla and Titan cards were produced in this generation, aiming only for the most expensive end of the marketintroduced. They were the first cards to launch with Tensor coresoffer FP16xFP16+FP32, supporting 4x4 FP16 matrix multiplications -multiplication-accumulate-units, used to accelerate convolutional 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>
* Tesla V100==AMD Matrix Cores==* Titan V: 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).
=== Pascal Architecture =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].
=Host-Device Latencies= 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 a measurable latency for null-kernels of 5 microseconds <ref>[https://imagesdevtalk.nvidia.com/contentdefault/pdftopic/tesla1047965/whitepapercuda-programming-and-performance/pascalhost-device-latencies-architecture/post/5318041/#5318041 host-whitepaperdevice latencies?] by [[Srdja Matovic]], Nvidia CUDA ZONE, Feb 28, 2019</ref> up to 100s of microseconds <ref>[https://community.amd.pdf Architecture Whitepapercom/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>.
Pascal cards =Deep Learning=GPUs are much more suited than CPUs to implement and train [[Neural Networks#Convolutional|Convolutional Neural Networks]] (CNN), and were first released 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 2016[[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.
* Tesla P100= Architectures =* Titan Xp* GTX 1080 Ti* GTX 1080* GTX 1070 Ti* GTX 1060* GTX 1050* GTX 1030The 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 CAD/CGI programs and server 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 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 Wikipedia]
 
=== CDNA3 ===
CDNA3 HPC architecture was unveiled in December, 2023. With MI300A APU model (CPU+GPU+HBM) and MI300X GPU model, both with multi-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 (sparsity). Supported by AMD's ROCm open software stack for AMD Instinct accelerators.
=== Navi RDNA 1* [https://www.0 === 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-set-architecture.pdf AMD Instinct MI300/CDNA3 Instruction Set Architecture]* [https://www.amd.com/en/developer/resources/rocm-hub.html AMD ROCm Developer Hub]
[https://gpuopen=== Navi 3x RDNA3 === RDNA3 architecture in Radeon RX 7000 series was announced on November 3, 2022, featuring dedicated AI tensor operation acceleration.com/wp-content/uploads/2019/08/RDNA_Architecture_public.pdf Architecture Slide Deck]
RDNA cards were first released in 2019* [https://en.wikipedia. RDNA is a major change for org/wiki/Radeon_RX_7000_series AMD cardsRadeon RX 7000 on Wikipedia]* [https: the underlying hardware supports both Wave32 and Wave64 gangs of threads//developer. Compute Units have 2x32 wide SIMD units, each of which executes 32 threads per clock tickamd. A Wave64 workgroup will execute on a single SIMD unit, but over two clock tickscom/wp-content/resources/RDNA3_Shader_ISA_December2022.pdf RDNA3 Instruction Set Architecture]
=== CDNA2 === CDNA2 architecture in MI200 HPC-GPU with optimized FP64 throughput (matrix and vector), multi-chip-module design and Infinity Fabric was unveiled in November, 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.pdf CDNA2 Instruction Set Architecture] === CDNA === CDNA architecture in MI100 HPC-GPU with Matrix Cores was unveiled in November, 2020. * [https://www.amd.com/system/files/documents/amd-cdna-whitepaper.pdf AMD CDNA Whitepaper]* [https://developer.amd.com/wp-content/resources/CDNA1_Shader_ISA_14December2020.pdf CDNA Instruction Set Architecture] === Navi 2x RDNA2 === [https://en.wikipedia.org/wiki/RDNA_(microarchitecture)#RDNA_2 RDNA2] cards were unveiled on October 28, 2020. * [https://en.wikipedia.org/wiki/Radeon_RX_6000_series AMD Radeon 5700 XTRX 6000 on Wikipedia]* Radeon 5700[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.
Vega cards were first released in 2017* [https://www.amd.com/system/files/documents/polaris-whitepaper. Vega is the last in the line of the GCN pdf ArchitectureWhitepaper]* [https: 64 threads per wavefront//developer.amd. Each compute unit contains 4x SIMD units, supporting a total of 40 wavefronts per compute unit (a queue of 10-wavefronts per SIMD Unit)com/wordpress/media/2013/12/AMD_GCN3_Instruction_Set_Architecture_rev1. 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 pdf GCN3/4 clock ticks == 64 threads per Wavefront).Instruction Set Architecture]
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.=== Southern Islands GCN 1st gen ===
* Radeon VII* Vega64* Vega56Southern 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 = Polaris GCN 4th gen = == = M series === Apple released its M series SoC (system on a chip) with integrated GPU for desktops and 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. * [https://en.wikipedia.org/wiki/Mali_(GPU)#Variants 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 (high-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 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 ([[ARM|ARM v9]]) 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://wwwen.amdwikipedia.comorg/systemwiki/files/documents/polarisAda_Lovelace_(microarchitecture) Ada Lovelace microarchitecture] was announced on September 20, 2022, featuring 4th-whitepapergeneration Tensor Cores with FP8, FP16, BF16, TF32 and sparsity acceleration.pdf Architecture Whitepaper]
* RX 580[https://images.nvidia.com/aem-dam/Solutions/geforce/ada/nvidia-ada-gpu-architecture.pdf Ada GPU Whitepaper]* RX 570* RX 560[https://docs.nvidia.com/cuda/ada-tuning-guide/index.html Ada Tuning Guide]
=Inside== Hopper Architecture === Modern GPUs consist of up to hundreds of [[SIMD and SWAR Techniques|SIMD]] or [https://en.wikipedia.org/wiki/Vector_processor Vector] units, coupled to compute units. Each compute unit processes multiple The [https://en.wikipedia.org/wiki/Thread_block#Warps Warps] Hopper_(Nvidia termmicroarchitecture) resp. Wavefronts ([[AMDHopper GPU Datacenter microarchitecture]] term) in [https://en.wikipedia.org/wiki/Single_instructionwas announced on March 22, 2022,_multiple_threads SIMT] fashion. Each Warp resp. Wavefront runs n (32 or 64) [[Thread|threads]] simultaneouslyfeaturing Transfomer Engines for large language models.
The Nvidia * [https://enresources.wikipedianvidia.orgcom/wiki/List_of_Nvidia_graphics_processing_units#GeForce_500_Series GeForce GTX 580en-us-tensor-core Hopper H100 Whitepaper], 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>CUDA C Programming Guide v7.0, Appendix G. COMPUTE CAPABILITIES, Table 12 Technical Specifications per Compute Capability</ref>The AMD * [https://endocs.wikipedianvidia.orgcom/wikicuda/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 coreshopper-tuning-guide/index. <ref>AMD Accelerated Parallel Processing OpenCL Programming html Hopper Tuning 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.]
=Memory== Ampere Architecture === The memory hierarchy of an GPU consists in main of private memory [https://en.wikipedia.org/wiki/Ampere_(registers accessed by an single thread resp. work-itemmicroarchitecture)Ampere microarchitecture] was announced on May 14, local memory (shared by threads of an block resp2020 <ref>[https://devblogs.nvidia. workcom/nvidia-ampere-architecture-in-items of an workdepth/ NVIDIA Ampere Architecture In-group )Depth | NVIDIA Developer Blog] by [https://people.csail.mit.edu/ronny/ Ronny Krashinsky], constant memory[https://cppcast.com/guest/ogiroux/ Olivier Giroux], different types of cache [https://blogs.nvidia.com/blog/author/stephenjones/ Stephen Jones], [https://blogs.nvidia.com/blog/author/nick-stam/ Nick Stam] and global memory[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>[https://devblogs.nvidia.com/cuda-11-features-revealed/ CUDA 11 Features Revealed | NVIDIA Developer Blog] by [https://devblogs.nvidia. Sizecom/author/pramarao/ Pramod Ramarao], May 14, latency and bandwidth vary between vendors and architectures2020</ref>.
Here the data for the Nvidia GeForce GTX 580 (* [https://enwww.wikipedianvidia.orgcom/content/dam/en-zz/Solutions/wikiData-Center/Fermi_%28microarchitecture%29 Fermi)nvidia-ampere-architecture-whitepaper.pdf Ampere GA100 Whitepaper] as an example* [https: <ref>CUDA C Programming Guide v7//www.0, Appendix Gnvidia.COMPUTE CAPABILITIES<com/content/PDF/ref>* 128 KiB private memory per compute unit* 48 KiB (16 KiB) local memory per compute unit (configurable)* 64 KiB constant memory* 8 KiB constant cache per compute unit* 16 KiB (48 KiB) L1 cache per compute unit (configurable)* 768 KiB L2 cachenvidia-ampere-ga-102-gpu-architecture-whitepaper-v2.pdf Ampere GA102 Whitepaper]* 1.5 GiB to 3 GiB global memoryHere the data for the AMD Radeon HD 7970 ([https://endocs.wikipedianvidia.orgcom/cuda/wikiampere-tuning-guide/Graphics_Core_Next GCNindex.html Ampere GPU Architecture Tuning Guide]) 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>* 256 KiB private memory per compute unit* 64 KiB local memory per compute unit* 64 KiB constant memory* 16 KiB constant cache per four compute units* 16 KiB L1 cache per compute unit* 768 KiB L2 cache* 3 GiB to 6 GiB global memory
=Instruction Throughput== Turing Architecture === GPUs are used in [https://en.wikipedia.org/wiki/High-performance_computing HPCTuring_(microarchitecture) Turing] environments because of their good [https://encards were first released in 2018.wikipedia.org/wiki/FLOP FLOP]/Watt ratio. The instruction throughput in general depends on They are the architecture (like Nvidia's [https://en.wikipedia.org/wiki/Tesla_%28microarchitecture%29 Tesla]first consumer cores to launch with RTX, for [https://en.wikipedia.org/wiki/Fermi_%28microarchitecture%29 FermiRay_tracing_(graphics) raytracing], [https://enfeatures.wikipedia.org/wiki/Kepler_%28microarchitecture%29 Kepler], These are also the first consumer cards to launch with TensorCores used for matrix multiplications to accelerate [https://en.wikipedia.org/wiki/Maxwell_%28microarchitecture%29 Maxwell] or AMD's [https://en.wikipedia.org/wiki/TeraScale_%28microarchitecture%29 TerascaleNeural Networks#Convolutional|convolutional neural networks], [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] The Turing GTX line of chips do not offer RTX 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 modelTensorCores.
* 32 bit Integer Performance [https://www.nvidia.com/content/dam/en-zz/Solutions/design-visualization/technologies/turing-architecture/NVIDIA-Turing-Architecture-Whitepaper.pdf Turing Architecture Whitepaper]* [https: The 32 bit integer performance can be architecture and operation depended less than 32 bit FLOP or 24 bit integer performance//docs.nvidia.com/cuda/turing-tuning-guide/index.html Turing Tuning Guide]
* 64 bit Integer Performance=== Volta Architecture === : Current GPU [https://en.wikipedia.org/wiki/Processor_register registersVolta_(microarchitecture) Volta] and Vector-[https://encards were released in 2017.wikipedia.org/wiki/Arithmetic_logic_unit ALUs] are 32 bit wide and have They were the first cards to launch with TensorCores, supporting matrix multiplications to emulate 64 bit integer operations.<ref>accelerate [[https://en.wikichip.org/w/images/a/a1/vega-whitepaper.pdf AMD Vega White PaperNeural Networks#Convolutional|convolutional neural networks]]</ref> <ref>[https://www.nvidia.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://enimages.wikipedianvidia.orgcom/wikicontent/Turing_(microarchitecture) Turing] and AMD [https:volta-architecture/pdf/envolta-architecture-whitepaper.wikipedia.org/wiki/AMD_RX_Vega_series Vegapdf Volta Architecture Whitepaper] have mixed precision support. Vega doubles the * [https://endocs.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.anandtechnvidia.com/showcuda/13282/nvidia-turing-architecturevolta-deeptuning-diveguide/4 AnandTech - Nvidia Turing Deep Dive page 4index.html Volta Tuning Guide]</ref>
* TensorCores=== Pascal Architecture ===: With Nvidia [https://en.wikipedia.org/wiki/Volta_Pascal_(microarchitecture) VoltaPascal] series TensorCores cards were introduced. They offer fp16*fp16+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> Turings 2nd gen TensorCores add FP16, INT8, INT4 optimized computation.<ref>[https://wwwfirst released in 2016.anandtech.com/show/13282/nvidia-turing-architecture-deep-dive/6 AnandTech - Nvidia Turing Deep Dive page 6]</ref>
==Throughput Examples==* [https://images.nvidia.com/content/pdf/tesla/whitepaper/pascal-architecture-whitepaper.pdf Pascal Architecture Whitepaper]* [https://docs.nvidia.com/cuda/pascal-tuning-guide/index.html Pascal Tuning Guide]
Nvidia GeForce GTX 580 (Fermi, CC 2.0) - 32 bit integer operations=== Maxwell Architecture ===[https://clock cycle per compute unit <ref>CUDA C Programming Guide v7en.0, Chapter 5wikipedia.4org/wiki/Maxwell(microarchitecture) Maxwell] cards were first released in 2014.1. Arithmetic Instructions</ref>
MAD 16* [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] MUL 16 ADD 32 Bit* [https://docs.nvidia.com/cuda/maxwell-tuning-shift 16 Bitwise XOR 32guide/index.html Maxwell Tuning Guide]
Max theoretic ADD operation throughput: 32 Ops * 16 CUs * 1544 MHz = 790= 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.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>=== PowerVR ===
MAD 1* [https:/4 MUL 1/4 ADD 1 Bit-shift 1 Bitwise XOR 1en.wikipedia.org/wiki/PowerVR#PowerVR_Graphics PowerVR series on Wikipedia]
Max theoretic ADD operation throughput: 1 Op * 2048 PEs * 925 MHz = 1894.4 GigaOps/sec == IMG ===
=Host-Device Latencies= 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 measurable latency for null-kernels of 5 microseconds <ref>* [https://devtalken.nvidiawikipedia.com/default/topic/1047965/cuda-programming-and-performanceorg/host-device-latencies-/post/5318041wiki/PowerVR#5318041 hostIMG_A-device latencies?Series_(Albiorix) IMG A series on Wikipedia] 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://wwwen.talkchesswikipedia.comorg/forum3wiki/viewtopic.php?f=7&t=67347PowerVR#p761239 Re: GPU ANN, how to deal with hostIMG_B-device latencies?] by [[Milos Stanisavljevic]], [[CCCSeries IMG B series on Wikipedia]], May 06, 2018</ref>.
=Deep Learning=Qualcomm ==Qualcomm offers Adreno GPUs in various types as a component of their Snapdragon SoCs. Since Adreno 300 series OpenCL support is offered.
GPUs were originally intended to process matrix multiplications for graphical transformations and rendering=== Adreno ===* [https://en.wikipedia. [[Neural Networksorg/wiki/Adreno#Convolutional|Convolutional Neural NetworksVariants Adreno variants on Wikipedia]] 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== Vivante Corporation ==Vivante licenses IP to third parties for embedded systems, 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 GC series offers optional OpenCL 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. === GC-Series ===
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 TPUsVivante_Corporation#Products GC series on Wikipedia], and the open source projects [[Leela Zero]] headed by [[Gian-Carlo Pascutto]] for [[Go]] and its [[Leela Chess Zero]] adaption.
=See also=
* [[Deep Learning]]
** [[AlphaGo]]
** [[AlphaZero]]
** [[Neural Networks#Convolutional|Convolutional Neural Networks]]
** [[Leela Zero]]
** [[Leela Chess Zero]]
* [[FPGA]]
* [[Graphics Programming]]
* [[SIMD and SWAR Techniques]]
* [[Thread]]
* [[Zeta]]
=Publications=
 ==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>
* [[Ren Wu]], [http://www.cedar.buffalo.edu/~binzhang/ Bin Zhang], [http://www.hpl.hp.com/people/meichun_hsu/ Meichun Hsu] ('''2009'''). ''[http://portal.acm.org/citation.cfm?id=1531668 Clustering billions of data points using GPUs]''. [http://www.computingfrontiers.org/2009/ ACM International Conference on Computing Frontiers]
* [https://github.com/markgovett Mark Govett], [https://www.linkedin.com/in/craig-tierney-9568545 Craig Tierney], [[Jacques Middlecoff]], [https://www.researchgate.net/profile/Tom_Henderson4 Tom Henderson] ('''2009'''). ''Using Graphical Processing Units (GPUs) for Next Generation Weather and Climate Prediction Models''. [http://www.cisl.ucar.edu/dir/CAS2K9/ CAS2K9 Workshop]
* [[Hank Dietz]], [https://dblp.uni-trier.de/pers/hd/y/Young:Bobby_Dalton Bobby Dalton Young] ('''2009'''). ''[https://link.springer.com/chapter/10.1007/978-3-642-13374-9_5 MIMD Interpretation on a GPU]''. [https://dblp.uni-trier.de/db/conf/lcpc/lcpc2009.html LCPC 2009], [http://aggregate.ee.engr.uky.edu/EXHIBITS/SC09/mogsimlcpc09final.pdf pdf], [http://aggregate.org/GPUMC/mogsimlcpc09slides.pdf slides.pdf]
* [https://dblp.uni-trier.de/pid/28/7183.html Sander van der Maar], [[Joost Batenburg]], [https://scholar.google.com/citations?user=TtXZhj8AAAAJ&hl=en Jan Sijbers] ('''2009'''). ''[https://link.springer.com/chapter/10.1007/978-3-642-03138-0_33 Experiences with Cell-BE and GPU for Tomography]''. [https://dblp.uni-trier.de/db/conf/samos/samos2009.html#MaarBS09 SAMOS 2009] <ref>[https://en.wikipedia.org/wiki/Cell_(microprocessor) Cell (microprocessor) from Wikipedia]</ref>
==2010...==
* [https://www.linkedin.com/in/avi-bleiweiss-456a5644 Avi Bleiweiss] ('''2010'''). ''Playing Zero-Sum Games on the GPU''. [https://en.wikipedia.org/wiki/Nvidia NVIDIA Corporation], [http://www.nvidia.com/object/io_1269574709099.html GPU Technology Conference 2010], [http://www.nvidia.com/content/gtc-2010/pdfs/2207_gtc2010.pdf slides as pdf]
* [[Damian Sulewski]] ('''2011'''). ''Large-Scale Parallel State Space Search Utilizing Graphics Processing Units and Solid State Disks''. Ph.D. thesis, [[University of Dortmund]], [https://eldorado.tu-dortmund.de/dspace/bitstream/2003/29418/1/Dissertation.pdf pdf]
* [[Damjan Strnad]], [[Nikola Guid]] ('''2011'''). ''[http://cit.fer.hr/index.php/CIT/article/view/2029 Parallel Alpha-Beta Algorithm on the GPU]''. [http://cit.fer.hr/index.php/CIT CIT. Journal of Computing and Information Technology], Vol. 19, No. 4 » [[Parallel Search]], [[Othello|Reversi]]
* [[Balázs Jako|Balázs Jákó]] ('''2011'''). ''Fast Hydraulic and Thermal Erosion on GPU''. M.Sc. thesis, Supervisor [https://hu.linkedin.com/in/bal%C3%A1zs-t%C3%B3th-1b151329 Balázs Tóth], [http://eg2011.bangor.ac.uk/ Eurographics 2011], [http://old.cescg.org/CESCG-2011/papers/TUBudapest-Jako-Balazs.pdf pdf]
'''2012'''
* [[Liang Li]], [[Hong Liu]], [[Peiyu Liu]], [[Taoying Liu]], [[Wei Li]], [[Hao Wang]] ('''2012'''). ''[https://www.semanticscholar.org/paper/A-Node-based-Parallel-Game-Tree-Algorithm-Using-Li-Liu/be21d7b9b91957b700aab4ce002e6753b826ff54 A Node-based Parallel Game Tree Algorithm Using GPUs]''. CLUSTER 2012 » [[Parallel Search]]
* [https://dblp.uni-trier.de/pers/hd/k/Karami:Ali Ali Karami], [[S. Ali Mirsoleimani]], [https://dblp.uni-trier.de/pers/hd/k/Khunjush:Farshad Farshad Khunjush] ('''2013'''). ''[https://ieeexplore.ieee.org/document/6714232 A statistical performance prediction model for OpenCL kernels on NVIDIA GPUs]''. [https://ieeexplore.ieee.org/xpl/mostRecentIssue.jsp?punumber=6708586 CADS 2013]
* [[Diego Rodríguez-Losada]], [[Pablo San Segundo]], [[Miguel Hernando]], [https://dblp.uni-trier.de/pers/hd/p/Puente:Paloma_de_la Paloma de la Puente], [https://dblp.uni-trier.de/pers/hd/v/Valero=Gomez:Alberto Alberto Valero-Gomez] ('''2013'''). ''GPU-Mapping: Robotic Map Building with Graphical Multiprocessors''. [https://dblp.uni-trier.de/db/journals/ram/ram20.html IEEE Robotics & Automation Magazine, Vol. 20, No. 2], [https://www.acin.tuwien.ac.at/fileadmin/acin/v4r/v4r/GPUMap_RAM2013.pdf pdf]
* [https://dblp.org/pid/28/977-2.html David Williams], [[Valeriu Codreanu]], [https://dblp.org/pid/88/5343-1.html Po Yang], [https://dblp.org/pid/54/784.html Baoquan Liu], [https://www.strath.ac.uk/staff/dongfengprofessor/ Feng Dong], [https://dblp.org/pid/136/5430.html Burhan Yasar], [https://scholar.google.com/citations?user=FZVGYiQAAAAJ&hl=en Babak Mahdian], [https://scholar.google.com/citations?user=8WO6cVUAAAAJ&hl=en Alessandro Chiarini], [https://zhaoxiahust.github.io/ Xia Zhao], [https://scholar.google.com/citations?user=jCFYHlkAAAAJ&hl=en Jos Roerdink] ('''2013'''). ''[https://link.springer.com/chapter/10.1007/978-3-642-55224-3_42 Evaluation of Autoparallelization Toolkits for Commodity GPUs]''. [https://dblp.org/db/conf/ppam/ppam2013-1.html#WilliamsCYLDYMCZR13 PPAM 2013]
'''2014'''
* [https://dblp.uni-trier.de/pers/hd/d/Dang:Qingqing Qingqing Dang], [https://dblp.uni-trier.de/pers/hd/y/Yan:Shengen Shengen Yan], [[Ren Wu]] ('''2014'''). ''[https://ieeexplore.ieee.org/document/7097862 A fast integral image generation algorithm on GPUs]''. [https://dblp.uni-trier.de/db/conf/icpads/icpads2014.html ICPADS 2014]
* [[S. Ali Mirsoleimani]], [https://dblp.uni-trier.de/pers/hd/k/Karami:Ali Ali Karami Ali Karami], [https://dblp.uni-trier.de/pers/hd/k/Khunjush:Farshad Farshad Khunjush] ('''2014'''). ''[https://link.springer.com/chapter/10.1007/978-3-319-04891-8_12 A Two-Tier Design Space Exploration Algorithm to Construct a GPU Performance Predictor]''. [https://dblp.uni-trier.de/db/conf/arcs/arcs2014.html ARCS 2014], [https://en.wikipedia.org/wiki/Lecture_Notes_in_Computer_Science Lecture Notes in Computer Science], Vol. 8350, [https://en.wikipedia.org/wiki/Springer_Science%2BBusiness_Media Springer]
* [[Steinar H. Gunderson]] ('''2014'''). ''[https://archive.fosdem.org/2014/schedule/event/movit/ Movit: High-speed, high-quality video filters on the GPU]''. [https://en.wikipedia.org/wiki/FOSDEM FOSDEM] [https://archive.fosdem.org/2014/ 2014], [https://movit.sesse.net/movit-fosdem2014.pdf pdf]* [https://dblp.org/pid/54/784.html Baoquan Liu], [https://scholar.google.com/citations?user=VspO6ZUAAAAJ&hl=en Alexandru Telea], [https://scholar.google.com/citations?user=jCFYHlkAAAAJ&hl=en Jos Roerdink], [https://dblp.org/pid/87/6797.html Gordon Clapworthy], [https://dblp.org/pid/28/977-2.html David Williams], [https://dblp.org/pid/88/5343-1.html Po Yang], [https://www.strath.ac.uk/staff/dongfengprofessor/ Feng Dong], [[Valeriu Codreanu]], [https://scholar.google.com/citations?user=8WO6cVUAAAAJ&hl=en Alessandro Chiarini] ('''2018'''). ''Parallel centerline extraction on the GPU''. [https://www.journals.elsevier.com/computers-and-graphics Computers & Graphics], Vol. 41, [https://strathprints.strath.ac.uk/70614/1/Liu_etal_CG2014_Parallel_centerline_extraction_GPU.pdf pdf]
==2015 ...==
* [[Peter H. Jin]], [[Kurt Keutzer]] ('''2015'''). ''Convolutional Monte Carlo Rollouts in Go''. [http://arxiv.org/abs/1512.03375 arXiv:1512.03375] » [[Deep Learning]], [[Go]], [[Monte-Carlo Tree Search|MCTS]]
* [[Liang Li]], [[Hong Liu]], [[Hao Wang]], [[Taoying Liu]], [[Wei Li]] ('''2015'''). ''[https://ieeexplore.ieee.org/document/6868996 A Parallel Algorithm for Game Tree Search Using GPGPU]''. [[IEEE#TPDS|IEEE Transactions on Parallel and Distributed Systems]], Vol. 26, No. 8 » [[Parallel Search]]
* [[Simon Portegies Zwart]], [https://github.com/jbedorf Jeroen Bédorf] ('''2015'''). ''[https://www.computer.org/csdl/magazine/co/2015/11/mco2015110050/13rRUx0Pqwe Using GPUs to Enable Simulation with Computational Gravitational Dynamics in Astrophysics]''. [[IEEE #Computer|IEEE Computer]], Vol. 48, No. 11
'''2016'''
* <span id="Astro"></span>[https://www.linkedin.com/in/sean-sheen-b99aba89 Sean Sheen] ('''2016'''). ''[https://digitalcommons.calpoly.edu/theses/1567/ Astro - A Low-Cost, Low-Power Cluster for CPU-GPU Hybrid Computing using the Jetson TK1]''. Master's thesis, [https://en.wikipedia.org/wiki/California_Polytechnic_State_University California Polytechnic State University], [https://digitalcommons.calpoly.edu/cgi/viewcontent.cgi?referer=&httpsredir=1&article=2723&context=theses pdf] <ref>[http://www.nvidia.com/object/jetson-tk1-embedded-dev-kit.html Jetson TK1 Embedded Development Kit | NVIDIA]</ref> <ref>[http://www.talkchess.com/forum/viewtopic.php?t=61761 Jetson GPU architecture] by [[Dann Corbit]], [[CCC]], October 18, 2016</ref>
* [https://scholar.google.com/citations?user=YyD7mwcAAAAJ&hl=en Jingyue Wu], [https://scholar.google.com/citations?user=EJcIByYAAAAJ&hl=en Artem Belevich], [https://scholar.google.com/citations?user=X5WAGdEAAAAJ&hl=en Eli Bendersky], [https://www.linkedin.com/in/mark-heffernan-873b663/ Mark Heffernan], [https://scholar.google.com/citations?user=Guehv9sAAAAJ&hl=en Chris Leary], [https://scholar.google.com/citations?user=fAmfZAYAAAAJ&hl=en Jacques Pienaar], [http://www.broune.com/ Bjarke Roune], [https://scholar.google.com/citations?user=Der7mNMAAAAJ&hl=en Rob Springer], [https://scholar.google.com/citations?user=zvfOH0wAAAAJ&hl=en Xuetian Weng], [https://scholar.google.com/citations?user=s7VCtl8AAAAJ&hl=en Robert Hundt] ('''2016'''). ''[https://dl.acm.org/citation.cfm?id=2854041 gpucc: an open-source GPGPU compiler]''. [https://cgo.org/cgo2016/ CGO 2016]
* [[David Silver]], [[Shih-Chieh Huang|Aja Huang]], [[Chris J. Maddison]], [[Arthur Guez]], [[Laurent Sifre]], [[George van den Driessche]], [[Julian Schrittwieser]], [[Ioannis Antonoglou]], [[Veda Panneershelvam]], [[Marc Lanctot]], [[Sander Dieleman]], [[Dominik Grewe]], [[John Nham]], [[Nal Kalchbrenner]], [[Ilya Sutskever]], [[Timothy Lillicrap]], [[Madeleine Leach]], [[Koray Kavukcuoglu]], [[Thore Graepel]], [[Demis Hassabis]] ('''2016'''). ''[http://www.nature.com/nature/journal/v529/n7587/full/nature16961.html Mastering the game of Go with deep neural networks and tree search]''. [https://en.wikipedia.org/wiki/Nature_%28journal%29 Nature], Vol. 529 » [[AlphaGo]]
* [[Balázs Jako|Balázs Jákó]] ('''2016'''). ''[https://www.semanticscholar.org/paper/Hardware-accelerated-hybrid-rendering-on-PowerVR-J%C3%A1k%C3%B3/d9d7f5784263c5abdcd6c1bf93267e334468b9b2 Hardware accelerated hybrid rendering on PowerVR GPUs]''. <ref>[https://en.wikipedia.org/wiki/PowerVR PowerVR from Wikipedia]</ref> [[IEEE]] [https://ieeexplore.ieee.org/xpl/conhome/7547434/proceeding 20th Jubilee International Conference on Intelligent Engineering Systems]
* [[Diogo R. Ferreira]], [https://dblp.uni-trier.de/pers/hd/s/Santos:Rui_M= Rui M. Santos] ('''2016'''). ''[https://github.com/diogoff/transition-counting-gpu Parallelization of Transition Counting for Process Mining on Multi-core CPUs and GPUs]''. [https://dblp.uni-trier.de/db/conf/bpm/bpmw2016.html BPM 2016]
* [https://dblp.org/pers/hd/s/Sch=uuml=tt:Ole Ole Schütt], [https://developer.nvidia.com/blog/author/peter-messmer/ Peter Messmer], [https://scholar.google.ch/citations?user=ajbBWN0AAAAJ&hl=en Jürg Hutter], [[Joost VandeVondele]] ('''2016'''). ''[https://onlinelibrary.wiley.com/doi/10.1002/9781118670712.ch8 GPU Accelerated Sparse Matrix–Matrix Multiplication for Linear Scaling Density Functional Theory]''. [https://www.cp2k.org/_media/gpu_book_chapter_submitted.pdf pdf] <ref>[https://en.wikipedia.org/wiki/Density_functional_theory Density functional theory from Wikipedia]</ref>
: Chapter 8 in [https://scholar.google.com/citations?user=AV307ZUAAAAJ&hl=en Ross C. Walker], [https://scholar.google.com/citations?user=PJusscIAAAAJ&hl=en Andreas W. Götz] ('''2016'''). ''[https://onlinelibrary.wiley.com/doi/book/10.1002/9781118670712 Electronic Structure Calculations on Graphics Processing Units: From Quantum Chemistry to Condensed Matter Physics]''. [https://en.wikipedia.org/wiki/Wiley_(publisher) John Wiley & Sons]
'''2017'''
* [[David Silver]], [[Thomas Hubert]], [[Julian Schrittwieser]], [[Ioannis Antonoglou]], [[Matthew Lai]], [[Arthur Guez]], [[Marc Lanctot]], [[Laurent Sifre]], [[Dharshan Kumaran]], [[Thore Graepel]], [[Timothy Lillicrap]], [[Karen Simonyan]], [[Demis Hassabis]] ('''2017'''). ''Mastering Chess and Shogi by Self-Play with a General Reinforcement Learning Algorithm''. [https://arxiv.org/abs/1712.01815 arXiv:1712.01815] » [[AlphaZero]]
* [[Tristan Cazenave]] ('''2017'''). ''[http://ieeexplore.ieee.org/document/7875402/ Residual Networks for Computer Go]''. [[IEEE#TOCIAIGAMES|IEEE Transactions on Computational Intelligence and AI in Games]], Vol. PP, No. 99, [http://www.lamsade.dauphine.fr/~cazenave/papers/resnet.pdf pdf]
* [https://scholar.google.com/citations?user=zLksndkAAAAJ&hl=en Jayvant Anantpur], [https://dblp.org/pid/09/10702.html Nagendra Gulur Dwarakanath], [https://dblp.org/pid/16/4410.html Shivaram Kalyanakrishnan], [[Shalabh Bhatnagar]], [https://dblp.org/pid/45/3592.html R. Govindarajan] ('''2017'''). ''RLWS: A Reinforcement Learning based GPU Warp Scheduler''. [https://arxiv.org/abs/1712.04303 arXiv:1712.04303]
'''2018'''
* [[David Silver]], [[Thomas Hubert]], [[Julian Schrittwieser]], [[Ioannis Antonoglou]], [[Matthew Lai]], [[Arthur Guez]], [[Marc Lanctot]], [[Laurent Sifre]], [[Dharshan Kumaran]], [[Thore Graepel]], [[Timothy Lillicrap]], [[Karen Simonyan]], [[Demis Hassabis]] ('''2018'''). ''[http://science.sciencemag.org/content/362/6419/1140 A general reinforcement learning algorithm that masters chess, shogi, and Go through self-play]''. [https://en.wikipedia.org/wiki/Science_(journal) Science], Vol. 362, No. 6419
* [http://www.talkchess.com/forum/viewtopic.php?t=66280 Announcing lczero] by [[Gary Linscott|Gary]], [[CCC]], January 09, 2018 » [[Leela Chess Zero]]
* [http://www.talkchess.com/forum3/viewtopic.php?f=7&t=67347 GPU ANN, how to deal with host-device latencies?] by [[Srdja Matovic]], [[CCC]], May 06, 2018 » [[Neural Networks]]
* [http://www.talkchess.com/forum3/viewtopic.php?f=2&t=67357 GPU contention] by [[Ian Kennedy]], [[CCC]], May 07, 2018 » [[Leela Chess Zero]]
* [http://www.talkchess.com/forum3/viewtopic.php?f=2&t=68448 How good is the RTX 2080 Ti for Leela?] by Hai, September 15, 2018 » [[Leela Chess Zero]] <ref>[https://en.wikipedia.org/wiki/GeForce_20_series GeForce 20 series from Wikipedia]</ref>
: [http://www.talkchess.com/forum3/viewtopic.php?f=2&t=68448&start=2 Re: How good is the RTX 2080 Ti for Leela?] by [[Ankan Banerjee]], [[CCC]], September 16, 2018
* [http://www.talkchess.com/forum3/viewtopic.php?f=7&t=69447 Generate EGTB with graphics cards?] by [[Pham Hong Nguyen|Nguyen Pham]], [[CCC]], January 01, 2019 » [[Endgame Tablebases]]
* [http://www.talkchess.com/forum3/viewtopic.php?f=2&t=69478 LCZero FAQ is missing one important fact] by [[Jouni Uski]], [[CCC]], January 01, 2019 » [[Leela Chess Zero]]
* [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++]]
* [http://www.talkchess.com/forum3/viewtopic.php?f=7&t=71058 Lazy-evaluation of futures for parallel work-efficient Alpha-Beta search] by Percival Tiglao, [[CCC]], June 06, 2019
* [https://www.game-ai-forum.org/viewtopic.php?f=21&t=694 My home-made CUDA kernel for convolutions] by [[Rémi Coulom]], [[Computer Chess Forums|Game-AI Forum]], November 09, 2019 » [[Deep Learning]]
* [http://www.talkchess.com/forum3/viewtopic.php?f=2&t=72320 GPU rumors 2020] by [[Srdja Matovic]], [[CCC]], November 13, 2019
==2020 ...==
* [http://www.talkchess.com/forum3/viewtopic.php?f=7&t=74771 AB search with NN on GPU...] by [[Srdja Matovic]], [[CCC]], August 13, 2020 » [[Neural Networks]] <ref>[https://forums.developer.nvidia.com/t/kernel-launch-latency/62455 kernel launch latency - CUDA / CUDA Programming and Performance - NVIDIA Developer Forums] by LukeCuda, June 18, 2018</ref>
* [http://www.talkchess.com/forum3/viewtopic.php?f=2&t=75073 I stumbled upon this article on the new Nvidia RTX GPUs] by [[Kai Laskos]], [[CCC]], September 10, 2020
* [http://www.talkchess.com/forum3/viewtopic.php?f=2&t=75639 Will AMD RDNA2 based Radeon RX 6000 series kick butt with Lc0?] by [[Srdja Matovic]], [[CCC]], November 01, 2020
* [http://www.talkchess.com/forum3/viewtopic.php?f=7&t=76986 Zeta with NNUE on GPU?] by [[Srdja Matovic]], [[CCC]], March 31, 2021 » [[Zeta]], [[NNUE]]
* [https://talkchess.com/forum3/viewtopic.php?f=2&t=77097 GPU rumors 2021] by [[Srdja Matovic]], [[CCC]], April 16, 2021
* [https://www.talkchess.com/forum3/viewtopic.php?f=7&t=79078 Comparison of all known Sliding lookup algorithms <nowiki>[CUDA]</nowiki>] by [[Daniel Infuehr]], [[CCC]], January 08, 2022 » [[Sliding Piece Attacks]]
* [https://talkchess.com/forum3/viewtopic.php?f=7&t=72566&p=955538#p955538 Re: China boosts in silicon...] by [[Srdja Matovic]], [[CCC]], January 13, 2024
=External Links=
* [https://en.wikipedia.org/wiki/General-purpose_computing_on_graphics_processing_units General-purpose computing on graphics processing units (GPGPU) from Wikipedia]
* [https://en.wikipedia.org/wiki/List_of_AMD_graphics_processing_units List of AMD graphics processing units from Wikipedia]
* [https://en.wikipedia.org/wiki/List_of_Intel_graphics_processing_units List of Intel graphics processing units from Wikipedia]
* [https://en.wikipedia.org/wiki/List_of_Nvidia_graphics_processing_units List of Nvidia graphics processing units from Wikipedia]
* [https://developer.nvidia.com/ NVIDIA Developer]
422
edits

Navigation menu