Chapter 2
23 “Data parallelism refers to the phenomenon in which the computation work to be performed on different parts of the dataset can be done independently of each other and thus can be done in parallel with each other. ”
27 “All the threads that are launched by a kernel call are collectively called a grid.”
31 “In practice, such a “transparent” outsourcing model can be very inefficient because of all the copying of data back and forth. One would often keep large and important data structures on the device and simply invoke device functions on them from the host code”
It is pretty intuitive that a lot of data copying from the cpu to the device is expensive and not efficient
“hardware cards that come with their own dynamic random-access memory called device global memory, or simply global memory”
It is called the global memory to distinguish from other types of memory in the GPU.
“The first parameter to the cudaMalloc function is the address of a pointer vari
able that will be set to point to the allocated object. The address of the pointer variable should be cast to (void **)”
33 “The addresses in A_d, B_d, and C_d point to locations in the device global memory. These addresses should not be dereferenced in the host code. They should be used in calling API functions and kernel functions. Dereferencing a device global memory pointer in host code can cause exceptions or other types of runtime errors”
This is pretty straightforward since these are pointers pointing to the GPU memory
“The addresses in A_d, B_d, and C_d point to locations in the device global memory. These addresses should not be dereferenced in the host code. They should be used in calling API functions and kernel functions. Dereferencing a device global memory pointer in host code can cause exceptions or other types of runtime errors”
34
“The vecAdd function, outlined in Fig. 2.5, allocates space in device global memory, requests data transfers, and calls the kernel that performs the actual vector addition. We refer to this type of host code as a stub for calling a kernel.”
It is called a stub because the client calls this function, he doesn’t really know about all the stuff happening in the background from memory allocation to data transfer.
36
“For a given grid of threads, the number of threads in a block is available in a built-in variable named blockDim. The “blockDim” variable is a struct with three unsigned integer fields (x, y, and z) that help the programmer to organize the threads into a one-, two-, or three-dimensional array”
39
“The “device” keyword indicates that the function being declared is a CUDA device function. A device function executes on a CUDA device and can be called only from a kernel function or another device function. The device function is executed by the device thread that calls it and does not result in any new device threads being launched.”
“Note that one can use both “host” and “device” in a function declaration. This combination tells the compilation system to generate two versions of object code for the same function.
Chapter 3
Ideas
- When solving a parallel algorithm problem, and can’t think about how to parallelize, think of a single part of the output array, how is it calculated ? Usually a single thread will be doing the work to get this single output
Questions
- Difference between thread-to-output and thread-to-data mapping
Chapter 4
P.70
“It is organized into an array of highly threaded streaming multiprocessors (SMs). Each SM has several processing units called streaming processors or CUDA cores, show an as small tiles that share control logic and memory resources, the SMs also come with different on-chip memory structures collectively labeled as “Memory”, GPUs also come with gigabytes of off-chip device memory, referred to as “Global Memory””
“These threads are assigned to SMs on a block-by-block basis. That is, all threads in a block are simultaneously assigned to the same SM”
71 ” This guarantee makes it possible for threads in the same block to interact with each other in ways that threads across different blocks cannot.”
This is talking about the threads that are in the same block always run on the SMs simultaneously
72
“When a thread calls __syncthreads(), it will be held at the program location of the call until every thread in the same block reaches that location”
74 “the CUDA runtime system can execute blocks in any order relative to each other, since none of them need to wait for each other”
This happens due to not allowing threads in different blocks to perform, barrier synchronization, which leads to transparent scalability, i.e the ability to run the application on multiple devices that have different speed sine the threads in different blocks don’t rely on each other
75 “In most implementations to date, once a block has been assigned to an SM, it is further divided into 32-thread units called warps.”
“For a block whose size is not a multiple of 32, the last warp will be padded with inactive threads to fill up the 32 thread positions”
77 “An SM is designed to execute all threads in a warp following the single-instruction, multiple-data (SIMD) model. That is, at any instant in time, one instruction is fetched and executed for all threads in the warp”
“Fig. 4.8 shows how the cores in an SM are grouped into processing blocks in which every 8 cores form a processing block”
“Because the SIMD hardware effectively restricts all threads in a warp to execute the same instruction at any point in time, the execution behavior of a warp is often referred to as single instruction, multiple-thread”
GPU → SMs → processing blocks Same threads run on the same processing blocks
80 “When threads in the same warp follow different execution paths, we say that these threads exhibit control divergence,”
“However, when threads within a warp take different control flow paths, the SIMD hardware will take multiple passes through these paths, one pass for each path, for example for an if-else construct, if some threads in a warp follow the if-path while others follow the else path, the hardware will take two passes. One pass executes the threads that follow the if-path, and the other executes the threads that follow the else-path. During each pass, the threads that follow the other path are not allowed to take effect”
What is meant by a pass here is an execution cycle. When there is a an if-else conditional, the cores don’t run the statement in parallel, instead, it first runt the if-then, masking the cores that will run the else statement, then mask the if-then and run the else.
There is a performance penalty here due to not all warp threads are running in parallel because of the if-conditional
In some architectures when there is different passes, they can be executed concurrently
85 “he execution resources in an SM include registers, shared memory (discussed in Chapter 5, Memory Architecture and Data Locality), thread block slots, and thread slots. These resources are dynamically partitioned across threads to support their execution. For example, an Ampere A100 GPU can support a maximum of 32 blocks per SM, 64 warps (2048 threads) per SM, and 1024 threads per block. If a grid is launched with a block size of 1024 threads (the maximum allowed), the 2048 thread slots in each SM are partitioned and assigned to 2 blocks.”
Chapter 5
95 “Thus the ratio of floating-point operations (FLOP) to bytes (B) accessed from global memory is 2 FLOP to 8 B, or 0.25 FLOP/B. We will refer to this ratio as the compute to global memory access ratio, defined as the number of FLOPs performed for each byte access from the global memory within a region of a program. This ratio is sometimes also referred to as arithmetic intensity or computational intensity in the literature.”
“The compute to global memory access ratio has major implications forthe performance of a CUDA kernel. For example, the Ampere A100 GPU has a peak global memory bandwidth of 1555 GB/second. Since the matrix multiplication kernel performs 0.25 OP/B, the global memory bandwidth limits the throughput of single-precision FLOPs that can be performed by the kernel to 389 giga FLOPs per second (GFLOPS), obtained by multiplying 1555 GB/ second with 0.25 FLOP/B”
This is a very important consideration I should take while coding
To illustrate more, let’s say I can get 500MB of data per second, If i only do 0.25 FLOP per bytes I will only a small amount of throughput which is
“. We refer to programs whose execution speed is limited by memory bandwidth as memory-bound programs.”
96 “At the bottom of the figure, we see global memory and constant memory. Both these types of memory can be written (W) and read (R) by the host. The global memory can also be written and read by the device, whereas the constant memory supports short-latency, high-bandwidth read-only access by the device. ”
“Another type of memory is the local memory, which can also be read and written. The local memory is actually placed in global memory and has similar access latency, but it is not shared across threads. Each thread has its own section”
97 “Registers and shared memory in Fig. 5.2 are on-chip memories. Variables that reside in these types of memory can be accessed at very high speed in a highly parallel manner. Registers are allocated to individual threads; each thread can access only its own register”
“Shared memory is allocated to thread blocks; all threads in a block can access shared memory variables declared for the block.”
“Furthermore, whenever a variable is stored in a register, its accesses no longer consume off-chip global memory bandwidth. This will be reflected as an increased compute to global memory access ratio.”
I.e. It is no longer used in the compute to global memory access ratio
101
” If a variable’s scope is a single thread, a private version of the variable will be created for every thread; each thread can access only its private version of the variable. For example, if a kernel declares a variable whose scope is a thread and it is launched with one million threads, one million versions of the variable will be created so that each thread initializes and uses its own version of the variable.”
” If a variable’s lifetime is within a grid’s execution, it must be declared within the kernel function body and will be available for use only by the kernel’s code. If the kernel is invoked several times, the value of the variable is not maintained across these invocations”
“if a variable’s lifetime is throughout the entire application, it must be declared outside of any function body.”
“A private version of the shared variable is created for and used by each block during kernel execution”
There is emphasis here on each different block
“Declaration of constant variables must be outside any function body. ”
“Currently, the total size of constant variables in an application is limited to 65,536 bytes. One may need to break up the input data volume to fit within this limitation.”
“Accesses to a global variable are slow. Latency and throughput of accessing global variables have been improved with caches in more recent devices”
110
“The two __syncthreads() calls in lines 21 and 26 demonstrate two different types of data dependence that parallel programmers often have to reason about when they are coordinating between threads. The first is called a read-after-write dependence because threads must wait for data to be written to the proper place by other threads before they try to read it. The second is called a write-after-read dependence because a thread must wait for the data to be read by all threads that need it before overwriting it. Other names for read-after-write and write-after-read dependences are true and false dependences, respectively. A read-after-write dependence is a true dependence because the reading thread truly needs the data supplied by the writing thread, so it has no choice but to wait for it. A write-after-read dependence is a false dependence because the writing thread does not need any data from the reading thread. The dependence is caused by the fact that they are reusing the same memory location and would not exist if they used different
locations.”
116
“We can enable such adjustment with a different style of declaration in CUDA by adding a C extern keyword in front of the shared memory declaration and omitting the size of the array in the declaration. Based on this style, the declarations for Mds and Nds need to be merged into one dynamically allocated array”
The adjustment here is creating a dynamically allocated shared memory in the GPU, during that we use the extern keyword. Example code snippet
extern __shared__ float arr[];Then we calculate the size of the array in bytes, and pass it in the kernel configuration parameters and it is a size_t type to hold a size
size_t is architecture specific, and is the largest an array can be in your system
example kernel configuration
kernel<<<dimGrid, dimBlock, size in bytes>>>();