Parallel Thread Execution

The programming guide to using PTX (Parallel Thread Execution) and ISA (Instruction Set Architecture). This document describes PTX, a low-level parallel thread execution virtual machine and instruction set architecture (ISA). PTX exposes the GPU as a data-parallel computing device . Driven by the insatiable market demand for real-time, high-definition 3D graphics, the programmable GPU has evolved into a highly parallel, multithreaded, many-core processor with tremendous computational horsepower and very high memory bandwidth. The GPU is especially well-suited to address problems that can be expressed as data-parallel computations - the same program is executed on many data elements in parallel - with high arithmetic intensity - the ratio of arithmetic operations to memory operations. Because the same program is executed for each data element, there is a lower requirement for sophisticated flow control; and because it is executed on many data elements and has high arithmetic intensity, the memory access latency can be hidden with calculations instead of big data caches. Data-parallel processing maps data elements to parallel processing threads. Many applications that process large data sets can use a data-parallel programming model to speed up the computations. In 3D rendering large sets of pixels and vertices are mapped to parallel threads. Similarly, image and media processing applications such as post-processing of rendered images, video encoding and decoding, image scaling, stereo vision, and pattern recognition can map image blocks and pixels to parallel processing threads. In fact, many algorithms outside the field of image rendering and processing are accelerated by data-parallel processing, from general signal processing or physics simulation to computational finance or computational biology. PTX defines a virtual machine and ISA for general purpose parallel thread execution. PTX programs are translated at install time to the target hardware instruction set. The PTX-to-GPU translator and driver enable NVIDIA GPUs to be used as programmable parallel computers. PTX provides a stable programming model and instruction set for general purpose parallel programming. It is designed to be efficient on NVIDIA GPUs supporting the computation features defined by the NVIDIA Tesla architecture. High level language compilers for languages such as CUDA and C/C++ generate PTX instructions, which are optimized for and translated to native target-architecture instructions. Provide a common source-level ISA for optimizing code generators and translators, which map PTX to specific target machines. Provide a scalable programming model that spans GPU sizes from a single unit to many parallel units. Adds support for .phase type:: qualifier for mbarrier.test wait and mbarrier.try wait instructions. Adds support for reportPredicate and reportValue operands for mbarrier.test wait and mbarrier.try wait instructions. Adds support for .layout qualifier for mbarrier instructions and support for mbarrier.check layout instruction. Adds support for .sem and .scope qualifiers for cp.async.bulk , cp.reduce.async.bulk , multimem.cp.async.bulk and multimem.cp.reduce.async.bulk instructions. PTX Machine Model gives an overview of the PTX virtual machine model. State Spaces, Types, and Variables describes state spaces, types, and variable declarations. Abstracting the ABI describes the function and call syntax, calling convention, and PTX support for abstracting the Application Binary Interface (ABI) . Release Notes provides release notes for PTX ISA versions 2.x and beyond. Threads within a CTA can communicate with each other. To coordinate the communication of the threads within the CTA, one can specify synchronization points where threads wait until all threads in the CTA have arrived. Each thread has a unique thread identifier within the CTA. Programs use a data parallel decomposition to partition inputs, work, and results across the thre