GPU Computing with CUDA - Memory Hierarchy

University:
Boston University
Course:
GRS CS 910 | Directed Study: Systems
Academic year:

2022
Views:

452

Pages:

23
Author:

Priemosobbilimluj

Memory hierarchy ‣ CUDA works in both the CPU and GPU - One has to keep track of which memory is operating on (host - device) - Within the GPU there are also different memory spaces 13 Memory hierarchy ‣ Global memory - Main GPU memory - Communicates with host - Can be seen by all threads - Order of GB - Off chip, slow (~400 cycles) __device__ float variable; 14 Memory hierarchy ‣ Shared memory - Per SM - Seen by threads of the same thread block - Order of kB - On chip, fast (~4 cycles) __shared__ float variable; ‣ Registers - Private to each thread - On chip, fast float variable; 15 Memory hierarchy ‣ Local memory - Private to each thread - Off chip, slow - Register overflows float variable [10]; ‣ Constant memory - Read only - Off chip, but fast (cached) - Seen by all threads __constant__ float variable; - 64kB with 8kB cache 16 Memory hierarchy ‣ Texture memory - Seen by all threads - Read only - Off chip, but fast (cached) if cache hit - Cache optimized for 2D locality - Binds to global texture tex_var; cudaChannelFormatDesc(); cudaBindTexture2D(...); tex2D(tex_var, x_index, y_index); 17 Memory hierarchy ‣ Texture memory - Seen by all threads - Read only - Off chip, but fast (cached) if cache hit - Cache optimized for 2D locality - Binds to global Initialize texture tex_var; cudaChannelFormatDesc(); cudaBindTexture2D(...); tex2D(tex_var, x_index, y_index); 17 Memory hierarchy ‣ Texture memory - Seen by all threads - Read only - Off chip, but fast (cached) if cache hit - Cache optimized for 2D locality - Binds to global Initialize texture tex_var; Options cudaChannelFormatDesc(); cudaBindTexture2D(...); tex2D(tex_var, x_index, y_index); 17 Memory hierarchy ‣ Texture memory - Seen by all threads - Read only - Off chip, but fast (cached) if cache hit - Cache optimized for 2D locality - Binds to global Initialize texture tex_var; Options cudaChannelFormatDesc(); Bind cudaBindTexture2D(...); tex2D(tex_var, x_index, y_index); 17 Memory hierarchy ‣ Texture memory - Seen by all threads - Read only - Off chip, but fast (cached) if cache hit - Cache optimized for 2D locality - Binds to global Initialize texture tex_var; Options cudaChannelFormatDesc(); Bind cudaBindTexture2D(...); tex2D(tex_var, x_index, y_index); Fetch 17 Memory hierarchy 18 Memory hierarchy ‣ Case of Fermi - Added an L1 cache to each SM - Shared + cache = 64kB: ‣ Shared = 48kB, cache = 16kB ‣ Shared = 16kB, cache = 48kB 19 Memory hierarchy ‣ Use you memory strategically - Read only: __constant__ (fast) - Read/write and communicate within a block: __shared__ (fast and communication) - Read/write inside thread: registers (fast) - Data locality: texture 20 Resource limits ‣ Number of thread blocks per SM at the same time is limited by - Shared memory usage - Registers - No more than 8 thread blocks per SM - Number of threads 21 Resource limits - Examples Number of blocks How big should my blocks be? 8x8, 16x16 or 64x64? Max threads per block ‣ 8x8 8*8 = 64 threads per block, 1024/64 = 16 blocks. An SM can have up to 8 blocks: only 512 threads will be active at the same time ‣ 16x16 16*16 = 256 threads per block, 1024/256 = 4 blocks. An SM can take all blocks, then all 1024 threads will be active and achieve full capacity unless other resource overrule ‣ 64x64 64*64 = 4096 threads per block: doesn’t fit in a SM ! 22 Resource limits - Examples Registers We have a kernel that uses 10 registers. With 16x16 block, how many blocks can run in G80 (max 8192 registers per SM, 768 threads per SM)? 10*16*16 = 2560. SM can hold 8192/2560 = 3 blocks, meaning we will use 3*16*16 = 768 threads, which is within limits. If we add one more register, the number of registers grows to 11*16*16 = 2816. SM can hold 8192/2816 = 2 blocks, meaning we will use 2*16*15 = 512 threads. Now as less threads are running per SM is more difficult to have enough warps to have the GPU always busy! Introduction to Finite Difference Method ‣ Throughout the course many examples will be taken from Finite Difference Method ‣ Also, next lab involves implementing a stencil code ‣ Just want to make sure we’re all in the same page! 24 Introduction to Finite Difference Method ‣ Finite Difference Method (FDM) is a numerical method for solution of differential equations ‣ Domain is discretized with a mesh ‣ Derivative at a node are approximated by the linear combination of the values of points nearby (including itself) yi − yi−1 ∂y ≈ ∂x i xi − xi−1 Example using first order one sided difference 25 FDM - Accuracy and order ‣ From Taylor’s polynomial, we get f (x0 + h) = f (x0 ) + f ! (x0 ) · h + R1 (x) f (a + h) − f (a) f (x0 ) = − R(x) h ! ‣ Sources of error - Roundoff (Machine) - Truncation (R(x)): gives the order of the method 26 FDM - Stability ‣ Stability of the algorithm may be conditional 2 ∂u ∂ u = ∂t ∂x2 ‣ Explicit method n n n n un+1 − u u − 2u + u j j j+1 j j−1 = k h2 Conditionally stable k < 0.5 2 h ‣ Implicit n+1 n+1 n+1 n un+1 − u u − 2u + u j j j+1 j j−1 Unconditionally stable = 2 k h ‣ Crank Nicolson un+1 j − k unj 1 = 2 ! unj+1 − 2unj h2 + unj−1 + un+1 j+1 − 2un+1 j h2 + un+1 j−1 " Unconditionally stable 27 FDM - 2D example ‣ We’re going to be working with this example in the labs ‣ Explicit diffusion 2 ∂u ∂ u =α 2 ∂t ∂x un+1 i,j = uni,j αk n + 2 (ui,j+1 + uni,j−1 + uni+1,j + uni−1,j − 4uni,j ) h 28 FDM on the GPU ‣ Mapping the physical problem to the software and hardware ‣ Each thread will operate on one node: node indexing groups them naturally! un+1 i,j = uni,j αk n + 2 (ui,j+1 + uni,j−1 + uni+1,j + uni−1,j − 4uni,j ) h 29 FDM on the GPU ‣ Mapping the physical problem to the software and hardware ‣ Each thread will operate on one node: node indexing groups them naturally! Thread = node un+1 i,j = uni,j Thread block αk n + 2 (ui,j+1 + uni,j−1 + uni+1,j + uni−1,j − 4uni,j ) h 29 FDM on the GPU ‣ Mapping the physical problem to the software and hardware ‣ Each thread will operate on one node: node indexing groups them naturally! un+1 i,j = uni,j αk n + 2 (ui,j+1 + uni,j−1 + uni+1,j + uni−1,j − 4uni,j ) h 29

GPU Computing with CUDA - Memory Hierarchy

Related Documents

Recommended Documents

Description

Related Documents

Free up your schedule!

Take 5 seconds to unlock