| ⭅ Previous (OpenCL 101) |
Introducing Gouda: Write CUDA Run Anywhere
Welcome back to our series on GPU and parallel programming. In our last two posts, we compared the code required to run a simple CUDA parallel program on an NVIDIA GPU, vs an OpenCL program on any other GPU or CPU.
After writing a few CUDA and OpenCL programs by hand, I had an idea. Perhaps we could generate the required OpenCL code in order to make OpenCL easy to write?
This experiment is called Gouda.
While still in the very early stages, it is now able to compile some non-trivial CUDA programs. In the rest of the article, I’ll talk about how it works, show some demos, and then share how you can try it yourself.
How it works: Not a compatibility layer
Gouda is not the first attempt to open up parallel programming. Though I believe it is unique in its approach. A few other notable approaches:
-
OpenGL (1.0 in 1992, 4.6 in 2017, Khronos) : A graphics-focused API. Cross platform, in contrast to Microsoft’s Direct-X graphics APIS. Later versions provide a shader language, which let you write parallel code to control pieces of the graphics pipeline.
-
OpenCL(2008, Khronis) : An open standard for running compute on parallel devices, including CPUs and GPUS. Unfortunately, its API is very verbose, and tooling is lacking compared to those for NVIDIA devices.
-
Vulkan(2016, Khronos) : Another standard, more focused on graphics. While it allows much lower level access to hardware of the GPU, this comes at a cost. While allowing for potentially higher performance, it is even harder to program than OpenGL.
-
ZLUDA(2020, Andrzej Janik) : Binary translation layer for CUDA. To allow compiled CUDA code to run on non-NVIDIA devices, this replaces the runtime library files, and substitutes its own version of the CUDA API, on-the-fly, before communicating with the real hardware.
-
Triton(2021, OpenAI) : A new, python-inspired language for writing GPU code. It is higher level, and aims to allow machine learning researchers to quickly write high performance code for the GPU.
Gouda’s approach is very simple. It neither tries to be a compatibility layer, nor invent a new language. The C language used by CUDA is essentially CPP with a few extra primitives to handle task distribution, and moving memory from host to device. We support the same API, and our compiler will rewrite your code for a variety of backends we support. C code in, C code out. Then you can integrated the generated code with the rest of your build system.
This combines the ease-of-use of CUDA, with the wide hardware support of OpenCL (and more backends to come).
Lets take a look at a simple but representative demo program.
Conway’s Game of Life
John Conway, the late mathemtician, studied many things but among them were cellular automata. These are simulated universes, where the value of each element in a grid depends on its neighbors. Cellular automata(CA) are an interesting model for exploring computation.
In Conway’s CA, called the Game of Life or sometimes just Life, every cell is either dead or alive. The cells evolve according to these rules, based on the 8 cells neighboring it.
- A live cell with less than 2 live neighbors dies (loneliness).
- A live cell greater than 3 live neighbors dies (overpopulation)
- A dead cell with exactly 3 live neighbors comes to life (reproduction).
These simple rules lead to very interesting behaviors. This setup is actually universal, so it is possible to build general purpose computers within the Game of Life.
We are interested in implementing the Game of Life, since this type of simulation is both very simple, but represents many types of parallel simulation of interest in science and engineering.
Life in C
(All source code can be found in the github here)
The main logic for life can be summarized in these two functions.
const int WIDTH=300;
const int HEIGHT=300;
// yellow for alive, black for dead
#define COLOR_ALIVE 0xFFFF00FF
#define COLOR_DEAD 0x000000FF
int countNeighbors(int x, int y, const vector<uint32_t> grid) {
int living = 0;
for (int dy = -1; dy <= 1; dy++) {
for (int dx = -1; dx <= 1; dx++) {
if (dx == 0 && dy == 0) continue; // exclude self.
int ny = y + dy;
int nx = x + dx;
if (ny >= 0 && ny < HEIGHT &&
nx >= 0 && nx < WIDTH) {
living += grid[ny*WIDTH + nx] == COLOR_ALIVE;
}
}
}
return living;
}
// Computes one step of life. Pixels from previous generation are inputs
// to the next.
static void stepLife(
const std::vector<uint32_t>& in_pixels,
std::vector<uint32_t>* out_pixels) {
for (int py = 0; py < HEIGHT; ++py) {
for (int px = 0; px < WIDTH; ++px) {
int neighbors = countNeighbors(px, py, in_pixels);
const bool wasLive = in_pixels[py*WIDTH + px] == COLOR_ALIVE;
bool alive = false;
if (wasLive) {
// underpopulation: <2 living neighbors dies
// survival: 2-3 neighbors
// overpopulation: >3 live neighbors dies
alive = neighbors >= 2 && neighbors <= 3;
} else {
// reproduction: exactly 3 neighbors: dead->live
alive = neighbors == 3;
}
uint32_t pixel = alive ? COLOR_ALIVE : COLOR_DEAD;
(*out_pixels)[py * WIDTH + px] = pixel;
}
}
}
We use an image for both the input and the output, to more easily integrate with our visualization / frontend. For each step, we run through each cell in the input, and produce one cell of output.
We instrument our application, so that every 60 steps we log how long the last one took to run. This gives us a simple way of estimating performance.
When running the CPU version above on our development laptop (Lenovo X1 carbon, 8th gen intel core i7 cpu), we observed the following:
$ ./life_cpu
Grid size: 200x200 cells.
Last frame time: 0.149863 sec (6.672766 fps)
Last frame time: 0.209113 sec (4.782110 fps)
Last frame time: 0.267670 sec (3.735948 fps)
Last frame time: 0.335950 sec (2.976632 fps)
Last frame time: 0.381501 sec (2.621226 fps)
Last frame time: 0.413070 sec (2.420898 fps)
Last frame time: 0.389354 sec (2.568360 fps)
Last frame time: 0.421950 sec (2.369949 fps)
Last frame time: 0.410780 sec (2.434392 fps)
Porting to Gouda
When writing for Gouda, we express the work needed as a function of some grid index. We eliminate the loop that runs over x and y, and instead spawn our kernel function to compute over all the pixels in the input.
// Kernel is almost unchanged, besides the px=,py= lines to figure out our assigned pixel.
__global__ void lifeCell(const uint32_t* in_pixels, uint32_t* out_pixels, int WIDTH, int HEIGHT) {
int px = blockDim.x*blockIdx.x + threadIdx.x;
int py = blockDim.y*blockIdx.y + threadIdx.y;
if (px >= WIDTH || py >= HEIGHT) {
return;
}
int neighbors = countNeighbors(px, py, in_pixels, WIDTH, HEIGHT);
const bool wasLive = in_pixels[py*WIDTH + px] == COLOR_ALIVE;
bool alive = false;
if (wasLive) {
// underpopulation: <2 living neighbors dies
// survival: 2-3 neighbors
// overpopulation: >3 live neighbors dies
alive = neighbors >= 2 && neighbors <= 3;
} else {
// reproduction: exactly 3 neighbors: dead->live
alive = neighbors == 3;
}
uint32_t pixel = alive ? COLOR_ALIVE : COLOR_DEAD;
out_pixels[py * WIDTH + px] = pixel;
}
// misc code omitted
// Spawn the kernel across the input:
dim3 tile(16, 16, 1);
dim3 tiles(
div_ceil(WIDTH, 16),
div_ceil(HEIGHT, 16),
1);
lifeCell<<<tiles, tile>>>(device_pixels_in, device_pixels_out, WIDTH, HEIGHT);
And when we run it, we observe quite a nice speedup:
# full makefile in the repo.
# https://github.com/goudalang/gouda/blob/main/demos/life/Makefile
gouda life.cu --backend opencl --transpile # outputs tmp.cc
g++ tmp.cc -lsdl2 -lOpenCL -o life_cu
./life_cu
Last frame time: 0.000191 sec (5228.512207 fps) (209140496.000000 cells/sec)
Last frame time: 0.000254 sec (3929.643555 fps) (157185744.000000 cells/sec)
Last frame time: 0.000192 sec (5204.240723 fps) (208169616.000000 cells/sec)
Our time-per frame went from 0.149863 sec to 0.000191 sec, or a speedup of ~784.
Getting started with Gouda
The latest version of Gouda, which was used to compile the above demo, is available via releases on Github. Many things are not yet working, but as we saw above it can already compile some interesting simulation code.
If you encounter issues, feel free to open an issue or send patches :)
| ⭅ Previous (OpenCL 101) |