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
Top Related