.. title: PAR Class 6, Wed 2018-02-21 .. slug: class06 .. date: 2018-02-15 .. tags: class .. category: .. link: .. description: .. type: text .. raw:: html .. role:: red .. role:: blue .. sectnum:: .. contents:: Table of contents Narayanaswami Chandrasekhar talk on Blockchains ----------------------------------------------- Class will be short today because of this. Optional Homework - bring your answers and discuss next week ------------------------------------------------------------ Paper questions =============== #. Research and then describe the main changes from NVidia Maxwell to Pascal. #. Although a thread can use 255 registers, that might be bad for performance. Why? #. Give a common way that the various threads in a block can share data with each other. #. Reading a word from global memory might take 400 cycles. Does that mean that a thread that reads many words from global memory will always take hundreds of times longer to complete? #. Since the threads in a warp are executed in a SIMD fashion, how can an if-then-else block be executed? #. What is unified virtual addressing and how does it make CUDA programming easier? Programming questions ===================== #. Repeat homework 2's matrix multiplication problem, this time in CUDA. Report how much parallel speedup you get. #. Look at the dataset **/parallel-class/data/bunny**. It contains 35947 points for the Stanford bunny. Assuming that each point has a mass of 1, and is gravitationally attracted to the others, compute the potential energy of the system. The formula is this: :math:`U = - \sum_{i=1}^{N-1} \sum_{j=i+1}^N \frac{1}{r_{ij}}` where :math:`r_{ij}` is the distance between points :math:`i` and :math:`j` . (This assumes that G=1). #. Now look at the dataset **/parallel-class/data/blade**, which contains 882954 points for a turbine blade. Can you process it? Stanford lectures ----------------- #. `Lecture 5 performance considerations <../../stanford/lectures/lecture_5/performance_considerations.pdf>`_ shows how to fine tune your program once it's already working, if you need the extra speed. #. `Lecture 6 parallel patterns 1 <../../stanford/lectures/lecture_6/parallel_patterns_1.pdf>`_ presents some paradigms of parallel programming. These are generally useful building blocks for parallel algorithms. Misc CUDA --------- #. The demo programs are in **/local/cuda/samples/** . Their coding style is suboptimal. However, in **/local/cuda/samples/1_Utilities/** , **bandwidthTest** and **deviceQuery** are interesting. For your convenience, **/parallel-class/deviceQuery** is a link. Run it to see the GPU's capabilities. #. The program **nvidia-smi** shows the current load on the GPU. #. My `web copy <../../stanford/tutorials>`_ of the tutorial programs from Stanford's parallel course notes is also on parallel at **/parallel-class/stanford/tutorials/** . a. I've edited some of them, and put the originals in **orig/** , and created new ones. #. To compile them, you need **/local/cuda/bin** in your **PATH** and **/local/cuda/lib64** in your **LD_LIBRARY_PATH** . #. Name your source program **foo.cu** for some **foo** . #. Compile it thus: **nvcc foo.cu -o foo** . #. **hello_world.cu** shows a simple CUDA program and uses a hack to print from a device function. #. **hello_world2.cu** shows printing from several threads. #. **global_functions.cu** shows some basic CUDA stuff. #. **device_functions.cu** extends it. #. **vector_addition.cu** does (you figure it out). #. **vector_addition2.cu** is my modification to use unified memory, per http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html . I also cleaned up the code and shrank the number of lines for better display. IMO, unified memory makes programming a lot easier. Notes: i. In linux, what's the easiest way to find the smallest prime larger than a given number? #. To find the number of blocks needed for N threads, you can do it the Stanford way: grid_size = num_elements / block_size; if(num_elements % block_size) ++grid_size; or you can do it the RPI (i.e., my) way: grid_size = (num_elements + block_size - 1) / block_size; Managed Variables ----------------- Last time we saw 2 ways to create managed variables. They can be accessed by either the host or the device and are paged automatically. This makes programming much easier. #. Create static variables with **__device__ __managed__**. See **/parallel-class/stanford/tutorials/vector_addition2.cu** on parallel. #. Use **cudaMallocManaged**. See **/parallel-class/stanford/tutorials/vector_addition3.cu** on parallel. #. In either case, you need to call **cudaDeviceSynchronize();** on the host after starting a parallel kernel before reading the data on the host. The reason is that the kernel is started asynchonously and control returns while it is still executing. #. When the linux kernel gets HMM (heterogeneous memory management), all data on the heap will automatically be managed. #. The reason is that virtual addresses are long enough to contain a tag saying what device they are on. The VM page mapper will read and write pages to various devices, not just swap files. #. Any CUDA example using **cudaMemcpy** is now obsolete (on Pascal GPUs). Doc --- Nvidia's `CUDA programming guide `_ is excellent, albeit obsolescent in places. The Pascal info looks like it's been tacked onto an older document. The whitepaper `NVIDIA GeForce GTX 1080 `_ describes, from a gaming point of view, the P104 GPU, which is in the GTX 1080, the card in parallel.ecse. NVIDIA now has a higher level GPU, the P100, described in the `P100 whitepaper `_ and `P100 technical overview `_. Note that the P100 is a Tesla (scientific computing) not a GeForce (gaming). This description is much more technical. Managed memory issues --------------------- I'm sometimes seeing a 2.5x speed reduction using managed memory on the host, compared to using unmanaged memory. Dunno what's happening. Misc hints ---------- Vim === To get **vim** to show line numbers, create a file **~/.exrc** containing this line: :se nu It will be read everytime vim starts, and will set line number mode.