Summary and Schedule

Introduction

---

This is an advanced course in GPU programming using Python. You will learn about the types of problems that are suitable for parallelisation on the GPU and how to reformulate a problem using data parallelism. We will show you how to use the high-level GPU API provided by CuPy and then how to write your own, highly optimised kernels.

This course can be separated into two days:

Day 1

Day 2

• An intro to parallelism
• GPU fundamentals
• High-level GPU programming with CuPy
• Writing your first kernel

• Floating point precision
• Shared memory optimisations
• Thread coarsening
• Explicit FMA instructions

• Kernel optimisation, including:
  • Floating point precision
  • Shared memory optimisations
  • Thread coarsening
  • Explicit FMA instructions

Where: The workshop will be held at Swinburne university in the Engineering building.

When: The workshop will run on the 21st-22nd of July, 10am - 4pm each day. Arrive as early as 9am if you need help setting up your computer/environment.

Why use GPU acceleration?

It is increasingly common for science to make use of massive datasets—such as high-resolution sky surveys, complex N-body simulations, or signal processing over terabytes of data. Processing these vast datasets can be computationally prohibitive on standard CPUs. However, it is our good fortune that many of these tasks are trivially parallelisable, meaning that the work can be progressed by many processing cores at once.

GPU acceleration is the next step beyond CPU parallelisation, and allows you to leverage literally thousands of parallel cores to accelerate these tasks. GPUs are built specifically with high multi-threaded workloads in mind, and are efficient at these kinds of tasks in a way that CPUs are simply not. GPU programming is not easy, but with a little care (and experimentation) it is possible to achieve speedups of several orders of magnitude in processing time.

Or, let Mythbusters convince you:

Source: Nvidia on Youtube

Why Python?

This course is taught using Python and a number of libraries that allow writing CUDA kernels directly using Python. Python has been chosen to make this course as accessible as possible, which allows us to focus on the underlying concepts without getting snagged on unfamiliar syntax. Under the hood, the Python kernels are compiled to CUDA kernels and run just as fast as if we had written them in C.

And rest assured: the concepts learned here can be applied wholesale to CUDA programming in C or C++ (or Julia, Rust, …).

	Setup Instructions	Download files required for the lesson
Duration: 00h 00m	1. An introduction to parallelism	Why does adding more cores not always make a program faster? What distinguishes data parallelism from task parallelism, and why do GPUs rely on the former? Why does a parallel program sometimes give a different answer each time you run it?
Duration: 00h 40m	2. A GPU Deep Dive	How is a GPU organised internally, and how does that differ from a CPU? Why does the GPU have multiple types of memory with very different speeds and sizes? How does a GPU avoid sitting idle while waiting for slow memory operations?
Duration: 00h 50m	3. CuPy: A Numpy-like GPU experience	If NumPy arrays live in CPU memory, where do CuPy arrays live and what does that mean for your code? Why does benchmarking GPU code give misleadingly fast results if you aren’t careful? How can you overlap independent computations and data transfers on the GPU?
Duration: 01h 00m	4. Writing your first GPU kernels	How does a single function body know which piece of data to process when launched thousands of times? How do you configure a kernel launch so every element of an array is covered, regardless of its size? How can threads safely combine their results into a single value?
Duration: 01h 10m	5. Optimisation	Why is accessing global memory in a kernel’s hot loop a performance bottleneck? How can threads within a block cooperate to share the cost of reading data from global memory? When does using lower-precision floating point numbers make your kernel faster?
Duration: 01h 20m	Finish

The actual schedule may vary slightly depending on the topics and exercises chosen by the instructor.

Expected background

---

This course sets a reasonably high bar for what you are expected to already know, including:

• Python
• Primitive data types: integers, floats, doubles, etc. and their machine representation
• numpy, including broadcasting rules
• Some experience with numba and its parallelisation functions

Additionally, the examples given here are aimed towards people coming from a physics, astronomy, or engineering background. It will help if you are familiar with:

• Basic mathematical sum notation, e.g. for discrete Fourier transform
• Complex numbers (and, for example, Euler’s formula)

Setup

---

It is expected that you will be working on OzStar for this workshop, but if you have a CUDA based GPU you can work on your local machine (not recommended).

To work on ozstar you will need an account, and you will need to have joined the group oz983.

Making an account and joining the group takes time as there is a human approval process required for each, so please make sure that you do this before attending the workshop.

Connecting to Ozstar

Log into ozstar using ssh via:

BASH

ssh <USERNAME>@ozstar.swin.edu.au

This should connect you to either farnarkle1 or farnarkle2 as these are the login nodes.

See the documentation here for tips on how to setup easy access via ssh.

Installing the environment on HPC

Installing the environment on a HPC like OzStar will require loading the gcc and CUDA modules and installing uv.

Installing uv should require no change to the Linux installation instructions as the bash script that runs should install uv to your home directory on the login node of the HPC. Using uv to install the environment on OzStar:

BASH

# Work in your home directory
cd ~
# Load the required software modules
module load gcc/13.3.0 cuda/12.8.0 python/3.12.3
# Install uv
curl -LsSf https://astral.sh/uv/install.sh | sh
# Now make a new directory and environment for us to work in
uv init gpu-tutorial
cd gpu-tutorial
# Note that the following may take up to 1/2 hour to complete (thanks cupy!)
uv add numpy \
       numba \
       numba-cuda \
       cupy \
       matplotlib \
       pytest
# Activate this new environment
source .venv/bin/activate
# Test that everything is working
python -c 'import numpy, numba, numba.cuda, cupy, pytest, matplotlib'

If you can run all of the above without generating errors then you are ready to go. If you find errors then either email the organisers of the workshop, or be sure to arrive early on the first day to get in-person help.

Running Python on a HPC

Running the code in this tutorial, will require you to interact with SLURM on OzStar, you can do this via sinteractive (recommended) or via sbatch (optional).

sinteractive

sbatch

Start an interactive session with OzStar:

BASH

# You may want to make the time longer or shorter depending on what you need
sinteractive --gres=gpu:1 --mem 16G --cpus-per-task=8 --time=04:00:00 --reservation=curtin_training 

You will then see something like the following

OUTPUT

srun: job 14313053 queued and waiting for resources
srun: job 14313053 has been allocated resources
[phancock@john26 gpu-tutorial]$ 

There may be a delay between the “waiting” and “allocated” stage, though with the reservation in place this may only a be a few seconds. You will know that your job has been allocated because your command prompt should now be [<user>@<node> <working directory>], where <node> is not farnarkle[12].

Once you’re in the interactive shell, you’ll need to reload the modules and the environment before running any code:

BASH

module load gcc/13.3.0 cuda/12.8.0 python/3.12.3
cd ~/gpu-tutorial
source .venv/bin/activate

The usual way to run scripts on an HPC is to submit a job script which contains all the information about the job you are running, and how it should run.

In a file called run_python_script.sh you would include:

BASH

#!/bin/bash 
#SBATCH --job-name=myjob 
#SBATCH --out=myjob-%j.out
#SBATCH --gres=gpu:1 
#SBATCH --cpus-per-task=8
#SBATCH --mem=16G 
#SBATCH --time=00:15:00 
#SBATCH --reservation=curtin_training

module load gcc/13.3.0 cuda/12.8.0 python/3.12.3
cd ~/gpu-tutorial
source .venv/bin/activate

python some_python_file_here.py

Then submit this script to the job scheduler using sbatch:

$ sbatch run_python_script.sh

When the job runs it will put any output (STDERR + STDOUT) into the file myjob-<jobid>.out in the current directory. This file will include a short summary from the job scheduler about the resources used.

Editing your files on an HPC

There are many options available for the edit/run/fix loop that we will be engaged in during this workshop. Depending on your prefered editor (and patience) we recommend the following:

VSCode

vim/emacs/nano

Other

1. Download and install VSCode on your local machine.
2. Install the extension remote-ssh, plus whatever other extensions you love.
3. Press the >< looking button on the bottom left of your window.
4. Choose “Connect to host” 1. If you set up your .ssh/config file you should see “ozstar” as an option 1. Otherwise choose “Add new SSH host” and enter ssh <username>@ozstar.swin.edu.au and then your password if/when requested.
5. Choose File -> Open workspace from Folder”
6. Enter /home/<username>>/gpu-tutorial/
7. You should see the files on the remote system in your file explorer tab
8. Open the built in terminal (default is <Ctrl>+~)
9. This terminal is already logged into ozstar so you can use the sinteractive commands above to join an interactive session
10. From here, you can edit/save files in VSCode, and run them in the terminal without having to swap between them.

The workflow here is:

1. ssh into ozstar
2. sinteractive to get a node with a gpu
3. use vim/emacs/nano to create/edit files
4. save file, quit editor
5. run file using python myfile.py
6. observe success or errors
7. if errors GOTO 3
8. success

This works just fine for short scripts but the continually open/edit/save/run loop can become tedious, especially when you are making typo level errors.

Whatever your prefered editor is. We’ll try to help if you have problems, but consider this a self supported option!

Imports

---

We will frequently omit imports from the code snippets listed throughout this tutorial. The following imports should be assumed:

PYTHON

# Python imports
import math

# Package imports
import cupy
import cupyx
import matplotlib.pyplot as plt
from numba import cuda, njit, prange
import numpy as np

AI Declaration

---

With the exception of the following, no AI or LLM tools were used to prepare these notes in any capacity (including planning, writing, or otherwise):

• The shared memory optimisation animation (using Gemini 3).
• Final proof and code checking (using Qwen 3.6)