University:
Boston University
Course:
MET CS 201 | Introduction to Programming
Academic year:

2021
Views:

105

Pages:

20
Author:

otkopaokyd0

Advanced Algorithmic Techniques for GPUs Lecture 1: Introduction and Computational Thinking Course Objective • To master the most commonly used algorithm techniques and computational thinking skills needed for many-core GPU programming – Especially the simple ones! • In particular, to understand – Many-core hardware limitations and constraints – Desirable and undesirable computation patterns – Commonly used algorithm techniques to convert undesirable computation patterns into desirable ones Performance Advantage of GPUs • An enlarging peak performance advantage: – – Calculation: 1 TFLOPS vs. 100 GFLOPS Memory Bandwidth: 100-150 GB/s vs. 32-64 GB/s Courtesy: John Owens – GPU in every PC and workstation – massive volume and potential impact CPUs and GPUs have fundamentally different design philosophies. ALU ALU ALU ALU Control CPU GPU Cache DRAM DRAM UIUC/NCSA AC Cluster • 32 nodes – 4-GPU (GTX280, Tesla) nodes – GPUs donated by NVIDIA – Host boxes funded by NSF • Coulomb Summation: – 1.78 TFLOPS/node – 271x speedup vs. one Intel QX6700 CPU core w/ SSE UIUC/NCSA AC Cluster http://iacat.uiuc.edu/resources/cluster/ A partnership between NCSA and UIUC academic departments. EcoG - One of the Most Energy Efficient Supercomputers in the World • #3 of the Nov 2010 Green 500 list • 128 nodes • One Fermi GPU per node • 934 MFLOPS/Watt • 33.6 TFLOPS DP Linpack • Built by Illinois students and NVIDIA researchers GPU computing is catching on. Financial Analysis Scientific Simulation Engineering Simulation Data Intensive Analytics Medical Imaging Digital Audio Processing Digital Video Processing Computer Vision Biomedical Informatics Electronic Design Automation Statistical Modeling Ray Tracing Rendering Interactive Physics Numerical Methods • 280 submissions to GPU Computing Gems – 110 articles included in two volumes A Common GPU Usage Pattern • A desirable approach considered impractical – Due to excessive computational requirement – But demonstrated to achieve domain benefit – Convolution filtering (e.g. bilateral Gaussian filters), De Novo gene assembly, etc. • Use GPUs to accelerate the most time-consuming aspects of the approach – Kernels in CUDA or OpenCL – Refactor host code to better support kernels • Rethink the domain problem CUDA /OpenCL – Execution Model • Integrated host+device app C program – Serial or modestly parallel parts in host C code – Highly parallel parts in device SPMD kernel C code Serial Code (host) Parallel Kernel (device) KernelA<<< nBlk, nTid >>>(args); ... Serial Code (host) Parallel Kernel (device) KernelB<<< nBlk, nTid >>>(args); ... CUDA Devices and Threads • A compute device – – – – • • Is a coprocessor to the CPU or host Has its own DRAM (device memory) Runs many threads (work elements for OpenCL) in parallel Is typically a GPU but can also be another type of parallel processing device Data-parallel portions of an application are expressed as device kernels which run on many threads Differences between GPU and CPU threads – GPU threads are extremely lightweight • – Very little creation overhead GPU needs 1000s of threads for full efficiency • Multi-core CPU needs only a few Arrays of Parallel Threads • A CUDA kernel is executed by an array of threads – All threads run the same code (SPMD) – Each thread has an index that it uses to compute memory addresses and make control decisions threads 0 1 2 3 4 5 6 7 … float a = input[threadIdx]; float b = func(a); output[threadIdx] = b; … Thread Blocks: Scalable Cooperation • Divide monolithic thread array into multiple blocks – Threads within a block cooperate via shared memory, atomic operations and barrier synchronization – Threads in different blocks cannot cooperate Thread Block 1 Thread Block 0 threads 0 1 2 3 4 5 6 7 … float a = input[threadIdx]; float b = func(a); output[threadIdx] = b; … 0 1 2 3 4 5 6 Thread Block N - 1 7 … float a = input[threadIdx]; float b = func(a); output[threadIdx] = b; … 0 … 1 2 3 4 5 6 7 … float a = input[threadIdx]; float b = func(a); output[threadIdx] = b; … blockIdx and threadIdx Host • Each thread uses indices to decide what data to work on – – • blockIdx: 1D or 2D threadIdx: 1D, 2D, or 3D Simplifies memory addressing when processing multidimensional data – – – Image processing Solving PDEs on volumes … Device Grid 1 Kernel 1 Block (0, 0) Block (1, 0) Block (0, 1) Block (1, 1) Grid 2 Kernel 2 Block (1, 1) (0,0,1) (1,0,1) (2,0,1) (3,0,1) Thread Thread Thread Thread (0,0,0) (1,0,0) (2,0,0) (3,0,0) Thread Thread Thread Thread (0,1,0) (1,1,0) (2,1,0) (3,1,0) Courtesy: NDVIA Example: Vector Addition Kernel Device Code // Compute vector sum C = A+B // Each thread performs one pair-wise addition __global__ void vecAdd(float* A, float* B, float* C, int n) { int i = threadIdx.x + blockDim.x * blockIdx.x; if(i>>(d_A, d_B, d_C, n); } Example: Vector Addition Kernel // Compute vector sum C = A+B // Each thread performs one pair-wise addition __global__ void vecAdd(float* A, float* B, float* C, int n) { int i = threadIdx.x + blockDim.x * blockIdx.x; if(i>>(d_A, d_B, d_C, N); } Kernel execution in a nutshell __global__ __host__ void example() { int B = 128, P = ceil(n/B); vecAdd<<>>(n,a,x,y); } void saxpy(int n, float a, float *x, float *y) { int i = blockIdx.x * blockDim.x + threadIdx.x; if( i