Proseminar

GPU-Computing Cuda vs. OpenCL

SS 2013 Alexander Stepanov

Inhaltsverzeichnis • 1. Einführung: Warum GPU Computing?

• CPU vs. GPU • GPU Architektur

• 2. CUDA • Architektur • Beispiel Matrix Multiplikation

• 3. OpenCL • Architektur • Beispiel Matrix Multiplikation

• 4. CUDA vs. OpenCL • 5. Fazit

CPU vs. GPU • Top CPUs:

• Top GPUs:

HD 7870 GHz HD 7990 GTX Titan GTX 690

Takt 1.050 MHz 1.000 MHz 876 MHz 1.019 MHz

Shader 2.048 2 × 2.048 2.688 2 × 1.536

GFLOPs (SP|DP) 4.096 | 1024 2 × 4.096 | 2 × 1024 4.500 | 1.300 2 × 2.810 | 2 × 117

Speicher (GDDR5) 3 GB 2 × 3 GB 6 GB 2 × 2 GB

AMD FX-8350 Opteron 6386 SE i7-3970X Xeon E5-2690

Takt 4,2 GHz 3,2 GHz 4,0 GHz 3,8 GHz

Kerne | Threads 8 | 8 16 | 16 6 | 12 8 | 16

GFLOPs (SP|DP) 256 | 64 179 158 | 109 348 | 224

CPU vs. GPU

http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html

http://en.wikipedia.org/wiki/Flynn's_taxonomy

GPU Architektur (GK110 – GTX Titan)

http://www.nvidia.de/content/PDF/kepler/NVIDIA-Kepler-GK110-Architecture-Whitepaper.pdf

GPU Architektur (GK110 – GTX Titan)

• L1-Cache: 64 KB / SMX • Data-Cache: 48 KB /SMX • L2-Cache: 1.536 KB

• 196 SP Cores / SMX • 64 DPU / SMX • 32 SFU / SMX • 16 Texture Units / SMX • 4 Warp Scheduler / SMX

http://www.nvidia.de/content/PDF/kepler/NVIDIA-Kepler-GK110-Architecture-Whitepaper.pdf

CPU vs. GPU

• GPU • Massiv Parallelisiert • Wenig Cache • Langsam bei Verzweigungen

• CPU • Großer Befehlssatz • Branch Prediction

CUDA (Compute Unified Device Architecture) • 2007 von NVIDIA released • Wird nur von NVIDIA gepflegt und läuft nur auf deren Hardware • Aktuellste Version: 5.0 • Sprache: CUDA-C und PTX Assembler • Läuft auf Windows, Linux und Mac OS X

CUDA Ablauf: • CPU kopiert kompilierten Kernel-Code und Daten

auf die GPU. • GPU führt den Code in Warps (32 Threads) aus und

überträgt das Ergebnis zurück zum Host

Warp-Scheduler: • Partitioniert mehrere Blöcke zu Warps • Ein Warp führt eine Instruktion gleichzeitig aus (SIMD)

http://sidkashyap.wordpress.com/2013/05/02/cuda-faq/

CUDA // Kernel Code

__global__ void mmult ( const float *A, const float *B, float *C,

int width)

{

int i = blockIdx .x * blockDim .x + threadIdx .x;

int j = blockIdx .y * blockDim .y + threadIdx .y;

float val = 0;

for( int k = 0; k < width; k++ )

val += A[width *j+k] * B[width *k+i];

C[width *j+i] = val ;

} • __global__ zeichnet Kernelcode aus. Wird zur Laufzeit vom Nvcc (NVIDIA CUDA

Compiler) kompiliert.

• __device__ (global), __constant__ (constant), __share__ (shared) sind Schlüsselwörter für die Speicherhierarchie in CUDA

• Nachteil: Matrixgröße beschränkt auf Blockgröße => Kernel mehrere Zellen berechnen lassen => Mehr Blöcke zur Problemlösung beauftragen

http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html

CUDA // Host Code __host__ int main (int argc , char ** argv ) { // Matrizengrößen A (1024 x 512), B (512 x 2048) => C (512 x 512) size_t sizeA = 1024 * 512 * sizeof(float); size_t sizeB = 512 * 2048 * sizeof(float); size_t sizeC = 512 * 512 * sizeof(float); // Matrizen im Host anlegen float *h_A = (float *) malloc(sizeA); float *h_B = (float *) malloc(sizeB); float *h_C = (float *) malloc(sizeC);

// Matrizenspeicher auf der GPU (Device) reservieren float* d_A, d_B, d_C; cudaMalloc (( void **) &d_A , sizeA ); cudaMalloc (( void **) &d_B , sizeB ); cudaMalloc (( void **) &d_C , sizeC ); // Matrizen vom Host zur GPU kopieren cudaMemcpy(d_A , h_A , sizeA , cudaMemcpyHostToDevice); cudaMemcpy(d_B , h_B , sizeB , cudaMemcpyHostToDevice); // Kernel vorbereiten und ausführen dim3 threads(16 , 16) ; //Blockgröße dim3 grid(512 / threads.x, 512 / threads.y ); //Gridgröße mmult<<<grid , threads>>>(d_A , d_B , d_C , 512); // Ergebnis Matrix d_C von der GPU zum Host h_C kopieren cudaMemcpy (h_C , d_C , sizeC , cudaMemcpyDeviceToHost ); }

OpenCL (Open Computing Language) • Spezifikation Ende 2008 veröffentlich • Aktuelle Version: 1.2 • Wurde ursprünglich von Apple entwickelt, jetzt in den Händen der Kronos Group • Offener Standard (wie OpenGL) • Sprache: Kernel in OpenCL C und Host in C/C++ • Plattform- und Geräteunabhängige

• Gerätehersteller kümmern sich selbst um den OpenCL-Support in ihren Chips

OpenCL Ablauf (Ähnlich wie bei CUDA): • CPU kopiert kompilierten Kernel-Code und Daten auf die GPU. • GPU führt den Code in den Compute Units aus und überträgt das

Ergebnis zurück zum Host

http://www.khronos.org/assets/uploads/developers/library/overview/opencl_overview.pdf

OpenCL const char kernel_src [] = " __kernel void mmult ( __global const float *A, " " __global const float *B, " " __global float *C, " " int wA , int wB) " "{ " " int i = get_global_id (0); " " int j = get_global_id (1); " " " " float val = 0; " " for( int k = 0; k < wA; k++ ) " " val += A[wA*j+k] * B[wB*k+i]; " " C[wA*j+i] = val; " "} ";

• Kernelcode wird als String übergeben und dann zur Laufzeit kompiliert.

• Schlüsselwörter für Speicherhierarchie: __global (global), __constant (constant), __local (share) und __private (privat)

http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html

OpenCL // Matrizen erstellen float * h_A = new float[sizeA]; float * h_B = new float[sizeB]; float * h_C = new float[sizeC]; cl_mem d_A = clCreateBuffer ( context , CL_MEM_READ_ONLY , sizeA , 0, 0); cl_mem d_B = clCreateBuffer ( context , CL_MEM_READ_ONLY , sizeB , 0, 0); cl_mem d_C = clCreateBuffer ( context , CL_MEM_WRITE_ONLY , sizeC , 0, 0); // Speicher reservieren clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_A); clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_B); clSetKernelArg(kernel, 2, sizeof(cl_mem), &d_C); // work size berechnen size_t ws_global [] = {512, 512}; size_t ws_local [] = {16 , 8}; // 128 items per group // Eingabematrizen zur GPU kopieren, Kernel starten und Ergebnis zurückkopieren clEnqueueWriteBuffer(cmd_queue, d_A, CL_FALSE, 0, sizeA, h_A, 0, 0, 0); clEnqueueWriteBuffer(cmd_queue, d_B, CL_FALSE, 0, sizeB, h_B, 0, 0, 0); clEnqueueNDRangeKernel(cmd_queue, kernel, 2, 0, ws_global, ws_local, 0, 0, 0); clEnqueueReadBuffer(cmd_queue, d_C, CL_FALSE, 0, sizeC, h_C, 0, 0, 0); clFinish(cmd_queue); }

int main (int argc , const char * argv []) { // Matrizengrößen A (1024 x 512), B (512 x 2048) => C (512 x 512) size_t sizeA = 1024 * 512 * sizeof(float); size_t sizeB = 512 * 2048 * sizeof(float); size_t sizeC = 512 * 512 * sizeof(float); // Matrizen im Host anlegen float *h_A = (float *) malloc(sizeA); float *h_B = (float *) malloc(sizeB); float *h_C = (float *) malloc(sizeC); // Auf welcher Platform soll das laufen cl_uint num_platforms ; cl_platform_id platform ; cl_int err = clGetPlatformIDs (1, & platform , & num_platforms ); // Auf welchem Gerät soll das laufen cl_device_id device ; clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device , 0); // Erstelle context , command queue , program und kernel cl_context context = clCreateContext(0, 1, &device , 0, 0, &err); cl_command_queue cmd_queue = clCreateCommandQueue(context, device , 0, 0); cl_program program = clCreateProgramWithSource(context , 1, &kernel_src, 0, &err); clBuildProgram(program , 0, 0, 0, 0, 0); cl_kernel kernel = clCreateKernel (program, " mmult ", &err);

CUDA vs. OpenCL • Wegen CUDA nur auf NVIDIA GPUs vergleichbar

• OpenCL ist von vornerein benachteiligt -> Wird mit CUDA umgesetzt

• Vergleich ist somit wegen dem

„Performanceverlust“ von Interesse

CUDA OpenCL

Global Memory Global Memory

Constant Memory Constant Memory

Shared Memory Local Memory

Local Memory Private Memory

Thread Work-item

Thread-block Work-group http://www.evga.com/FORUMS/tm.aspx?m=91863&mpage=1

CUDA vs. OpenCL • Test aus „A Performance Comparison of CUDA and OpenCL” 2010 von Kamran Karimi, Neil G. Dickson und

Firas Hamze von D-Wave Systems Inc. in Canada

• Quanten-Spin-System Simulation mit AQUA (Adiabatic QUantum Algorthms)

• GPU: GTX 260 mit 192 Kernen Chiptakt: 576 MHz

• CUDA-Version: 2.3 OpenCL: 1.0

CUDA vs. OpenCL

Qubits Kernel Laufzeit Datentransfer Gesamtlaufzeit

8 1,12 0,94 649,05 Kb 1,51

16 1,23 1,24 1.633,32 Kb 1,38

32 1,19 1,4 3.553,44 Kb 1,26

48 1,45 1,32 8.210,22 Kb 1,5

72 1,63 1,58 15.338,77 Kb 1,68

96 1,17 1,36 33.124,49 Kb 1,22

128 1,12 1,34 49.541,04 Kb 1,16

• CUDA -> OpenCL: • Syntaktische Anpassung (siehe Terminologie) • keine Referenzen in OpenCL

CUDA vs. OpenCL • CUDA unterstützt spezifische Hardware-Features (z.b. Textur-Speicher)

• CUDA-Compiler besser optimiert

• => Selbst Kernel optimieren

z.b. Schleifen „entrollen“

“A Comprehensive Performance Comparison of CUDA and OpenCL“ - 2011 International Conference on Parallel Processing - Jianbin Fang, Ana Lucia Varbanescu and Henk Sips

Fazit

• Cuda: • Hardwarenäher (PTX Assembler) => Performanter • Von nur einem Hersteller betreut: schnellere Updates • Großer Umfang an Tools

• OpenCL: • Universeller