direkt zum Inhalt springen

direkt zum Hauptnavigationsmenü

Sie sind hier

TU Berlin

Inhalt des Dokuments

Überblick über das System

Ein neuer Trend im High Performance Computing ist der Einsatz von Grafikkarten als leistungsstarke, günstige Co-Prozessoren der CPU. Man spricht hier vom sogenannten General-purpose computing on graphics processing units (GPGPU).

Der nVidia TESLA Cluster der TU-Berlin besitzt 16 Knoten mit je zwei Intel Xeon X5550 (2.66 GHz) Prozessoren und 24 GB RAM.

Die Knoten sind wie folgt partitioniert:

  • 14 Knoten sind mit je zwei nVidia TESLA C1060 Karten ausgestattet. Jede C1060 besitzt 240 Recheneinheiten, die parallel arbeiten. Damit stehen pro Knoten theoretische 1866 GFLOPS für Single Precision Arithemtik und 156 GFLOPS für Double Precision Arithmetik durch die GPUs bereit. Jede Karte ist mit 4096 MB RAM ausgestattet.
  • 2 Knoten sind mit je vier nVidia GTX 295 Karten ausgestattet. Damit stehen jedem Knoten theoretische 7152 GFLOPS für Single Precision Arithmetik und 596 GFLOPS für Double Precision Arithmetik durch die GPUs bereit. Jede Karte ist mit 1792 MB RAM ausgestattet.

Alle Knoten verfügen über eine 300 GB @ 10.000 RPM Festplatte und sind per Infiniband QDR vernetzt.

Nutzung

Der Zugriff auf den Cluster erfolgt über die üblichen Login-Knoten (cluster.math.tu-berlin.de) oder den headnode des GPU-Clusters (cluster-g.math.tu-berlin.de).

Die hardware von cluster-g ist identisch mit den GPU-Knoten. Insbesondere ist es dort  möglich, nach importieren des CUDA Moduls durch Aufruf von
$ gpurun program
ein Programm zu debugging-Zwecken testen, ohne das batch-System nutzen zu müssen.

Auf den Login-Knoten muss zunächst die Umgebung für die Entwicklung mit Cuda vorbereitet werden, indem

module add cuda-3.0

ausgeführt wird. Danach steht die CUDA Entwicklungsumgebung zur Verfügung.

Jobs werden über "qsub -l gpus=X ..."  übermittelt, wobei X angibt, wie viele GPU Prozessoren verlangt werden. Auf den C1060-Knoten stehen 2 GPU Prozessoren zur Verfügung, auf den GTX 295-Knoten stehen 8 GPU Prozessoren zur Verfügung.

Entwicklung

Zur Entwicklung können CUDA sowie OpenCL eingesetzt werden. Des weiteren existieren Bibliotheken, die aus C und Fortran heraus benutzt werden können. Die Portland-Compiler können Teile des Codes in die GPU auslagern.
Näheres ist unten unter dem jeweiligen Stichwort zu finden.

CUDA

Die CUDA API wurde von nVidia entwickelt und erlaubt eine sehr einfache Entwicklung für GPUs mittels C und Fortran.

Die CUDA Webseite enthält wichtige Dokumentation wie


Das CUDA SDK enthält Beispiele für die folgenden Domänen:

  • Parallel bitonic sort
  • Matrix multiplication
  • Matrix transpose
  • Performance profiling using timers
  • Parallel prefix sum (scan) of large arrays
  • Image convolution
  • 1D DWT using Haar wavelet
  • OpenGL and Direct3D graphics interoperation examples
  • CUDA BLAS and FFT library usage examples
  • CPU-GPU C- and C++-code integration
  • Binomial Option Pricing
  • Black-Scholes Option Pricing
  • Monte-Carlo Option Pricing
  • Parallel Mersenne Twister (random number generation)
  • Parallel Histogram
  • Image Denoising
  • Sobel Edge Detection Filter



Thrust ist eine highlevel C++ API im STL Stil, welche die Entwicklung mit CUDA nocheinmal stark vereinfacht.

OpenCL

OpenCL stellt eine plattformunabhängige Alternative zu CUDA dar, die sich noch in der Entwicklung befindet.
Eine erste Version der OpenCL-Entwicklertools von nVidia sind nach Laden des entsprechenden Modules verfügbar:

module add opencl-1.0

Erste Schritte mit CUDA

Die TESLA und GTX 295 besitzen eine hohe Zahl von Recheneinheiten, die in der Lage sind, Daten parallel zu verarbeiten. CUDA erlaubt dabei dem Programmierer, C-Funktionen zu schreiben, sogenannte Kernels, die parallel in N sogenannten CUDA threads augeführt werden. Um eine C-Funktion zu einem Kernel zu machen, genügt es zunächst einmal, diese Funktion mit dem __global__ Attribut zu annotieren. Der CUDA Compiler interpretiert dieses Attribut und sorgt für die Ausführung auf der GPU.

Da der gleiche Kernel vielfach parallel ausgeführt wird, dient eine threadID dazu, die Threads zu unterscheiden und den Zuständigkeitsbereich der einzelnen Threads zu definieren. Dies erlaubt, einen Eingabe-Vektor oder eine Eingabe-Matrix zu partitionieren und die Partitionen eindeutigen Threads zuzuordnen.

Die Eingabedaten werden in ein ein- oder zweidimensionales Raster (Grid) partitioniert. Jeder Block dieses Rasters wird durch ein- oder zweidimensionale Koordinaten, den sogenannte Block-Index, identifiziert und einem Prozessorkern der GPU zugeordnet. Die Eingabedaten eines Blocks werden wiederum in viele ein-, zwei- oder dreidimensionale Threadblöcke partitioniert, die je einem Thread entsprechen. Ein Block kann bis zu 512 Threads enthalten.

Das Speichermodell in CUDA sieht eine Unterscheidung zwischen Host-Speicher, also normalem RAM des Computers, und Device-Speicher, also Speicher der GPU, vor. Das Hauptprogramm hat die Aufgabe, Eingabedaten aus dem Host-Speicher in den Device-Speicher zu kopieren und Ergebnisse aus dem Device-Speicher zurück in den Hauptspeicher zu transferieren. Dazu stehen sehr einfache Kopieroperationen zur Verfügung. Auf weitere Speicherarten wie Konstanter-Speicher, Textur-Speicher, Thread-lokale und Thread-Block-lokale Speicher soll an dieser Stelle nicht weiter eingegangen werden. Sie werden im CUDA Tutorial genau erläutert.

Das folgende Beispiel addiert zwei Nx1 Vektoren. Dabei wird der Vektor in viele Blöcke zerlegt, die jeweils 256 Threads enthalten. Ein Thread hat lediglich die Aufgabe, die Zahlen der beiden Vektoren an einer bestimmten Koordinate zu addieren. Die Koordinate ergibt sich aus dem Block-Index und dem Thread-Index innerhalb des Blocks.

 

addvectors.cu

#include <cuda.h>
#include <stdio.h>
 
// Device code (executed on TESLA card, AKA "kernel")
__global__ void VecAdd(float* A, float* B, float* C, int N) {
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < N)
C[i] = A[i] + B[i];
}
 
// Host code
int main() {
int N = 10000;
size_t size = N * sizeof(float);
 
// Allocate input vectors h_A and h_B in host memory
float * h_A = (float *)malloc(size);
float * h_B = (float *)malloc(size);
float * h_C = (float *)malloc(size);
int i;
 
// Fill input vectors with content
for (i=0; i<N; ++i)
h_A[i] = h_B[i] = 1.0;
 
// Allocate vectors in device memory
float * d_A;
cudaMalloc((void**) &d_A, size);
float * d_B;
cudaMalloc((void**) &d_B, size);
float * d_C;
cudaMalloc((void**) &d_C, size);
 
// Copy vectors from host memory to device memory
cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
 
// Invoke kernel
int threadsPerBlock = 256;
int blocksPerGrid =
(N + threadsPerBlock - 1) / threadsPerBlock;
VecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
 
// Copy result from device memory to host memory
// -> h_C conains the result in host memory
cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
 
// Free device memory
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
 
// Print result
for (i=0; i<N; ++i)
printf("%d %f\n", i, h_C[i]);
 
// Free host memory
free(h_A);
free(h_B);
free(h_C);
return 0;
}

Das Programm wird mit der Endung .cu abgespeichert und über "nvcc -o addvectors addvectors.cu" kompiliert. Danach kann es über "gpurun ./addvectors" auf dem Frontend-Knoten gestartet werden. Der Wrapper 'gpurun' setzt dabei die nötigen Rechte der GPU.

Achtung: Bei direkter Ausführung ohne 'gpurun' läuft das Programm möglicherweise ohne Fehlermeldung durch. Berechnungen werden aber nicht ausgeführt!

Zur Übergabe des Jobs an das Queueing-System, kann über "qsub -l gpus=1 ./addvectors" erfolgen.

Für weitere Details zur Programmierung wird auf das ausführliche CUDA Programming Guide verwiesen.

Thrust

Mit der Thrust-Abstraktion lässt sich obiges Beispiel wie folgt vereinfachen:

addvectors.cu

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/transform.h>
#include <thrust/functional.h>
#include <thrust/copy.h>
#include <cstdlib>
 
int main(void)
{
int N = 10000;
thrust::host_vector<float> h_A(N);
thrust::host_vector<float> h_B(N);
thrust::host_vector<float> h_C(N);
for (int i=0; i<N; ++i)
h_A[i] = h_B[i] = 1.0;
 
// transfer data to the device
thrust::device_vector<float> d_A = h_A;
thrust::device_vector<float> d_B = h_B;
thrust::device_vector<float> d_C(N);
 
// apply + operator
thrust::transform(
d_A.begin(), d_A.end(), // input 1
d_B.begin(), // input 2
d_C.begin(), // output
thrust::plus<float>());
 
// transfer data back to host
h_C = d_C;
 
for (int i=0; i<N; ++i)
printf("%d %f\n", i, h_C[i]);
 
return 0;
}

Thrust bietet diverse Funktionen zum Suchen, Kopieren, Reduzieren, Sortieren und Transformieren von Vektoren.

cuBLAS und cuFFT

Unter

/afs/math/software/nvidia/cuda/2.3/opensuse11.1.64/sdk/C/src

findet sich ein Beispiel, wie man sehr einfach mit der cuBLAS Bibliothek für Lineare Algebra arbeiten kann. Diese Bibliothek versteckt die Nutzung der GPUs weitgehend transparent für den Programmierer. Die cuBLAS Referenz sowie eine spezielle Seite zum Thema Lineare Algebra von nVidia bieten weitere Details. CuBLAS wird sowohl für C++ als auch für Fortran 77 unterstützt.

Eine ähnliche Abstraktion für FFT findet man unter

/afs/math/software/nvidia/cuda/2.3/opensuse11.1.64/sdk/C/src/simpleCUFFT.

Die cuFFT Referenz bietet diverse Details zur Nutzung.

PGI GPU Accelerator

Die Portland Compiler bieten die Möglichkeit, Fragmente des Quellcodes zu annotieren, sodass diese automatisch an die GPU Karte ausgelagert werden.

Das folgende Beispiel illustriert die Einfachheit:
#pragma acc region
{
for( i = 0; i < n; ++i ) r[i] = a[i]*2.0f;
}
Details finden sich auf den PGI Webseiten.

Zur Nutzung muss mittels "module add pgi-10.0" die PGI Umgebung geladen werden. Im Anschluss kann das Programm mit "pgcc program.c -ta=nvidia" kompiliert werden. Für Details sei auf die Tutorials verwiesen.

Links

Zusatzinformationen / Extras

Direktzugang

Schnellnavigation zur Seite über Nummerneingabe