CUDA Parallel Programming Tutorial
Richard Membarth
[Link]@[Link]
Hardware-Software-Co-Design
University of Erlangen-Nuremberg
19.03.2009
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
1
Outline
Tasks for CUDA
CUDA programming model
Getting started
Example codes
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
2
Tasks for CUDA
Provide ability to run code on GPU
Manage resources
Partition data to t on cores
Schedule blocks to cores
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
3
Data Partitioning
Partition data in smaller
blocks that can be processed
by one core
Up to 512 threads in one
block
All blocks dene the grid
All blocks execute same
program (kernel)
Independent blocks
Only ONE kernel at a time
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
4
Memory Hierarchy
Memory types (fastest memory
rst):
Registers
Shared memory
Device memory (texture,
constant, local, global)
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
5
Tesla Architecture
30 cores, 240 ALUs (1 mul-add)
(1 mul-add + 1 mul): 240 * (2+1) * 1.3 GHz = 936 GFLOPS
4.0 GB GDDR3, 102 GB/s Mem BW, 4GB/s PCIe BW to CPU
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
6
CUDA: Extended C
Function qualiers
Variable qualiers
Built-in keywords
Intrinsics
Function calls
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
7
Function Qualiers
Functions: device , global , host
__global__ void filter(int *in, int *out) {
...
}
Default: host
No function pointers
No recursion
No static variables
No variable number of arguments
No return value
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
8
Variable Qualiers
Variables: device , constant , shared
__constant__ float matrix[10] = {1.0f, ...};
__shared__ int [32][2];
Default: Variables reside in registers
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
9
Built-In Variables
Available inside of kernel code
Thread index within current block:
threadIdx.x , threadIdx.y ,
threadIdx.z
Block index within grid:
blockIdx.x , blockIdx.y
Dimension of grid, block:
gridDim.x , gridDim.y
blockDim.x , blockDim.y ,
blockDim.z
Warp size: warpSize
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
10
Intrinsics
void __syncthreads();
Synchronizes in all thread of current block
Use in conditional code may lead to deadlocks
Intrinsics for most mathematical functions exists, e.g.
__sinf(x), __cosf(x), __expf(x), ...
Texture functions
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
11
Function Calls
Launch parameters:
Grid dimension (up to 2D)
Block dimension (up to 3D)
Optional: stream ID
Optional: shared memory size
kernel<<<grid, block, stream, shared_mem>>>();
__global__ void filter(int *in, int *out);
...
dim3 grid(16, 16);
dim3 block(16, 16);
filter <<< grid, block, 0, 0 >>> (in, out);
filter <<< grid, block >>> (in, out);
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
12
Getting Started
Compiler path
Sample Makele
Debugging
Memory management
Time measurement
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
13
Compiler Path
gcc/g++ compiler for host
code
nvcc compiler for device code
gcc/g++ for linking
icc/icpc works as well
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
14
Simple Project Makele
Use different les for host and device code
Compile device/host code with nvcc
Compile additional code with gcc
Adjust Makele from SDK:
# Add source files here
EXECUTABLE := vector_add
# CUDA source files (compiled with cudacc)
CUFILES := vector_add_host.cu
# CUDA dependency files
CU_DEPS := \
vector_add_device.cu \
defines.h
# C/C++ source files (compiled with gcc / c++)
CCFILES := \
vector_add_cpu.cpp
#set directory for [Link]
CUDA_SDK_PATH ?= /opt/cuda/sdk
ROOTDIR := $(CUDA_SDK_PATH )/projects
ROOTBINDIR := bin
ROOTOBJDIR := obj
include $(CUDA_SDK_PATH )/common/[Link]
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
15
Building the Program
Makele offers different options:
Production mode: make
Debug mode: make dbg=1
Emulation mode: make emu=1
Debug+Emulation mode: make dbg=1 emu=1
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
16
Debugging
SDK offers wrappers for function calls:
For CUDA function calls: cutilSafeCall(function);
For kernel launches (calls internally cudaThreadSynchronize()):
cutilCheckMsg(function);
For SDK functions: cutilCheckError(function);
Additional tools (recommended):
CudaVisualProler
valgrind in emulation mode only, there is no MMU on the GPU!
gdb in emulation mode: #ifdef __DEVICE_EMULATION__
real (!) gdb support, for GNU Linux unfortunately 32bit only :(
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
17
Memory Management
Host manages GPU memory
cudaMalloc(void **pointer, size_t size);
cudaMemset(void *pointer, int value, size_t count);
cudaFree(void *pointer);
Memcopy for GPU:
cudaMemcpy(void *dst, void *src, size_t size, cudaMemcpyKind
direction
cudaMemcpyKind:
cudaMemcpyHostToDevice
cudaMemcpyDeviceToHost
cudaMemcpyDeviceToDevice
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
18
Time Measurement
Initialization biases execution time
Dont measure rst kernel launch!
SDK provides timer:
int timer=0;
cutCreateTimer (&timer);
cutStartTimer (timer);
...
cutStopTimer (timer);
cutGetTimerValue (timer);
cutDeleteTimer (timer);
Use events for asynchronous functions:
cudaEvent_t start_event, stop_event ;
cutilSafeCall (cudaEventCreate (&start_event ));
cutilSafeCall (cudaEventCreate (&stop_event ));
cudaEventRecord (start_event, 0); // record in stream-0, to ensure that all
previous CUDA calls have completed
...
cudaEventRecord (stop_event, 0);
cudaEventSynchronize (stop_event ); // block until the event is actually
recorded
cudaEventElapsedTime (&time_memcpy, start_event, stop_event );
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
19
Example
Vector addition:
CPU Implementation
Host code
Device code
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
20
Vector Addition - CPU Implementation
void vector_add (float *iA, float *iB, float* oC, int width) {
int i;
for (i=0; i<width; i++) {
oC[i] = iA[i] + iB[i];
}
}
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
21
Vector Addition - GPU Initialization
// include CUDA and SDK headers - CUDA 2.1
#include <cutil_inline.h>
// include CUDA and SDK headers - CUDA 2.0
#include <cuda.h>
#include <cutil.h>
// include kernels
#include "vector_add_kernel.cu"
int main( int argc, char** argv) {
int dev;
// CUDA 2.1
dev = cutGetMaxGflopsDeviceId();
cudaSetDevice(dev);
// CUDA 2.0
CUT_DEVICE_INIT(argc, argv);
}
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
22
Vector Addition - Memory Management
// allocate device memory
int *device_idata_A, *device_idata_B, *device_odata_C;
cudaMalloc ((void**) &device_idata_A, mem_size);
cudaMalloc ((void**) &device_idata_B, mem_size);
cudaMalloc ((void**) &device_odata_C, mem_size);
// copy host memory to device
cudaMemcpy (device_idata_A, host_idata_A, mem_size,
cudaMemcpyHostToDevice);
cudaMemcpy (device_idata_B, host_idata_B, mem_size,
cudaMemcpyHostToDevice);
...
// copy result from device to host
cudaMemcpy (host_odata_C, device_odata_C, mem_size,
cudaMemcpyDeviceToHost);
// free memory
cudaFree(device_idata_A);
cudaFree(device_idata_B);
cudaFree(device_odata_C);
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
23
Vector Addition - Launch Kernel
// setup execution parameters
dim3 grid(1, 1);
dim3 threads(num_elements, 1);
// execute the kernel
vec_add<<< grid, threads >>>(device_idata_A, device_idata_B,
device_odata_C);
cudaThreadSynchronize();
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
24
Vector Addition - Kernel Function
__global__ void vector_add (float *iA, float *iB, float* oC) {
int idx = threadIdx.x + blockDim.x * blockId.x ;
oC[idx] = iA[idx] + iB[idx];
}
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
25
Questions?
Krakow, Pontical Residency
Courtesy of Robert Grimm
Friedrich-Alexander University of Erlangen-Nuremberg
Richard Membarth
26