10-605 / 10-805

Machine Learning from Large Datasets

Recap: the course so far

• Tool: Scalable computation on clusters
• MapReduce, Spark and iteration, Spark workflows
• Tool: Learning as optimization
• MLE for binomials, linear regression, logistic regression, k-means clustering
• Tools: exact solutions, gradient descent (batch, parallel, stochastic), EM and EM-like methods (for k-means)
• “Craft” of ML
• Tool: feature extraction / augmentation/ kernels
• Tool: regularization and sparsity hyperparameter tuning
• Tool: the hash trick and efficient SGD
• Tool: distributed ML – interaction with sparsity, AllReduce
• Tool: Randomized algorithms
• Bloom filter (a randomized set)
• Count-Min sketch (a randomized mapping x 🡪R+ )
• Locality sensitive hashing (LSH)

2

Many of these are different with GPUs

• Tool: Scalable computation on clusters
• MapReduce, Spark and iteration, Spark workflows
• Tool: Learning as optimization
• MLE for binomials, linear regression, logistic regression, k-means clustering
• Tools: exact solutions, gradient descent (batch, parallel, stochastic), EM and EM-like methods (for k-means)
• “Craft” of ML
• Tool: feature extraction / augmentation/ kernels
• Tool: regularization and sparsity hyperparameter tuning
• Tool: the hash trick and efficient SGD
• Tool: distributed ML – interaction with sparsity, AllReduce
• Tool: Randomized algorithms
• Bloom filter (a randomized set)
• Count-Min sketch (a randomized mapping x 🡪R+ )
• Locality sensitive hashing (LSH)

3

Many of these are different with GPUs

• Tool: Scalable computation on clusters
• MapReduce, Spark and iteration, Spark workflows
• Tool: Learning as optimization
• MLE for binomials, linear regression, logistic regression, k-means clustering
• Tools: exact solutions, gradient descent (batch, parallel, stochastic), EM and EM-like methods (for k-means)
• Today and Wed:
• How GPUs work and how this changes optimization for ML
• “Craft” of ML
• Tool: feature extraction / augmentation/ kernels
• Tool: regularization and sparsity hyperparameter tuning
• Tool: the hash trick and efficient SGD
• Tool: distributed ML – interaction with sparsity, AllReduce
• Tool: Randomized algorithms
• Bloom filter, Count-Min sketch, Locality sensitive hashing (LSH)

4

WHAT ARE GPUS?

5

What is a GPU?

6

A graphics processing unit (GPU) is a specialized electronic circuit designed to rapidly manipulate and alter memory to accelerate the creation of images in a frame buffer intended for output to a display device. [wikipedia]

The term GPU was popularized by Nvidia in 1999, who marketed the GeForce 256 as "the world's first …Graphics Processing Unit.^” It was presented as a "single-chip processor with integrated transform, lighting, triangle setup/clipping, and rendering engines".^[3] 

Rendering and lighting require fast parallel rotations, translations, and scaling

GPUs do linear algebra in parallel

Started to be used in mid-2000’s for ML – Theano (2007) was first DL framework based on GPUs

7

Summary of GPUs vs CPUs

​

• less total memory
• more cores and more parallelism

​

• Multicore CPUs are mostly multiple-instruction multiple-data (MIMD)
• GPUs are mostly single-instruction multiple-data (SIMD)

​

8

https://www.youtube.com/watch?v=Ge-g3xZ5bb8

​

9

https://www.karlrupp.net/2013/06/cpu-gpu-and-mic-hardware-characteristics-over-time/

GPUs are faster than CPUs:

maximum FLOPS/clock cycle

old slide

10

https://www.karlrupp.net/2013/06/cpu-gpu-and-mic-hardware-characteristics-over-time/

GPUs have more cores than CPUs

hi-end CPU: 32

hi-end GPU: 2000

GPUs have less memory than CPUs

hi-end CPU: 256 Gb

hi-end GPU: 8 Gb

old slide

11

https://www.karlrupp.net/2013/06/cpu-gpu-and-mic-hardware-characteristics-over-time/

GPUs compute faster because they are parallel

old slide

12

https://www.karlrupp.net/2013/06/cpu-gpu-and-mic-hardware-characteristics-over-time/

GPUs also read memory faster

(memory bandwidth: bits/sec)

because they are parallel

old slide

13

newer slide

Newer Accelerators

newer slide

15

newer slide

NVIDIA GB200 Superchip Incl. Two Blackwell GPUs and One Grace CPU

GraceHopper H100

Roughly speaking performance doubles every 2.5-3 years

GPUs vs CPUs

• Highly parallel floating point operations (FLOPs)
• Harder to program
• in general
• For ML use linear algebra abstractions
• Memory access
• Registers and caches
• On-board memory (HBM – connected to CPU memory)

​

​

​

• Easier to program
• in general
• Many programming abstractions
• Memory access
• Registers and caches
• On-board memory
• Disk
• Network, …

17

newer slide

GPUs vs TPUs/NPUs

• Highly parallel floating point operations (FLOPs)
• Harder to program
• in general
• For ML use linear algebra abstractions
• Memory access
• Registers and caches
• On-board memory (HBM – connected to CPU memory)

​

• Specialized to certain parallel operations
• Even harder to program
• more specialized h/w
• Only fast for vectorized operations and matmul
• Memory access
• Registers and caches
• eg for matmul units
• On-board memory (HBM – connected to CPU memory)

​

History of Google’s TPU’s

Bigger models need more memory

Faster clocks mean faster training

Will discuss HBM later

More bandwidth 🡪 Faster Communication

HOW DO YOU USE A GPU?

22

http://developer.download.nvidia.com/compute/developertrainingmaterials/presentations/cuda_language/Introduction_to_CUDA_C.pptx

Using GPUs for ML

• Programming parallel machines is complicated
• To use the parallelism of a GPU in an ML algorithm we almost always use matrix algebra as an abstraction layer – i.e. vectorize the algorithm
• Then you can do an entire mini-batch in parallel
• Adjust the size of the mini-batch to suite the GPU that you have available
• More on this later…

23

Aside: Vectorizing logistic regression

• Computation we’d like to parallelize:
• For each x in the minibatch X_batch, compute

​

​

​

24

Note : x・w can be partially computed in parallel

Aside: Vectorizing logistic regression

• Computation we’d like to parallelize:
• For each x in the minibatch X_batch, compute

​

​

​

25

Note: X_batch ・w can be computed in parallel…

Aside: Vectorizing logistic regression

• Computation we’d like to parallelize:
• For each x in the minibatch X_batch, compute

​

​

​

26

Similarly for the rest of the algorithm…

28

Streaming SGD:

​

• Iterative
• Sequential
• Fast
• Scale up by bounding memory

​

• You can handle very large datasets … but slowly

​

​

x1

x2

x3

w

x4

x5

x6

Single examples x, y

29

Streaming SGD:

​

• Iterative
• Sequential
• Fast
• Scale up by bounding memory

​

• You can handle very large datasets … but slowly

​

• You can speed it up by making the tasks in the stream bigger and doing them in parallel
• Minibatch SGD with a GPU is a good way of doing that

X1

X2

X3

w

X4

X5

X6

Minibatch examples X, y

Limitations are

• communication to GPU
• GPU memory

Parameters being learned are stored on the GPU to reduce communication costs

Copy them out occasionally to checkpoint and test the model

Using GPUs for ML

30

#include <iostream>

#include <algorithm>

​

using namespace std;

​

#define N 1024

#define RADIUS 3

#define BLOCK_SIZE 16

​

__global__ void stencil_1d(int *in, int *out) {

__shared__ int temp[BLOCK_SIZE + 2 * RADIUS];

int gindex = threadIdx.x + blockIdx.x * blockDim.x;

int lindex = threadIdx.x + RADIUS;

​

// Read input elements into shared memory

temp[lindex] = in[gindex];

if (threadIdx.x < RADIUS) {

temp[lindex - RADIUS] = in[gindex - RADIUS];

temp[lindex + BLOCK_SIZE] = in[gindex + BLOCK_SIZE];

}

​

// Synchronize (ensure all the data is available)

__syncthreads();

​

// Apply the stencil

int result = 0;

for (int offset = -RADIUS ; offset <= RADIUS ; offset++)

result += temp[lindex + offset];

​

// Store the result

out[gindex] = result;

}

​

void fill_ints(int *x, int n) {

fill_n(x, n, 1);

}

​

int main(void) {

int *in, *out; // host copies of a, b, c

int *d_in, *d_out; // device copies of a, b, c

int size = (N + 2*RADIUS) * sizeof(int);

​

// Alloc space for host copies and setup values

in = (int *)malloc(size); fill_ints(in, N + 2*RADIUS);

out = (int *)malloc(size); fill_ints(out, N + 2*RADIUS);

// Alloc space for device copies

cudaMalloc((void **)&d_in, size);

cudaMalloc((void **)&d_out, size);

​

// Copy to device

cudaMemcpy(d_in, in, size, cudaMemcpyHostToDevice);

cudaMemcpy(d_out, out, size, cudaMemcpyHostToDevice);

​

// Launch stencil_1d() kernel on GPU

stencil_1d<<<N/BLOCK_SIZE,BLOCK_SIZE>>>(d_in + RADIUS, d_out + RADIUS);

​

// Copy result back to host

cudaMemcpy(out, d_out, size, cudaMemcpyDeviceToHost);

​

// Cleanup

free(in); free(out);

cudaFree(d_in); cudaFree(d_out);

return 0;

}

​

serial code

parallel code

serial code

parallel fn

Simple Processing Flow

31

1. Copy input data from CPU memory to GPU memory

PCI Bus

Simple Processing Flow

32

1. Copy input data from CPU memory to GPU memory
2. Load GPU program and execute,�caching data on chip for performance

PCI Bus

Simple Processing Flow

33

1. Copy input data from CPU memory to GPU memory
2. Load GPU program and execute,�caching data on chip for performance
3. Copy results from GPU memory to CPU memory

PCI Bus

SOME EXAMPLE CODE

34

Hello World!

35

int main(void) {

printf("Hello World!\n");

return 0;

}

​

• Standard C that runs on the host (the CPU)

​

• NVIDIA compiler (nvcc)

​

• We can also write code for the device (the GPU)

​

​

​

​

​

Output:

​

$ nvcc hello_world.cu

$ a.out

Hello World!

$

Hello World! with Device Code

36

__global__ void mykernel(void) {

}

​

int main(void) {

mykernel<<<1,1>>>();

printf("Hello World!\n");

return 0;

}

​

• Two new things here…

​

Hello World! with Device Code

__global__ void mykernel(void) {

}

​

• CUDA C/C++ keyword __global__ indicates a function that:
• Runs on the device
• Is called from host code

​

• nvcc separates source code into host and device components
• Device functions (e.g. mykernel()) processed by NVIDIA compiler
• Host functions (e.g. main()) processed by standard host compiler
• gcc, cl.exe

​

37

Hello World! with Device COde

mykernel<<<1,1>>>();

​

• Triple angle brackets mark a call from host code to device code
• Also called a “kernel launch”
• We’ll get to the parameters (1,1) soon

​

38

Hello World! with Device Code

39

__global__ void mykernel(void){

}

​

int main(void) {

mykernel<<<1,1>>>();

printf("Hello World!\n");

return 0;

}

​

• mykernel() does nothing at all in this example …. so let’s fix that.

​

​

​

​

Output:

​

$ nvcc hello.cu

$ a.out

Hello World!

$

Addition on the Device

• A simple kernel to add two integers (coming up: adding two arrays)

​

__global__ void add(int *a, int *b, int *c) {

*c = *a + *b;

}

​

​

• As before __global__ is a CUDA C/C++ keyword meaning
• add() will execute on the device
• add() will be called from the host

​

40

Addition on the Device

• Note that we use pointers for the variables

​

__global__ void add(int *a, int *b, int *c) {

*c = *a + *b;

}

​

• add() runs on the device, so a, b and c must point to device memory

​

• We need to allocate memory on the GPU

​

​

41

Memory Management

• Host and device memory are separate entities
• Device pointers point to GPU memory

May be passed to/from host code

May not be dereferenced in host code

• Host pointers point to CPU memory

May be passed to/from device code

May not be dereferenced in device code

​

• Simple CUDA API for handling device memory
• cudaMalloc(), cudaFree(), cudaMemcpy()
• Similar to the C equivalents malloc(), free(), memcpy()

​

42

Addition on the Device: add()

• Returning to our add() kernel

​

__global__ void add(int *a, int *b, int *c) {

*c = *a + *b;

}

​

​

• Let’s take a look at main()…

​

43

Addition on the Device: main()

int main(void) {

int a, b, c; // host copies of a, b, c

int *d_a, *d_b, *d_c; // device copies of a, b, c

int size = sizeof(int);

// Allocate space for device copies of a, b, c

cudaMalloc((void **)&d_a, size);

cudaMalloc((void **)&d_b, size);

cudaMalloc((void **)&d_c, size);

​

// Setup input values

a = 2;

b = 7;

​

44

We’re getting ready to do this…

45

1. Copy input data from CPU memory to GPU memory is coming up next!

PCI Bus

Addition on the Device: main()

// Copy inputs to device

cudaMemcpy(d_a, &a, size, cudaMemcpyHostToDevice);

cudaMemcpy(d_b, &b, size, cudaMemcpyHostToDevice);

​

// Launch add() kernel on GPU

add<<<1,1>>>(d_a, d_b, d_c);

​

// Copy result back to host

cudaMemcpy(&c, d_c, size, cudaMemcpyDeviceToHost);

​

// Cleanup

cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);

return 0;

}

46

Next: Vector Addition on the Device

• With add() running in parallel we can do vector addition

​

• Terminology: each parallel invocation of add() is referred to as a block
• The set of blocks is referred to as a grid
• Each invocation can refer to its block index using blockIdx.x

​

__global__ void add(int *a, int *b, int *c) {

c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];

}

​

• By using blockIdx.x to index into the array, each block handles a different index

​

47

Spoiler: blockIdx.y is also a thing (and probably blockIdx.z)

Vector Addition on the Device

__global__ void add(int *a, int *b, int *c) {

c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];

}

​

• On the device, each block can execute in parallel:

​

48

c[0] = a[0] + b[0];

c[1] = a[1] + b[1];

c[2] = a[2] + b[2];

c[3] = a[3] + b[3];

Block 0

Block 1

Block 2

Block 3

Vector Addition on the Device: add()

• Returning to our parallelized add() kernel

​

__global__ void add(int *a, int *b, int *c) {

c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];

}

​

​

• Let’s take a look at main()…

​

49

Vector Addition on the Device: main()

50

#define N 512

int main(void) {

int *a, *b, *c; // host copies of a, b, c

int *d_a, *d_b, *d_c; // device copies of a, b, c

int size = N * sizeof(int);

// Alloc space for device copies of a, b, c

cudaMalloc((void **)&d_a, size);

cudaMalloc((void **)&d_b, size);

cudaMalloc((void **)&d_c, size);

​

// Alloc space for host copies of a, b, c and setup input values

a = (int *)malloc(size); random_ints(a, N);

b = (int *)malloc(size); random_ints(b, N);

c = (int *)malloc(size);

Vector Addition on the Device: main()

51

// Copy inputs to device

cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);

cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);

​

// Launch add() kernel on GPU with N blocks

add<<<N,1>>>(d_a, d_b, d_c);

​

// Copy result back to host

cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost);

​

// Cleanup

free(a); free(b); free(c);

cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);

return 0;

}

a little more magic…

Coordinating Host & Device

• Kernel launches are asynchronous
• Control returns to the CPU immediately
• CPU needs to synchronize before consuming the results
• “Pipelineing” operations is useful

52

ML PLATFORMS AND GPUS

53

Some sample code

54

Demo

55

MORE DETAILS ON GPU PROGRAMMING

56

http://developer.download.nvidia.com/compute/developertrainingmaterials/presentations/cuda_language/Introduction_to_CUDA_C.pptx

Threads and Cores in GPUs

57

58

Simplified summary of GPUs vs CPUs

​

• less total memory
• more cores and more parallelism

​

• Multicore CPUs are mostly multiple-instruction multiple-data (MIMD)
• GPUs are mostly single-instruction multiple-data (SIMD)

​

Blocks, Grids, Threads, Warps

• Recall blocks are the things that work in parallel, and blocks are arranged in grids

​

• c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];

​

59

• That would be SIMD (single instruction multiple data)
• It’s actually more complicated than that….

​

Comparison

60

SIMT: Single Instruction Multiple Threads

61

A thread can access its own block id and also thread id. Blocks and threads are in a grid, which is 2D or 3D (there’s a .x and a .y part)

What’s in a GPU?

62

Threads (SIMT, synchronous threads) are grouped into cores (which are decoupled, like a MIMD machine)

IDs and Dimensions

• A kernel is launched as a grid of blocks of threads
• blockIdx and threadIdx are 3D
• We showed only one dimension (x)

​

• Built-in variables:
• threadIdx
• blockIdx
• blockDim
• gridDim

​

63

Device

Grid 1

Block�(0,0,0)

Block�(1,0,0)

Block�(2,0,0)

Block�(1,1,0)

Block�(2,1,0)

Block�(0,1,0)

Block (1,1,0)

Other kinds of accelerators

64

GPUs vs NPUs and TPUs

• Neural processing units: specifically designed for AI/ML (not graphics)
• Tensor processing units: Google’s brand of TPU
• Special memory (SparseCore) for embeddings
• Special hardware for matrix multiplication tasks
• High-bandwidth connections between TPU units in the same “pod”

​

65

• Other ML-oriented GPUs often designed to be clustered in a mesh or “torus” topology

​

Memory types

66

Memory Hierarchy in GPUs

• Slowest: fetching/storing on the CPU
• Faster:
• HBM (high-bandwidth memory): stored on or close to GPU processors, very wide bus
• eg 4k vs 32 bytes for GDDR
• SparseCore memory (sometimes)
• Even faster: GPU “global” memory (SRAM)
• Even faster: thread-local memory, caches, registers, …

67

Memory Hierarchy Inside A GPU

History of Google’s TPU’s

Bigger models need more memory

Will discuss HBM vs on-chip memory later