PAR Class 6, Wed 2018-02-21

1   Narayanaswami Chandrasekhar talk on Blockchains

Class will be short today because of this.

2   Optional Homework - bring your answers and discuss next week

3   Stanford lectures

1. Lecture 5 performance considerations shows how to fine tune your program once it's already working, if you need the extra speed.
2. Lecture 6 parallel patterns 1 presents some paradigms of parallel programming. These are generally useful building blocks for parallel algorithms.

4   Misc CUDA

1. 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.

2. The program nvidia-smi shows the current load on the GPU.

3. My web copy of the tutorial programs from Stanford's parallel course notes is also on parallel at /parallel-class/stanford/tutorials/ .

   1. I've edited some of them, and put the originals in orig/ , and created new ones.

   2. To compile them, you need /local/cuda/bin in your PATH and /local/cuda/lib64 in your LD_LIBRARY_PATH .

   3. Name your source program foo.cu for some foo .

   4. Compile it thus: nvcc foo.cu -o foo .

   5. hello_world.cu shows a simple CUDA program and uses a hack to print from a device function.

   6. hello_world2.cu shows printing from several threads.

   7. global_functions.cu shows some basic CUDA stuff.

   8. device_functions.cu extends it.

   9. vector_addition.cu does (you figure it out).

   10. 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:

       1. In linux, what's the easiest way to find the smallest prime larger than a given number?

       2. 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;

5   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.

1. Create static variables with __device__ __managed__. See /parallel-class/stanford/tutorials/vector_addition2.cu on parallel.
2. Use cudaMallocManaged. See /parallel-class/stanford/tutorials/vector_addition3.cu on parallel.
3. 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.
4. When the linux kernel gets HMM (heterogeneous memory management), all data on the heap will automatically be managed.
5. 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.
6. Any CUDA example using cudaMemcpy is now obsolete (on Pascal GPUs).

6   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.

7   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.

8   Misc hints