Table of contents
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.
I'm open to questions and discussions about any legal ethical topic. Even after you graduate.
Table of contents
Try to use the IEEE conference format . It allows either latex or MS word. Submit the PDF paper to gradescope.
"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.
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.
Update: However, with CUDA, maybe managed variables must be on the heap.
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
Go Behind-the-Scenes of a Quantum Experiment (2:10) https://quantumexperience.ng.bluemix.net/qx/community/question?questionId=5ae975690f020500399ed39a&channel=videos
or
https://www.youtube.com/watch?v=tfZpJLdkzRU&list=PLOFEBzvs-VvpzQnlazij7cL1mjKvJTAwk&index=7
A Qubit in the Making (2:01)
https://www.youtube.com/watch?v=2pB87H3_F_c&list=PLOFEBzvs-VvpzQnlazij7cL1mjKvJTAwk&index=10&t=0s
Behold the Mighty Qubit (2:51) https://www.youtube.com/watch?v=_P7K8jUbLU0&list=PLOFEBzvs-VvpzQnlazij7cL1mjKvJTAwk&index=10
Classical and Quantum Randomness (3:39) https://www.youtube.com/watch?v=8kyJfAC4VAo&list=PLOFEBzvs-VvpzQnlazij7cL1mjKvJTAwk&index=6
Quantum Entanglement (2:21) https://www.youtube.com/watch?v=RmXasxLm43k&list=PLOFEBzvs-VvpzQnlazij7cL1mjKvJTAwk&index=5
Benchmarking Quantum Systems (1:58) https://www.youtube.com/watch?v=-7L5o-mzLqU&list=PLOFEBzvs-VvpzQnlazij7cL1mjKvJTAwk&index=8
SC18 Invited Talk: Matthias Troyer https://www.youtube.com/watch?v=97s3FEhG14g (45:46)
This is very nice. I'll show the start and then a highlight, starting at 19:00.
SC18: NVIDIA CEO Jensen Huang on the New HPC (1:44:46) https://www.youtube.com/watch?v=PQbhxfRH2H4
We'll watch the wrapup,
SCinet Must See: How to Build the World’s Fastest Temporary Computer Network Time Lapse (0:57) https://sc18.supercomputing.org/scinet-must-see-how-to-build-the-worlds-fastest-temporary-computer-network-time-lapse/
He has visited RPI more than once.
Algorithms for future emerging technologies (2:55:45) https://www.youtube.com/watch?v=TCgHNMezmZ8
We'll watch the start. You're encouraged to watch it all. It gets more technical later.
Stony Brook talk https://www.youtube.com/watch?v=D1hfrtoVZDo
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.
Table of contents
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.
Sava C & Davi P & Emil V
27 Quantum Mechanics - Measurement 8:07
Shor on, what is Shor's factoring algorithm? (2:09) https://www.youtube.com/watch?v=hOlOY7NyMfs
Hacking at Quantum Speed with Shor's Algorithm (16:35) https://www.pbs.org/video/hacking-at-quantum-speed-with-shors-algorithm-8jrjkq/
43 Quantum Mechanics - Quantum factoring Period finding (19:27) https://www.youtube.com/watch?v=crMM0tCboZU
44 Quantum Mechanics - Quantum factoring Shor's factoring algorithm (25:42) https://www.youtube.com/watch?v=YhjKWAMFBUU
Table of contents
Due last class day, Wed April 29.
See syllabus.
To the 2 students in the grad version: do more work.
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.
An overview of Grover's Algorithm. An unstructured search algorithm that can find an item in a list much faster than a classical computer can. Several sources are listed.
Can we make quantum technology work? | Leo Kouwenhoven | TEDxAmsterdam (18:19)
Introduction to Quantum Computing (18) - Grover's Algorithm: Quantum Problem Statement (4:24)
Introduction to Quantum Computing (19) - Grover's Algorithm: Outline (16:09)
Introduction to Quantum Computing (20) - Grover's Algorithm: Subspace (7:42)
Introduction to Quantum Computing (21) - Grover's Algorithm: Geometric Interpretation
Introduction to Quantum Computing (23) - Grover's Algorithm: Implementation (9:33)
Introduction to Quantum Computing (24) - Grover's Algorithm: IBM Quantum Experience (11:39)
Table of contents
Anyone who did not sign up is signed up for Thurs April 23.
I had to reinstall linux and the packages. If something is missing, please tell me. /opt, /local and your accounts should be the same. Cuda is now up to version 10.2.
The problem was a failed update that resulted in the boot image being damaged.
Table of contents
Remember to sign up.
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.
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.
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:
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.
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
Table of contents
GTC_2010_Part_2_Thrust_By_Example.pdf
Start at slide 27.
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 /parallel-class/thrust/examples-1.8/tiled_range.cu into /parallel-class/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());
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));
auto data = make_transform_iterator(make_counting_iterator(0), _1*_1);
tiled_range5.cu shows how to use a lambda instead of the _1 notation:
auto output=make_permutation_iterator(data, make_transform_iterator(make_counting_iterator(0), [](const int i){return i%N;} ));
You have to compile with --std c++11 .
This can be rewritten thus: ..
auto f = [](const int i){return i%N;}; auto output = make_permutation_iterator(data, make_transform_iterator(make_counting_iterator(0), f));
The shortest lambda is this:
auto f = [](){};
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));
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.
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
This homework is a status report on the term project.
Please include specific info about the parallel computing component.
Total: 10 pts.
Table of contents
https://wrf.ecse.rpi.edu/Teaching/parallel-s2020/posts/quantum.html
though that might change.
I don't care how qbits are implemented. That's important, even necessary, but it's someone else's problem.
This recapitulates the development of quantum computing, where the theory was worked out starting in the 1980s, long before actual quantum computers were built.
(Analogously, the theory for classical computers was worked out in the 1930s, long before real computers were built. Alan Turing contributed to both the theory and the practice.)
This presentation may have many errors, but I hope that the general tenor is not too misleading.
warmup
Classical metaphor for entanglement:
Technical details later.
Qbit, $q$.
Its state is a linear combo of two basis states, $|0>$ and $|1>$:
$q = a|0> + b|1>$ ,
where $a$ and $b$ are complex numbers, and $ | a | ^2 + | b | ^2 = 1$.
IOW, its state is a superposition of those two basis states, with those weights.
It is wrong to think that $q$ is really in one of the two states, but you don't know which one. This is the hidden variable theory. It has been proved experimentally to be false.
$q$ is really in both states simultaneously.
Alice laughed. "There's no use trying," she said: "one can't believe impossible things." "I daresay you haven't had much practice," said the Queen. "When I was your age, I always did it for half-an-hour a day. Why, sometimes I've believed as many as six impossible things before breakfast." - Through the Looking-Glass, and What Alice Found There (1871), by Lewis Carroll (Charles Lutwidge Dodgson).
You cannot observe its state, unless it is $|0>$ and $|1>$, in which case you observe $0$ or $1$. This is the classical case.
Otherwise you observe it with a measurement operator that transforms it to either $|0>$ and $|1>$, with probabilities
$| a | ^2$ and $| b | ^2 = 1$, respectively.
$a$ and $b$ are complex.
That measurement changes $q$; it no longer has its old value.
$q$, that is, $q$ 's value, can be considered to be a vector of length one: $$\begin{pmatrix} a | 0> \\ b | 1> \end{pmatrix} $$ or simply $$\begin{pmatrix}a\\b\end{pmatrix}$$.
You operate on $q$ with a matrix multiplication: $q_2 = M q$.
This changes $q$; the old value is no longer available.
No cloning: You cannot copy a qbit, but can move it.
The life cycle of a qbit:
So far, not very powerful.
Now, let $q$ be a system with two qbits, i.e., a 2-vector of qbits.
$q$ is now a linear combo of 4 basis values, $ | 00>$, $ | 01>$, $ | 10>$, $ | 11>$.
$q = a_0 | 00> + a_1 | 01> + a_2 | 10> + a_3 | 11> $
where $a_i$ are complex and $ \sum | a_i | ^2 = 1$.
$q$ exists in all 4 states simultaneously.
If $q$ is a vector with n component qbits, then it exists in $2^n$ states simultaneously.
This is part of the reason that quantum computation is powerful.
Measuring $q$ will collapse it to one of {00, 01, 10, 11} with probabilities $ | a_i | ^2$.
You operate on $q$ by multiplying it by a 4x4 matrix operator.
The matrices are all invertible, and all leave $ | q | = 1$.
You set the initial value of $q$ by setting its two qbits each to 0 or 1.
How this is done depends on the particular hw.
I.e., initially, $q_1 = \begin{pmatrix}a_1 | 0> \\b_1 | 1> \end{pmatrix}$ and $q_2 = \begin{pmatrix}a_2 | 0> \\b_2 | 1> \end{pmatrix}$, and so
$$q = \begin{pmatrix} q_1 \\ q_2 \end{pmatrix} = \begin{pmatrix} a_1 a_2 | 00 > \\ a_1 b_2 | 01 > \\ a_2 b_1 | 10 > \\ b_1 b_2 | 11 > \end{pmatrix}$$.
Here, the combined state is the tensor product of the individual qbits.
For $n$ qbits, the tensor product is a vector with $2^n$ elements, one element for each possible value of each qbit.
Each element of the tensor product has a complex weight.
You transform a state by multiplying it by a matrix.
The matrix is invertible.
The transformation doesn't destroy information.
When you measure a state, it collapses into one of the component states. (This may be inaccurate.)
You don't need to bring in consciousness etc. The collapse happens because the measurement causes the state to interact with the outside world.
The probability of collapsing into a particular state is the squared magnitude of its complex weight.
For some sets of weights, particularly after a transformation, the combined state cannot be separated into a tensor product of individual qbits. In this case, the individual qbits are entangled.
That is the next part of why quantum computation is powerful.
Entanglement means that if you measure one qbit then what you observe restricts what would be observed when you measure the other qbit.
However that does not let you communicate.
From page 171 of
Quantum Computing for Computer Scientists 1st Edition
All quantum algorithms work with the following basic framework:
- The system will start with the qubits in a particular classical state.
- From there the system is put into a superposition of many states.
- This is followed by acting on this superposition with several unitary operations.
- And finally, a measurement of the qubits.
More from
Quantum Computing for Computer Scientists 1st Edition
The current limitations are that IBM does only a few qbits and that the operation is noisy.
https://en.m.wikipedia.org/wiki/Bloch_sphere:
Common operations (aka gates):
https://en.wikipedia.org/wiki/Quantum_logic_gate
Hadamard.
Creates a superposition.
https://cs.stackexchange.com/questions/63859/intuition-behind-the-hadamard-gate
Toffoli, aka CCNOT.
Universal for classical boolean functions.
(a,b,c) -> (a,b, c xor (a and b))
Fredkin, aka CSWAP.
https://en.wikipedia.org/wiki/Fredkin_gate
3 inputs. Swaps last 2 if first is true.
sample app: 5 gates make a 3-bit full adder.
Algorithms:
I've created ECSE-4964 Quantum Computer Programming, CRN 30195. Preliminary syllabus.
This will force me to learn the material.
As of 3/25/2020, 17 students have preregistered.
They have several quantum computers.
The older ones are freely available on the web.
Those have 5 and 14 qbits.
You submit a batch job and get emailed when it runs.
IBM github site: https://github.com/Qiskit with
a free simulator.
It doesn't match all the physical complexity of the real computer, but it's a good start.
and tutorials and presentations:
and a SW development framework. https://qiskit.org/
You can create a quantum computation program either by
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.
An overview of Grover's Algorithm. An unstructured search algorithm that can find an item in a list much faster than a classical computer can. Several sources are listed.
Can we make quantum technology work? | Leo Kouwenhoven | TEDxAmsterdam (18:19)
IBM Introduces First Integrated Quantum Computing System for Commercial Use
The World’s First Integrated Quantum Computing System (1:06)
Shor on, what is Shor's factoring algorithm? (2:09) https://www.youtube.com/watch?v=hOlOY7NyMfs
Hacking at Quantum Speed with Shor's Algorithm (16:35) https://kcts9.org/programs/digital-studios-infinite-series/episodes/0121
43 Quantum Mechanics - Quantum factoring Period finding (19:27) https://www.youtube.com/watch?v=crMM0tCboZU
44 Quantum Mechanics - Quantum factoring Shor's factoring algorithm (25:42) https://www.youtube.com/watch?v=YhjKWAMFBUU
These are available on IBM's site (you need to create an account) and on youtube.
Go Behind-the-Scenes of a Quantum Experiment (2:10) https://quantumexperience.ng.bluemix.net/qx/community/question?questionId=5ae975690f020500399ed39a&channel=videos
or
https://www.youtube.com/watch?v=tfZpJLdkzRU&list=PLOFEBzvs-VvpzQnlazij7cL1mjKvJTAwk&index=7
A Qubit in the Making (2:01)
https://www.youtube.com/watch?v=2pB87H3_F_c&list=PLOFEBzvs-VvpzQnlazij7cL1mjKvJTAwk&index=10&t=0s
Behold the Mighty Qubit (2:51) https://www.youtube.com/watch?v=_P7K8jUbLU0&list=PLOFEBzvs-VvpzQnlazij7cL1mjKvJTAwk&index=10
Classical and Quantum Randomness (3:39) https://www.youtube.com/watch?v=8kyJfAC4VAo&list=PLOFEBzvs-VvpzQnlazij7cL1mjKvJTAwk&index=6
Quantum Entanglement (2:21) https://www.youtube.com/watch?v=RmXasxLm43k&list=PLOFEBzvs-VvpzQnlazij7cL1mjKvJTAwk&index=5
Benchmarking Quantum Systems (1:58) https://www.youtube.com/watch?v=-7L5o-mzLqU&list=PLOFEBzvs-VvpzQnlazij7cL1mjKvJTAwk&index=8
Experiment with Basic Quantum Algorithms (Ali Javadi-Abhari, ISCA 2018) (19:05) https://www.youtube.com/watch?v=M1UHi9UXTWI&list=PLOFEBzvs-VvruANdBhTb-9YRDes07pZWQ&index=2
Table of contents
To the several people who haven't uploaded it to gradescope: please do.
is online.
Thrust is an API that looks like STL. Its backend can be CUDA, OpenMP, or sequential host-based code.
The online Thrust directory structure is a mess. Three main sites appear to be these:
This points to the above site.
https://developer.nvidia.com/thrust
This has links to other Nvidia docs, some of which are obsolete.
http://docs.nvidia.com/cuda/thrust/index.html
easy-to-read, thorough, obsolete, doc
https://code.google.com/ - no longer exists.
Functional-programming philosophy.
Many possible backends: host, GPU, OpenMP, TBB...
Easier programming, once you get used to it.
Code is efficient.
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.
On parallel in /parallel-class/thrust/ are many little demo programs from the thrust distribution, with my additions.
CUDACast videos on Thrust:
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:
Sometimes, e.g., in saxpy.cu, you see saxpy_functor(A).
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:
Continue Stanford's parallel course notes.
Lecture 8: Thrust ctd, starting at slides 24.
There's a lot of stuff here.
Table of contents
I use this for quick announcements that are not so much of permanent interest, and for things that don't necessarily need to be on the global internet forever.
I think I added everyone who wasn't already in it, but for complicated reasons might have missed someone.
So, if you're not in it, please add yourself or email me to add you.
I'll send a test message. If you don't get it, check your spam filter. RPI's spam filter recently blocked my own messages to piazza being forwarded to myself.
I prepared it for an internal RPI research discussion last week. RPI faculty in several departments are considering how to do research on this. A group of RPI people, including two vice presidents, visited IBM before the break to talk about this. I'll work this into this course later, but you're welcome to read it now.
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?
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.
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., 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.
Table of contents
The videos of our classes will be in my Mediatesite channel ECSE-4740 Applied Parallel Computing for Engineers.
Please report any problems. I have no easy way to see the system from a student's perspective.
Stanford's parallel course notes.
Starting with lecture 5, which shows us a little CUDA code and, more important, optimizing techniques and parallel paradigms.
I showed starting from lecture 5 through lecture 8, slide 19.
The github repository, with demo is https://github.com/thrust/thrust.git .
Nvidia's proprietary version is slightly newer.
The most comprehensive doc is online at http://thrust.github.io/doc/index.html
It is badly written and slightly obsolete.
There are various tutorials online, most obsolescent. E.g., they don't use C++-11 lambdas, which are a big help.
Look at some Thrust programs in /parallel-class/cuda/thrust
One Nvidia-sponsored alternative is agency, at https://github.com/agency-library/ .
There are other alternatives that I'll mention later.
The alternatives are lower-level (= faster and harder to use) and newer (= possibly less debugged, fewer users).
However OpenACC now looks competitive.
The biggest problem with Thrust is that it appears that Nvidia has de-emphasized it, and appears to be making it proprietary. The two latest versions of Thrust do not allow Intel as a backend. That is a bad sign, and may be a reason to stop using it. The Thrust developers say that this is temporary but they haven't fixed it in the release version.
Table of contents
RPI wants to continue to give you a solid education.
We recognize that new problems may arise and will try to be humane.
You will still have to study.
Anything listed here is only a best attempt and might be changed if there's a good reason.
Your feedback is always welcome.
For the rest of the semester, the idea is that I'll still lecture at the appointed times, Mon and Thurs at noon eastern time.
However, now I'll try to lecture with Webex. You can access it via your favorite browser. Or, you can use an app on some platforms like the Ipad.
The lectures will be recorded so that you can rewatch them later.
I'll post the link on piazza with email to you before the lecture.
I'll use piazza---
This blog will continue to be sort of a permanent record of the class.
Gradescope will continue to be used for grading.
parallel.ecse will continue to be available remotely.
Two remaining students will present your talks.
If possible, post your slides in advance or email them to me for posting.
Quickly finish the 3 slide sets.
From https://www.openacc.org/events/openacc-online-course-2018.
My local copy on parallel: /parallel-class/openacc/online-course
I've mentioned this many times. Now we'll get into details.
- Lecture 5: Finish CUDA,
- Lectures 6, 7: Some parallel patterns.
Table of contents
10
Table of contents
To fill in the remaining time.
From https://www.openacc.org/events/openacc-online-course-2018
My local copy on parallel: /parallel-class/openacc/online-course
This homework is the 2nd round of student talks.
Total: 10 pts.
Table of contents
Table of contents
A group of us will be visiting IBM's Quantum Computing group.
In preparation for quantum computing, which I'll start in a few weeks, read Quantum Computing for Computer Scientists slides and video
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.
I put the invoice on parallel.ecse in /parallel-class/ . It gives the hardware specifics.
Since the initial purchase, parallel has been modified:
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.
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:
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:
Generational changes:
Refs:
(I'll keep adding to this. Suggestions are welcome.)
CUDA function qualifiers:
CUDA variable qualifiers:
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.
This material accompanies Programming Massively Parallel Processors A Hands-on Approach, Third Edition, David B. Kirk Wen-mei W. Hwu. I recommend it. (The slides etc are free but the book isn't.)
Finishing /parallel-class/GPU-Teaching-Kit-Accelerated-Computing, Module 21, OpenACC.
Adds pragmas to C++ code, like OpenMP.
Unlike OpenMP, targets accelerators, e.g., GPUs.
Tries to be higher level than OpenMP.
E.g., kernel pragma says to try to parallelize everything you can.
Nevertheless, the more specific parallel pragma probably leads to better code.
G++ does OpenACC badly.
Use PGI compiler, which Nvidia bought.
In parallel, I've installed it both directly and in a docker file.
Good books:
OpenACC for Programmers: Concepts and Strategies By: Sunita Chandrasekaran, Guido Juckeland kindle $31.19
Parallel Programming with OpenACC By: Rob Farber kindle: $37.74
Table of contents
A group of us will be visiting IBM's Quantum Computing group.
On Monday I'll make suggestions for how to occupy your time.
Choose a topic (or group of topics) from the 589 on-demand sessions at the GPU Tech Conference.
Summarize it in class. OK to show the whole presentation (or parts of it) with your commentary.
Groups of two people are allowed. However, this time, please keep your talks to 6 min.
Presentation dates: March 2 and 5.
No need to sign up for your topic since there'll probably be few overlaps. However here's a signup site for your presentation date.
https://doodle.com/poll/hwqrmeyxgsvk2mzk
I've allowed up to 11 talks per day.
Please enter your name and rcsid.
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.
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.
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.
This material accompanies Programming Massively Parallel Processors A Hands-on Approach, Third Edition, David B. Kirk Wen-mei W. Hwu. I recommend it. (The slides etc are free but the book isn't.)
Continuing /parallel-class/GPU-Teaching-Kit-Accelerated-Computing, starting at Modules 17 slide 14.
Table of contents
I've installed pgc++ directly. Run it thus:
/local/pgi/linux86-64/19.10/bin/pgc++ foo.cc -o foo
Here's a set of good switches:
/local/pgi/linux86-64/19.10/bin/pgc++ -fast -mp -Msafeptr -O3 -Minfo=all -Mconcur=allcores -ta:tesla -acc foo.cc -o foo
In parallel.ecse:/parallel-class/matmul, I experiment with gcc and pgi with openmp running on the Xeon, and openacc running on the GPU, multiplying matrices stored as global data, on the local stack, as STL vectors of vectors, and as a heap array. Some of my conclusions:
This material accompanies Programming Massively Parallel Processors A Hands-on Approach, Third Edition, David B. Kirk Wen-mei W. Hwu. I recommend it. (The slides etc are free but the book isn't.)
Continuing /parallel-class/GPU-Teaching-Kit-Accelerated-Computing, Modules 13 to 17 slide 14.
Since it has 256GB of main memory, there's no paging, and pinning memory is not a problem.
Linux now has Heterogeneous Memory Management (HMM) .
Assume that a kernel is launched with 1000 thread blocks each of which has 512 threads. If a variable is declared as a shared memory variable, how many versions of the variable will be created through the lifetime of the execution of the kernel? (module 4)
Assume that each atomic operation in a DRAM system has a total latency of 100ns. What is the maximal throughput we can get for atomic operations on the same global memory variable? (module 7)
For a processor that supports atomic operations in L2 cache, assume that each atomic operation takes 4ns to complete in L2 cache and 100ns to complete in DRAM. Assume that 90% of the atomic operations hit in L2 cache. What is the approximate throughput for atomic operations on the same global memory variable?
For the following basic reduction kernel code fragment, if the block size is 1024 and warp size is 32, how many warps in a block will have divergence during the iteration where stride is equal to 1? (module 9):
unsigned int t = threadIdx.x;
Unsigned unsigned int start = 2*blockIdx.x*blockDim.x;
partialSum[t] = input[start + t];
partialSum[blockDim.x+t] = input[start+ blockDim.x+t];
for (unsigned int stride = 1; stride <= blockDim.x; stride *= 2)
{
__syncthreads();
if (t % stride == 0) {partialSum[2*t]+= partialSum[2*t+stride];}
}
In the previous question, how many warps in a block will have divergence during the iteration where stride is equal to 16?
For the work inefficient scan kernel based on reduction trees, assume that we have 1024 elements, which of the following gives the closest approximation of the number of add operations performed? (module 10)
Which of the following statements is true? (module 14)
What is the CUDA API call that makes sure that all previous kernel executions and memory copies in a device have been completed?
Total: 90 pts.
Table of contents
This material accompanies Programming Massively Parallel Processors A Hands-on Approach, Third Edition, David B. Kirk Wen-mei W. Hwu. I recommend it. (The slides etc are free but the book isn't.)
Continuing /parallel-class/GPU-Teaching-Kit-Accelerated-Computing, Modules 9 to 12.
Table of contents
Docker is worth learning, apart from its use by Nvidia for parallel computing. You might also look up Kubernetes.
On parallel, I think that docker is probably secure, but am not certain. In return for my allowing access to docker on parallel, I expect you to report, and not to exploit, holes that you might discover.
Compiling:
This material accompanies Programming Massively Parallel Processors A Hands-on Approach, Third Edition, David B. Kirk Wen-mei W. Hwu. I recommend it. (The slides etc are free but the book isn't.)
Continuing /parallel-class/GPU-Teaching-Kit-Accelerated-Computing, Modules 4 to 8.
These may require some research.
Total: 40 pts.
Table of contents
I've installed docker, a popular lightweight virtualization system, on parallel, because Nvidia uses it to distribute SW.
Docker runs images that define virtual machines.
Docker images share resources with the host, in a controlled manner.
You can install private copies of images, or to see what images I've installed, do: docker images .
Run the hello-world image thus: docker run hello-world
Here's a more complicated example:
docker run -it --mount type=bind,source=/parallel-class,destination=/parallel-class --mount type=bind,source=$HOME,destination=/home --gpus=all nvidia/cuda:10.1-devel
This interactively runs a virtual machine with
E.g., go into /parallel-class/openmp/rpi and run some programs.
Copy some .cc files to your home dir, compile, and run them.
There are ways to make the image's contents persistent. E.g., you can customize and save an image.
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. More later. However, e.g.,
docker run --gpus=all nvidia/cuda:10.1-devel nvidia-smi
parallel has CUDA sample programs in /local/cuda/samples. To make them available in a docker image, include -v local/cuda/samples:/samples
Continuing /parallel-class/GPU-Teaching-Kit-Accelerated-Computing at Module 3.
Table of contents
Note: in a month, we'll have another round.
This is from https://developer.nvidia.com/teaching-kits-downloads
My local copy of what I'm using is in /parallel-class/GPU-Teaching-Kit-Accelerated-Computing
Today we did through module 2.
For students who've looked at last year's course, here are some upcoming changes.
Table of contents
is online here , due in a week.
The goal is to measure whether OpenMP actually makes matrix multiplication faster, with and w/o SIMD.
You may use anything in /parallel-class/openmp that seems useful.
Write a C++ program on parallel.ecse to initialize pseudorandomly and multiply two 100x100 float matrices. One possible initialization:
a[i][j] = i*3 + (j*j)%5; b[i][j] = i*2 + (j*j)%7;
(10 points) Report the elapsed time. Include the program listing.
Add an OpenMP pragma to do the work in parallel.
(10 points) Report the elapsed time, varying the number of threads thus: 1, 2, 4, 8, 16, 32, 64.
What do you conclude?
(5 points) Repeat that two more times to see how consistent the times are.
(10 points) Modify the pragma to use SIMD.
Report the elapsed time, varying the number of threads thus: 1, 2, 4, 8, 16, 32, 64.
What do you conclude?
(10 points) Compile and run your program with two different levels of compiler optimization: O1 and O3, reporting the elapsed time. Modify your program to prevent the optimizer from optimizing the program away to nothing. E.g., print a few values.
(5 points) What do you conclude about everything?
Total: 40 pts.
Table of contents
continuing for several classes.
Table of contents
| Title: | ECSE-4740-01 Applied Parallel Computing for Engineers, CRN 94829 |
|---|---|
| Semesters: | Spring term annually |
| Credits: | 3 credit hours |
| Time and place: | Mon and Thurs noon-1:20pm, JONSSN 4107 |
| Multicore Intel Xeon: | |
|---|---|
| universally available and inexpensive, comparatively easy to program, powerful | |
| Nvidia GPU accelerator: | |
| 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. | |
| IBM quantum computer: | |
| It's new and hot and might some day be useful. | |
| OpenMP C++ extension: | |
|---|---|
| widely used, easy to use if your algorithm is parallelizable, backend is primarily multicore Xeon but also GPUs. | |
| Thrust C++ functional programming library: | |
| FP is nice, hides low level details, backend can be any major parallel platform. | |
| MATLAB, Mathematica: | |
| easy to use parallelism for operations that they have implemented in parallel, etc. | |
| CUDA C++ extension and library for Nvidia: | |
| low level access to Nvidia GPUs. | |
ECSE-2660 CANOS or equivalent, knowledge of C++.
W. Randolph Franklin. BSc (Toronto), AM, PhD (Harvard)
| Office: |
Jonsson Engineering Center (JEC) 6026 |
|---|---|
| Phone: |
+1 (518) 276-6077 (forwards) |
| Email: |
Email is my preferred communication medium. Sending from non-RPI accounts are fine, but please show your name, at least in the comment field. A subject prefix of #Prob is helpful. GPG encryption is fine. |
| Web: |
A quick way to get there is to google RPIWRF. |
| Office hours: |
After each lecture, usually as long as anyone wants to talk. Also by appointment. |
| Informal meetings: | |
|
If you would like to lunch with me, either individually or in a group, just mention it. We can then talk about most anything legal and ethical. |
|
Elkin Cruz, cruzce@THEUSUALDOMAIN
Office hour: Tuesday 6-8pm in the flip flop Lounge.
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.
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. If interested, you might start here:
This course will primarily use (remotely via ssh) parallel.ecse.rpi.edu.
Parallel has:
Material for the class is stored in /parallel-class/ .
We might also use a parallel virtual machine on the Amazon EC2. If so, you will be expected to establish an account. I expect the usage to be in the free category.
Piazza will be available for discussion and questions.
Gradescope will be used for you to submit homeworks and for us to distribute grades.
The entry code for this course is 9DYRB2.
Please add yourself.
There will be some homeworks.
You may do homeworks in teams of 2 students. Create a gradescope team and make one submission with both names.
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 up to 3 people. 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:
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.
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.
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
is online, due in a week.
This is homework 1.
Having talks today is too rushed, so let's have all the talks next week (unless anyone wants to talk today).
As of 1/21, 10 people have picked a topic. I assigned them all to talk on Thur 1/30. The other people have been assigned to talk on Mon 1/27.
Grad students who haven't registered for a reading class and who also haven't doodled are also not listed. Add yourselves to Mon.
For privacy purposes, I've abbreved the names.
Monday speakers: please tell me your topics.
If anyone is part of a group, please tell me.
If any Thursday speaker would like to speak Mon, please tell me. You'll get first choice for dates on the next presentation.
The machine is parallel.ecse.rpi.edu . You must be on campus or use the VPN to connect. Connect with ssh.
I've created accounts for everyone in ECSE-4740 but not assigned passwords (and so you cannot login). In class today I'll let you type in a password. I'll create accounts for the reading students.
Its accounts are completely separate from RCS; I used the same userid for convenience. Changing your password on parallel does not affect your RCS password, and vv.
I recommend that you create key pairs to connect to parallel w/o typing your password each time.
To avoid retyping your ssh private key passphrase in the same shell, do this
ssh-add
One advantage of ssh is that you can mount your parallel directory on your local machine. In the files browser on my laptop, I connect to the server sftp://parallel.ecse.rpi.edu . Any other program to mount network shares would work as well.
Where the share is actually mounted varies. On my laptop, it's here: /var/run/user/1000/gvfs/sftp:host=parallel.ecse.rpi.edu,user=wrf
1000 is my uid on my laptop. Yours is not 1000.
As with any network mount, fancy filesystem things like attributes, simultaneous reading and writing from different machines, etc., probably won't work.
With ssh, you can also copy files back and forth w/o mounting or typing passwords:
It even does filename completion on remote files.
You can also run single commands:
ssh parallel.ecse.rpi.edu hostname
In emacs, use tramp mode: Access file as /HOST:/FILE
On parallel, most filesystems use zfs.
Versions: OpenMP is a living project, that is being updated every few years.
We'll review some examples running on parallel, at /parallel-class/openmp/rpi/ .
Now that we've seen some examples, look at the LLNL tutorial. https://computing.llnl.gov/tutorials/openMP/
We'll see some more examples running on parallel, at /parallel-class/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.
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.
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.
stacksize.cc - get and set stack size. Can also do ulimit -s.
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.
http://stackoverflow.com/questions/13788638/difference-between-section-and-task-openmp
Each is 5 points.
Total: 50 pts.
Table of contents
We'll quickly finish it. https://computing.llnl.gov/tutorials/parallel_comp/
https://doodle.com/poll/68axuciyyp3avwz5. Pick a topic by entering your name and rcsid (because sometimes students have similar names and sometimes students use different names from what's on the classlist). The topics are:
- Describe the fastest 2 computers on the top500 list. https://www.top500.org/lists/2018/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.
- Describe an application where astronomy 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.
- More topics TBA.
- Other related topic of your choice. Tell me and I'll add it to the list.
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.
We'll see some examples running on parallel, at /parallel-class/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.
(10 points) This is the in-class presentation on Jan 27 and 30 mentioned in class.
Table of contents
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.
Table of contents
| Title: | ECSE-6xxx Independent study: Parallel, GPU, quantum comput. |
|---|---|
| Semesters: | Spring, on demand |
| Credits: | 3 credit hours |
| Time and place: | Mon and Thurs noon-1:20pm, JONSSN 4107 |
| Multicore Intel Xeon: | |
|---|---|
| universally available and inexpensive, comparatively easy to program, powerful | |
| Nvidia GPU accelerator: | |
| 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. | |
| IBM quantum computer: | |
| It's new and hot and might some day be useful. | |
| OpenMP C++ extension: | |
|---|---|
| widely used, easy to use if your algorithm is parallelizable, backend is primarily multicore Xeon but also GPUs. | |
| Thrust C++ functional programming library: | |
| FP is nice, hides low level details, backend can be any major parallel platform. | |
| MATLAB, Mathematica: | |
| easy to use parallelism for operations that they have implemented in parallel, etc. | |
| CUDA C++ extension and library for Nvidia: | |
| low level access to Nvidia GPUs. | |
ECSE-2660 CANOS or equivalent, knowledge of C++.
W. Randolph Franklin. BSc (Toronto), AM, PhD (Harvard)
| Office: |
Jonsson Engineering Center (JEC) 6026 |
|---|---|
| Phone: |
+1 (518) 276-6077 (forwards) |
| Email: |
Email is my preferred communication medium. Sending from non-RPI accounts are fine, but please show your name, at least in the comment field. A subject prefix of #Prob is helpful. GPG encryption is fine. |
| Web: |
A quick way to get there is to google RPIWRF. |
| Office hours: |
After each lecture, usually as long as anyone wants to talk. Also by appointment. |
| Informal meetings: | |
|
If you would like to lunch with me, either individually or in a group, just mention it. We can then talk about most anything legal and ethical. |
|
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.
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. If interested, you might start here:
This course will primarily use (remotely via ssh) parallel.ecse.rpi.edu.
Parallel has:
Material for the class is stored in /parallel-class/ .
We might also use a parallel virtual machine on the Amazon EC2. If so, you will be expected to establish an account. I expect the usage to be in the free category.
Piazza will be available for discussion and questions.
Gradescope will be used for you to submit homeworks and for us to distribute grades.
The entry code for this course is 9DYRB2.
Please add yourself.
Students will be required to perform individual research into parallel computing principles and to record their findings as weekly entries in a personal blog, a sort of online lab book.
Possible platforms include:
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 up to 3 people. 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:
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.
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.
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.