More on GPU Programming Ghufran Baig https devblogs

More on GPU Programming Ghufran Baig https: //devblogs. nvidia. com/how-overlap-data-transfers-cuda-cc/ http: //on-demand. gputechconf. com/gtc/2014/presentations/S 4158 -cuda-streams-best-practices-common-pitfalls. pdf https: //www. nvidia. com/docs/IO/116711/sc 11 -cuda-c-basics. pdf

Heterogeneous Computing #include <iostream> #include <algorithm> using namespace std; #define N 1024 #define RADIUS 3 #define BLOCK_SIZE 16 __global__ void stencil_1 d(int *in, int *out) { __shared__ int temp[BLOCK_SIZE + 2 * RADIUS]; int gindex = thread. Idx. x + block. Idx. x * block. Dim. x; int lindex = thread. Idx. x + RADIUS; // Read input elements into shared memory temp[lindex] = in[gindex]; if (thread. Idx. x < RADIUS) { temp[lindex - RADIUS] = in[gindex - RADIUS]; temp[lindex + BLOCK_SIZE] = in[gindex + BLOCK_SIZE]; } // Synchronize (ensure all the data is available) __syncthreads(); // Apply the stencil int result = 0; for (int offset = -RADIUS ; offset <= RADIUS ; offset++) result += temp[lindex + offset]; parallel fn // Store the result out[gindex] = result; } void fill_ints(int *x, int n) { fill_n(x, n, 1); } int main(void) { int *in, *out; // host copies of a, b, c int *d_in, *d_out; // device copies of a, b, c int size = (N + 2*RADIUS) * sizeof(int); // Alloc space for host copies and setup values in = (int *)malloc(size); fill_ints(in, N + 2*RADIUS); out = (int *)malloc(size); fill_ints(out, N + 2*RADIUS); // Alloc space for device copies cuda. Malloc((void **)&d_in, size); cuda. Malloc((void **)&d_out, size); // Copy to device cuda. Memcpy(d_in, size, cuda. Memcpy. Host. To. Device); cuda. Memcpy(d_out, size, cuda. Memcpy. Host. To. Device); serial code // Launch stencil_1 d() kernel on GPU stencil_1 d<<<N/BLOCK_SIZE, BLOCK_SIZE>>>(d_in + RADIUS, d_out + RADIUS); // Copy result back to host cuda. Memcpy(out, d_out, size, cuda. Memcpy. Device. To. Host); // Cleanup free(in); free(out); cuda. Free(d_in); cuda. Free(d_out); return 0; } parallel code serial code

Simple Processing Flow PCI Bus 1. Copy input data from CPU memory to GPU memory © NVIDIA 2013

Simple Processing Flow PCI Bus 1. Copy input data from CPU memory to GPU memory 2. Load GPU program and execute, caching data on chip for performance

Simple Processing Flow PCI Bus 1. Copy input data from CPU memory to GPU memory 2. Load GPU program and execute, caching data on chip for performance 3. Copy results from GPU memory to CPU memory

Opportunity for More Concurrency • Different kinds of action overlap are possible 1. Overlapped host computation and device computation 2. Overlapped host computation and host-device data transfer 3. Overlapped host-device data transfer and device computation 4. Concurrent device computation

Concurrency Serial (1 x) Overlap D 2 H with kernel (upto 2 x) Overlap both memcpy with Kernel (upto 3 x) Overlap with CPU (upto 3 x+)

CUDA Streams • CUDA Stream: a FIFO queue of CUDA actions to be performed – Every action (kernel launch, cuda. Memcpy, etc) is enqueued in a stream – No operation in the stream will begin until all previously issued operations complete • Operations in different streams are unordered and can overlap head CUDA Application Kernel tail CUDA Stream cuda. Memcpy CUDA Runtime & GPU

CUDA Streams for Overlap • Two types of streams in a CUDA program – The implicitly declared stream (NULL stream) – Explicitly declared streams (non-NULL streams) • Up until now, all code has been using the NULL stream by default cuda. Memcpy(. . . ); kernel<<<. . . >>>(. . . ); cuda. Memcpy(. . . ); • Non-NULL streams require manual allocation and management by the CUDA programmer

Synchronicity In CUDA • All CUDA calls are either synchronous or asynchronous w. r. t the host – Synchronous: enqueue work and wait for completion – Asynchronous: enqueue work and return immediately • Kernel Launches are asynchronous Automatic overlap with host

Asynchronous Operations for Overlap • cuda. Memcpy. Async: Asynchronous memcpy • cuda. Memcpy. Async does the same as cuda. Memcpy, but may return before the transfer is actually complete

Asynchronous Operations for Overlap • Performing a cuda. Memcpy. Async: page-locked memory allocation int *h_arr, *d_arr; cuda. Stream_t stream; cuda. Malloc((void **)&d_arr, nbytes); cuda. Malloc. Host((void **)&h_arr, nbytes); cuda. Stream. Create(&stream); Call return before transfer complete cuda. Memcpy. Async(d_arr, h_arr, nbytes, cuda. Memcpy. Host. To. Device, stream); . . . Do something while data is being moved cuda. Stream. Synchronize(stream); cuda. Free(d_arr); cuda. Free. Host(h_arr); cuda. Stream. Destroy(stream); Sync to make sure operations complete

CUDA Streams • Associate kernel launches with a non-NULL stream – Note that kernels are always asynchronous kernel<<<nblocks, threads_per_block, smem_size, stream>>>(. . . ); • The effects of cuda. Memcpy. Async and kernel launching – Operations added to stream queue for execution – Actually operations may not happen yet

CUDA Streams • Vector sum example, A + B = C NULL stream Copy A Copy B vector_sum<<<. . . >>> Copy C • Partition the vectors and use CUDA streams to overlap copy and compute Stream A Stream B Stream C Stream D A B A v_s B A C v_s B C v_s C

Implementation • Asynchronous implementation 1 – loop over all the operations for each chunk for (int i = 0; i < n. Streams; ++i) { int offset = i * stream. Size; cuda. Memcpy. Async(&d_a[offset], &a[offset], stream. Bytes, cuda. Memcpy. Host. To. Device, stream[i]); kernel<<<stream. Size/block. Size, 0, stream[i]>>>(d_a, offset); cuda. Memcpy. Async(&a[offset], &d_a[offset], stream. Bytes, cuda. Memcpy. Device. To. Host, stream[i]); }

Implementation • Asynchronous implementation 2 – batch similar operations together for (int i = 0; i < n. Streams; ++i) { int offset = i * stream. Size; cuda. Memcpy. Async(&d_a[offset], &a[offset], stream. Bytes, cuda. Memcpy. Host. To. Device, stream[i]); } for (int i = 0; i < n. Streams; ++i) { int offset = i * stream. Size; kernel<<<stream. Size/block. Size, 0, stream[i]>>>(d_a, offset); } for (int i = 0; i < n. Streams; ++i) { int offset = i * stream. Size; cuda. Memcpy. Async(&a[offset], &d_a[offset], stream. Bytes, cuda. Memcpy. Device. To. Host, stream[i]); }

Execution over C 1060 • One copy engine • One kernel engine

Execution over C 2050 • Two copy engines (H 2 D + D 2 H) • One kernel engine • Multiple kernels complete together

Overlap Data Transfers • Optimized implementation must be tailored for GPU architecture • Latest GPUs provide support to workaround this tailoring

PRIORITY STREAMS • You can give streams priority – High priority streams will preempt lower priority streams. – Currently executing blocks will complete but new blocks will only be scheduled after higher priority work has been scheduled. – Not applicable to memory transfers • Query available priorities: – cuda. Device. Get. Stream. Priority. Range(&low, &high) • Kepler: low: -1, high: 0 • Create using special API: – cuda. Stream. Create. With. Priority(&stream, flags, priority)

Implicit and Explicit Synchronization • Two types of host-device synchronization: – Implicit synchronization causes the host to wait on the GPU, but as a side effect of other CUDA actions – Explicit synchronization causes the host to wait on the GPU because the programmer has asked for that behavior 21

Implicit and Explicit Synchronization • CUDA operations that include implicit synchronization: 1. A device memset (cuda. Memset) 2. A memory copy between two addresses on the same device (cuda. Memcpy(. . . , cuda. Memcpy. Device. To. Device)) 22

Implicit and Explicit Synchronization • Four ways to explicitly synchronize in CUDA: 1. Synchronize on a device cuda. Error_t cuda. Device. Synchronize(); 2. Synchronize on a stream cuda. Error_t cuda. Stream. Synchronize(); 3. Synchronize on an event cuda. Error_t cuda. Event. Synchronize(); 4. Synchronize across streams using an event cuda. Error_t cuda. Stream. Wait. Event(cuda. Stream_t stream, cuda. Event_t event); 23

Implicit and Explicit Synchronization • cuda. Stream. Wait. Event adds inter-stream dependencies – Causes the specified stream to wait on the specified event before executing any further actions – event does not need to be an event recorded in stream cuda. Event. Record(event, stream 1); . . . cuda. Stream. Wait. Event(stream 2, event); . . . – No actions added to stream 2 after the call to cuda. Stream. Wait. Event will execute until event is satisfied

Cooperating GPU Threads

Cooperating GPU Threads • Although GPU threads SIMD archicture – Not all threads are being executed at a given instant – Cooperating threads need some synchronization • No mutexes or semaphores for explicit synchronization between threads

1 D Stencil • Consider applying a 1 D stencil to a 1 D array of elements – Each output element is the sum of input elements within a radius • If radius is 3, then each output element is the sum of 7 input elements: radius

Implementing Within a Block • Each thread processes one output element – block. Dim. x elements per block • Input elements are read several times – With radius 3, each input element is read seven times

Sharing Data Between Threads • Terminology: within a block, threads share data via shared memory • Extremely fast on-chip memory, user-managed • Declare using __shared__, allocated per block • Data is not visible to threads in other blocks

Implementing With Shared Memory • Cache data in shared memory – Read (block. Dim. x + 2 * radius) input elements from global memory to shared memory – Compute block. Dim. x output elements – Write block. Dim. x output elements to global memory – Each block needs a halo of radius elements at each boundary halo on right halo on left block. Dim. x output elements

Stencil Kernel __global__ void stencil_1 d(int *in, int *out) { __shared__ int temp[BLOCK_SIZE + 2 * RADIUS]; int gindex = thread. Idx. x + block. Idx. x * block. Dim. x; int lindex = thread. Idx. x + RADIUS; // Read input elements into shared memory temp[lindex] = in[gindex]; if (thread. Idx. x < RADIUS) { temp[lindex - RADIUS] = in[gindex - RADIUS]; temp[lindex + BLOCK_SIZE] = in[gindex + BLOCK_SIZE]; }

Stencil Kernel // Apply the stencil int result = 0; for (int offset = -RADIUS ; offset <= RADIUS ; offset++) result += temp[lindex + offset]; // Store the result out[gindex] = result; }

Data Race! § The stencil example will not work… § Suppose thread 15 reads the halo before thread 0 has fetched it… temp[lindex] = in[gindex]; Store at temp[18] if (thread. Idx. x < RADIUS) { Skipped, thread. Idx > RADIUS temp[lindex – RADIUS = in[gindex – RADIUS]; temp[lindex + BLOCK_SIZE] = in[gindex + BLOCK_SIZE]; } int result = 0; result += temp[lindex + 1]; Load from temp[19]

__syncthreads() • void __syncthreads(); • Synchronizes all threads within a block – Used to prevent RAW / WAR / WAW hazards • All threads must reach the barrier – In conditional code, the condition must be uniform across the block

Stencil Kernel __global__ void stencil_1 d(int *in, int *out) { __shared__ int temp[BLOCK_SIZE + 2 * RADIUS]; int gindex = thread. Idx. x + block. Idx. x * block. Dim. x; int lindex = thread. Idx. x + radius; // Read input elements into shared memory temp[lindex] = in[gindex]; if (thread. Idx. x < RADIUS) { temp[lindex – RADIUS] = in[gindex – RADIUS]; temp[lindex + BLOCK_SIZE] = in[gindex + BLOCK_SIZE]; } // Synchronize (ensure all the data is available) __syncthreads();

Zero Copy Memory • Zero copy Access host memory directly from device code – Transfers implicitly performed as needed by device code • Zero copy will be faster if data is only read/written from/to global memory once – Copy input data to GPU memory – Run one kernel – Copy output data back to CPU memory