Your First Kernel, Called from Go — The Programming Model

The idea in one minute#

A real GPU program has two halves. The kernel is written in CUDA C and compiled for the GPU. The host program — ours is in Go — allocates GPU memory, copies data in, launches the kernel, and copies results out. Go reaches the CUDA half through cgo, Go’s built-in bridge to C.

Every GPU program, in any language, performs the same five steps you are about to write.

A picture#

sequenceDiagram
  participant Go as Go program (host)
  participant C as CUDA C wrapper
  participant GPU as GPU (device)
  Go->>C: saxpy(a, x, y, n) via cgo
  C->>GPU: 1. allocate device memory
  C->>GPU: 2. copy x and y to the device
  C->>GPU: 3. launch kernel on the grid
  Note over GPU: thousands of threads run
  C->>GPU: 4. copy result back
  GPU-->>C: bytes
  C->>GPU: 5. free device memory
  C-->>Go: return

How it really works#

The kernel (CUDA C)#

CUDA C++

// saxpy.cu
#include <cuda_runtime.h>

// __global__ marks a function that runs on the GPU and is launched from the host.
__global__ void saxpy_kernel(int n, float a, const float *x, float *y) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;   // which element am I?
    if (i >= n) return;                              // bounds check
    y[i] = a * x[i] + y[i];
}

// A plain C function Go can call. It performs the five steps.
extern "C" int saxpy(int n, float a, const float *host_x, float *host_y) {
    float *x, *y;
    size_t bytes = n * sizeof(float);

    if (cudaMalloc(&x, bytes) != cudaSuccess) return 1;         // 1. allocate
    if (cudaMalloc(&y, bytes) != cudaSuccess) return 1;

    cudaMemcpy(x, host_x, bytes, cudaMemcpyHostToDevice);       // 2. copy in
    cudaMemcpy(y, host_y, bytes, cudaMemcpyHostToDevice);

    int block = 256;
    int grid  = (n + block - 1) / block;
    saxpy_kernel<<<grid, block>>>(n, a, x, y);                  // 3. launch

    cudaMemcpy(host_y, y, bytes, cudaMemcpyDeviceToHost);       // 4. copy out (waits for the kernel)

    cudaFree(x); cudaFree(y);                                   // 5. free
    return cudaGetLastError() == cudaSuccess ? 0 : 2;
}

The only unusual syntax is <<<grid, block>>>: it is how CUDA says “launch this function on a grid of this many blocks of this many threads”.

Compile it into a shared library with NVIDIA’s compiler:

Shell

nvcc -O2 -shared -fPIC -o libsaxpy.so saxpy.cu

The host program (Go)#

// main.go — build with: CGO_ENABLED=1 go build && LD_LIBRARY_PATH=. ./first-kernel
package main

/*
#cgo LDFLAGS: -L. -lsaxpy
int saxpy(int n, float a, const float *x, float *y);
*/
import "C"

import (
	"fmt"
	"unsafe"
)

func main() {
	const n = 1 << 20
	x, y := make([]float32, n), make([]float32, n)
	for i := range x {
		x[i], y[i] = float32(i), 1
	}

	// Hand C a pointer to each slice's backing array. cgo keeps the slices
	// alive and in place for the duration of the call.
	rc := C.saxpy(C.int(n), C.float(2),
		(*C.float)(unsafe.Pointer(&x[0])),
		(*C.float)(unsafe.Pointer(&y[0])))
	if rc != 0 {
		panic(fmt.Sprint("saxpy failed: ", rc))
	}
	fmt.Println(y[0], y[1], y[n-1]) // 1 3 2.097151e+06
}

Three things to notice#

The kernel is tiny; the plumbing is not. Three lines of real work, fifteen of memory management. That ratio is normal.
This program is slow on purpose. It copies 8 MB over PCIe, does a trivial amount of arithmetic, and copies it back. The copies take far longer than the kernel. A CPU would win. GPUs pay off when data stays on the device across many kernels.
cgo has a cost of its own. Each Go→C call costs on the order of 50–100 nanoseconds and pins the calling goroutine to an OS thread while it runs. That is nothing compared with a kernel launch, but it means you make few, large calls — never one cgo call per element.

Other ways to reach a GPU from Go#

Approach	What it is	When to use
cgo + your own CUDA C	What you just did	Custom kernels, full control
cgo + a library (cuBLAS, ONNX Runtime, llama.cpp)	Call existing optimized code	Almost always the practical answer
Driver API without cgo (`purego`-style bindings)	Load `libcuda` at run time, no C compiler in the build	Simpler builds and cross-compiling; kernels shipped precompiled
A separate inference server (vLLM, Triton, Ollama) over HTTP/gRPC	The GPU lives in another process	Production serving; the usual choice for Go services
`go-nvml`	Management only: read state, no compute	Monitoring and scheduling (IV.05)

Most production Go code never launches a kernel. It talks to a server that does. Knowing what that server is doing underneath is what this course is for.

Code (no GPU? do this)#

If you cannot run the CUDA version, extend the simulator from lesson 01 so that it has a separate “device memory” and forces you through the same five steps:

// device.go — a pretend GPU that enforces the five-step discipline.
package main

import "fmt"

type Device struct {
	mem    map[int][]float32 // device allocations, by handle
	next   int
	copies int // how many times data crossed the "PCIe link"
}

func NewDevice() *Device { return &Device{mem: map[int][]float32{}} }

func (d *Device) Malloc(n int) int              { d.next++; d.mem[d.next] = make([]float32, n); return d.next }
func (d *Device) Free(h int)                    { delete(d.mem, h) }
func (d *Device) CopyIn(h int, host []float32)  { copy(d.mem[h], host); d.copies++ }
func (d *Device) CopyOut(host []float32, h int) { copy(host, d.mem[h]); d.copies++ }

// Launch runs kernel for every index. The kernel can only see device memory.
func (d *Device) Launch(n int, kernel func(i int, mem map[int][]float32)) {
	for i := 0; i < n; i++ {
		kernel(i, d.mem)
	}
}

func main() {
	const n = 8
	hostX, hostY := make([]float32, n), make([]float32, n)
	for i := range hostX {
		hostX[i], hostY[i] = float32(i), 1
	}

	dev := NewDevice()
	x, y := dev.Malloc(n), dev.Malloc(n) // 1. allocate
	dev.CopyIn(x, hostX)                 // 2. copy in
	dev.CopyIn(y, hostY)
	dev.Launch(n, func(i int, m map[int][]float32) { // 3. launch
		m[y][i] = 2*m[x][i] + m[y][i]
	})
	dev.CopyOut(hostY, y) // 4. copy out
	dev.Free(x)           // 5. free
	dev.Free(y)

	fmt.Println(hostY, "link crossings:", dev.copies)
}

Remember this#

Five steps: allocate, copy in, launch, copy out, free.
Kernels are CUDA C; the host can be Go via cgo.
Copies across PCIe usually cost more than a simple kernel. Keep data on the device.
Make few, large cgo calls.

Try it#

(GPU) Build and run the CUDA version. Time the whole saxpy call for n = 2^10, 2^20 and 2^26. Then time a plain Go loop doing the same. Where does the GPU start to win, if at all?
(GPU) Split saxpy into upload, run and download functions so run can be called 100 times without copying. Time 100 runs. This is how real programs are structured.
(No GPU) In device.go, run the kernel 100 times. Count link crossings when you copy in and out every time versus once. Which design would you ship?

Check yourself#

What are the five steps of every GPU program?
Why is the SAXPY example slower on a GPU than on a CPU?
Why should cgo calls be few and large?