GPU

Home * Hardware * GPU



GPU (Graphics Processing Unit), a specialized processor primarily intended to fast image processing. GPUs may have more raw computing power than general purpose CPUs but need a specialized and massive parallelized way of programming. Leela Chess Zero has proven that a Best-first Monte-Carlo Tree Search (MCTS) with deep learning methodology will work with GPU architectures.

=GPGPU=

The traditional job of a GPU is to take the x,y,z coordinates of triangles, and map these triangles to screen space through a 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 OpenGL or 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 vertex shaders or pixel shaders. GPUs evolved to accelerate general purpose compute from pixel shader and vertex shader programmers, and even merged the functionality into "universal" shaders (which can perform either vertex shading or pixel shading).

Today, these universal shaders are flexible enough to provide General Purpose compute for GPUs (GPGPU). GPGPU languages, such as OpenCL or CUDA, is how the programmer can access this capability.

Khronos OpenCL
The Khronos group is a committee formed to oversee the OpenGL, OpenCL, and Vulkan standards. Although compute shaders exist in all languages, OpenCL is the designated general purpose compute language.

OpenCL 1.2 is widely supported by AMD, 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.


 * OpenCL 1.2 Specification
 * OpenCL 1.2 Reference


 * OpenCL 2.0 Specification
 * OpenCL 2.0 C Language Specification
 * OpenCL 2.0 Reference

Nvidia Software overview
Nvidia CUDA is their general purpose compute framework. CUDA has a C++ compiler based on LLVM / clang, which compiles into an assembly-like language called 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.


 * Nvidia CUDA Zone
 * Nvidia PTX ISA
 * Nvidia CUDA Toolkit Documentation

AMD Software Overview
AMD's original software stack, called AMDGPU-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 ROCm. ROCm is AMD's open source compiler and device driver stack intended for general purpose compute. ROCm supports two languages: 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 PCIe 3.0 and relatively recent GPUs (such as the RX 580, and Vega GPUs).

AMD regularly publishes the assembly language details of their architectures. Their "GCN Assembly" changes slightly from generation to generation, but the fundamental principles have remained the same.

AMD's OpenCL documentation, especially the "OpenCL Programming Guide" and the "Optimization Guide" are good places to start for beginners looking to program their GPUs. For Linux developers, the ROCm environment is under active development and has enough features to get code working well.


 * ROCm Homepage
 * AMD OpenCL Programming Guide
 * AMD OpenCL Optimization Guide
 * RDNA Instruction Set
 * Vega Instruction Set

Other 3rd party tools

 * DirectCompute (GPGPU API by Microsoft)
 * OpenMP Device Offload
 * OpenACC Device Offload
 * Metal (GPU and GPGPU API by Apple)
 * oneAPI (Data Parallel C++ by Intel)

=The SIMT Programming Model=

CUDA, OpenCL, ROCm HIP, all have the same model of implicitly parallel programming. All threads are given an identifier: 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. In the above example code, 97% of the threads will be effectively idle during doA, while 3% of the threads will be idle during doB.

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
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 1920x1080 grid for convenience). Specifying these 2,073,600 work items is the purpose of a CUDA Grid or OpenCL NDRange.

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

These blocks and workgroups will execute with as much parallel processing as the underlying hardware can support. Roughly 150 CUDA Blocks or OpenCL Workgroups at a time on a typical midrange GPU circa from 2019 (such as a Nvidia 2060 Super or AMD 5700). The most important note is that blocks within a grid (or workgroups within an NDRange) may not execute concurrently with each other. Some degree of sequential processing may happen. If thread #0 creates a Spinlock waiting for thread #1000000 to communicate with it, modern hardware will probably never have the two threads executing concurrently with each other, and the code would likely timeout. In practice, the easiest mechanism for Grid or NDRange sized synchronization is to wait for the kernel to finish executing: to have the CPU wait and process the results in between Grid or NDRanges.

For example: LeelaZero will schedule an NDRange for each Convolve operation, as well as merge and other primitives. The convolve operation is over a 3-dimensional NDRange for . To build up a full CNN operation, the CPU will schedule different operations for the GPU: convolve, merge, transform and more.

Memory Model
OpenCL, CUDA, ROCM, and other GPU-languages all have a similar memory model.


 * __device__ (CUDA) or __global (OpenCL) memory -- OpenCL __global and CUDA __device__ memory exists on the GPU's VRAM. Any threads can access any part of __device__ or __global memory, although memory-ordering and caching details can get quite complicated if multiple threads simultaneously read and write to a particular memory location. Proper memory ordering with __threadfence (CUDA) or mem_fence (OpenCL) is essential to preventing memory-consistency issues.


 * __constant__ (CUDA) or __constant (OpenCL) memory -- Constants are not allowed to change during the execution of a particular kernel. Historically, this was used by Pixel Shaders as they read texture data. The texture-data could be computed and loaded onto the GPU, but the data was not allowed to change during the Pixel Shader's execution. Both Nvidia and AMD GPUs have special caches (and in AMD's case: special registers called sGPRs) which accelerate constant-data. The caches associated with this memory space is sometimes called K$ (Konstant-cache), and has to be independently flushed if its data ever changes. The main benefit in both AMD and Nvidia systems is that K$ values are broadcast extremely efficiently to all threads in a wavefront, but only if all threads in a wavefront are reading from the same memory location. Instead of haing 32-memory reads (Nvidia) or 64-memory reads (AMD GCN), a read from K$ can be optimized into a single-read, broadcast to all 32 or 64-threads of a Warp or Wavefront.


 * __shared__ (CUDA) or __local (OpenCL) memory -- This is highly-accelerated memory regions designed for threads to exchange data within a CUDA Block or OpenCL Workgroup. On AMD Systems, there is more Local "LDS" memory than even L1 Cache (GCN) or L0 Cache (RDNA).


 * Default (CUDA) or __private (OpenCL) Memory -- Private memory typically maps to a GPU-register, and is inaccessible to other threads. If a kernel requires more memory than what can exist in GPU-registers, the data will automatically spill over into global VRAM (with an associated performance penalty). In practice, this spillover is well interleaved, well-optimized, and reduced to as small a subset as possible through compiler optimizations.

Here the data for the Nvidia GeForce GTX 580 (Fermi) as an example: Here the data for the AMD Radeon HD 7970 (GCN) as an example:
 * 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
 * 1.5 GiB to 3 GiB global memory
 * 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

= Architectures and Physical Hardware =

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.

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.

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.

Nvidia
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".

Nvidia's "Titan" line of Geforce cards use consumer drivers, but use professional or server class chips. As such, the Titan line can cost anywhere from $1000 to $3000 per card.

Ampere Architecture
The Ampere microarchitecture was announced on May 14, 2020. The Nvidia A100 GPU based on the Ampere architecture delivers a generational leap in accelerated computing in conjunction with CUDA 11.


 * DGX A100
 * HGX A100

Turing Architecture
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-intersections with lists of bounding-boxes, accelerating raytracing performance. These are also the first consumer cards to launch with Tensor cores, 4x4 matrix multiplication FP16 instructions to accelerate convolutional neural networks.


 * RTX 2080 Ti
 * RTX 2080
 * RTX 2070 Ti
 * RTX 2070 Super
 * RTX 2070
 * RTX 2060 Super
 * RTX 2060
 * GTX 1660 -- Low-end GPU without Tensor cores or RTX Cores.

Volta Architecture
Architecture Whitepaper

Volta cards were released in 2017. Only Tesla and Titan cards were produced in this generation, aiming only for the most expensive end of the market. They were the first cards to launch with Tensor cores, supporting 4x4 FP16 matrix multiplications to accelerate convolutional neural networks.


 * Tesla V100
 * Titan V

Pascal Architecture
Architecture Whitepaper

Pascal cards were first released in 2016.


 * Tesla P100
 * Titan Xp
 * GTX 1080 Ti
 * GTX 1080
 * GTX 1070 Ti
 * GTX 1060
 * GTX 1050
 * GTX 1030

Navi 2X RDNA 2.0

 * RDNA 2 from Wikipedua

RDNA 2.0 cards were unveiled on October 28, 2020.


 * Radeon RX 6000 series from Wikipedua

Navi RDNA 1.0

 * RDNA Whitepaper
 * Architecture Slide Deck
 * RDNA (microarchitecture) from Wikipedua

RDNA cards were first released in 2019. RDNA is a major change for AMD cards: the underlying hardware supports both Wave32 and Wave64 gangs of threads. Compute Units have 2x32 wide SIMD units, each of which executes 32 threads per clock tick. A Wave64 workgroup will execute on a single SIMD unit, but over two clock ticks. It should be noted that these Wave32 still have 5 cycles of latency before registers can be reused, so a Wave64 executing over two clock ticks will have fewer stalls than a Wave32.


 * Radeon RX 5000 series from Wikipedua
 * Radeon 5700 XT
 * Radeon 5700

Vega GCN 5th gen
Architecture Whitepaper

Vega cards were first released in 2017. Vega is the last in the line of the GCN Architecture: 64 threads per wavefront. Each compute unit contains 4x SIMD units, supporting a total of 40 wavefronts per compute unit (a queue of 10-wavefronts per SIMD Unit). Each SIMD unit contains 16 vALUs for general compute + 1 sALU for branching and constant logic. Each SIMD unit executes the same instruction over four clock ticks (16 vALUs x 4 clock ticks == 64 threads per Wavefront).

Vega specifically added Packed FP16 instructions, such as dot-product and packed add and packed multiply. From a programming level, these packed FP16 instructions are SIMD-within-SIMD, each SIMD thread could operate its own SIMD FP16 instruction akin to AVX or SSE from the x64 architecture.


 * Radeon VII
 * Vega64
 * Vega56

Polaris GCN 4th gen
Polaris cards were first released in 2016 under the AMD Radeon 400 series name.

Architecture Whitepaper


 * RX 580
 * RX 570
 * RX 560

=Instruction Throughput= GPUs are used in HPC environments because of their good FLOP/Watt ratio. The instruction throughput in general depends on the architecture (like Nvidia's Tesla, Fermi, Kepler, Maxwell or AMD's Terascale, GCN, RDNA), the brand (like Nvidia GeForce, Quadro, Tesla or AMD Radeon, Radeon Pro, Radeon Instinct) and the specific model.

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
 * Current GPU registers and Vector-ALUs are 32 bit wide and have to emulate 64 bit integer operations.


 * INT8
 * Some architectures like AMD Vega or Intel Xe offer higher throughput with lower precision. They double the FP16 and quadruple the INT8 throughput.

Floating Point Instruction Throughput

 * FP32
 * Consumer GPU performance is measured usually in single-precision (32 bit) floating point FMA, fused-multiply-add, throughput.


 * FP64
 * Consumer GPUs have in general a lower ratio (FP32:FP64) for double-precision (64 bit) floating point operations than server brand GPUs, like 4:1 down to 32:1 compared to 2:1 to 4:1.


 * FP16
 * Some newer GPGPU architectures offer half-precision (16 bit) floating point operation throughput with an FP32:FP16 ratio of 1:2. Older architectures migth not support FP16 at all, at the same rate as FP32, or at very low rates.

Nvidia TensorCores

 * With Nvidia Volta series TensorCores were introduced. They offer FP16xFP16+FP32, matrix-multiplication-accumulate-units, used to accelerate neural networks. Turing's 2nd gen TensorCores add FP16, INT8, INT4 optimized computation. Amperes's 3rd gen adds support for BF16, TF32, FP64 and sparsity acceleration.

AMD Matrix Cores

 * AMD released 2020 its server-class CDNA architecture with Matrix Cores which support MFMA, matrix-fused-multiply-add, operations on various data types like INT8, FP16, BF16, FP32.

Intel XMX Cores

 * Intel plans XMX, Xe Matrix eXtensions, for its upcoming Xe discrete GPU series.

Throughput Examples
Nvidia GeForce GTX 580 (Fermi, CC 2.0) - 32 bit integer operations/clock cycle per compute unit

MAD 16 MUL 16 ADD 32 Bit-shift 16 Bitwise XOR 32

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

AMD Radeon HD 7970 (GCN 1.0) - 32 bit integer operations/clock cycle per processing element

MAD 1/4 MUL 1/4 ADD 1 Bit-shift 1 Bitwise XOR 1

Max theoretic ADD operation throughput: 1 Op * 2048 PEs * 925 MHz = 1894.4 GigaOps/sec

=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 up to 100s of microseconds. One solution to overcome this limitation is to couple tasks to batches to be executed in one run.

=Deep Learning= GPUs were originally intended to process matrix multiplications for graphical transformations and rendering. Convolutional Neural Networks can have their operations interpreted as a series of matrix multiplications. GPUs are therefore a natural fit to parallelize and process CNNs.

GPUs traditionally operated on 32-bit floating point numbers. However, CNNs can make due with 16-bit half floats (FP16), or even 8-bit or 4-bit numbers. One thousand single-precision floats will take up 4kB of space, while one-thousand FP16 will take up 2kB of space. A half-float uses half the memory, eats only half the memory bandwidth, and only half the space in caches. As such, GPUs such as AMD Vega or Nvidia Volta added support for FP16 processing.

Specialized units, such as Nvidia Volta's "Tensor cores", can perform an entire 4x4 block of FP16 matrix multiplications in just one PTX assembly language statement. It is with these instructions that CNN operations are accelerated.

GPUs are much more suited than CPUs to implement and train Convolutional Neural Networks (CNN), and were therefore also responsible for the deep learning boom, also affecting game playing programs combining CNN with MCTS, as pioneered by Google DeepMind's AlphaGo and AlphaZero entities in Go, Shogi and Chess using TPUs, and the open source projects Leela Zero headed by Gian-Carlo Pascutto for Go and its Leela Chess Zero adaption.

=History= In the 1970s and 1980s RAM was expensive and Home Computers used custom graphics chips to operate directly on registers/memory without a dedicated frame buffer, like TIAin the Atari VCS gaming system, GTIA+ANTIC in the Atari 400/800 series, or Denise+Agnus in the Commodore Amiga series. The 1990s would make 3D graphics and 3D modeling more popular, especially for video games. Cards specifically designed to accelerate 3D math, such as the 3dfx Voodoo2, were used by the video game community to play 3D graphics. Some game engines, such as Quake, could use instead the SIMD-capabilities of CPUs such as the Intel MMX instruction set or AMD's 3DNow!. Sony's 3D capable chip used in the PlayStation (1994) and Nvidia's 2D/3D combi chips like NV1 (1995) coined the term GPU for 3D graphics hardware acceleration. With the advent of the unified shader architecture, like in Nvidia Tesla (2006), ATI/AMD TeraScale (2007) or Intel GMA X3000 (2006), GPGPU frameworks like CUDA and OpenCL emerged and gained in popularity.

The large number of regular matrix multiplications led to natural SIMD-style algorithms. The 3D graphics community drew upon the rich history of vector-compute and SIMD-compute from 1980s and 1970s supercomputers. As such, many publications relating to Cray-vector supercomputers or the Connection Machine supercomputer easily apply to modern GPUs. For example, all the algorithms described in the 1986 publication Data Parallel Algorithms can be efficiently executed on a modern GPU workgroup (roughly ~256x GPU threads). The Data Parallel Algorithms paper is a beginner-level algorithms paper, demonstrating simple and efficient parallel-prefix sum, parallel-linked list traversal, parallel RegEx matching on the 4096x parallel Connection Machine-2 supercomputer.

Modern papers on GPUs, such as Nvidia's excellent Parallel Prefix Sum (Scan) with CUDA (GPU Gems 3), are built on top of these papers from the 1980s or 1990s. As such, the beginner will find it far easier to read the papers from the 1980s or 90s before attempting to read a modern piece like GPU Gems 3.

=Chess Engines=
 * Category:GPU

=See also=
 * Deep Learning
 * FPGA
 * Graphics Programming
 * Monte-Carlo Tree Search
 * MCαβ
 * UCT
 * Parallel Search
 * Perft(15)
 * SIMD and SWAR Techniques
 * Thread

=Publications=

1986

 * W. Daniel Hillis, Guy L. Steele, Jr. (1986). Data parallel algorithms. Communications of the ACM, Vol. 29, No. 12, Special Issue on Parallelism

1990

 * Guy E. Blelloch (1990). Vector Models for Data-Parallel Computing. MIT Press, pdf

2008 ...

 * Vlad Stamate (2008). Real Time Photon Mapping Approximation on the GPU. in ShaderX6 - Advanced Rendering Techniques
 * Ren Wu, Bin Zhang, Meichun Hsu (2009). Clustering billions of data points using GPUs. ACM International Conference on Computing Frontiers
 * Mark Govett, Craig Tierney, Jacques Middlecoff, Tom Henderson (2009). Using Graphical Processing Units (GPUs) for Next Generation Weather and Climate Prediction Models. CAS2K9 Workshop
 * Hank Dietz, Bobby Dalton Young (2009). MIMD Interpretation on a GPU. LCPC 2009, pdf, slides.pdf
 * Sander van der Maar, Joost Batenburg, Jan Sijbers (2009). Experiences with Cell-BE and GPU for Tomography. SAMOS 2009

2010...
2011 2012 2013 2014
 * Avi Bleiweiss (2010). Playing Zero-Sum Games on the GPU. NVIDIA Corporation, GPU Technology Conference 2010, slides as pdf
 * Mark Govett, Jacques Middlecoff, Tom Henderson (2010). Running the NIM Next-Generation Weather Model on GPUs. CCGRID 2010
 * John Nickolls, William J. Dally (2010). The GPU Computing Era. IEEE Micro.
 * Mark Govett, Jacques Middlecoff, Tom Henderson, Jim Rosinski, Craig Tierney (2011). Parallelization of the NIM Dynamical Core for GPUs. slides as pdf
 * Ľubomír Lackovič (2011). Parallel Game Tree Search Using GPU. Institute of Informatics and Software Engineering, Faculty of Informatics and Information Technologies, Slovak University of Technology in Bratislava, pdf
 * Dan Anthony Feliciano Alcantara (2011). Efficient Hash Tables on the GPU. Ph. D. thesis, University of California, Davis, pdf » Hash Table
 * Damian Sulewski (2011). Large-Scale Parallel State Space Search Utilizing Graphics Processing Units and Solid State Disks. Ph.D. thesis, University of Dortmund, pdf
 * Damjan Strnad, Nikola Guid (2011). Parallel Alpha-Beta Algorithm on the GPU. CIT. Journal of Computing and Information Technology, Vol. 19, No. 4 » Parallel Search, Reversi
 * Balázs Jákó (2011). Fast Hydraulic and Thermal Erosion on GPU. M.Sc. thesis, Supervisor Balázs Tóth, Eurographics 2011, pdf
 * Liang Li, Hong Liu, Peiyu Liu, Taoying Liu, Wei Li, Hao Wang (2012). A Node-based Parallel Game Tree Algorithm Using GPUs. CLUSTER 2012 » Parallel Search
 * S. Ali Mirsoleimani, Ali Karami Ali Karami, Farshad Khunjush (2013). A parallel memetic algorithm on GPU to solve the task scheduling problem in heterogeneous environments. GECCO '13
 * Ali Karami, S. Ali Mirsoleimani, Farshad Khunjush (2013). A statistical performance prediction model for OpenCL kernels on NVIDIA GPUs. CADS 2013
 * Diego Rodríguez-Losada, Pablo San Segundo, Miguel Hernando, Paloma de la Puente, Alberto Valero-Gomez (2013). GPU-Mapping: Robotic Map Building with Graphical Multiprocessors. IEEE Robotics & Automation Magazine, Vol. 20, No. 2, pdf
 * Qingqing Dang, Shengen Yan, Ren Wu (2014). A fast integral image generation algorithm on GPUs. ICPADS 2014
 * S. Ali Mirsoleimani, Ali Karami Ali Karami, Farshad Khunjush (2014). A Two-Tier Design Space Exploration Algorithm to Construct a GPU Performance Predictor. ARCS 2014, Lecture Notes in Computer Science, Vol. 8350, Springer
 * Steinar H. Gunderson (2014). Movit: High-speed, high-quality video filters on the GPU. FOSDEM 2014, pdf

2015 ...
2016
 * Peter H. Jin, Kurt Keutzer (2015). Convolutional Monte Carlo Rollouts in Go. arXiv:1512.03375 » Deep Learning, Go, MCTS
 * Liang Li, Hong Liu, Hao Wang, Taoying Liu, Wei Li (2015). A Parallel Algorithm for Game Tree Search Using GPGPU. IEEE Transactions on Parallel and Distributed Systems, Vol. 26, No. 8 » Parallel Search
 * Sean Sheen (2016). Astro - A Low-Cost, Low-Power Cluster for CPU-GPU Hybrid Computing using the Jetson TK1. Master's thesis, California Polytechnic State University, pdf
 * Jingyue Wu, Artem Belevich, Eli Bendersky, Mark Heffernan, Chris Leary, Jacques Pienaar, Bjarke Roune, Rob Springer, Xuetian Weng, Robert Hundt (2016). gpucc: an open-source GPGPU compiler. CGO 2016
 * David Silver, 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). Mastering the game of Go with deep neural networks and tree search. Nature, Vol. 529 » AlphaGo
 * Balázs Jákó (2016). Hardware accelerated hybrid rendering on PowerVR GPUs. IEEE 20th Jubilee International Conference on Intelligent Engineering Systems
 * Diogo R. Ferreira, Rui M. Santos (2016). Parallelization of Transition Counting for Process Mining on Multi-core CPUs and GPUs. BPM 2016
 * Ole Schütt, Peter Messmer, Jürg Hutter, Joost VandeVondele (2016). GPU Accelerated Sparse Matrix–Matrix Multiplication for Linear Scaling Density Functional Theory. pdf
 * Chapter 8 in Ross C. Walker, Andreas W. Götz (2016). Electronic Structure Calculations on Graphics Processing Units: From Quantum Chemistry to Condensed Matter Physics. John Wiley & Sons

2017 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 (2017). Mastering Chess and Shogi by Self-Play with a General Reinforcement Learning Algorithm. arXiv:1712.01815 » AlphaZero
 * Tristan Cazenave (2017). Residual Networks for Computer Go. IEEE Transactions on Computational Intelligence and AI in Games, Vol. PP, No. 99, pdf
 * Jayvant Anantpur, Nagendra Gulur Dwarakanath, Shivaram Kalyanakrishnan, Shalabh Bhatnagar, R. Govindarajan (2017). RLWS: A Reinforcement Learning based GPU Warp Scheduler. arXiv:1712.04303
 * 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). A general reinforcement learning algorithm that masters chess, shogi, and Go through self-play. Science, Vol. 362, No. 6419

=Forum Posts=

2005 ...

 * Hardware assist by Nicolai Czempin, Winboard Forum, August 27, 2006
 * Monte carlo on a NVIDIA GPU ? by Marco Costalba, CCC, August 01, 2008

2010 ...
2011
 * Using the GPU by Louis Zulli, CCC, February 19, 2010
 * GPGPU and computer chess by Wim Sjoho, CCC, February 09, 2011
 * Possible Board Presentation and Move Generation for GPUs? by Srdja Matovic, CCC, March 19, 2011
 * Re: Possible Board Presentation and Move Generation for GPUs by Steffan Westcott, CCC, March 20, 2011

2012 2013
 * Zeta plays chess on a gpu by Srdja Matovic, CCC, June 23, 2011 » Zeta
 * GPU Search Methods by Joshua Haglund, CCC, July 04, 2011
 * Possible Search Algorithms for GPUs? by Srdja Matovic, CCC, January 07, 2012
 * uct on gpu by Daniel Shawul, CCC, February 24, 2012 » UCT
 * Is there such a thing as branchless move generation? by John Hamlen, CCC, June 07, 2012 » Move Generation
 * Choosing a GPU platform: AMD and Nvidia by John Hamlen, CCC, June 10, 2012
 * Nvidias K20 with Recursion by Srdja Matovic, CCC, December 04, 2012
 * Kogge Stone, Vector Based by Srdja Matovic, CCC, January 22, 2013 » Kogge-Stone Algorithm
 * GPU chess engine by Samuel Siltanen, CCC, February 27, 2013
 * Fast perft on GPU (upto 20 Billion nps w/o hashing) by Ankan Banerjee, CCC, June 22, 2013 » Perft, Kogge-Stone Algorithm

2015 ...
2017 2018
 * GPU chess update, local memory... by Srdja Matovic, CCC, June 06, 2016
 * Jetson GPU architecture by Dann Corbit, CCC, October 18, 2016 » Astro
 * Pigeon is now running on the GPU by Stuart Riffle, CCC, November 02, 2016 » Pigeon
 * Back to the basics, generating moves on gpu in parallel... by Srdja Matovic, CCC, March 05, 2017 » Move Generation
 * Re: Perft(15): comparison of estimates with Ankan's result by Ankan Banerjee, CCC, August 26, 2017 » Perft(15)
 * Chess Engine and GPU by Fishpov, Rybka Forum, October 09, 2017
 * To TPU or not to TPU... by Srdja Matovic, CCC, December 16, 2017 » Deep Learning
 * Announcing lczero by Gary, CCC, January 09, 2018 » Leela Chess Zero
 * GPU ANN, how to deal with host-device latencies? by Srdja Matovic, CCC, May 06, 2018 » Neural Networks
 * GPU contention by Ian Kennedy, CCC, May 07, 2018 » Leela Chess Zero
 * How good is the RTX 2080 Ti for Leela? by Hai, September 15, 2018 » Leela Chess Zero
 * Re: How good is the RTX 2080 Ti for Leela? by Ankan Banerjee, CCC, September 16, 2018

2019
 * My non-OC RTX 2070 is very fast with Lc0 by Kai Laskos, CCC, November 19, 2018 » Leela Chess Zero
 * LC0 using 4 x 2080 Ti GPU's on Chess.com tourney? by M. Ansari, CCC, December 28, 2018 » Leela Chess Zero
 * Generate EGTB with graphics cards? by Nguyen Pham, CCC, January 01, 2019 » Endgame Tablebases
 * LCZero FAQ is missing one important fact by Jouni Uski, CCC, January 01, 2019 » Leela Chess Zero
 * Michael Larabel benches lc0 on various GPUs by Warren D. Smith, LCZero Forum, January 14, 2019 » Lc0
 * Using LC0 with one or two GPUs - a guide by Srdja Matovic, CCC, March 30, 2019 » Lc0
 * Wouldn't it be nice if C++ GPU by Chris Whittington, CCC, April 25, 2019 » C++
 * Lazy-evaluation of futures for parallel work-efficient Alpha-Beta search by Percival Tiglao, CCC, June 06, 2019
 * My home-made CUDA kernel for convolutions by Rémi Coulom, Game-AI Forum, November 09, 2019 » Deep Learning
 * GPU rumors 2020 by Srdja Matovic, CCC, November 13, 2019

2020 ...

 * AB search with NN on GPU... by Srdja Matovic, CCC, August 13, 2020 » Neural Networks
 * I stumbled upon this article on the new Nvidia RTX GPUs by Kai Laskos, CCC, September 10, 2020
 * Will AMD RDNA2 based Radeon RX 6000 series kick butt with Lc0? by Srdja Matovic, CCC, November 01, 2020
 * Zeta with NNUE on GPU? by Srdja Matovic, CCC, March 31, 2021 » Zeta, NNUE

=External Links=
 * Graphics processing unit from Wikipedia
 * Video card from Wikipedia
 * Heterogeneous System Architecture from Wikipedia
 * Tensor processing unit from Wikipedia
 * General-purpose computing on graphics processing units (GPGPU) from Wikipedia
 * List of AMD graphics processing units from Wikipedia
 * List of Nvidia graphics processing units from Wikipedia
 * NVIDIA Developer
 * NVIDIA GPU Programming Guide

OpenCL

 * OpenCL from Wikipedia
 * Part 1: OpenCL™ – Portable Parallelism - CodeProject
 * Part 2: OpenCL™ – Memory Spaces - CodeProject

CUDA

 * CUDA from Wikipedia
 * CUDA Zone | NVIDIA Developer
 * Nvidia CUDA Compiler (NVCC) from Wikipedia
 * Compiling CUDA with clang — LLVM Clang documentation
 * CppCon 2016: “Bringing Clang and C++ to GPUs: An Open-Source, CUDA-Compatible GPU C++ Compiler" by Justin Lebar, YouTube Video

Deep Learning

 * Deep Learning | NVIDIA Developer » Deep Learning
 * NVIDIA cuDNN | NVIDIA Developer
 * Efficient mapping of the training of Convolutional Neural Networks to a CUDA-based cluster
 * Deep Learning in a Nutshell: Core Concepts by Tim Dettmers, Parallel Forall, November 3, 2015
 * Deep Learning in a Nutshell: History and Training by Tim Dettmers, Parallel Forall, December 16, 2015
 * Deep Learning in a Nutshell: Sequence Learning by Tim Dettmers, Parallel Forall, March 7, 2016
 * Deep Learning in a Nutshell: Reinforcement Learning by Tim Dettmers, Parallel Forall, September 8, 2016
 * Faster deep learning with GPUs and Theano
 * Theano (software) from Wikipedia
 * TensorFlow from Wikipedia

Game Programming

 * Advanced game programming | Session 5 - GPGPU programming by Andy Thomason
 * Leela Zero by Gian-Carlo Pascutto » Leela Zero
 * GitHub - gcp/leela-zero: Go engine with no human-provided knowledge, modeled after the AlphaGo Zero paper

Chess Programming

 * Chess on a GPGPU
 * GPU Chess Blog
 * ankan-ban/perft_gpu · GitHub » Perft
 * LCZero · GitHub » Leela Chess Zero
 * GitHub - StuartRiffle/Jaglavak: Corvid Chess Engine » Jaglavak
 * Zeta OpenCL Chess » Zeta

=References= Up one Level