PCA 21: The CUDA Memory Model

In order to successfully complete this assignment you must do the required reading, watch the provided videos and complete all instructions. The embedded survey form must be entirely filled out and submitted on or before 11:59pm. Students must come to class the next day prepared to discuss the material covered in this assignment.

PCA 21: The CUDA Memory Model#

Goals for today’s pre-class assignment#

  1. Quick Overview of the CUDA Memory Model

  2. Coalescing global memory accesses

  3. Using shared memory

  4. CUDA Memory Example by doing Reduction

  5. Assignment wrap-up

1. Quick Overview of the CUDA Memory Model#

The following video shows you how to schedule basic CUDA jobs on the HPCC.

from IPython.display import YouTubeVideo
YouTubeVideo("HQejUtJtBlg",width=640,height=360)

2. Coalescing global memory accesses#

from IPython.display import YouTubeVideo
YouTubeVideo("mLxZyWOI340",width=640,height=360)

✅ QUESTION: Why does a large stride lower performance?

Put your answer to the above question here.

The following is a great reference on how to access global memory effectively:

3. Using shared memory#

The following is a great reference on using Shared memory on cuda:

The basic syntac can be found here:

Static Shared memory:

__global__ void staticReverse(int *d, int n)
{
  __shared__ int s[64];
  int tid = threadIdx.x;
  int i = tid + blockIdx.x*blockDim.x;
  
  int t = tid;
  int tr = n-t-1;
  s[t] = d[t];
  __syncthreads();
  d[t] = s[tr];
}

Dynamic Shared Memory:

dynamicReverse<<<1, n, n*sizeof(int)>>>(d_d, n);

4. CUDA Memory Example by doing Reduction#

from IPython.display import YouTubeVideo
YouTubeVideo("RzPDlnZhxtQ",width=640,height=360)

✅ DO THIS: I found the following code on Github which does a reduction similar to the one shown in the above video. Download the code and get it working on the HPCC.

You will need to do the following

  • Comment out the CPU print statements (lines 341 and 342)

  • Remove the <512> syntax error on line 217

✅ DO THIS: Change the kernel function pointer on line 233 to point to the five different kernel options. Try each one and record the timing differences as compared to the cpu.

5. Assignment wrap-up#

Please fill out the form that appears when you run the code below. You must completely fill this out in order to receive credits for the assignment!

Direct Link to Survey Form

✅ Assignment-Specific QUESTION: Were you able to get the CUDA reduction example working on the HPCC? If not, where did you get stuck?

Put your answer to the above question here

✅ QUESTION: Summarize what you did in this assignment.

Put your answer to the above question here

✅ QUESTION: What questions do you have, if any, about any of the topics discussed in this assignment after working through the jupyter notebook?

Put your answer to the above question here

✅ QUESTION: How well do you feel this assignment helped you to achieve a better understanding of the above mentioned topic(s)?

Put your answer to the above question here

✅ QUESTION: What was the most challenging part of this assignment for you?

Put your answer to the above question here

✅ QUESTION: What was the least challenging part of this assignment for you?

Put your answer to the above question here

✅ QUESTION: What kind of additional questions or support, if any, do you feel you need to have a better understanding of the content in this assignment?

Put your answer to the above question here

✅ QUESTION: Do you have any further questions or comments about this material, or anything else that’s going on in class?

Put your answer to the above question here

✅ QUESTION: Approximately how long did this pre-class assignment take?

Put your answer to the above question here

from IPython.display import HTML
HTML(
"""
<iframe 
	src="https://cmse.msu.edu/cmse401-pc-survey" 
	width="100%" 
	height="500px" 
	frameborder="0" 
	marginheight="0" 
	marginwidth="0">
	Loading...
</iframe>
"""
)

Congratulations, we’re done!#

To get credit for this assignment you must fill out and submit the above survey from on or before the assignment due date.

Written by Dr. Dirk Colbry, Michigan State University Creative Commons License
This work is licensed under a Creative Commons Attribution-NonCommercial 4.0 International License.