Table of contents
Adrian-James G, Kevin L, Ryan W
Cohen D, Zach O
Amanda M
Reagan W, Ean S
Tom P, Allan N
Saaif A
Next fall I'm teaching Quantum Computer Programming.
Table of contents
Watch some videos that I listed in class 26 but didn't show.
Table of contents
Phu TT, Mark B
Ivan H
Luoyan Z
Sissi J
If your final project is building on, or sharing with, another course or project (say on GitHub), then you must give the details, and say what's new for this course.
Try to use the IEEE conference format . It allows either latex or MS word. Submit the PDF paper to gradescope.
from the syllabus (modified).
An implementation showing parallel computing.
An extended abstract or paper on your project, written up like a paper. You should follow the style guide for some major conference (I don't care which, but can point you to one).
A more detailed manual, showing how to use it.
A 10-minute project presentation.
A write-up uploaded on the last class day. This will contain an academic paper, code and perhaps video or user manual.
I'll assign some good things to watch.
Git is good to simultaneously keep various versions. Here's a git intro:
Create a dir for the project:
mkdir PROJECT; cd PROJECT
Initialize:
git init
Create a branch (you can do this several times):
git branch MYBRANCHNAME
Go to a branch:
git checkout MYBRANCHNAME
Do things:
vi, make, ....
Save it:
git add .; git commit -mCOMMENT
Repeat
At times I've used this to modify a program for class while keeping a copy of the original.
One of my rules is to push design decisions to take effect as early in the process execution as possible. Constructing variables at compile time is best, at function call time is 2nd, and on the heap is worst.
If I have to construct variables on the heap, I construct few and large variables, never many small ones.
Often I compile the max dataset size into the program, which permits constructing the arrays at compile time. Recompiling for a larger dataset is quick (unless you're using CUDA).
Accessing this type of variable uses one less level of pointer than accessing a variable on the heap. I don't know whether this is faster with a good optimizing compiler, but it's probably not slower.
If the data will require a dataset with unpredictably sized components, such as a ragged array, then I may do the following.
Read the data once to accumulate the necessary statistics.
Construct the required ragged array.
Reread the data and populate the array.
Update: However, with CUDA, maybe managed variables must be on the heap.
https://www.nextplatform.com/2019/12/04/openacc-cozies-up-to-c-c-and-fortran-standards/
https://www.nextplatform.com/2019/01/09/two-thirds-of-the-way-home-with-exascale-programming/
They agree with me that fewer bigger parallel processors are better than many smaller processors connected with MPI. I.e., they like my priorities in this course.
CppCon is the annual, week-long face-to-face gathering for the entire C++ community. https://cppcon.org/
CppCon 2014: Herb Sutter "Paying for Lunch: C++ in the ManyCore Age" https://www.youtube.com/watch?v=AfI_0GzLWQ8
CppCon 2016: Combine Lambdas and weak_ptrs to make concurrency easy (4min) https://www.youtube.com/watch?v=fEnnmpdZllQ
A Pragmatic Introduction to Multicore Synchronization by Samy Al Bahra. https://www.youtube.com/watch?v=LX4ugnzwggg
CppCon 2018: Tsung-Wei Huang “Fast Parallel Programming using Modern C++” https://www.youtube.com/watch?v=ho9bqIJkvkc&list=PLHTh1InhhwT7GoW7bjOEe2EjyOTkm6Mgd&index=13
CppCon 2018: Anny Gakhokidze “Workflow hacks for developers” https://www.youtube.com/watch?v=K4XxeB1Duyo&list=PLHTh1InhhwT7GoW7bjOEe2EjyOTkm6Mgd&index=33
CppCon 2018: Bjarne Stroustrup “Concepts: The Future of Generic Programming (the future is here)” https://www.youtube.com/watch?v=HddFGPTAmtU
CppCon 2018: Herb Sutter “Thoughts on a more powerful and simpler C++ (5 of N)” https://www.youtube.com/watch?v=80BZxujhY38
CppCon 2018: Jefferson Amstutz “Compute More in Less Time Using C++ Simd Wrapper Libraries” https://www.youtube.com/watch?v=80BZxujhY38
CppCon 2018: Geoffrey Romer “What do you mean "thread-safe"?” https://www.youtube.com/watch?v=s5PCh_FaMfM
"The International Conference for High Performance Computing, Networking, Storage, and Analysis"
He has visited RPI more than once.
An Overview of High Performance Computing and Future Requirements Jul 17, 2021, (35:03)
Abstract: In this talk we examine how high performance computing has changed over the last 10-year and look toward the future in terms of trends. These changes have had and will continue to have a major impact on our numerical scientific software. A new generation of software libraries and algorithms are needed for the effective and reliable use of (wide area) dynamic, distributed and parallel environments. Some of the software and algorithm challenges have already been encountered, such as management of communication and memory hierarchies through a combination of compile--time and run--time techniques, but the increased scale of computation, depth of memory hierarchies, range of latencies, and increased run--time environment variability will make these problems much harder.
About the Speaker: Jack Dongarra holds an appointment at the University of Tennessee, Oak Ridge National Laboratory, and the University of Manchester. He specializes in numerical algorithms in linear algebra, parallel computing, use of advanced-computer architectures, programming methodology, and tools for parallel computers. He was awarded the IEEE Sid Fernbach Award in 2004; in 2008 he was the recipient of the first IEEE Medal of Excellence in Scalable Computing; in 2010 he was the first recipient of the SIAM Special Interest Group on Supercomputing's award for Career Achievement; in 2011 he was the recipient of the IEEE Charles Babbage Award; in 2013 he received the ACM/IEEE Ken Kennedy Award; in 2019 he received the ACM/SIAM Computational Science and Engineering Prize, and in 2020 he received the IEEE Computer Pioneer Award. He is a Fellow of the AAAS, ACM, IEEE, and SIAM and a foreign member of the Russian Academy of Science, a foreign member of the British Royal Society, and a member of the US National Academy of Engineering.
Lorena Barba IPDPS21 keynote May 20, 2021 (37:11)
Keynote at the IEEE International Parallel and Distributed Processing Symposium, May 19, 2021
Abstract Thirty years ago, David Bailey published a humorous piece in the Supercomputing Review magazine, listing 12 ways of presenting results to artificially boost performance claims. That was at a time when the debate was between Cray "two-oxen" machines versus parallel "thousand-chickens" systems, when parallel standards (like MPI) were still unavailable, and the Top500 list didn't yet exist. In the years since, David and others updated the list of tricks a few times, notably in 2010–11 (when the marketing departments of Intel and Nvidia were really going at each other) Georg Hager in his blog and Scott Pakin in HPC Wire. Heterogeneity of computing systems has only escalated in the last decade, and many remiss reporting tactics continue unabated. Alas, two new ingredients have entered into the mix: wide adoption of machine learning techniques both in the science applications and systems research; and a swell of concern over reproducibility and replicability. My talk will be a new twist on the 12 ways to fool the masses, focusing on how researchers in computational science and high-performance computing miss the mark when conducting or reporting their results with poor reproducibility. By showcasing in a lighthearted manner a set of anti-patterns, I aim to encourage us to see the value and commit to adapting our practice to achieve more trustworthy scientific evidence with high-performance computing.
There's a link to the slides there.
I see my role as a curator, selecting the best stuff to present to you. Since I like the topic, and asked to be given permission to create this course, I pick things I like since you might like them too.
So, if you liked the course, then please officially tell RPI by completing the survey and saying what you think. Thanks.
I'm open to questions and discussions about any legal ethical topic. Even after you graduate.
Table of contents
RENSSELAER CENTER FOR MATERIALS, DEVICES & INTEGRATED SYSTEMS (cMDIS), SCHOOL OF ENGINEERING, AND SCHOOL OF SCIENCE
In Partnership With
LOS ALAMOS NATIONAL LABORATORY INSTITUE FOR MATERIALS SCIENCE
Research Nucleation Workshop
Quantum Materials and Devices
04/29/2022, 2.30-5.00 pm EST, to be held on Webex
It's not certain whether students are invited. However, if interested, ask me.
Only the part from 40:45 to 41:20
Ean S
Everyone who requested a specific data got it. I assigned the others randomly.
Phu TT, Mark B
Ivan H
Luoyan Z
Sissi J
Adrian-James G, Kevin L, Ryan W
Cohen D, Zach O
Amanda M
Reagan W, Ean S
Tom P, Allan N
Saaif A
Quantum Machine Learning - 37 - Overview of the HHL Algorithm 5:48.
quick, deep, intro.
Quantum algorithm for solving linear equations 36:31.
quite understandable, but no time.
This is in Huawei HiQ, an open-source software framework for quantum computing.
https://en.wikipedia.org/wiki/Quantum_algorithm_for_linear_systems_of_equations
What is an Ion Trap Quantum Computer? Simon Benjamin (2:30).
What is Quantum Annealing?, D-Wave Systems (6:14).
Elementary but sets the stage.
How The Quantum Annealing Process Works, D-Wave Systems (6:09)
Smart: Guy Creates An AI Clone Of Himself To Sit In On Zoom Meetings So He Doesn't Have To
"The goal of the contest is to write code that is as readable, clear, innocent and straightforward as possible, and yet it must fail to perform at its apparent function. To be more specific, it should do something subtly evil. Every year, we will propose a challenge to coders to solve a simple data processing problem, but with covert malicious behavior. Examples include miscounting votes, shaving money from financial transactions, or leaking information to an eavesdropper. The main goal, however, is to write source code that easily passes visual inspection by other programmers."
https://www.awesomestories.com/asset/view/Space-Race-American-Rocket-Failures
Moral: After early disasters, sometimes you can eventually get things to work.
Upload a 100 word project report to Gradescope.
Table of contents
Reagan W, Ean S, Tom P, and Allan N: computer vision
Reading/viewing material will be assigned.
April 18 or 25. Email me your preferred date. FCFS.
We'll watch some in class.
Shor's Factoring Algorithm by Umesh Vazirani (25:42)
This is an excellent video, which there's no time to show.
This video is difficult because the topic is difficult, but important. At least you'll get a sense of it. This is what potentially breaks a lot of cryptosystems.
7. Shor's Algorithm I: Understanding Quantum Fourier Transform, Quantum Phase Estimation - Part 1 by Abraham Asfaw, IBM. (44:05).
This is from IBM's Introduction to Quantum Computing and Quantum Hardware. Lecture Notes and Labs: https://qiskit.org/learn/intro-qc-qh
This is the 1st of 5 videos describing the algorithm. There's no time for the rest, but this will give you an idea.
RPI's quantum group, including me, visited Asfaw last spring just before everything shut down.
Quantum Machine Learning - 37 - Overview of the HHL Algorithm 5:48.
quick, deep, intro.
Quantum algorithm for solving linear equations 36:31.
quite understandable, but no time.
This is in Huawei HiQ, an open-source software framework for quantum computing.
https://en.wikipedia.org/wiki/Quantum_algorithm_for_linear_systems_of_equations
No lecture today.
Finish watching Quantum Computing for Computer Scientists. We stopped around 18 minutes.
Browse around Amazon Braket .
Next Monday, I'll invite you to talk about anything interesting you saw.
Table of contents
THE WORLD'S SMALLEST PC?! "The Dot" Coming Soon From Newegg (3:01)
Adrian-James G: IoT/5G
Saaif A: Accelerating Distributed Reinforcement Learning
Ryan W: Accelerated Vehicle Routing Algorithms [S41708]
Cohen D: Accelerating Automobile Noise and Vibration Analysis Using GPUs
Kevin L: Accelerate Your AI and HPC Journey on Google Cloud (Presented by Google Cloud) [S42583].
Table of contents
Homeworks up through 6 have been graded.
A course grade up to Mar 30 was calculated, weighting all homeworks the same, and uploaded to LMS.
You may convert it to a letter grade as follows:
90: A, 85: A-, 80: B+, 75: B, 70: B-, 65: C+, 60: C, 55: C-, 50: D+, 45: D, 0: F
Mark B and Phu T.T: GTC Cybersecurity
Zach O: Advanced Technologies and Techniques for Debugging CUDA GPU Applications [S41252].
Amanda M: Autonomous Vehicles
Adrian-James G: IoT/5G
Saaif A: Accelerating Distributed Reinforcement Learning
Ryan W: Accelerated Vehicle Routing Algorithms [S41708]
Cohen D: Accelerating Automobile Noise and Vibration Analysis Using GPUs
Reagan W, Ean S, Tom P, and Allan N: computer vision
continuing this, the 3rd major topic listed in the syllabus (after Multicore Intel Xeon and Nvidia GPU accelerator).
investigate some other widely used parallel tool, like Kokkos or Intel TBB, or parallel STL or the parallel primitives in C++20.
try NVIDIA's ML tools.
try a parallel VM on Amazon ECC and compare it to parallel.ecse.
Table of contents
My quantum computing summary .
We'll do parts of it from time to time.
I present quantum computing for computer scientists, mostly on a level above the physics.
Videos. They present the subject better than I can. They show you leaders in the field. My value is in selecting the best ones for you.
Quantum Computing for Computer Scientists (1:28:30), the 1st 1/2 hour or so.
Which problems can quantum computers solve exponentially faster than classical computers? David Gosset, IBM quantum computing research scientist, explains why algorithms are key to finding out.
Quantum Computing 2021 Update (13:12)
https://www.telegraph.co.uk/technology/2020/09/02/britain-must-act-fast-prevent-brain-drain-quantum-computing/ - good summary of programs around the world.
The Map of Quantum Physics (21:17), Jul 31, 2020.
Table of contents
was caused by my updating some packages but not rebooting, causing new libraries to try to use the old kernel driver.
Examples
bucket_sort2d.cu overlays a grid on a set of 2D points and finds the points in each grid cell (bucket).
The tuple is an efficient class for a short vector of fixed length.
Note how random numbers are generated. You combine an engine that produces random output with a distribution.
However you might need more complicated coding to make the numbers good when executing in parallel. See monte_carlo_disjoint_sequences.cu.
The problem is that the number of points in each cell is unpredictable.
The cell containing each point is computed and that and the points are sorted to bring together the points in each cell.
Then lower_bound and upper_bound are used to find each bucket in that sorted vector of points.
See this lower_bound description.
mode.cu shows:
Counting the number of unique keys in a vector.
Sort the vector.
Do an inner_product. However, instead of the operators being times and plus, they are not equal to the next element and plus.
Counting their multiplicity.
Construct vectors, sized at the number of unique keys, to hold the unique keys and counts.
Do a reduce_by_keys on a constant_iterator using the sorted vector as the keys. For each range of identical keys, it sums the constant_iterator. That is, it counts the number of identical keys.
Write a vector of unique keys and a vector of the counts.
Finding the most used key (the mode).
Do max_element on the counts vector.
repeated_range.cu repeats each element of an N-vector K times: repeated_range([0, 1, 2, 3], 2) -> [0, 0, 1, 1, 2, 2, 3, 3]. It's a lite version of expand.cu, but uses a different technique.
Here, N=4 and K=2.
The idea is to construct a new iterator, repeated_range, that, when read and incremented, will return the proper output elements.
The construction stores the relevant info in structure components of the variable.
Treating its value like a subscript in the range [0,N*K), it divides that value by K and returns that element of its input.
See also strided_range.cu and tiled_range.cu.
The Thrust device can be CUDA, OpenMP, TBB, etc.
You can spec it in 2 ways:
by adding an extra arg at the start of a function's arg list.
with an envar
This is hardware support to make programming devices like GPUs easier. It took years to get into the official kernel, presumably because you don't want bugs in memory management.
This is also a nice intro into current operating systems concerns. Do our operating systems courses talk about this?
See the syllabus.
April 4: team and title.
April 11: 100 word project report.
April 18, 25: 10 minute presentations. Email me with preferred date.
Wed Apr 27: report due.
April 7, 21. If you wish, I can assign some excellent quantum computing videos to watch.
Rules:
Any number of students may work together to present one talk.
Prepare and give 5-10 minute talk on Mon 3/28, Thurs 3/30, Mon 4/4, or Mon 4/11. A team would jointly give one 5-10 minute talk. I'll pack six talks per day, filling the earliest days first, and fill out unused time with new material.
Email me with your team, topic, and preferred dates.
Your topic may be anything from GTC 2022 or a comparable conference.
Table of contents
Today is the spring equinox. However, today in Albany, sunrise is at 6:56am, sunset 7:08pm. That makes the day 12 hours and 12 minutes long, not 12 hours. Why?
See saxpy2 and saxpy3 in files/nvidia/thrust/rpi/
Continue examples from last class.
Table of contents
I'll lecture from .../files/nvidia/thrust/doc/GTC_2010_Part_2_Thrust_By_Example.pdf on parallel.ecse starting around slide 23.
Nvidia's thrust examples are now in /parclass/2022/files/nvidia/thrust/examples .
The github repository is here: https://github.com/NVIDIA/thrust
It references https://godbolt.org a web-based compiler explorer.
I won't cover these in class; they're presented in case you're interested.
The point of these examples is to show you some techniques for parallel programming that are not obvious. I.e., you probably wouldn't think of them if you did not already know them.
I rewrote thrust/examples-1.8/tiled_range.cu into thrust/rpi/tiled_range2.cu .
It is now much shorter and much clearer. All the work is done here: ..
gather(make_transform_iterator(make_counting_iterator(0), _1%N), make_transform_iterator(make_counting_iterator(N*C), _1%N), data.begin(), V.begin());
make_counting_iterator(0) returns pointers to the sequence 0, 1, 2, ...
_1%N is a function computing modulo N.
make_transform_iterator(make_counting_iterator(0), _1%N) returns pointers to the sequence 0%N, 1%N, ...
gather populates V. The i-th element of V gets make_transform_iterator...+i element of data, i.e., the i%N-th element of data.
tiled_range3.cu is even shorter. Instead of writing an output vector, it constructs an iterator for a virtual output vector: ..
auto output=make_permutation_iterator(data, make_transform_iterator(make_counting_iterator(0), _1%N));
*(output+i) is *(data+(i%N)).
You can get as many tiles as you want by iterating.
repeated_range2.cu is my improvement on repeated_range.cu: ..
auto output=make_permutation_iterator(data.begin(), make_transform_iterator(make_counting_iterator(0), _1/3));
make_transform_iterator(make_counting_iterator(0), _1/3)) returns pointers to the sequence 0,0,0,1,1,1,2,2,2, ...
Unmodified thrust examples:
expand.cu takes a vector like V= [0, 10, 20, 30, 40] and a vector of repetition counts, like C= [2, 1, 0, 3, 1]. Expand repeats each element of V the appropriate number of times, giving [0, 0, 10, 30, 30, 30, 40]. The process is as follows.
Since the output vector will be longer than the input, the main program computes the output size, by reduce summing C, and constructs a vector to hold the output.
Exclusive_scan C to obtain output offsets for each input element: C2 = [0, 2, 3, 3, 6].
Scatter_if the nonzero counts into their corresponding output positions. A counting iterator, [0, 1, 2, 3, 4] is mapped with C2, using C as the stencil, giving C3 = [0, 0, 1, 3, 0, 0, 4].
An inclusive_scan with max fills in the holes in C3, to give C4 = [0, 0, 1, 3, 3, 3, 4].
Gather uses C4 to gather elements of V: [0, 0, 10, 30, 30, 30, 40].
set_operations.cu. This shows methods of handling an operation whose output is of unpredictable size. The question is, is space or time more important?
If the maximum possible output size is reasonable, then construct an output vector of that size, use it, and then erase it down to its actual size.
Or, run the operation twice. The 1st time, write to a discard_iterator, and remember only the size of the written data. Then, construct an output vector of exactly the right size, and run the operation again.
I use this technique a lot with ragged arrays in sequential programs.
sparse_vector.cu represents and sums sparse vectors.
A sparse vector has mostly 0s.
The representation is a vector of element indices and another vector of values.
Adding two sparse vectors goes as follows.
Allocate temporary index and element vectors of the max possible size (the sum of the sizes of the two inputs).
Catenate the input vectors.
Sort by index.
Find the number of unique indices by applying inner_product with addition and not-equal-to-next-element to the indices, then adding one.
E.g., applied to these indices: [0, 3, 3, 4, 5, 5, 5, 8], it gives 5.
Allocate exactly enough space for the output.
Apply reduce_by_key to the indices and elements to add elements with the same keys.
The size of the output is the number of unique keys.
What's the best way to sort 16000 sets of 1000 numbers each? E.g., sort the rows of a 16000x1000 array? On geoxeon, /pc/thrust/rpi/comparesorts.cu, which I copied from http://stackoverflow.com/questions/28150098/how-to-use-thrust-to-sort-the-rows-of-a-matrix|stackoverflow, compares three methods.
Call the thrust sort 16000 times, once per set. That took 10 secs.
Sort the whole list of 16,000,000 numbers together. Then sort it again by key, with the keys being the set number, to bring the elements of each set together. Since the sort is stable, this maintains the order within each set. (This is also how radix sort works.) That took 0.04 secs.
Call a thrust function (to sort a set) within another thrust function (that applies to each set). This is new in Thrust 1.8. That took 0.3 secs.
This is a surprising and useful paradigm. It works because
There's an overhead to starting each thrust function, and
Radix sort, which thrust uses for ints, takes linear time.
The Thrust examples teach several non-intuitive paradigms. As I figure them out, I'll describe a few. My descriptions are modified and expanded versions of the comments in the programs. This is not a list of all the useful programs, but only of some where I am adding to their comments.
arbitrary_transformation.cu and dot_products_with_zip.cu. show the very useful zip_iterator. Using it is a 2-step process.
Combine the separate iterators into a tuple.
Construct a zip iterator from the tuple.
Note that operator() is now a template.
boundingbox.cu finds the bounding box around a set of 2D points.
The main idea is to do a reduce. However, the combining operation, instead of addition, is to combine two bounding boxes to find the box around them.
The combining op can be any associative op.
is online.
Modify some Thrust example programs use lambdas and placeholder notation. Good choices might be saxpy, repeated_range, bounding_box, or simple_moving_average.
Teams are ok.
Table of contents
I'll lecture from .../files/nvidia/thrust/doc/GTC_2010_Part_2_Thrust_By_Example.pdf on parallel.ecse
Why something from 2010? It's well presented, hasn't been updated by NVidia, and is current, apart from not using lambdas.
The deeper meaning of this stuff is the set of parallel programming and algorithm paradigms.
On parallel in .../files/nvidia/thrust are many little demo programs from the thrust distribution, with my additions.
Thrust is fast because the functions that look like they would need linear time really take only log time in parallel.
In functions like reduce and transform, you often see an argument like thrust::multiplies<float>(). The syntax is as follows:
thrust::multiplies<float> is a class.
It overloads operator().
However, in the call to reduce, thrust::multiplies<float>() is calling the default constructor to construct a variable of class thrust::multiplies<float>, and passing it to reduce.
reduce will treat its argument as a function name and call it with an argument, triggering operator().
You may also create your own variable of that class, e.g., thrust::multiplies<float> foo. Then you just say foo in the argument list, not foo().
The optimizing compiler will replace the operator() function call with the defining expression and then continue optimizing. So, there is no overhead, unlike if you passed in a pointer to a function.
Sometimes, e.g., in saxpy.cu, you see saxpy_functor(A).
The class saxpy_functor has a constructor taking one argument.
saxpy_functor(A) constructs and returns a variable of class saxpy_functor and stores A in the variable.
The class also overloads operator().
(Let's call the new variable foo). foo() calls operator() for foo; its execution uses the stored A.
Effectively, we did a closure of saxpy_functor; this is, we bound a property and returned a new, more restricted, variable or class.
I found and reported a bug in version 100904. This version does not work with OpenMP. It was immediately closed because they already knew about it. They just hadn't told us users.
Awhile ago I reported a bug in nvcc, where it went into an infinite loop for a certain array size. The next minor version of CUDA was released the next day.
Two observations:
I'm good at breaking SW by expecting it to meet its specs.
Nvidia is responsive, which I like.
Table of contents
March 21-24.
mostly free.
I encourage you to browse around.
There'll probably be a homework or 2 based on this.
There are many possible levels of virtualization.
At a low level, one might emulate the HW. This is quite flexible but too slow.
At a higher level, a basic OS runs separate virtual machines, each with its own file system.
Harmless machine instructions execute normally.
Powerful ones are trapped and emulated.
This requires a properly designed instruction set.
IBM has been doing this commercially for 40 years, with something originally called CP/CMS.
I think that IBM lucked out with their instruction set design, and didn't plan it.
Well-behaved clients might have problematic code edited before running, to speed the execution.
I think that Vmware does that.
It seems that compute-intensive clients might have almost no overhead.
However, the emulated file system can be pretty slow.
With Vmware, several clients can all be different OSs, and the host can be any compatible OS.
E.g., I've had a linux vmware host simultaneously running both linux and windows clients.
In linux, root no longer has infinite power.
The next level of virtualization has an nontrivial host OS, but separates the clients from each other.
They see a private view of the process space, file system, and other resources.
This is lighter weight, e.g., quicker to start a VM and less overhead.
The host and client must be the same OS.
This might be called paravirtualization.
Linux supports this with things like namespace isolation and control groups (cgroups). Wikipedia et al describe this.
Ubuntu snaps do something like this.
E.g., firefox is distributed as a snap to increase isolation and security.
However starting firefox now takes 15 sec.
The next level up is the normal linux security.
You can see all the processes and similar resources.
The file system has the usual protections.
This is hard to make secure when doing something complicated.
How do I protect myself from firefox going bad?
It's easy to describe what it should be allowed to do, but almost impossible to implement.
That includes using apparmor etc.
Who guards the guards? I get spammed at a unique address that I used only to register with apparmor.
In theory, packaging an app in a virtual machine has fewer dependencies and is more secure.
You can run the vm w/o changes on different hosts.
A Vmware client can run w/o change on both linux and windows hosts.
You can run a client on your own hardware, then spill over to a commercial cloudy platform when necessary.
Docker is a popular lightweight virtualization system, which Nvidia uses to distribute SW.
Docker runs images that define virtual machines.
Docker images share resources with the host, in a controlled manner.
For simple images, which is not nvidia/cuda, starting the image is so cheap that you can do it to run one command, and encapsulate the whole process in a shell function.
Docker is worth learning, apart from its use by Nvidia for parallel computing. You might also look up Kubernetes.
More info:
I installed docker on parallel to run nvidia images like pgc++. Then I removed it because it wasn't necessary, was complicated, and it was insecure.
Traditional top level function
auto add(int a, int b) { return a+b;}
You can pass this to a function. This really passes a pointer to the function. It doesn't optimize across the call.
These have global scope.
Note auto. It's underused.
Overload operator() in a new class
Each different variable of the class is a different function. The function can use the variable's value. This is a closure.
This is local to the containing block.
This form optimizes well.
Lambda, or anon function.
auto add = [](int a, int b) { return a+b;};
This is local to the containing block.
This form optimizes well.
Placeholder notation.
As an argument in, e.g., thrust transform, you can do this:
transform(..., _1+_2);
This is nice and short.
As this is implemented by overloading the operators, the syntax of the expression is limited to what was overloaded.
Thrust is an API that looks like STL. Its backend can be GPU-CUDA, OpenMP, TBB, or sequential host-based code.
Functional-programming philosophy.
Easier programming, once you get used to it.
Source code is compact.
Compiled code runs efficiently.
Uses some unusual C++ techniques, like overloading operator().
Since the Stanford slides were created, Thrust has adopted unified addressing, so that pointers know whether they are host or device.
Links:
https://developer.nvidia.com/thrust Developer doc
Introduction to GPU Programming with CUDA and Thrust 1:18:19
All the preceding contains other interesting links.
You just want to run a big application, and figure a parallel computer might be useful.
You don't care for the theory.
Then try to implement your application on using existing parallel apps like Matlab or Mathematica.
Parallel matrix ops and Fourier transforms are already available.
Due Mon 2022-03-14.
Submit the answers to Gradescope.
You may do homeworks in teams of 2 students. Create a gradescope team and make one submission with both names.
Read the NVIDIA Ampere info linked on Class 13, and other places should you choose to look around.
Write a page on what's new, interesting, and different about Ampere.
Table of contents
This is from https://developer.nvidia.com/teaching-kits-downloads
My local copy of what I'm using is in /parclass/2022/files/nvidia/GPU-Teaching-Kit-Accelerated-Computing.zip
Start at Module 18 - MPI, then 23 - Dynamic parallelism.
Module 20.
Apple's competition to CUDA.
is largely CUDA but they changed the names and made it look like OpenGL and clunkier.
not interesting so long as Nvidia is dominant.
Generally more up-to-date and accurate, but drier. A little disorganized because it keeps growing. The root is here: https://docs.nvidia.com/
Two major relevant sets are:
We've seen almost everything, except:
Texture and surface maps.
ML HW like A=BC+D for 4x4 matrices.
Ray tracing HW, to compute a ray's intersections with boxes.
Cooperative groups: with Ampere, subsets of a warp can synchronize.
Subsets of a GPU can be defined as virtual GPUS, which are walled off from each other.
Memory can be compressed when stored, making a space-time tradeff.
The terminology CUDA core is obsolete. Now, they say that an SM has, perhaps 32 single float units, 32 integer units, 32 CUDA instruction dispatchers, and 16 double float units, etc. Each unit operates independently.
As always, this is as I understand it, and could be wrong. Nvidia uses their own terminology inconsistently. They may use one name for two things (E.g., Tesla and GPU), and may use two names for one thing (e.g., module and accelerator). As time progresses, they change their terminology.
At the bottom is the hardware micro-architecture. This is an API that defines things like the available operations. The last several Nvidia micro-architecture generations are, in order, Tesla (which introduced unified shaders), Fermi, Kepler, Maxwell (introduced in 2014), Pascal (2016), and Volta (2018).
Each micro-architecture is implemented in several different microprocessors. E.g., the Kepler micro-architecture is embodied in the GK107, GK110, etc. Pascal is GP104 etc. The second letter describes the micro-architecture. Different microprocessors with the same micro-architecture may have different amounts of various resources, like the number of processors and clock rate.
To be used, microprocessors are embedded in graphics cards, aka modules or accelerators, which are grouped into series such as GeForce, Quadro, etc. Confusingly, there is a Tesla computing module that may use any of the Tesla, Fermi, or Kepler micro-architectures. Two different modules using the same microprocessor may have different amounts of memory and other resources. These are the components that you buy and insert into a computer. A typical name is GeForce GTX1080.
There are many slightly different accelerators with the same architecture, but different clock speeds and memory, e.g. 1080, 1070, 1060, ...
The same accelerator may be manufactured by different vendors, as well as by Nvidia. These different versions may have slightly different parameters. Nvidia's reference version may be relatively low performance.
The term GPU sometimes refers to the microprocessor and sometimes to the module.
There are at least four families of modules: GeForce for gamers, Quadro for professionals, Tesla for computation, and Tegra for mobility.
Nvidia uses the term Tesla in two unrelated ways. It is an obsolete architecture generation and a module family.
Geoxeon has a (Maxwell) GeForce GTX Titan and a (Kepler) Tesla K20xm. Parallel has a (Volta) RTX 8000 and (Pascal) GeForce GTX 1080. We also have an unused (Kepler) Quadro K5000.
Since the highest-end (Tesla) modules don't have video out, they are also called something like compute modules.
Here is an example of the wide range of Nvidia GPU speeds; all times are +-20%.
The Quadro RTX 8000 has 4608 CUDA cores @ 1.77GHz and 48GB of memory. matrixMulCUBLAS runs at 5310 GFlops. The specs claim 16 TFlops. However those numbers understate its capabilities because it also has 576 Tensor cores and 72 ray tracing cores to cast 11G rays/sec.
The GeForce GTX 1080 has 2560 CUDA cores @ 1.73GHz and 8GB of memory. matrixMulCUBLAS runs at 3136 GFlops. However the reported time (0.063 msec) is so small that it may be inaccurate. The quoted speed of the 1080 is about triple that. I'm impressed that the measured performance is so close.
The Quadro K2100M in my Lenovo W540 laptop has 576 CUDA cores @ 0.67 GHz and 2GB of memory. matrixMulCUBLAS runs at 320 GFlops. The time on the GPU was about .7 msec, and on the CPU 600 msec.
It's nice that the performance almost scaled with the number of cores and clock speed.
CUDA has a capability version, whose major number corresponds to the micro-architecture generation. Kepler is 3.x. The K20xm is 3.5. The GTX 1080 is 6.1. The RTX 8000 is 7.5. Here is a table of the properties of different compute capabilities. However, that table is not completely consistent with what deviceQuery shows, e.g., the shared memory size.
nvcc, the CUDA compiler, can be told which capabilities (aka architectures) to compile for. They can be given as a real architecture, e.g., sm_61, or a virtual architecture. e.g., compute_61.
The CUDA driver and runtime also have a software version, defining things like available C++ functions. The latest is 10.1. This is unrelated to the capability version.
With CUDA, the dominant problem in program optimization is optimizing the data flow. Getting the data quickly to the cores is harder than processing it. It helps big to have regular arrays, where each core reads or writes a successive entry.
This is analogous to the hardware fact that wires are bigger (hence, more expensive) than gates.
That is the opposite optimization to OpenMP, where having different threads writing to adjacent addresses will cause the false sharing problem.
has links to other Nvidia docs.
can be a little old.
Here's a brief overview of my understanding of the various places that you can assign memory in a program.
Static. Define a fixed-size array global array. The variable is constructed at compile time, so accesses might perhaps be faster. Global vars with non default initial values increase the executable file size. If they're large enough, you need to use the compiler option -mcmodel=medium or -mcmodel=large. They cause the compiler to generate wider addresses. I don't know the effect on the program's size or speed, but suspect that it's small.
Stack. Define local arrays, that are created and freed as the routine is entered and exited. Their addresses relative to the base of this call frame may be constant. The default stack size is 8MB. You can increase this with the command ulimit or in the program as shown in stacksize.cc. I believe that in OpenMP, the max stacksize may be allocated when each thread is created. Then, a really big stackssize might have a penalty.
Heap. You use new and destroy. Variables are constructed whenever you want. The more objects on the heap, the more time that each new or destroy takes. If you have lots of objects consider using placement new or creating an array of them.
For CUDA, some variables must be on the heap.
I like to use static, then stack, and heap only when necessary. However, allocating few but large, blocks on the heap is also fast.
Google's allocator is noticeably better than the default one. To use it, link your programs with -ltcmalloc. You can often use it on an existing executable foo thus:
LD_PRELOAD="/usr/lib/libtcmalloc.so" foo
I found it to save 15% to 30% in time.
Another memory concern is speed. Parallel has a NUMA (Non Uniform Memory Architecture). It has two 14-core Xeons. Each core has 128GB of main memory. Although all 256GB are in a common address space, accessing memory on same core as the thread is running on is faster.
The following is what I think based on some research, but may be wrong: A 4KB page of memory is assigned to a specific core when it is first written (not when it is reserved). So, each page of a large array may be on a different core. This can be used to optimize things. This gets more fun with 8-processor systems.
All that is separate from cache issues.
You can also assign your OpenMP threads to specific cores. This affects speed in ways I don't understand. The issues are resource sharing vs conflicts.
Here's a summary of the Nvidia Pascal GP104 GPU architecture as I understand it. It's more compact than I've found elsewhere. I'll add to it from time to time. Some numbers are probably wrong.
The host is the CPU.
The device is the GPU.
The device contains 20 streaming multiprocessors (SMs).
Different GPU generations have used the terms SMX or SMM.
A thread is a sequential program with private and shared memory, program counter, etc.
Threads are grouped, 32 at a time, into warps.
Warps of threads are grouped into blocks.
Often the warps are only implicit, and we consider that the threads are grouped directly into blocks.
That abstract hides details that may be important; see below.
Blocks of threads are grouped into a grid, which is all the threads in the kernel.
A kernel is a parallel program executing on the device.
The kernel runs potentially thousands of threads.
A kernel can create other kernels and wait for their completion.
There may be a limit, e.g., 5 seconds, on a kernel's run time.
Thread-level resources:
Each thread can use up to 255 fast registers. Registers are private to the thread.
All the threads in one block have their registers allocated from a fixed pool of 65536 registers. The more registers that each thread uses, the fewer warps in the block can run simultaneously.
Each thread has 512KB slow local memory, allocated from the global memory.
Local memory is used when not enough registers are available, and to store thread-local arrays.
Warp-level resources:
Threads are grouped, 32 at a time, into warps.
Each warp executes as a SIMD, with one instruction register. At each cycle, every thread in a warp is either executing the same instruction, or is disabled. If the 32 threads want to execute 32 different instructions, then they will execute one after the other, sequentially.
If you read in some NVidia doc that threads in a warp run independently, then continue reading the next page to get the info mentioned in the previous paragraph.
If successive instructions in a warp do not depend on each other, then, if there are enough warp schedulers available, they may be executed in parallel. This is called Instruction Level Parallelism (ILP).
For an array in local memory, which means that each thread will have its private copy, the elements for all the threads in a warp are interleaved to potentially increase the I/O rate.
Therefore your program should try to have successive threads read successive words of arrays.
A thread can read variables from other threads in the same warp, with the shuffle instruction. Typical operation are to read from the K-th next thread, to do a butterfly permutation, or to do an indexed read. This happens in parallel for the whole warp, and does not use shared memory.
A warp vote combines a bit computed by each thread to report results like all or any.
Block-level resources:
A block may contain up to 1024 threads.
Each block has access to 65536 fast 32-bit registers, for the use of its threads.
Each block can use up to 49152 bytes of the SM's fast shared memory. The block's shared memory is shared by all the threads in the block, but is hidden from other blocks.
Shared memory is basically a user-controllable cache of some global data. The saving comes from reusing that shared data several times after you loaded it from global memory once.
Shared memory is interleaved in banks so that some access patterns are faster than others.
Warps in a block run asynchronously and run different instructions. They are scheduled and executed as resources are available.
However they are all running the same instruction sequence, perhaps at different points in it.
That is call SPMD, single program multiple data.
The threads in a block can be synchonized with __syncthreads().
Because of how warps are scheduled, that can be slow.
The threads in a block can be arranged into a 3D array, up to 1024x1024x64.
That is for convenience, and does not increase performance (I think).
I'll talk about textures later.
Streaming Multiprocessor (SM) - level resources:
Each SM has 128 single-precision CUDA cores, 64 double-precision units, 32 special function units, and 32 load/store units.
In total, the GPU has 2560 CUDA cores.
A CUDA core is akin to an ALU. The cores, and all the units, are pipelined.
A CUDA core is much less powerful than one core of an Intel Xeon. My guess is 1/20th.
Beware that, in the CUDA C Programming Guide, NVidia sometimes calls an SM a core.
The limited number of, e.g., double precision units means that an DP instruction will need to be scheduled several times for all the threads to execute it. That's why DP is slower.
Each SM has 4 warp schedulers and 8 instruction dispatch units.
64 warps can simultaneously reside in an SM.
Therefore up to 32x64=2048 threads can be executed in parallel by an SM.
Up to 16 blocks that can simultaneously be resident in an SM.
However, if each block uses too many resources, like shared memory, then this number is reduced.
Each block sits on only one SM; no block is split. However a block's warps are executed asynchronously (until synced).
Each SM has 64KiB (?) fast memory to be divided between shared memory and an L1 cache. Typically, 48KiB (96?) is used for the shared memory, to be divided among its resident blocks, but that can be changed.
The 48KB L1 cache can cache local or global memory.
Each SM has a read-only data cache of 48KB to cache the global constant memory.
Each SM has 8 texture units, and many other graphics capabilities.
Each SM has 256KB of L2 cache.
Grid-level resources:
The blocks in a grid can be arranged into a 3D array. up to \((2^{31}-1, 2^{16}-1, 2^{16}-1)\).
Blocks in a grid might run on different SMs.
Blocks in a grid are queued and executed as resources are available, in an unpredictable parallel or serial order. Therefore they should be independent of each other.
The number of instructions in a kernel is limited.
Any thread can stop the kernel by calling assert.
Device-level resources:
There is a large and slow 48GB global memory, which persists from kernel to kernel.
Transactions to global memory are 128 bytes.
Host memory can also be memory-mapped into global memory, although the I/O rate will be lower.
Reading from global memory can take hundreds of cycles. A warp that does this will be paused and another warp started. Such context switching is very efficient. Therefore device throughput stays high, although there is a latency. This is called Thread Level Parallelism (TLP) and is a major reason for GPU performance.
That assumes that an SM has enough active warps that there is always another warp available for execution. That is a reason for having warps that do not use all the resources (registers etc) that they're allowed to.
There is a 2MB L2 cache, for sharing data between SMs.
There is a 64KiB Small and fast global constant memory, , which also persists from kernel to kernel. It is implemented as a piece of the global memory, made fast with caches.
(Again, I'm still resolving this apparent contradiction).
Grid Management Unit (GMU) schedules (pauses, executes, etc) grids on the device. This is more important because grids can start other grids (Dynamic Parallelism).
Hyper-Q: 32 simultaneous CPU tasks can launch kernels into the queue; they don't block each other. If one kernel is waiting, another runs.
CUDA Work Distributor (CWD) dispatches 32 active grids at a time to the SMs. There may be 1000s of grids queued and waiting.
GPU Direct: Other devices can DMA the GPU memory.
The base clock is 1607MHz.
GFLOPS: 8873.
Memory bandwidth: 320GB/s
GPU-level resources:
Being a Geforce product, there are many graphics facilities that we're not using.
There are 4 Graphics processing clusters (GPCs) to do graphics stuff.
Several perspective projections can be computed in parallel, for systems with several displays.
There's HW for texture processing.
Generational changes:
With each new version, Nvidia tweaks the numbers. Some get higher, others get lower.
E.g., Maxwell had little HW for double precision, and so that was slow.
Pascal's clock speed is much higher.
Refs:
The CUDA program deviceDrv.
http://developer.download.nvidia.com/compute/cuda/compute-docs/cuda-performance-report.pdf
Better Performance at Lower Occupancy, Vasily Volkov, UC Berkeley, 2010.
https://www.pgroup.com/lit/articles/insider/v2n1a5.htm - well written but old.
(I'll keep adding to this. Suggestions are welcome.)
CUDA function qualifiers:
__global__ device function called from host, starting a kernel.
__device__ device function called from device function.
__host__ (default) host function called from host function.
CUDA variable qualifiers:
__shared__
__device__ global
__device__ __managed__ automatically paged between host and device.
__constant__
(nothing) register if scalar, or local if array or if no more registers available.
If installing CUDA on your machine, this repository seems best:
http://developer.download.nvidia.com/compute/cuda/repos/ubuntu1604/x86_64
That includes the Thrust headers but not example programs.
aka overlay FS, translucent FS.
If a, b are directories, and m is an empty directory, then
unionfs -o cow a=RW:b m
makes m to be a combo of a and b, with a being higher priority
Writing a file into m writes it in a.
Changing a file in b writes the new version into a
Deleting a file in b causes a white-out note to be stored in a.
Unmount it thus:
fusermount -u m
None of this requires superuser.
Application: making a read-only directory into a read-write directory.
Note: IBM had a commercial version of this idea in its CP/CMS OS in the 1960s.
On parallel.ecse in /parclass/2022/files/stanford/
Very well done.
Duplicative of the nvidia course, so I'm not spending time on them this year.
However, worth your time to scan them.
Next topic.
Table of contents
This is from https://developer.nvidia.com/teaching-kits-downloads
My local copy of what I'm using is in /parclass/2022/files/nvidia/GPU-Teaching-Kit-Accelerated-Computing.zip
Start at Module 16.
We'll cover this in detail later.
Table of contents
This is from https://developer.nvidia.com/teaching-kits-downloads
My local copy of what I'm using is in /parclass/2022/files/nvidia/GPU-Teaching-Kit-Accelerated-Computing.zip
Do Module 12 – Floating-Point Considerations: Lecture 12.2 - Numerical Stability.
Then start at Module 14 – Efficient Host-Device Data Transfer: Lecture 14.1 - Pinned Host Memory
Table of contents
This is from https://developer.nvidia.com/teaching-kits-downloads
My local copy of what I'm using is in /parclass/2022/files/nvidia/GPU-Teaching-Kit-Accelerated-Computing.zip
Starting at Module 9 - Reduction.
Due next Tues 2022-02-22.
Submit the answers to Gradescope.
You may do homeworks in teams of 2 students. Create a gradescope team and make one submission with both names.
Use the ideas we've seen to implement a parallel text histogram program, both on the multicore CPU and on the manycore GPU. Do it with and without private subtotal arrays.
Use current tools, like managed memory on the device.
Compare the performance processing a big file. Separate the time to read the data into memory from the time to process it.
Try to use debugging tools to find hotspots.
To reduce file reading time, before starting your program, put the data file into the in-core filesystem in /dev/shm/, aka /ds .
Do Chapter 9 of the free online Nvidia text, Question 4.
Do Chapter 9, Question 5.
Table of contents
This is from https://developer.nvidia.com/teaching-kits-downloads
My local copy of what I'm using is in /parclass/2022/files/nvidia/GPU-Teaching-Kit-Accelerated-Computing.zip
Continuing from Module 7 - Parallel Computation Patterns (Histogram)
Read Chapters 5 to 9.
Table of contents
This is from https://developer.nvidia.com/teaching-kits-downloads
My local copy of what I'm using is in /parclass/2022/files/nvidia/GPU-Teaching-Kit-Accelerated-Computing.zip
Continuing from Lecture 3.5 – CUDA Parallelism Model.
Read Chapter 3:
Module 3 - CUDA Parallelism Model/EBook Chapters/3rd-Edition-Chapter03-scalable-parallel-execution.pdf
Exercise from p66:
Assume that a CUDA device allows up to 8 blocks and 1024 threads per SM, whichever becomes a limitation first. Furthermore, it allows up to 512 threads in each block. For image blur, should we use 8 × 8, 16 × 16, or 32 × 32 thread blocks?
I want to project stuff and run programs during the class. Maybe even record the lectures.
If I project my ipad, I can project anything on its screen, like goodnotes and firefox. However it's hard to ssh to parallel, and webex doesn't want to record the screen. The ipad is not a full multiprocessing OS, perhaps to avoid taking customers from the macbook.
If I project my linux thinkpad x12, I can also project anything on its screen, including ssh to parallel. However the projector forces the thinkpad's display to a different resolution that the touchscreen doesn't understand. Where I touch the screen with my stylus is not where the mark is made on the display. That makes writing legibly impossible.
Webex can also broadcast and (usually) record the x12 screen.
When lecturing from home, I use a p73 thinkpad with an NVidia GPU that can remotely display the ipad as another window with uxplay. Then I write on the ipad, and firefox and ssh from the thinkpad, while webex broadcasts the thinkpad. That all works until something crashes (usually in the middle of a lecture).
Uxplay doesn't work on the x12 probably because it has integrated graphics.
Long time ago, RPI's classrooms had blackboards on every available wall, so the prof could move from blackboard to blackboard. The only tech problem was that the chalk was messy. Also recording was by the students writing in their notebooks.
At one time, the classrooms had transparency projectors with rolls of plastic. We wrote on them and rolled them up. For several years, I had the TAs xerox the marked up roll so I had a copy.
At another time I wrote on individual plastic sheets on the transparency projector. Those were easier to xerox.
I'm still thinking of other solutions. Suggestions are welcome.
Table of contents
This is from https://developer.nvidia.com/teaching-kits-downloads
My local copy of what I'm using is in /parclass/2022/files/nvidia/GPU-Teaching-Kit-Accelerated-Computing.zip
Continuing from Module 2 - Introduction to CUDA C/Slides/Lecture-2-2-cuda-data-allocation-API.pdf
There is more information in the included book chapter.
There are some version skew problems with some of the debugging examples.
Part of the class will be a chance for you to play with some examples, some from a Stanford course. I'll walk around and help.
/parclass/2022/files/stanford/tutorials
Due next Mon 2022-02-14
Submit the answers to Gradescope.
You may do homeworks in teams of 2 students. Create a gradescope team and make one submission with both names.
Do all 9 exercises at the end of chapter 2 of
Module 2 - Introduction to CUDA C/3rd-Edition-Chapter02-data-parallel-computing.pdf
They're in gradescope.
Table of contents
I recommend OpenACC for Programmers: Concepts and Strategies By Sunita Chandrasekaran, Guido Juckeland. Informit, Amazon. Github.
This is from https://developer.nvidia.com/teaching-kits-downloads
My local copy of what I'm using is in /parclass/2022/files/nvidia/GPU-Teaching-Kit-Accelerated-Computing.zip
We'll spend several classes on this topic, speed reading through the slides.
This is today's programmer productivity tip, unrelated to parallel computing but generally useful.
Rather than explicitly extracting large zip archives or tarballs in order to read files in them, I use archivemount to create a virtual file system. This saves disk space. It doesn't stress git as much (fewer files). When reading, the I/O time is insignificantly increased. For some formats, you can even write. You can have more confidence that the zip file wasn't changed, than in a directory with perhaps hundreds of files.
Table of contents
Would you like to point my phone at me lecturing and attempt broadcast and record it on webex?
We'll continue the lab to help you get started.
To compile foo.cc into foo, you can say make foo .
Envars, e.g.
OMP_NUM_THREADS=2
My file common has useful stuff.
ulimit -s 100000000 can be useful.
In paged virtual memory, stacks are more flexible than apparent.
Dunno if OMP_STACKSIZE does anything.
OpenMP on GPU:
https://www.nersc.gov/users/training/events/openmp-device-offload-sept-2021/
I've just given you an intro.
Newer than OpenMP.
More abstract; compiler tries to determine what to parallelize.
Your algorithm needs to be parallelizable. That's the hardest part.
I'll go quickly through the tutorials from https://www.openacc.org/events/openacc-online-course-2018 .
You can read more on your own.
Leave you wishing that I'd taught you more.
Table of contents
Amanda. OpenACC
Ivan. Portability issues in parallel designs
Ryan. Machine learning
Saaif. Gaming
This class is partly a lab to help you get started.
Due next Thu 2022-02-03
Submit the answers to Gradescope.
You may do homeworks in teams of 2 students. Create a gradescope team and make one submission with both names.
(10 pts) Investigate the effects of varying numbers of threads on elapsed time and errors. Test sum, sum_atomic, sum_crit, sum_reduc with 1, 2, 8, 16, 32, 64 threads. Make appropriate mods so that the programs run slowly enough for you to measure the elapsed times.
(10 pts) Look at https://docs.nvidia.com/hpc-sdk/compilers/hpc-compilers-user-guide/index.html#openmp-use
Try to modify sum_atomic and sum_crit to run on the GPU (using OpenMP). Full points for a good try, whether or not you succeed.
(5 pts) Why is matmul4 segfaulting on n=1000? Try to fix it.
Total: 25
Table of contents
https://hpc.llnl.gov/documentation/tutorials/introduction-parallel-computing-tutorial
Starting at SPMD and MPMD. You can study the rest on your own. We'll see what it teaches by example.
The full name is parallel.ecse.rpi.edu .
I've set up accounts for you.
Your user name is your RCSID. (However your UID is unrelated to your RCS UID).
Its accounts are completely separate from RCS; I will use the same userid for convenience. Changing your password on parallel does not affect your RCS password, and vv.
In class, I'll tell you how I created the passwords. Change yours today or I'll disable your account.
From a linux machine, connect with ssh. I recommend setting up a key pair.
From outside RPI, you need to use the vpn.
Your default shell is bash. Feel free to change it. I use zsh. If changing your shell, look at these last lines of /etc/bash.bashrc and adapt them accordingly.:
# nvidia; from https://docs.nvidia.com/hpc-sdk/hpc-sdk-install-guide/index.html NVARCH=`uname -s`_`uname -m`; export NVARCH NVCOMPILERS=/opt/nvidia/hpc_sdk MANPATH=$MANPATH:$NVCOMPILERS/$NVARCH/22.1/compilers/man PATH=$PATH:$NVCOMPILERS/$NVARCH/22.1/compilers/bin CUDA_PATH=$NVCOMPILERS/$NVARCH/22.1/cuda LD_LIBRARY_PATH=$LD_LIBRARY_PATH:$CUDA_PATH/lib64 CUDA_INCLUDE=/opt/nvidia/hpc_sdk/Linux_x86_64/22.1/math_libs/11.5/targets/x86_64-linux/include
Files used in the course will generally be put into /parclass2022/files/ .
Play with your account, e.g., compiling a small program.
Use NVIDIA's C++ compiler nvc++. You're welcome to download it from NVIDIA and install it on your own machine. That requires a free NVIDIA developer account.
nvidia-smi will describe the GPUs.
Versions: OpenMP is a living project, that is being updated every few years.
There is a lot of free online documentation, examples, and tutorials.
New features are regularly added, such as support for GPU backends.
However the documentation and implementations lags.
I'll use obsolete documentation when it's good enough. It's all still correct and fine for most applications.
Then I'll present recent some additions.
We'll see some examples running on parallel, at /parclass2022/files/openmp/rpi
We saw that running even a short OpenMP program could be unpredictable. The reason for that was a big lesson.
There are no guarantees about the scheduling of the various threads. There are no guarantees about fairness of resource allocation between the threads. Perhaps thread 0 finishes before thread 1 starts. Perhaps thread 1 finishes before thread 2 starts. Perhaps thread 0 executes 3 machine instructions, then thread 1 executes 1 instruction, then thread 0 executes 100 more, etc. Each time the process runs, something different may happen. One thing may happen almost all the time, to trick you.
This happened during the first space shuttle launch in 1981. A 1-in-70 chance event prevented the computers from synchonizing. This had been observed once before during testing. However when they couldn't get it to happen again, they ignored it.
This interlacing of different threads can happen at the machine instruction level. The C statement K=K+3 can translate into the machine instructions
LOAD K; ADD 3; STORE K.
Let's color the threads red and blue.
If the threads interlace thus:
LOAD K; ADD 3; STORE K; LOAD K; ADD 3; STORE K
then K increases by 6. If they interlace thus:
LOAD K; ADD 3; LOAD K; ADD 3; STORE K; STORE K;
then K increases by 3.
This can be really hard to debug, particularly when you don't know where this is happening.
One solution is to serialize the offending code, so that only one thread executes it at a time. The limit would be to serialize the whole program and not to use parallel processing.
OpenMP has several mechanisms to help.
A critical pragma serializes the following block. There are two considerations.
You lose the benefits of parallelization on that block.
There is an overhead to set up a critical block that I estimate might be 100,000 instructions.
An atomic pragma serializes one short C statement, which must be one of a small set of allowed operations, such as increment. It matches certain atomic machine instructions, like test-and-set. The overhead is much smaller.
If every thread is summing into a common total variable, the reduction pragma causes each thread to sum into a private subtotal, and then sum the subtotals. This is very fast.
Another lesson is that sometimes you can check your program's correctness with an independent computation. For the trivial example of summing \(i^2\), use the formula
\(\sum_{i=1}^N i^2 = \frac{N(N+1)(2N+1)}{6}\).
There is a lot of useful algebra if you have the time to learn it. I, ahem, learned this formula in high school.
Another lesson is that even when the program gives the same answer every time, it might still be consistently wrong.
Another is that just including OpenMP facilities, like -fopenmp, into your program slows it down even if you don't use them.
Another is that the only meaningful time metric is elapsed real time. One reason that CPU time is meaningless is the OpenMP sometimes pauses a thread's progress by making it do a CPU-bound loop. (That is also a common technique in HW.)
Also note that CPU times can vary considerably with successive executions.
Also, using too many threads might increase the real time.
Finally, floating point numbers have their own problems. They are an approximation of the mathematical concept called the real number field. That is defined by 11 axioms that state obvious properties, like
A+B=B+A (commutativity) and
A+(B+C)=(A+B+C) (associativity).
This is covered in courses on modern algebra or abstract algebra.
The problem is that most of these are violated, at least somewhat, with floats. E.g.,
\(\left(10^{20}-10^{20}\right)+1=0+1=1\) but
\(10^{20}+\left(-10^{20}+1\right)=10^{20}-10^{20}=0\).
Therefore when threads execute in a different order, floating results might be different.
There is no perfect solution, though using double is a start.
On modern CPUs, double is just as fast as float.
However it's slower on GPUs.
How much slower can vary. Nvidia's Maxwell line spent very little real estate on double precision, so it was very slow. In contrast, Maxwell added a half-precision float, for implementing neural nets. Apparently neural nets use very few significant digits.
Nvidia's newer Pascal line reverted this design choice and spends more real estate on implementing double. parallel.ecse's GPU is Pascal.
Nvidia's Volta generation also does doubles well;
It also requires moving twice the data, and data movement is often more important than CPU time.
The large field of numerical analysis is devoted to finding solutions, with more robust and stable algorithms.
Summing an array by first sorting, and then summing the absolutely smaller elements first is one technique. Inverting a matrix by pivoting on the absolutely largest element, instead of on \(a_{11}\) is another.
Now that we've seen some examples, look at the LLNL tutorial. https://hpc-tutorials.llnl.gov/openmp/
Another nice, albeit obsolescent and incomplete, ref: Wikibooks OpenMP.
Its tasks examples are interesting.
OpenMP tasks
While inside a pragma parallel, you queue up lots of tasks - they're executed as threads are available.
My example is tasks.cc with some variants.
taskfib.cc is modified from an example in the OpenMP spec.
It shows how tasks can recursively spawn more tasks.
This is only an example; you would never implement fibonacci that way. (The fastest way is the following closed form. In spite of the sqrts, the expression always gives an integer. )
\(F_n = \frac{\left( \frac{1+\sqrt{5}}{2} \right) ^n - \left( \frac{1-\sqrt{5}}{2} \right) ^n }{\sqrt{5}}\)
Spawning a task is expensive; we can calculate the cost.
omps.cc - read assorted OMP variables.
Since I/O is not thread-safe, in hello_single_barrier.cc it has to be protected by a critical and also followed by a barrier.
OpenMP sections - my example is sections.cc with some variants.
Note my cool way to print an expression's name followed by its value.
Note the 3 required levels of pragmas: parallel, sections, section.
The assignment of sections to threads is nondeterministic.
IMNSHO, OpenMP considerably easier than pthreads, fork, etc.
http://stackoverflow.com/questions/13788638/difference-between-section-and-task-openmp
OpenMP is weak on compiling for GPUs. So....
Due next Thu 2022-01-27.
Submit the answers to Gradescope.
You may do homeworks in teams of 2 students. Create a gradescope team and make one submission with both names.
Short but informative answers are fine.
Each is 5 points.
Why is one CUDA core less powerful than one Intel Xeon core?
If 10% of a program cannot be parallelized, then what is the max speedup even with infinite parallelism?
"For short running parallel programs, there can actually be a decrease in performance compared to a similar serial implementation.". Why?
What is the difference between strong scaling and weak scaling?
What is cache coherency?
Define "embarrassingly parallel".
How many CUDA cores does the RTX 8000 have?
Why have machine cycle speeds stopped increasing?
Give an advantage and a disadvantage of shared memory versus distributed memory for parallel computation.
Although I am talking about parallel computing on NVIDIA, the AMD Radeon is a possible competitor. Write a paragraph (say, 100 words) comparing the 2.
Total: 50 pts.
Table of contents
Read the syllabus.
Read https://computing.llnl.gov/tutorials/parallel_comp/ for an intro to parallel computing.
Some points:
Parallel computing is decades old; there were commercial machines in the 1980s. I directed two PhD theses in parallel geometry then. However, then, clocks speeds were increasing so serial machines were more interesting.
Now: physics limits to processor speed.
History of Nvidia.
Curtis Priem had designed graphics HW for both IBM and Sun Microsystems. (For awhile Sun was THE Unix workstation company. They used open standards and had the best price / performance.)
Nvidia designed gaming graphics accelerators...
that just happened to be parallel coprocessors...
that started to be used for nongraphics parallel processing because of their value.
Nvidia noticed that and added more capability, e.g., double precision IEEE floats, to serve that market.
Currently, some of the highest performance Nvidia boards cannot even do graphics because they don't have video out ports.
Intel CPUs vs Nvidia CUDA cores.
Advantages and disadvantages of shared memory.
OpenMP vs CUDA.
Rate-limiting cost usually I/O not computation.
is online, due next Mon.
(10 points) Pick two topics from this list.
Describe the fastest 2 computers on the top500 list https://www.top500.org/lists/top500/2020/11/
Describe computers 3 and 4 on top500.
Describe computers 5 and 6 on top500.
Describe the Blue Gene.
Summarize python's parallel capabilities.
Summarize Matlab's parallel capabilities.
Summarize Mathematica's parallel capabilities.
Summarize MPI.
Describe an application where physics needs/uses parallel computing.
Describe an application where astronomy needs/uses parallel computing.
Describe an application where biology needs/uses parallel computing.
Give an application where machine learning needs/uses parallel computing.
Give an application where computational fluid dynamics needs/uses parallel computing.
Summarize OpenACC.
Summarize pthreads.
Other related topic of your choice.
Email me with your choices.
I'll assign you your first topic that wasn't already taken.
Two students may work together to present one talk.
Prepare and give 5-10 minute talk next Thu 1/20. We'll spread the talks out evenly. A team would jointly give one 3-5 minute talk.
I won't block your mike at 10 minutes, but will leave it to you to be responsible. There are 16 students in class. The class is 80 minutes long. If necessary, we'll run over to next Mon.
Table of contents
ECSE-4740-01 Applied Parallel Computing for Engineers
Spring term annually
3 credit hours
Mon and Thurs noon-1:20pm. Online for the 1st 2 weeks, then in class
This is intended to be a computer engineering course to provide students with knowledge and hands-on experience in developing applications software for affordable parallel processors. This course will mostly cover hardware that any lab can afford to purchase. It will cover the software that, in the prof's opinion, is the most useful. There will also be some theory.
The target audiences are ECSE seniors and grads and others with comparable background who wish to develop parallel software.
This course will have minimal overlap with parallel courses in Computer Science. We will not teach the IBM BlueGene, because it is so expensive, nor cloud computing and MPI, because most big data problems are in fact small enough to fit on our hardware.
You may usefully take all the parallel courses at RPI.
This unique features of this course are as follows:
For 2/3 of the course, use of only affordable hardware that any lab might purchase, such as Nvidia GPUs. This is currently the most widely used and least expensive parallel platform.
Emphasis on learning several programming packages, at the expense of theory. However you will learn a lot about parallel architecture.
Hardware taught, with reasons:
universally available and inexpensive, comparatively easy to program, powerful
widely available (Nvidia external graphics processors are on 1/3 of all PCs), very inexpensive, powerful, but harder to program. Good cards cost only a few hundred dollars.
It's new and hot and might some day be useful.
Software that might be taught, with reasons:
widely used, easy to use if your algorithm is parallelizable, backend is primarily multicore Xeon but also GPUs.
FP is nice, hides low level details, backend can be any major parallel platform.
easy to use parallelism for operations that they have implemented in parallel, etc.
low level access to Nvidia GPUs.
Nvidia's parallel C++ compiler.
The techniques learned here will also be applicable to larger parallel machines -- numbers 2, 3, 5, and 6 on the top 500 list use NVIDIA GPUs.
Effectively programming these processors will require in-depth knowledge about parallel programming principles, as well as the parallelism models, communication models, and resource limitations of these processors.
ECSE-2660 CANOS or equivalent, knowledge of C++,.
W. Randolph Franklin. BSc (Toronto), AM, PhD (Harvard)
online only
+1 (518) 276-6077 (forwards)
frankwr@YOUKNOWTHEDOMAIN or mailATwrfranklinDOTorg or wrfranklinATpmDOTme
Email is my preferred communication medium. I generally send email from my own domain.
Sending from non-RPI accounts are fine; I do that. But please show your name, at least in the comment field. A subject prefix of #Par is helpful. I support encrypted email, either through inline GPG, or via protonmail.
We can also use webex or facetime.
A quick way to get there is to google RPIWRF.
At 3:20 pm after my probability lecture, usually as long as anyone wants to talk. Also by appointment.
If you would like to talk with me, either individually or in a group, just mention it. We can then talk about most anything legal and ethical.
I asked to create and teach it, because I like the topic.
The homepage has lecture summaries, syllabus, homeworks, etc.
There is no required text, but the following inexpensive books may be used. I might mention others later.
Sanders and Kandrot, CUDA by example. It gets excellent reviews, although it is several years old. Amazon has many options, including Kindle and renting hardcopies.
Kirk and Hwu, 2nd edition, Programming massively parallel processors. It concentrates on CUDA.
One problem is that even recent books may be obsolete. For instance they may ignore the recent CUDA unified memory model, which simplifies CUDA programming at a performance cost. Even if the current edition of a book was published after unified memory was released, the author might not have updated the examples.
There is a lot of free material on the web, which I'll reference class by class. Because web pages are vanish so often (really!), I may cache some locally. We will start here:
This course will primarily use (remotely via ssh) parallel.ecse.rpi.edu.
Parallel has:
dual 14-core Intel Xeon E5-2660 2.0GHz
256GB of DDR4-2133 ECC Reg memory
Nvidia GPUs:
Quadro RTX 8000 with 48GB memory, 16 TFLOPS, and 4608 CUDA cores. This can do real time ray tracing.
GeForce GTX 1080 with 8GB memory and 2569 CUDA cores.
Samsung Pro 850 1TB SSD
WD Red 6TB 6GB/s hard drive
CUDA 10.1
OpenMP 4.5
Thrust
Ubuntu 21.10
Material for the class is stored in /parallel-class/ .
We might also use parallel virtual machines on the Amazon EC2. If so, you will be expected to establish an account. I expect the usage to be in the free category.
Gradescope will be used for you to submit homeworks and for us to distribute grades.
to deliver lectures, to discuss, etc.
This is only for distributing the computed (total) grades. (Gradescope has no way to upload computed grades.)
Delivering good remote courses is a work in progress. If I learn of a new tool that looks useful, I might use it.
There will be no exams.
Each student will make several in-class presentations summarizing some relevant topic.
There will be a term project.
There will be some homeworks.
You may do homeworks in teams of up to 3 students. Create a gradescope team and make one submission with all the names.
You may switch teams from homework to homework.
For the latter part of the course, most of your homework time will be spent on a term project.
You are encouraged do it in teams of any size. A team of 3 people would be expected to do twice as much work as 1 person.
You may combine this with work for another course, provided that both courses know about this and agree. I always agree.
If you are a grad student, you may combine this with your research, if your prof agrees, and you tell me.
You may build on existing work, either your own or others'. You have to say what's new, and have the right to use the other work. E.g., using any GPLed code or any code on my website is automatically allowable (because of my Creative Commons licence).
You will implement, demonstrate, and document something vaguely related to parallel computing.
Deliverables:
An implementation showing parallel computing.
An extended abstract or paper on your project, written up like a paper. You should follow the style guide for some major conference (I don't care which, but can point you to one).
A more detailed manual, showing how to use it.
A 2-minute project proposal given to the class around the middle of the semester.
A 10-minute project presentation and demo given to the class in the last week.
Some progress reports.
A write-up uploaded on the last class day. This will contain an academic paper, code and perhaps video or user manual.
Size
It's impossible to specify how many lines of code makes a good term project. E.g., I take pride in writing code that is can be simultaneously shorter, more robust, and faster than some others. See my 8-line program for testing whether a point is in a polygon: Pnpoly.
According to Big Blues, when Bill Gates was collaborating with around 1980, he once rewrote a code fragment to be shorter. However, according to the IBM metric, number of lines of code produced, he had just caused that unit to officially do negative work.
Parallel will start virtual and then switch to in-person when RPI tells me to. I intend to lecture live, with some supplementary videos. If the format changes unexpectedly, I'll email everyone.
Since this is a small class, students will be encouraged to ask questions, live and with chat during the class. I'll try to save the chat window, but it may be deleted after each class. If you cannot watch the class live, then you may miss any info presented in the chat window.
Try to share your video when talking.
If you would like, I can try to set up a webex teams space for discussions between class.
I will maintain a class blog that briefly summarizes each class. Important points will be written down.
I will attempt to record all virtual classes and upload them to mediasite for students to watch or download later. However, this fall, about 10% of classes failed to record properly.
Depending on the class composition, since it increases learning, there might be graded in-class quizzes. That means, that I would like to do this. However if it causes too much trouble for too many students then I won't.
As required by the Provost, we may post notes about you to EWS, for example, if you're having trouble doing homeworks on time, or miss an exam. E.g., if you tell me that you had to miss a class because of family problems, then I may forward that information to the Dean of Students office.
See the Student Handbook for the general policy. The summary is that students and faculty have to trust each other. After you graduate, your most important possession will be your reputation.
Specifics for this course are as follows.
You may collaborate on homeworks, but each team of 1 or 2 people must write up the solution separately (one writeup per team) using their own words. We willingly give hints to anyone who asks.
The penalty for two teams handing in identical work is a zero for both.
You may collaborate in teams of up to 3 people for the term project.
You may get help from anyone for the term project. You may build on a previous project, either your own or someone else's. However you must describe and acknowledge any other work you use, and have the other person's permission, which may be implicit. E.g., my web site gives a blanket permission to use it for nonprofit research or teaching. You must add something creative to the previous work. You must write up the project on your own.
However, writing assistance from the Writing Center and similar sources in allowed, if you acknowledge it.
The penalty for plagiarism is a zero grade.
Cheating will be reported to the Dean of Students Office.
You've seen them in your other classes. Assume that they're included here.
Since it's my desire to give you the best possible course in a topic I enjoy teaching, I welcome feedback during (and after) the semester. You may tell me or write me, or contact a third party, such as Prof James Lu, the ECSE undergrad head, or Prof John Wen, the ECSE Dept head.
Table of contents
Something went wrong; sorry.
My value to you lies on selecting what web material is worth showing to you, not in rewording it. I go through a lot of stuff and select a small fraction of it. Also, presenting the original material instead of my version of it is more accurate.
Reagan and Allan. Blue Gene
Ruisi. Physics
Ean and Tom. Top 2 computers
Kevin. 5th and 6th computers
Zach. Python's parallel capabilities
Mark & Phu. Pthreads
Adrian-James. Astronomy
Cohen. 3rd and 4th computers
Amanda. OpenACC
Ivan. Portability issues in parallel designs
Ryan. Machine learning
Saaif. Gaming
Table of contents
We'll do it. https://computing.llnl.gov/tutorials/parallel_comp/
Think about these questions.
Why have machine cycle speeds stopped increasing?
What architectures do the top 3 machines on the Top 500 list use?
Which one of the following 4 choices are most GPUs: SISD, SIMD, MISD, MIMD.
Which one of the following 4 choices are most current multicore CPUs: SISD, SIMD, MISD, MIMD.
Per Amdahl's law, if a program is 10% sequential and 90% parallelizable, what is the max speed up that can be obtained with an infinite number of parallel processors?
Intel Xeon Phi
IBM BlueGene
I do parallel geometry algorithms on large problems for CAD and GIS. See my home page. If this is interesting, talk to me. Maybe you can do something that leads to a jointly-authored paper.