Kaitlynware

Lectures and book found here.

The Problem

A problem like this could be solved by using a simple set. Before inserting a new element, we check for membership into the visited set and discard duplicates.

In the following examples, rust is used, and elements are Strings, coming from crawled web pages.

We want to ensure functionality and correctness on parallel processors.

Motivating GPU Computing

Recall Moore’s Law: Gordon Moore predicted that the number of transistors per unit area would double every 18-24 months. However, this trend has been fading as the transistors we make become too small to physically shrink any more. As transistors become smaller, it becomes easier to switch them on/off faster, leading to faster clock frequency => clock frequency also doubled every 18-24 months… until around 2005. This was because of power. When we switched transistors on and off faster, we burned more power, generating more heat. CPUs got so hot that our cooling technology was unable cool them adequately, leading to mostly stagnation in frequency. This led to a stagnation in single-thread performance, except for some improvements in architecture and compiler advancements. Around the time where frequency stagnated, the number of cores in mainstream parallel processors exploded. We went from a single core processor (one core/datapath per processor) to a large number of cores in the CPU. Stagnation in signle-thread performance made parallel computing mainstream as transistors were invested in adding more cores to processors. No “free lunch” in increased performance by running same program on a newer core two years later, software people now had to design and write code for parallel.

Design Approaches

To design a system that accomplishes some task, there are two approaches:

• Latency-oriented design Example: car, minimizes the time to takes to perform a single task (get from one place to another)

• Throughput-oriented design Example: bus, maximizes number of tasks that can be performed in a given time frame (maximizes people getting from one place to another, though each person will take more time)

Approaches to Processor Design

• CPU: Latency-oriented

  • A few powerful ALUs, take up a lot of space, designed to reduce latency of a single operation.
  • Large caches, convert long latency memory accesses to short latency cache accesses. Making caches larger reduces miss rate of CPU, allowing us to have short latency memory operations
  • Sophisticated control techniques, branch predication to reduce control hazards (to data hazards), data forwarding to reduce data hazards, out-of-order execution to improve latency on a single thread (?). Expensive hardware. Relative to the GPU, not as many stalls
  • High frequency

• GPU: Throughput-oriented

  • Much smaller, weaker ALUs, longer latency and high throughput.

  • Have MANY of them, heavily pipelined for better throughput, but each individual operation has higher latency

  • Small caches, less likely to hit cache, but we can dedicate more area to computation. More silicon for ALUs, memory accesses take long time

  • Simple control, more area dedicated to computation, pay price in latency to get more throughput.

  • How to minimize stalls? What can the hardware do to minimize number of empty cycles? Techniques: Reorder Buffer (R.O.B.), OOO concept. ROB helps when the CPU executes instructions out of order to keep our pipeline busy.

  • Moderate frequency

  • CPU: multithreading to hide latency, use same core to execute multiple threads, “asynchronous” multiprocessing. Use stalls from one thread to execute instructions from another thread. Modest multithreading on CPU, 2 is typical -> hyperthreading? Two different instruction streams.

  • On GPU, we expect massive number of threads to hide the very high latency. Many more cores, many more threads per core, may select 32-64 threads to fill an empty cycle.

  • What is ROB?

    • Stores results of instructions that have executed but can’t be committed (retired) as earlier instructions are still pending.
    • Lets later independent instructions continue executing even if earlier ones are stalled.
    • Ensures results are written back in order. That is, CPU can start and finish instructions in a different order than the original program sequence.
    • Paired with register renaming and reservation stations in superscalars
    • Example.

      I1: R1 = R2 + R3
      I2: R4 = R1 + 5
      I3: R5 = R6 + 7

      Copy

With OOO, we execute I3 first, I1, then finally I2 which has a dependency on R1. But even though I3 finished before I1, we can’t immediately write I3’s result to the register file R5. This is because if I1 or I2 causes an exception, the CPU would have already “committed” an instruction that comes after, which breaks program correctness. Thus I3’s result sits in the ROB. ROB tracks which instructions have finished execution but haven’t yet safely updated architectural state. When CPU confirms all prior instructions in the original program sequence have completed without exception, then it can commit (retire) I3, writing its result from ROB into register R5.

• Branch order, out of order prediction

Origins of GPU

• Originally built for graphics

• Graphics workloads are massively parallel, process a massive number of pixels in parallel. Often, those pixels are independent from each other.

• “Don’t care about one pixel appearing as fast as possible, care about thousands of pixels appearing within a certain framerate”. This is motivation behind GPU’s throughput-oriented design.

• People discovered that GPUs could be used for other parallel workloads

• Before 2007, only way to program GPUs was using graphics APIs vendors gave, e.g. OpenGL, Direct3D. People using GPUs for other than graphics had to reformulate computations as a function that paints a pixel. Take your function that had nothing to do with graphics, and rewrite it to relate to rendering pixels.

• When GPU vendors discovered this, in 2007 Nvidia released CUDA - a programming interface for using GPUs in a general purpose way. Rather than calling graphics to use GPU, we could write general-purpose programs for the GPU!

• This new interface also required hardware extensions to the GPU architecture.

• Why do GPUs take more power than CPUs? How can GPU be more energy-efficient tha computer with just CPUs? Energy is a function of power and time.

• Why GPUs? Why did they succeed?

  • Chips are expensive to build, require lots of engineering, lots of prototyping, testing, fabrication of chips has high fixed cost, requires a large volume of sales to amortize the cost. This makes the chip market very difficult to penetrate, due to the fixed cost -> have to make sure you have a large customer base ready to purchase your chip
  • When parallel computing became mainstream, GPUs already had (have) a large customer base from gaming sector. This gave them a headstart compared to other potentially massively parallel accelerators.
  • Some company may have had some chip designed to run scientific applications in parallel, highly customised. Problem? Small customer base.

GPU Market Sector Breakdown

• More than 50% for gaming!
• Datacentre around a quarter
  • Enterprise
  • Scientific computing
  • Hyperscale

Readings

• PMPP, Chapter 1, all sections

END OF LECTURE 1

Lecture 2

• CPU (host) -> main memory (host memory)
• GPU (device) -> GPU memory (global memory) CPU and GPU have separated memories and can’t access each other’s memories. (Unless with unified virtual memory)

What do we need to do add arrays in GPU? We need to transfer data from CPU to GPU. We would have some kind of interconnect from main memory to GPU memory (e.g. PCIE, NV Link).

1. Allocate memory on GPU
2. Copy data from CPU to GPU memory
3. Perform the computation on GPU
4. Copy data from GPU memory
5. Deallocate memory on the GPU

Cuda Memory Management API

The cudaMalloc() function allows us to allocate memory on the GPU:

1cudaError_t cudaMalloc(void** devPtr, size_t size)

• devPtr: pointer to pointer to allocated device memory
• size: requested allocation size in bytes

1cudaError_t cudaFree(void* devPtr)

• devPtr: pointer to device memory to free

How to copy memory?

1cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, enum cudaMemcpyKind kind)

• dst: destination memory address
• src: source memory address
• count: size in bytes to copy
• kind: kind of transfer
  • cudaMemcpyHostToHost
  • cudaMemcpyHostToDevice
  • cudaMemcpyDeviceToHost
  • cudaMemcpyDeviceToDevice

For parallel vector addition, we want to create a thread for each element in the array x/y. Have each thread responsible for loading two elements, adding them, and storing them in output vector z. Simple strategy: assign one GPU thread per vector element.

An array of threads on the GPU is called a grid. We say that we launch a grid on the GPU to perform some computation. Two-level hierarchy, threads are organised or grouped into thread blocks. Significance: threads in the same block can interact and collaborate in ways that threads in different blocks can’t.

Launching a Grid

• Threads in the same grid all execute the same function, known as a kernel.- The way to launch a grid is by calling the special function, the kernel, and configuring it with the appropriate grid and block sizes, i.e. number of threads in each block we want.

Implementing a Kernel

A kernel is similar to a C or C++ function, but is preceded by the keyword __global__ to indicate that it’s a GPU kernel. It will use some special keywords to distinguish threads from each other. For threads to execute, there needs to be a way for each thread to distinguish itself from other threads; thread 0 needs to know it’s thread 0, thread 1 needs to know it’s thread 1, and so on…

Grid Dimension

• gridDim.x: tells a thread what is the number of blocks in our grid.

Some sample code for vector addition on CPU vs. GPU:

 1#include "timer.h"
 2
 3void vecadd_cpu(float* x, float* y, float* z, int N) {
 4    for (unsigned int i = 0; i < N; ++i) {
 5        z[i] = x[i] + y[i];
 6    }
 7}
 8
 9// convention is add _kernel to the name
10__global__ void vecadd_kernel(float* x, float* y, float* z, int N) {
11
12}
13
14void vecadd_gpu(float* x, float* y, float* z, int N) {
15    // Allocate GPU memory
16    float *x_d, *y_d, *z_d; // _d convention to mark data on device
17    cudaMalloc((void**)&x_d, N*sizeof(float));
18    cudaMalloc((void**)&y_d, N*sizeof(float));
19    cudaMalloc((void**)&z_d, N*sizeof(float));
20
21    // Copy to the GPU
22    cudaMemcpy(x_d, x, N*sizeof(float), cudaMemcpyHostToDevice);
23    cudaMemcpy(y_d, y, N*sizeof(float), cudaMemcpyHostToDevice);
24
25    // Call a GPU kernel function (launch a grid of threads)
26    const unsigned int numThreadsPerBlock = 512;
27    const unsigned int numBlocks = N/512;
28    // Below creates a grid on GPU, each thread executes vecadd_kernel
29    vecadd_kernel<<<numBlocks, numThreadsPerBlock>>>(x_d, y_d, z_d, N);
30
31    // Copy from the GPU
32    cudaMemcpy(z, z_d, N*sizeof(float), cudaMemcpyDeviceToHost);
33
34    // Deallocate GPU memory
35    cudaFree(x_d);
36    cudaFree(y_d);
37    cudaFree(z_d);
38}
39
40
41int main(int argc, char** argv) {
42    cudaDeviceSyhchronize();
43
44    // Allocate memory and initialize data
45    Timer timer;
46    unsigned int N = (argc > 1) ? (atoi(argv[1])) : (1 << 25);
47    float* x = (float*) malloc(N*sizeof(float));
48    float* y = (float*) malloc(N*sizeof(float));
49    float* z = (float*) malloc(N*sizeof(float));
50    for (unsigned int i = 0; i < N; ++i) {
51        x[i] = rand();
52        y[i] = rand();
53    }
54
55    // Vector addition on CPU
56    startTime(&timer);
57    vecadd_cpu(x, y, z, N);
58    stopTime(&timer);
59    printElapsedTime(timer, "CPU time", CYAN);
60
61    // Vector addition on GPU
62    startTime(&timer);
63    vecadd_gpu(x, y, z, N);
64    stopTime(&timer);
65    printElapsedTime(timer, "GPU time", DGREEN);
66    
67    // Free memory
68    free(x);
69    free(y);
70    free(z);
71
72    return 0;
73}