Architecture Overview Introduction to CUDA Programming Andreas Moshovos

Architecture Overview Introduction to CUDA Programming Andreas Moshovos Winter 2009 Most slides/material from: UIUC course by Wen-Mei Hwu and David Kirk Real World Techonologies by David Kanter

System Architecture of a Typical PC / Intel

PCI-Express Programming Model • PCI device registers are mapped into the CPU’s physical address space – Accessed through loads/ stores (kernel mode) • Addresses assigned to the PCI devices at boot time – All devices listen for their addresses • That’s a reason why Windows XP cannot “see” 4 GB

PCI-E 1. x Architecture • Switched, point-to-point connection – Each card has a dedicated “link” to the central switch, no bus arbitration. – Packet switches messages form virtual channel – Prioritized packets for Qo. S • E. g. , real-time video streaming IO IO IO NB NB BUS: PCI or older PCI-E IO

PCI-E 1. x Architecture Contd. • Each link consists of one more lanes – Each lane is 1 -bit wide (4 wires, each 2 -wire pair can transmit 2. 5 Gb/s in one direction) • Upstream and downstream now simultaneous and symmetric • Differential signalling – Each Link can combine 1, 2, 4, 8, 12, 16 lanes- x 1, x 2, etc. – Each byte data is 8 b/10 b encoded into 10 bits with equal number of 1’s and 0’s; net data rate 2 Gb/s per lane each way. – Thus, the net data rates are 250 MB/s (x 1) 500 MB/s (x 2), 1 GB/s (x 4), 2 GB/s (x 8), 4 GB/s (x 16), each way

PCI-E 2. x and beyond Version Clock Speed Transfer Rate Overhead Data Rate 1. x 1. 25 Ghz 2. 5 GT/s 20% 250 MB/s 2. 0 2. 5 Ghz 5 GT/s 20% 500 MB/s 3. 0 4 Ghz 8 GT/s 0% 1 GB/s

Typical AMD System (for completeness) • AMD Hyper. Transport™ Technology bus replaces the Front-side Bus architecture • Hyper. Transport ™ similarities to PCIe: – Packet based, switching network – Dedicated links for both directions • Shown in 4 socket configuraton, 8 GB/sec per link • Northbridge/Hyper. Transport ™ is on die • Glueless logic – to DDR, DDR 2 memory – PCI-X/PCIe bridges (usually implemented in Southbridge)

A Typical Motherboad (m. ATX form factor)

CUDA Refresher • Grids of Blocks • Blocks of Threads Why? Realities of integrated circuits: need to cluster computation and storage to achieve high speeds

Thread Blocks Refresher • Programmer declares (Thread) Block: – Block size 1 to 512 concurrent threads – Block shape 1 D, 2 D, or 3 D – Block dimensions in threads • • All threads in a Block execute the same thread program Threads have thread id numbers within Block Threads share data and synchronize while doing their share of the work Thread program uses thread id to select work and address shared data Thread Id #: 0123… m Thread program

Architecture Goals • • Use multithreading to hide DRAM latency Support fine-grain parallel processing Virtualize the processors to achieve scalability Simplify programming. Develop program for one thread • Conventional Processors – Latency optimized – ILP – Caches 99% hit rate • GPU – Caches 90% or less. Not a good option – Throughput optimized – ILP + TLP

GT 200 Architecture Overview atomic

Terminology • SPA – Streaming Processor Array • TPC – Texture Processor Cluster • 3 SM + TEX • SM – Streaming Multiprocessor (8 SP) – Multi-threaded processor core – Fundamental processing unit for CUDA thread block • SP – Streaming Processor – Scalar ALU for a single CUDA thread

Thread Processing Cluster SM TEX SM SM

Stream Multiprocessor Overview • Streaming Multiprocessor (SM) – 8 Streaming Processors (SP) – 2 Super Function Units (SFU) – 1 Double-FP Unit (DPU) • Multi-threaded instruction dispatch – 1 to 512 threads active – Shared instruction fetch per 32 threads – Cover latency of texture/memory loads • • • 20+ GFLOPS 16 KB shared memory DRAM texture and memory access Streaming Multiprocessor Instruction L 1 Data L 1 Instruction Fetch/Dispatch Shared Memory SP SP SFU SP SP DPU

Thread Life • Grid is launched on the SPA Host • Thread Blocks are serially distributed to all the SM’s – Potentially >1 Thread Block per SM • • Grid 1 Kernel 1 Each SM launches Warps of Threads – 2 levels of parallelism • Device SM schedules and executes Warps that are ready to run As Warps and Thread Blocks complete, resources are freed – SPA can distribute more Thread Blocks Block (0, 0) Block (1, 0) Block (2, 0) Block (0, 1) Block (1, 1) Block (2, 1) Grid 2 Kernel 2 Block (1, 1) Thread Thread (0, 0) (1, 0) (2, 0) (3, 0) (4, 0) Thread Thread (0, 1) (1, 1) (2, 1) (3, 1) (4, 1) Thread Thread (0, 2) (1, 2) (2, 2) (3, 2) (4, 2)

Cooperative Thread Array • Break Blocks into warps • Allocate Resources – Registers, Shared Mem, Barriers • Then allocate for execution

Stream Multiprocessors Execute Blocks • Threads are assigned to SMs in Block granularity – Up to 8 Blocks to each SM as resource allows – SM in G 200 can take up to 1 K threads • • • Could be 256 (threads/block) * 4 blocks Or 128 (threads/block) * 8 blocks, etc. Threads run concurrently – SM assigns/maintains thread id #s – SM manages/schedules thread execution t 0 t 1 t 2 … tm SM 0 MT IU SP Blocks Shared Memory TF Texture L 1 L 2 Memory

Thread Scheduling and Execution • Each Thread Blocks is divided in 32 -thread Warps – This is an implementation decision, not part of the CUDA programming model • Warp: primitive scheduling unit Block 1 Warps … …Block 2 Warps t 0 t 1 t 2 … t 31 … … Streaming Multiprocessor Instruction L 1 Data L 1 Instruction Fetch/Dispatch • All threads in warp: – same instruction – control flow causes some to become inactive Shared Memory SP SP SFU SP SP DPU

Warp Scheduling • – Warps whose next instruction has its operands ready for consumption are eligible for execution – Eligible Warps are selected for execution on a prioritized scheduling policy – All threads in a Warp execute the same instruction when selected SM multithreaded Warp scheduler time warp 8 instruction 11 warp 1 instruction 42 warp 3 instruction 95. . . warp 8 instruction 12 warp 3 instruction 96 SM hardware implements zerooverhead Warp scheduling • 4 clock cycles needed to dispatch the same instruction for all threads in a Warp in G 200

How many warps are there? • If 3 blocks are assigned to an SM and each Block has 256 threads, how many Warps are there in an SM? – Each Block is divided into 256/32 = 8 Warps – There are 8 * 3 = 24 Warps – At any point in time, only one of the 24 Warps will be selected for instruction fetch and execution.

Warp Scheduling: Hiding Thread stalls

Warp Scheduling Ramifications • If one global memory access is needed for every 4 instructions • A minimal of 13 Warps are needed to fully tolerate a 200 -cycle memory latency • Why? – Need to hide 200 cycles every four instructions – Every Warp occupies 4 cycles during which the same instruction executes – Every 4 insts a thread stalls – Every 16 cycles a thread stalls – 200/16 =12. 5 or at least 13 warps

SM Instruction Buffer – Warp Scheduling • Fetch one warp instruction/cycle – from instruction L 1 cache – into any instruction buffer slot • Issue one “ready-to-go” warp instruction/cycle – from any warp - instruction buffer slot – operand scoreboarding used to prevent hazards • Issue selection based on roundrobin/age of warp • SM broadcasts the same instruction to 32 Threads of a Warp I$ L 1 Multithreaded Instruction Buffer R F C$ L 1 Shared Mem Operand Select MAD SFU

Scoreboarding • All register operands of all instructions in the Instruction Buffer are scoreboarded – Status becomes ready after the needed values are deposited – prevents hazards – cleared instructions are eligible for issue • Decoupled Memory/Processor pipelines – any thread can continue to issue instructions until scoreboarding prevents issue – allows Memory/Processor ops to proceed in shadow of Memory/Processor ops

Granularity Considerations • For Matrix Multiplication, should I use 8 X 8, 16 X 16 or 32 X 32 tiles? – For 8 X 8, we have 64 threads per Block. Since each SM can take up to 1024 threads, it can take up to 16 Blocks. However, each SM can only take up to 8 Blocks, only 512 threads will go into each SM – For 16 X 16, we have 256 threads per Block. Since each SM can take up to 1024 threads, it can take up to 4 Blocks and achieve full capacity unless other resource considerations overrule. – For 32 X 32, we have 1024 threads per Block. Not even one can fit into an SM.

Stream Multiprocessor Detail 64 entry

Scalar Units • • • 32 bit ALU and Multiply-Add IEEE Single-Precision Floating-Point Integer Latency is 4 cycles FP: Na. N, Denormals become signed 0. Round to nearest even

Special Function Units • Transcendental function evaluation and perpixel attribute interpolation • Function evaluator: – rcp, rsqrt, log 2, exp 2, sin, cos approximations – Uses quadratic interpolation based on Enhanced Minimax Approximation – 1 scalar result per cycle • Latency is 16 cycles • Some are synthesized: 32 cycles or so

Memory System Goals • • • High-Bandwidth As much parallelism as possible wide. 512 pins in G 200 / Many DRAM chips fast signalling. max data rate per pin. maximize utilization – Multiple bins of memory requests – Coalesce requests to get as wide as possible – Goal to use every cycle to transfer from/to memory • Compression: lossless and lossy • Caches where it makes sense. Small

DRAM considerations • multiple banks per chip – 4 -8 typical • 2^N rows. – 16 K typical • 2^M cols – 8 K typical • Timing contraints – 10~ cycles for row – 4 cycles within row • DDR – 1 Ghz --> 2 Gbit/pin – 32 -bit --> 8 bytes clock • GPU to memory: many traffic generators – no correlation if greedy scheduling – separate heaps / coalesce accesses • Longer latency

Parallelism in the Memory System • Thread Local Memory: – – • Shared Memory – • Private per thread Auto variables, register spill Shared Memory: – Block per-thread per-Block Shared by threads of the same block Inter-thread communication Global Memory: per-application – – Shared by all threads Inter-Grid communication Grid 0. . . Global Memory Grid 1. . . Sequential Grids in Time

SM Memory Architecture • Threads in a Block share data & results – In Memory and Shared Memory – Synchronize at barrier instruction • Per-Block Shared Memory Allocation – Keeps data close to processor – Minimize trips to global Memory – SM Shared Memory dynamically allocated to Blocks, one of the limiting resources

SM Register File • Register File (RF) – 64 KB – 16 K 32 -bit registers – Provides 4 operands/clock • TEX pipe can also read/write RF – 3 SMs share 1 TEX • Load/Store pipe can also read/write RF I$ L 1 Multithreaded Instruction Buffer R F C$ L 1 Shared Mem Operand Select MAD SFU

Programmer’s View of Register File • There are 16 K registers in each SM in G 200 – This is an implementation decision, not part of CUDA – Registers are dynamically partitioned across all Blocks assigned to the SM – Once assigned to a Block, the register is NOT accessible by threads in other Blocks – Each thread in the same Block only access registers assigned to itself 4 blocks 3 blocks

Register Use Implications Example • Matrix Multiplication • If each Block has 16 X 16 threads and each thread uses 10 registers, how many threads can run on each SM? – Each Block requires 10*16*16 = 2560 registers – 16384 = 6* 2560 + change – So, six blocks can run on an SM as far as registers are concerned • How about if each thread increases the use of registers by 1? – Each Block now requires 11*256 = 2816 registers – 16384 < 2816 *6 – Only five Blocks can run on an SM, 5/6 reduction of parallelism

Dynamic Partitioning • Dynamic partitioning gives more flexibility to compilers/programmers – One can run a smaller number of threads that require many registers each or a large number of threads that require few registers each • This allows for finer grain threading than traditional CPU threading models. – The compiler can tradeoff between instruction-level parallelism and thread level parallelism

Within or Across Thread Parallelism (ILP vs. TLP) • Assume: – – – kernel: 256 -thread Blocks 4 independent instructions for each global memory load, thread: 21 registers global loads: 200 cycles 6 Blocks can run on each SM • If a Compiler can use one more register to change the dependence pattern so that 8 independent instructions exist for each global memory load – Only three can run on each SM – However, one only needs 200/(8*4) = 7 Warps to tolerate the memory latency – Two Blocks have 16 Warps. – Conclusion: could be better

Constants • Immediate address constants • Indexed address constants • Constants stored in DRAM, and cached on chip – L 1 per SM – 64 KB total in DRAM • A constant value can be broadcast to all threads in a Warp – Extremely efficient way of accessing a value that is common for all threads in a Block! I$ L 1 Multithreaded Instruction Buffer R F C$ L 1 Shared Mem Operand Select MAD SFU

Shared Memory • Each SM has 16 KB of Shared Memory I$ L 1 – 16 banks of 32 bit words • CUDA uses Shared Memory as shared storage visible to all threads in a thread block – read and write access • Not used explicitly for pixel shader programs – we dislike pixels talking to each other • • • Key Performance Enhancement Move data in Shared memory Operate in there Multithreaded Instruction Buffer R F C$ L 1 Shared Mem Operand Select MAD SFU

Parallel Memory Architecture • In a parallel machine, many threads access memory – Therefore, memory is divided into banks – Essential to achieve high bandwidth • Each bank can service one address per cycle – A memory can service as many simultaneous accesses as it has banks • Multiple simultaneous accesses to a bank result in a bank conflict – Conflicting accesses are serialized Bank 0 Bank 1 Bank 2 Bank 3 Bank 4 Bank 5 Bank 6 Bank 7 Bank 15

Bank Addressing Examples • No Bank Conflicts – Linear addressing stride == 1 • No Bank Conflicts – Random 1: 1 Permutation Thread 0 Thread 1 Thread 2 Thread 3 Thread 4 Thread 5 Thread 6 Thread 7 Bank 0 Bank 1 Bank 2 Bank 3 Bank 4 Bank 5 Bank 6 Bank 7 Thread 15 Bank 15

Bank Addressing Examples • 2 -way Bank Conflicts – Linear addressing stride == 2 Thread 0 Thread 1 Thread 2 Thread 3 Thread 4 Thread 8 Thread 9 Thread 10 Thread 11 • 8 -way Bank Conflicts – Linear addressing stride == 8 x 8 Bank 0 Bank 1 Bank 2 Bank 3 Bank 4 Bank 5 Bank 6 Bank 7 Thread 0 Thread 1 Thread 2 Thread 3 Thread 4 Thread 5 Thread 6 Thread 7 Bank 15 Thread 15 Bank 0 Bank 1 Bank 2 x 8 Bank 7 Bank 8 Bank 9 Bank 15

How addresses map to banks on G 80 • Each bank has a bandwidth of 32 bits per clock cycle • Successive 32 -bit words are assigned to successive banks • G 80 has 16 banks – So bank = address % 16 – Same as the size of a half-warp • No bank conflicts between different half-warps, only within a single half-warp • G 200? Probably the same – Will find out

Shared memory bank conflicts • Shared memory is as fast as registers if there are no bank conflicts • The fast case: – If all threads of a half-warp access different banks, there is no bank conflict – If all threads of a half-warp access the identical address, there is no bank conflict (broadcast) • The slow case: – Bank Conflict: multiple threads in the same half-warp access the same bank – Must serialize the accesses – Cost = max # of simultaneous accesses to a single bank

Linear Addressing • Given: __shared__ float shared[256]; float foo = shared[base. Index + s * thread. Idx. x]; s=1 Thread 0 Thread 1 Bank 0 Bank 1 Thread 2 Thread 3 Bank 2 Bank 3 Thread 4 Bank 4 Thread 5 Thread 6 Bank 5 Bank 6 Thread 7 Bank 7 Thread 15 Bank 15 s=3 • This is only bank-conflict-free if s shares no common factors with the number of banks – 16 on G 200, so s must be odd Thread 0 Thread 1 Bank 0 Bank 1 Thread 2 Thread 3 Bank 2 Bank 3 Thread 4 Bank 4 Thread 5 Thread 6 Bank 5 Bank 6 Thread 7 Bank 7 Thread 15 Bank 15

Data types and bank conflicts • This has no conflicts if type of shared is 32 -bits: foo = shared[base. Index + thread. Idx. x] • But not if the data type is smaller – 4 -way bank conflicts: __shared__ char shared[]; foo = shared[base. Index + thread. Idx. x]; – 2 -way bank conflicts: __shared__ short shared[]; foo = shared[base. Index + thread. Idx. x]; Thread 0 Thread 1 Thread 2 Thread 3 Thread 4 Thread 5 Thread 6 Thread 7 Bank 0 Bank 1 Bank 2 Bank 3 Bank 4 Bank 5 Bank 6 Bank 7 Thread 15 Bank 15

Structs and Bank Conflicts • Struct assignments compile into as many memory accesses as there are struct members: struct vector { float x, y, z; }; struct my. Type { float f; int c; }; __shared__ struct vectors[64]; __shared__ struct my. Types[64]; • This has no bank conflicts for vector; struct size is 3 words – 3 accesses per thread, contiguous banks (no common factor with 16) struct vector v = vectors[base. Index + thread. Idx. x]; • This has 2 -way bank conflicts for my Type; (2 accesses per thread) struct my. Type m = my. Types[base. Index + thread. Idx. x]; Thread 0 Thread 1 Thread 2 Thread 3 Thread 4 Thread 5 Thread 6 Thread 7 Bank 0 Bank 1 Bank 2 Bank 3 Bank 4 Bank 5 Bank 6 Bank 7 Thread 15 Bank 15

Common Array Bank Conflict Patterns 1 D • Each thread loads 2 elements into shared mem: – 2 -way-interleaved loads result in 2 -way bank conflicts: int tid = thread. Idx. x; shared[2*tid] = global[2*tid]; shared[2*tid+1] = global[2*tid+1]; • This makes sense for traditional CPU threads, locality in cache line usage and reduced sharing traffice. – Not in shared memory usage where there is no cache line effects but banking effects Thread 0 Bank 0 Thread 1 Bank 1 Thread 2 Bank 2 Thread 3 Bank 3 Thread 4 Bank 5 Bank 6 Bank 7 Thread 8 Thread 9 Thread 10 Thread 11 Bank 15

A Better Array Access Pattern • Each thread loads one element in every consecutive group of block. Dim elements. shared[tid] = global[tid]; shared[tid + block. Dim. x] = global[tid + block. Dim. x]; Thread 0 Bank 0 Thread 1 Bank 1 Thread 2 Bank 2 Thread 3 Bank 3 Thread 4 Bank 4 Thread 5 Bank 5 Thread 6 Bank 6 Thread 7 Bank 7 Thread 15 Bank 15

Vector Reduction with Bank Conflicts Array elements 0 1 2 3 4 5 6 7 8 9 10 11

No Bank Conflicts 0 1 2 3 … 13 14 15 16 17 18 19

Common Bank Conflict Patterns (2 D) • Operating on 2 D array of floats in shared memory Bank Indices without Padding 0 1 2 3 4 5 6 7 15 0 1 2 3 4 5 6 7 15 – e. g. , image processing • Example: 16 x 16 block – Each thread processes a row – So threads in a block access the elements in each column simultaneously (example: row 1 in purple) – 16 -way bank conflicts: rows all start at bank 0 • Solution 1) pad the rows – Add one float to the end of each row • Solution 2) transpose before processing – Suffer bank conflicts during transpose – But possibly save them later 0 1 2 3 4 5 6 7 15 Bank Indices with Padding 1 2 3 4 5 6 7 15 0 2 3 4 5 6 7 8 0 1 3 4 5 6 7 8 9 1 2 4 5 6 7 8 9 10 2 3 5 6 7 8 9 10 11 3 4 6 7 8 9 10 11 12 4 5 7 8 9 10 11 12 13 5 6 8 9 10 11 12 13 14 7 8 15 0 1 2 3 4 5 6 14 15

Load/Store (Memory read/write) • Use LD to hide LD latency (non-dependent LD ops only) Clustering/Batching – Use same thread to help hide own latency • Instead of: – – LD 0 (long latency) Dependent MATH 0 LD 1 (long latency) Dependent MATH 1 • Do: – – LD 0 (long latency) LD 1 (long latency - hidden) MATH 0 MATH 1 • Compiler handles this! – But, you must have enough non-dependent LDs and Math