Multi-Core Architectures and Programming · Tesla-Architektur Graphikkarten Allgemeine Prozessoren...

Friedrich-Alexander-Universität Erlangen-Nürnberg

Axel Jena, Jürgen Pröll 1

CUDA

Axel Jena, Jürgen Pröll

Multi-Core Architectures and Programming



Warum Tesla?� Traditionelle Graphikkarten

� Getrennte Prozessoren für Vertex- / Pixelberechnungen- Nachteil: es werden mehr Pixel- als Vertexberechnungen benötigt

�Keine gute Lastbalance

� Unterstützen Pipelining

� Tesla-Architektur Graphikkarten

� Allgemeine Prozessoren (frei programmierbar mit CUDA)

� SIMT (Single Instruction Multiple Thread)-Architektur

� Hohe Skalierbarkeit



Geschichte und Leistungsdaten� Erste Tesla-Graphikkarte im November 2006 (GeForce

8800) mit 128 Streaming-Prozessoren

� Nächste Generation der Tesla-Architektur: Fermi

� Leistungsdatenvergleich:

Intel Core i7 980 Nvidia Tesla C2050

Flops 107,55 G 515 G

Taktrate 3,3 GHz 600 MHz

Leistung 130 W 238 W

Kerne 6 448



Befehlsabarbeitung

Skalierbar

-128 Streaming Prozessoren

- Organisiert in 16 StreamingMultiprozessoren

- Aufgeteilt in 8 unabhängige Textur/Prozessor Clusters



Texture/Processor Cluster� Die TPC rechnen

unabhängig voneinander.

� Der „Geometry Controller“bildet die Berechnungen auf die SMs (StreamingMultiprocessors) ab.

� SMC

� Verbindung der SMs mit der Textur-Unit zur Textur Erzeugung und Filterung (z.B. Anisotropie Filterung)

� Steuert Lastbalance



Streaming Multiprocessor� Besteht aus:

� 8 Streaming Prozessoren (SP)

� 2 Special Function Units (SFU)

� Multithreaded Instruction Fetch and Issue Unit (MT issue)

� Instruction Cache (I cache)

� Read only Cache

� 16KB read/write Shared Memory

� (evtl. DP-Prozessoren (Double Precision) )

� Die SMs unterstützen massives Multithreading (bis zu 768 Threads pro SM, ohne scheduling Overhead, SIMT-Architektur)

� Das „Shared Memory“ hält die Daten für die parallelen Berechnungen.

� Die SFUs werden für komplexe Berechnungen verwendet

� Die SPs werden für die fundamentalen Berechnungen verwendet (Add, Mult F)



Streaming Multiprocessor: SIMT� Die Streaming Multiprozessoren bilden Thread-Gruppen

von 32 Threads pro Gruppe (= 1 Warp)

� Jeder SM kann 24 Warps behandeln.

� Die SM arbeiten mit der SIMT-Architektur (Single Instruktion Multiple Thread) ähnlich SIMD.

� Jeder Befehl wird mit Hilfe von Pipelining berarbeitet.

� Jede konditionelle Abzweigung (Pfade) im Code (IF´s) werden getrennt parallel ausgewertet. Sobald klar ist, welcher Pfad die Richtige ist werden die anderen verworfen.

� Die SMs arbeiten ansonsten mit einem Register basierendem Befehlssatz.(floating-point, integer, bit, flow control, memory load/store, texture operations)



Raster Operations Processor� Können direkt auf dem Speicher arbeiten.

� Zu jedem Speicherbaustein gehört ein eigenes ROP.

� Können Daten von den TPCs erhalten.

� Kümmern sich u. A. um:

� Farbüberblendungen.

� Antialiasing.

� Interpolation.

� Können nicht mit CUDA verwendet werden.



Speicher und Kommunikation� Der Datenbus für die Kommunikation („Interconenction

Network“) ist 384 Pins breit und in 6 Partitionen a 64 Pins gegliedert. Jede Partition verwaltet 1/6 des gesamten physikalischen Adressraums.

� Das „Interconenction Network“ basiert auf einer Hub-Unitdie die die Anfragen zwischen den verschiedenen Komponenten (PCIe-Bus, TPCs, DRAM) routet.

� Die ROPs sind direkt an den Speicher angebunden und belasten damit nicht das „Interconnection Network“.

� Eine Memory Management Unit setzt virtuelle zu physikalischen Adressen um und kümmert sich um Paging.

� Als Speicher fungieren GDDR3 Module die mit ca. 1GHz getaktet sind.



CUDA Programming Model� Serieller Code in einem Thread des Hosts (CPU)

� Paralleler Code verteilt auf viele Threads im Device (GPU)

� Kernel: vom Host aufgerufene Funktion, die auf dem Device von vielen Threads ausgeführt wird



Thread Hierarchy� Ein Kernel wird von einem Grid ausgeführt

� Ein Grid besteht aus in Blöcken gruppierten Threads

� Alle Threads des Grids:

� Führen den selben Code aus

� Sind durch Thread ID und Block ID unterscheidbar

� Jeder Block wird, in SIMT Warps geteilt, von einem StreamingMultiprozessor (SM) ausgeführt



Memory Hierarchy� Local Memory:

� jeder Thread hat Register, um temporäre Variablen zu speichern

� sind diese nicht ausreichend, hat er zusätzlich noch einen LocalMemory für größere Dateien

� Shared Memory:

� gemeinsamer Speicher für Threads desselben Blockes

� Global Memory:

� zur Kommunikation unter sequentiell ablaufenden Grids



Transparente Skalierbarkeit� Thread Blöcke sind voneinander unabhängig, weshalb sie

auch in beliebiger Reihenfolge ausgeführt werden können

� Programmierer müssen nicht auf Hardware achten, da sich ein Grid beliebig auf vorhandene Ressourcen verteilt



Kompilierung

� NVCC teilt Code in seriellen und parallelen Teil

� PTX Code (Parallel Thread Execution):

� Pseudoassembler für Grafikkarten

� Host und Device Code werden anschließend zu

einer ausführbaren Datei zusammen gelinkt



CUDA API: Extended C

� Minimale Erweiterungen zu C/C++

� Durch Wrapper auch andere Programmiersprachen nutzbar

� Language Extensions

� Function Type Qualifiers

� Variable Type Qualifiers

� Built-in Variablen

� Runtime Library

� Common Component

� Device Component

� Host Component



Function Type Qualifiers� __global__ void KernelFunc():

� wird vom Host aufgerufen und auf dem Device ausgeführt

� muss void sein

� __device__ float DeviceFunc():� wird vom Device aufgerufen und ausgeführt

� __host__ float HostFunc():� wird auf dem Host aufgerufen und ausgeführt (optional)

� Funktionen auf dem Device:� keine Rekursion

� keine statischen Variablen

� keine variable Argumentenliste

� KernelFunc<<<grid,block>>>(args);� grid: Dimension des Grids (1D oder 2D)

� block: Dimension eines Blocks (1D, 2D oder 3D)



Variable Type Qualifiers� __device__ int GlobalVar;

� im Global Memory

� bleibt über die gesamte Programmausführung bestehen

� sichtbar für alle Threads und für den Host über die Runtime Library

� __constant__ int ConstantVar;� im Constant Memory

� Lebensdauer und Verfügbarkeit wie bei __device__

� __shared__ int SharedVar;� im Shared Memory eines Thread Blocks

� an Lebensdauer des Thread Blocks gebunden

� ist nur für Threads innerhalb des Blocks sichtbar

� int LocalVar; (innerhalb einer Device-Funktion)� in einem Register oder bei großen Daten im Local Memory



Built-in Variablen

� Verfügbar in Device-Funktionen

� dim3 threadIdx;

� Thread-ID innerhalb des Blocks

� threadIdx.x, threadIdx.y, threadIdx.z

� dim3 blockIdx;

� Block-ID innerhalb des Grids

� dim3 blockDim;

� Größe des Blocks in Threads

� dim3 gridDim;

� Größe des Grids in Blöcken

� int warpSize;

� Größe eines Warps in Threads



Runtime Library� Common Runtime Component

� Vektor Typen (dim3), Texture Typen

� Teile der C Runtime Library

� Host Runtime Component� Device Management

� Memory ManagementcudaMalloc(), cudaFree(), cudaMemcpy()

� Texture Management

� Kompatibilität mit OpenGL und Direct3D

� Event Management

� Error Handling

� Device Runtime Component� Mathematische Funktionen

� Atomic Funktionen

� Funktionen zur Synchronisation- __syncthreads() wartet, bis alle Threads eines Blocks die Barriere erreichen



Libraries & Development Tools

� CUDA Libraries:

� Thrust: Standard Template Library (STL) für CUDA

� CUBLAS: “CUDA Basic Linear Algebra Subprograms“

� CUFFT: “CUDA Fast Fourier Transformation“

� Development Tools:

� CUDA-gdb (Debugger)

� CUDA Visual Profiler (Performance Profiling Tool)

� CUDA-MemCheck



Beispiel - CPU



Beispiel – GPU



Noch Fragen?

Multi-Core Architectures and Programming · Tesla-Architektur Graphikkarten Allgemeine Prozessoren...

Documents

Transcript of Multi-Core Architectures and Programming · Tesla-Architektur Graphikkarten Allgemeine Prozessoren...