GPU Computing with CUDA
Lab 2 - FD in global and texture memory Objectives
‣ Implement a simple finite difference problem in CUDA
‣ Learn how to use texture memory Heat transfer with FD
‣ Heat diffusion on square 2D flat plate with Dirichlet boundary
conditions and 0 initial condition
T=0
T = 200
T=0
3.5
∂u
2
= α∇ u
∂t
α = 0.645
k = 1e-5
N = 128
T = 200
3.5
n+1
ui,j
=
uni,j
αk n
+ 2 (ui,j+1 + uni,j−1 + uni+1,j + uni−1,j − 4uni,j )
h Heat transfer with FD
‣ Each thread will do one mesh point
‣ Global memory data lives for whole execution!
‣ Experiment with other mesh sizes (N) that are not a multiple of
block size Heat transfer with FD
‣ Changes to global memory case: (start from pervious code!)
- Initialize texture (beginning of code)
texture tex_u;
- After allocation, bind texture to global memory
cudaChannelFormatDesc desc = cudaCreateChannelDesc();
cudaBindTexture2D(NULL, tex_u, u_d, desc, N_max, N_max,
sizeof(float)*N_max);
- In the kernel, fetch values from texture memory
int x = threadIdx.x + blockIdx.x*BSZ;
int y = threadIdx.y + blockIdx.y*BSZ;
c = tex2D(tex_u_prev, x, y);
- Unbind texture at the end
cudaUnbindTexture(tex_u); Heat transfer with FD
‣ Using texture memory:
- Should give good results as it is designed for 2D data locality
‣ Textures map multiples of 32 or 64:
- Need to pad or use cudaMallocPitch
‣ Texture memory is read only within a kernel
- Need to write separate kernel to update values in texture
GPU Computing with CUDA Lab 2 - FD in Global and Texture Memory
of 6
Report
Tell us what’s wrong with it:
Thanks, got it!
We will moderate it soon!
Free up your schedule!
Our EduBirdie Experts Are Here for You 24/7! Just fill out a form and let us know how we can assist you.
Take 5 seconds to unlock
Enter your email below and get instant access to your document