11. The GPU Host Side: Device Memory, Real Runtime Calls, and a Stub "Device"¶
What you will understand: what has to surround a GPU kernel on the host before it can run (allocating device memory, copying inputs in and the result out, loading the kernel module, launching it with a grid, and cleaning up) and how MLIR 18's real passes turn those steps into calls to a GPU runtime (mgpuMemAlloc, mgpuModuleLoadJIT, mgpuLaunchKernel, ...). This chapter closes the "what is still missing" list Chapter 10 ended with, as far as this sandbox allows, and it is explicit about the one thing it still cannot do: run the PTX. Every file, script and output below is embedded from the repository.
What you need to know first: Chapter 10's kernel (gpu.module @add_tensors_kernel, 2 blocks x 2 threads, 23 parameters) and Chapter 8's harness (the same Memref2D struct-return convention). New ground: the async form of gpu ops, and where in the pipeline gpu.launch_func actually becomes runtime calls.
Compile and run
Every command, listing and output on this page is reproduced by the commands in Reproducing this chapter at the bottom of the page, which also lists what must be built first. The chain of builds is in Getting Started.
A wrong or missing answer here is the point
This chapter's controls and reproductions are supposed to look broken: a wrong answer, a crash, or a non-zero exit status. They show what happens without the fix or without the real behavior. The correct build's output is shown beside them. Each script says so in its own header.
Primer: what the host does around a kernel¶
A GPU has its own memory. A kernel can only read and write memory on the device, so a host program that wants to use one must, in order:
- Allocate buffers on the device for the inputs and the output.
- Copy the inputs from host memory to those device buffers.
- Load the compiled kernel (here, the PTX text) as a module, and look up the kernel by name.
- Launch it with a grid size, a block size, and an array of pointers to its arguments.
- Copy the result back, and free the device buffers.
Operations on a GPU are normally queued on a stream and run asynchronously relative to the host: the host enqueues work and waits only when it needs the result. MLIR models this with async tokens: each async gpu operation takes the tokens of the operations it depends on and returns a token of its own, and a gpu.wait blocks the host on a token. Everything above is real runtime-library work (NVIDIA's driver wrappers, or MLIR's own wrapper library), which is why the compiled host code ends up as calls to functions with names like mgpuMemAlloc: the compiler emits the calls, and a library implements them.
What changed from Chapter 10's plan¶
Chapter 10 ended by saying the remaining steps were --gpu-to-llvm, a runtime library, and device-visible allocation. Two of those turned out differently when run:
--gpu-to-llvmalone does not produce the launch call in this toolchain (LLVM/MLIR 18.1.3, Ubuntu'smlir-opt-18); the launch is lowered later, at LLVM IR translation time (below). Chapter 10's sentence namingmgpuModuleLoadwas also slightly off: the real name ismgpuModuleLoadJIT.- There is still no GPU runtime library on this machine, so a stub one is written instead.
The host function, written by hand (and why)¶
Chapter 10's host code used plain memref.alloc, i.e. ordinary CPU memory, which a real device cannot read. This chapter's host function allocates device buffers with gpu.alloc, copies with gpu.memcpy, launches, copies the result back, and frees. This function is hand-written, not generated from mg.add: this book has no pass that inserts gpu.alloc/gpu.memcpy around a kernel, and writing one is not attempted here. The kernel module below it is Chapter 10's own add_gpu.mlir module, unchanged. The whole file:
#map = affine_map<(d0)[s0, s1] -> ((d0 - s0) ceildiv s1)>
#map1 = affine_map<(d0)[s0, s1] -> (d0 * s0 + s1)>
module attributes {gpu.container_module} {
func.func @add_tensors(%a: memref<2x2xf64>, %b: memref<2x2xf64>) -> memref<2x2xf64> {
%da = gpu.alloc () : memref<2x2xf64>
%db = gpu.alloc () : memref<2x2xf64>
%dc = gpu.alloc () : memref<2x2xf64>
gpu.memcpy %da, %a : memref<2x2xf64>, memref<2x2xf64>
gpu.memcpy %db, %b : memref<2x2xf64>, memref<2x2xf64>
%c0 = arith.constant 0 : index
%c1 = arith.constant 1 : index
%c2 = arith.constant 2 : index
gpu.launch_func @add_tensors_kernel::@add_tensors_kernel blocks in (%c2, %c1, %c1) threads in (%c2, %c1, %c1) args(%c1 : index, %c0 : index, %da : memref<2x2xf64>, %db : memref<2x2xf64>, %dc : memref<2x2xf64>)
%r = memref.alloc() : memref<2x2xf64>
gpu.memcpy %r, %dc : memref<2x2xf64>, memref<2x2xf64>
gpu.dealloc %da : memref<2x2xf64>
gpu.dealloc %db : memref<2x2xf64>
gpu.dealloc %dc : memref<2x2xf64>
return %r : memref<2x2xf64>
}
gpu.module @add_tensors_kernel {
gpu.func @add_tensors_kernel(%arg0: index, %arg1: index, %arg2: memref<2x2xf64>, %arg3: memref<2x2xf64>, %arg4: memref<2x2xf64>) kernel {
%0 = gpu.block_id x
%1 = gpu.block_id y
%2 = gpu.block_id z
%3 = gpu.thread_id x
%4 = gpu.thread_id y
%5 = gpu.thread_id z
%6 = gpu.grid_dim x
%7 = gpu.grid_dim y
%8 = gpu.grid_dim z
%9 = gpu.block_dim x
%10 = gpu.block_dim y
%11 = gpu.block_dim z
cf.br ^bb1
^bb1: // pred: ^bb0
%12 = affine.apply #map1(%0)[%arg0, %arg1]
%c0 = arith.constant 0 : index
%c2 = arith.constant 2 : index
%c1 = arith.constant 1 : index
%c1_0 = arith.constant 1 : index
%c0_1 = arith.constant 0 : index
%13 = affine.apply #map1(%3)[%c1_0, %c0_1]
%14 = memref.load %arg2[%12, %13] : memref<2x2xf64>
%15 = memref.load %arg3[%12, %13] : memref<2x2xf64>
%16 = arith.addf %14, %15 : f64
memref.store %16, %arg4[%12, %13] : memref<2x2xf64>
gpu.return
}
}
}
Host function first (lines 4 to 20): three gpu.allocs, two gpu.memcpys in, the gpu.launch_func with a 2x1x1 grid and a 2x1x1 block, a host-side memref.alloc for the result, a gpu.memcpy back, and three gpu.deallocs. The kernel module follows, exactly as in Chapter 10. mlir-opt-18 parses and verifies the file as is.
A real run of dead ends¶
Getting from this file to runtime calls took several real failures. They are recorded because each one cost time and none is obvious from the pass names. Chapter 12 later explained all three from MLIR's own source, and its symptoms.sh reproduces each one with saved output; the real outputs are shown here.
1. gpu-to-llvm after gpu-module-to-binary crashes mlir-opt-18. A pipeline of nvvm-attach-target, convert-gpu-to-nvvm, gpu-module-to-binary{format=isa}, gpu-to-llvm, then the usual llvm conversions died with a segmentation fault:
Stack dump:
0. Program arguments: mlir-opt-18 symptoms/async.mlir "--pass-pipeline=builtin.module(lower-affine,nvvm-attach-
#0 0x00007fd60f7a63bf llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/usr/lib/llvm-18/bin/../lib/libLLVM.so.18.1+0xda63bf)
#4 0x000055fed77ded5b mlir::gpu::GPUModuleOp::getTargetsAttr() (/usr/lib/llvm-18/bin/mlir-opt+0x14a0d5b)
(Update, added after Chapter 12: reading GPUToLLVMConversion.cpp at tag llvmorg-18.1.3 established the cause: the launch pattern looks the kernel module up as a gpu::GPUModuleOp, which gpu-module-to-binary has already replaced, and dereferences the null result. See Chapter 12. The explanation in the next sentence is what was believed when this chapter was written.) This chapter did not diagnose it beyond narrowing it down: with the gpu.launch_func line deleted the same gpu-to-llvm run does not crash, so the launch is involved. Treat "it crashes because the kernel gpu.module has become a gpu.binary" as a hypothesis consistent with the top frame, not a confirmed cause.
2. gpu-to-llvm silently left gpu.alloc and gpu.memcpy alone. With the launch removed, the pass "succeeded" but a grep for mgpu calls found none. The ops were still there:
%memref = gpu.alloc () : memref<2x2xf64>
%memref_0 = gpu.alloc () : memref<2x2xf64>
%memref_1 = gpu.alloc () : memref<2x2xf64>
gpu.memcpy %memref, %8 : memref<2x2xf64>, memref<2x2xf64>
gpu.memcpy %memref_0, %17 : memref<2x2xf64>, memref<2x2xf64>
gpu.memcpy %62, %memref_1 : memref<2x2xf64>, memref<2x2xf64>
gpu.dealloc %memref : memref<2x2xf64>
gpu.dealloc %memref_0 : memref<2x2xf64>
gpu.dealloc %memref_1 : memref<2x2xf64>
and a later pass failed with failed to legalize operation 'builtin.unrealized_conversion_cast'. The fix was gpu-async-region, which rewrites these ops into their async form with !gpu.async.tokens and stream handling. The evidence that async form is the requirement is that adding that one pass changed the output from no runtime calls to mgpuMemAlloc/mgpuMemcpy/mgpuStreamCreate (after it, the same listing is empty: 0 of these ops left). This chapter did not read the pass's source; Chapter 12 did.
3. Putting gpu-to-llvm before the kernel conversion got past the crash but failed verification:
symptoms/async.mlir:13:10: error: 'gpu.launch_func' op got 23 kernel operands but expected 5
%3 = gpu.launch_func async [%2] @add_tensors_kernel::@add_tensors_kernel blocks in (%c2, %c1, %c1) threads in (%c2, %c1, %c1) args(%c1 : index, %c0 : index, %memref : memref<2x2xf64>, %memref_0 :
The pass had already expanded the three memref launch arguments into 23 scalars at the call site while the kernel still took 5 memref arguments. It is the same 2 + 3 x 7 = 23 expansion Chapter 10 found on the kernel side, now seen from the host side.
What finally worked¶
MLIR's own packaged pipeline, run after gpu-async-region. The whole script (pipeline.sh):
#!/bin/sh
# Chapter 11: host + device lowering with the real runtime calls, then run against a CPU stub runtime.
# Input: add_host_device.mlir (Chapter 10's kernel module + a hand-written host function using gpu.alloc/memcpy).
set -e
# 1. Make the gpu.* host ops async -- LLVM 18's gpu-to-llvm patterns only match async ops.
mlir-opt-18 add_host_device.mlir --pass-pipeline='builtin.module(func.func(gpu-async-region))' -o async.mlir
# 2. MLIR's own pipeline: kernel -> nvvm -> PTX (gpu.binary), host -> llvm dialect.
mlir-opt-18 async.mlir --gpu-lower-to-nvvm-pipeline="cubin-chip=sm_70 cubin-features=+ptx60 cubin-format=isa" -o pipe_out.mlir
# 3. gpu.launch_func is turned into mgpuModuleLoadJIT / mgpuLaunchKernel calls at LLVM-IR translation time.
mlir-translate-18 --mlir-to-llvmir pipe_out.mlir -o host.ll
# 4. Link with Chapter 8's unmodified harness and the stub runtime, run.
clang-18 -c host.ll -o host.o 2>/dev/null
clang-18 harness.c host.o stub_gpu_runtime.c -o add_tensors_stub_demo
./add_tensors_stub_demo
--gpu-lower-to-nvvm-pipeline appeared in mlir-opt-18 --help while writing Chapter 10; this is its first use. Its options in this version are named cubin-*; Chapter 10's standalone passes used gpu-* and pass-parameter names.
Step 1: gpu-async-region¶
This is the real output of the first step (async.mlir, host function). Compare it with the hand-written version above:
#map = affine_map<(d0)[s0, s1] -> (d0 * s0 + s1)>
module attributes {gpu.container_module} {
func.func @add_tensors(%arg0: memref<2x2xf64>, %arg1: memref<2x2xf64>) -> memref<2x2xf64> {
%0 = gpu.wait async
%memref, %asyncToken = gpu.alloc async [%0] () : memref<2x2xf64>
%memref_0, %asyncToken_1 = gpu.alloc async [%asyncToken] () : memref<2x2xf64>
%memref_2, %asyncToken_3 = gpu.alloc async [%asyncToken_1] () : memref<2x2xf64>
%1 = gpu.memcpy async [%asyncToken_3] %memref, %arg0 : memref<2x2xf64>, memref<2x2xf64>
%2 = gpu.memcpy async [%1] %memref_0, %arg1 : memref<2x2xf64>, memref<2x2xf64>
%c0 = arith.constant 0 : index
%c1 = arith.constant 1 : index
%c2 = arith.constant 2 : index
%3 = gpu.launch_func async [%2] @add_tensors_kernel::@add_tensors_kernel blocks in (%c2, %c1, %c1) threads in (%c2, %c1, %c1) args(%c1 : index, %c0 : index, %memref : memref<2x2xf64>, %memref_0 : memref<2x2xf64>, %memref_2 : memref<2x2xf64>)
gpu.wait [%3]
%alloc = memref.alloc() : memref<2x2xf64>
%4 = gpu.wait async
%5 = gpu.memcpy async [%4] %alloc, %memref_2 : memref<2x2xf64>, memref<2x2xf64>
%6 = gpu.dealloc async [%5] %memref : memref<2x2xf64>
%7 = gpu.dealloc async [%6] %memref_0 : memref<2x2xf64>
%8 = gpu.dealloc async [%7] %memref_2 : memref<2x2xf64>
gpu.wait [%8]
return %alloc : memref<2x2xf64>
}
Read the token chain: %0 = gpu.wait async creates an initial token; the first gpu.alloc async [%0] depends on it and returns %asyncToken; the next alloc depends on that; and so on through the two copies ([%asyncToken_3], then [%1]) to the launch (gpu.launch_func async [%2]), whose token %3 the host then waits on (gpu.wait [%3]). After the result is copied back the three deallocations chain the same way, ending in gpu.wait [%8]. The pass did not change what happens, only made each step's ordering explicit, in exactly the form gpu-to-llvm's patterns accept.
Step 2: after the packaged pipeline¶
After that pipeline, the output (pipe_out.mlir) contains all the host-side runtime calls except the launch: mgpuMemAlloc x3, mgpuMemcpy x3, mgpuMemFree x3, mgpuStreamCreate/Synchronize/Destroy x2 each, one malloc. The kernel is a gpu.binary holding Chapter 10's PTX as a string, and gpu.launch_func is still an MLIR op, now with a stream operand and the 23 expanded operands (line 17):
gpu.launch_func <%4 : !llvm.ptr> @add_tensors_kernel::@add_tensors_kernel blocks in (%3, %2, %2) threads in (%3, %2, %2) : i64 args(%2 : i64, %1 : i64, %8 : !llvm.ptr, %8 : !llvm.ptr, %1 : i64, %3 : i64, %3 : i64, %3 : i64, %2 : i64, %9 : !llvm.ptr, %9 : !llvm.ptr, %1 : i64, %3 : i64, %3 : i64, %3 : i64, %2 : i64, %10 : !llvm.ptr, %10 : !llvm.ptr, %1 : i64, %3 : i64, %3 : i64, %3 : i64, %2 : i64)
Step 3: the launch is lowered at translation time¶
The remaining step is one Chapter 1 already used for a different purpose: mlir-translate-18 --mlir-to-llvmir pipe_out.mlir -o host.ll. The full translated file is 158 lines; here is what it does, in three pieces. First, create a stream, allocate three 32-byte device buffers (2 x 2 x 8 bytes) and copy the two inputs in:
%15 = call ptr @mgpuStreamCreate()
%16 = call ptr @mgpuMemAlloc(i64 ptrtoint (ptr getelementptr (double, ptr null, i32 4) to i64), ptr %15, i8 0)
%17 = call ptr @mgpuMemAlloc(i64 ptrtoint (ptr getelementptr (double, ptr null, i32 4) to i64), ptr %15, i8 0)
%18 = call ptr @mgpuMemAlloc(i64 ptrtoint (ptr getelementptr (double, ptr null, i32 4) to i64), ptr %15, i8 0)
call void @mgpuMemcpy(ptr %16, ptr %1, i64 ptrtoint (ptr getelementptr (double, ptr null, i32 4) to i64), ptr %15)
call void @mgpuMemcpy(ptr %17, ptr %8, i64 ptrtoint (ptr getelementptr (double, ptr null, i32 4) to i64), ptr %15)
Then the launch itself, the part gpu-to-llvm did not produce:
%67 = call ptr @mgpuModuleLoadJIT(ptr @add_tensors_kernel_bin_cst, i32 2)
%68 = call ptr @mgpuModuleGetFunction(ptr %67, ptr @add_tensors_kernel_add_tensors_kernel_kernel_name)
call void @mgpuLaunchKernel(ptr %68, i64 2, i64 1, i64 1, i64 2, i64 1, i64 1, i32 0, ptr %15, ptr %20, ptr null, i64 23)
call void @mgpuModuleUnload(ptr %67)
call void @mgpuStreamSynchronize(ptr %15)
call void @mgpuStreamDestroy(ptr %15)
Reading it: load the embedded PTX as a module (optimization level 2); look up the kernel by name; launch with grid x, y, z = 2, 1, 1, block x, y, z = 2, 1, 1, shared memory 0, on stream %15, with %20 as an array of 23 pointers, one per kernel parameter, a null extra, and the parameter count 23; unload the module; synchronize and destroy the stream. The geometry matches Chapter 10's kernel exactly. Finally, copy the result back and release everything:
%77 = call ptr @mgpuStreamCreate()
call void @mgpuMemcpy(ptr %69, ptr %18, i64 ptrtoint (ptr getelementptr (double, ptr null, i32 4) to i64), ptr %77)
call void @mgpuMemFree(ptr %16, ptr %77)
call void @mgpuMemFree(ptr %17, ptr %77)
call void @mgpuMemFree(ptr %18, ptr %77)
call void @mgpuStreamSynchronize(ptr %77)
call void @mgpuStreamDestroy(ptr %77)
The conclusion that the launch is lowered here rather than by gpu-to-llvm comes from this observation (the launch op is present in pipe_out.mlir, the runtime calls are present in host.ll), not from documentation.
A stub device, and what it can and cannot prove¶
host.ll is real LLVM IR that calls functions this machine does not have. A real build would link NVIDIA's runtime wrapper library. With none available, a stub implements the same entry points on the CPU: "device memory" is malloc, copies are memcpy, and mgpuLaunchKernel runs a small C loop that reads the parameter-pointer array and does what the PTX does. Its source, in full (this is the version after Chapter 12 extended it to also accept a 5-parameter launch; the 23-parameter behavior described here is unchanged):
// A CPU stand-in for MLIR's GPU runtime wrappers (the mgpu* entry points the lowered host code calls).
// There is no GPU here, so "device memory" is malloc and the "kernel" is a C function that reads the
// SAME launch parameters the PTX reads (params 0, 1, 3, 10, 17) and does the same arithmetic.
// This checks the HOST side (allocation, copies, launch geometry, 23-parameter packing). It does NOT
// execute the PTX, so it says nothing about whether the PTX itself is correct.
#include <stdint.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
void *mgpuStreamCreate(void) { fprintf(stderr, "[stub] stream create\n"); return malloc(1); }
void mgpuStreamDestroy(void *s) { free(s); }
void mgpuStreamSynchronize(void *s) { (void)s; }
void *mgpuMemAlloc(uint64_t n, void *s, int shared) { (void)s; (void)shared; fprintf(stderr, "[stub] alloc %llu bytes\n", (unsigned long long)n); return malloc(n); }
void mgpuMemFree(void *p, void *s) { (void)s; free(p); }
void mgpuMemcpy(void *dst, void *src, uint64_t n, void *s) { (void)s; fprintf(stderr, "[stub] memcpy %llu bytes\n", (unsigned long long)n); memcpy(dst, src, n); }
static const char *kernel_name_loaded;
void *mgpuModuleLoadJIT(void *ptx, int optLevel) {
const char *p = (const char *)ptx;
fprintf(stderr, "[stub] module load: optLevel=%d, PTX has .entry add_tensors_kernel: %s\n", optLevel,
strstr(p, ".entry add_tensors_kernel") ? "yes" : "NO");
return (void *)p;
}
void *mgpuModuleGetFunction(void *mod, const char *name) { (void)mod; kernel_name_loaded = name; return (void *)name; }
void mgpuModuleUnload(void *mod) { (void)mod; }
void mgpuLaunchKernel(void *fn, intptr_t gx, intptr_t gy, intptr_t gz, intptr_t bx, intptr_t by, intptr_t bz,
int32_t smem, void *stream, void **params, void **extra, size_t nparams) {
(void)gy; (void)gz; (void)by; (void)bz; (void)smem; (void)stream; (void)extra;
fprintf(stderr, "[stub] launch %s: grid=(%ld,1,1) block=(%ld,1,1) nparams=%zu\n", (const char *)fn, (long)gx, (long)bx, nparams);
// Two parameter layouts, depending on the kernel calling convention the host was compiled with:
// 23 = memref descriptors expanded (default); 5 = bare pointers (kernel-bare-ptr-calling-convention).
if (nparams != 23 && nparams != 5) { fprintf(stderr, "[stub] unexpected parameter count\n"); exit(2); }
int64_t step = *(int64_t *)params[0], lb = *(int64_t *)params[1];
double *a, *b, *c;
if (nparams == 23) { a = *(double **)params[3]; b = *(double **)params[10]; c = *(double **)params[17]; }
else { a = *(double **)params[2]; b = *(double **)params[3]; c = *(double **)params[4]; }
for (intptr_t blk = 0; blk < gx; blk++)
for (intptr_t t = 0; t < bx; t++) {
int64_t i = (blk * step + lb) * 2 + t;
c[i] = a[i] + b[i];
}
}
For the 23-parameter layout it reads parameters 0 and 1 (loop step and lower bound) and 3, 10 and 17 (the three data pointers): the exact five parameters Chapter 10 found the PTX body actually loads. It computes c[(block * step + lb) * 2 + thread] = a[i] + b[i], and it checks that the module handed to mgpuModuleLoadJIT really contains .entry add_tensors_kernel. It is linked with Chapter 8's unmodified harness:
#include <stdint.h>
#include <stdio.h>
typedef struct {
double *allocated;
double *aligned;
int64_t offset;
int64_t sizes[2];
int64_t strides[2];
} Memref2D;
extern Memref2D add_tensors(double *allocated0, double *aligned0, int64_t offset0,
int64_t size0_0, int64_t size0_1,
int64_t stride0_0, int64_t stride0_1,
double *allocated1, double *aligned1, int64_t offset1,
int64_t size1_0, int64_t size1_1,
int64_t stride1_0, int64_t stride1_1);
int main(void) {
double a[4] = {1.0, 2.0, 3.0, 4.0};
double b[4] = {5.0, 6.0, 7.0, 8.0};
Memref2D result = add_tensors(a, a, 0, 2, 2, 2, 1,
b, b, 0, 2, 2, 2, 1);
printf("add_tensors([[1,2],[3,4]], [[5,6],[7,8]]) =\n");
for (int64_t i = 0; i < result.sizes[0]; i++) {
for (int64_t j = 0; j < result.sizes[1]; j++) {
printf("%g ", result.aligned[i * result.strides[0] + j * result.strides[1]]);
}
printf("\n");
}
return 0;
}
clang-18 -c host.ll -o host.o
clang-18 harness.c host.o stub_gpu_runtime.c -o add_tensors_stub_demo
./add_tensors_stub_demo
Real output (the stub's trace on stderr first, then the program's output):
[stub] stream create
[stub] alloc 32 bytes
[stub] alloc 32 bytes
[stub] alloc 32 bytes
[stub] memcpy 32 bytes
[stub] memcpy 32 bytes
[stub] module load: optLevel=2, PTX has .entry add_tensors_kernel: yes
[stub] launch add_tensors_kernel: grid=(2,1,1) block=(2,1,1) nparams=23
[stub] stream create
[stub] memcpy 32 bytes
add_tensors([[1,2],[3,4]], [[5,6],[7,8]]) =
6 8
10 12
[[6, 8], [10, 12]] is the same answer Chapters 5 through 9 produced by other routes. The trace shows the sequence the host code is supposed to perform: three 32-byte device buffers, two copies in, a launch, one copy out.
A negative control, because a check that cannot fail proves nothing. negative_control.sh rebuilds the same program with the stub's kernel arithmetic line removed:
#!/bin/sh
# Negative control: the same host code and harness, linked with a stub whose kernel arithmetic is REMOVED.
# If the real stub's answer did not depend on the launch parameters being read correctly, this would still print 6 8 / 10 12.
# Needs ./host.o from pipeline.sh. Output: negative_control_out.txt
# NOTE: the output of this script is SUPPOSED to show a wrong answer (or no answer). That is the control: it proves the real stub's correct answer is not an accident.
cd "$(dirname "$0")"
sed 's/c\[i\] = a\[i\] + b\[i\];/(void)a;(void)b;(void)c;(void)i;/' stub_gpu_runtime.c > /tmp/stub_noop_ch11.c
clang-18 harness.c host.o /tmp/stub_noop_ch11.c -o /tmp/noop_demo_ch11 2>/dev/null && /tmp/noop_demo_ch11 2>/dev/null
rm -f /tmp/stub_noop_ch11.c /tmp/noop_demo_ch11
It printed zeros, the contents of a fresh malloc (those bytes are not guaranteed to be zero; they happened to be). So the answer above does depend on the launch parameters being packed and read correctly.
What this establishes: the host-side glue is correct end to end: device allocation, host-to-device and device-to-host copies, launch geometry, the 23-parameter packing, and the descriptor handling for the returned memref, all through Chapter 8's own calling convention.
What it does not establish: that the PTX is correct, or that it assembles, or that it would run on any real GPU. The stub's kernel is hand-written C that mirrors the PTX, not an execution of it. Anything wrong in the PTX itself, or any difference between how a real driver packs mgpuLaunchKernel's parameter array and how this stub reads it, is invisible to this check. A real run needs a GPU, the driver, the real runtime wrapper library, and ptxas; none exist in this sandbox.
Reproducing this chapter¶
pipeline.sh runs every step above and was verified by running it in an empty directory and byte-comparing pipe_out.mlir and host.ll with the files in code/. Tools: mlir-opt-18, mlir-translate-18, clang-18; no packages beyond Chapter 10's. run_out.txt is the unedited output, and negative_control.sh (run after pipeline.sh, which produces host.o) regenerates the control. The three symptom files are produced by Chapter 12's symptoms.sh.
What this chapter does not cite¶
As in Chapter 10, no llvm-project documentation was read for this chapter. Every claim about where a lowering step happens comes from running the pass and reading its output. In particular, the explanations for the three dead ends were, when this chapter was written, the most likely readings of what was observed, not verified root causes; Chapter 12 verified them from source.
Chapter summary¶
This chapter built the host side of Chapter 10's kernel as far as a GPU-less sandbox allows. A hand-written host function wraps the kernel in gpu.alloc, gpu.memcpy and gpu.dealloc; gpu-async-region rewrites those into the async form MLIR 18's lowering requires; MLIR's own --gpu-lower-to-nvvm-pipeline produces the kernel binary and the host runtime calls; and mlir-translate-18 turns the remaining gpu.launch_func into mgpuModuleLoadJIT, mgpuModuleGetFunction, mgpuLaunchKernel and mgpuModuleUnload calls with a 2x1x1 grid, a 2x1x1 block and 23 kernel parameters. Linked with Chapter 8's unmodified harness and a CPU stub runtime, the result prints [[6, 8], [10, 12]], and a negative control with the stub's arithmetic removed does not. Three real dead ends are recorded with their outputs.
Deliberately out of scope, stated explicitly: the host function is hand-written (no pass generates it), the stub's "kernel" is C code mirroring the PTX rather than an execution of it, so nothing here shows the PTX is correct or even assembles, and the three dead ends were explained from MLIR's source only afterwards, in Chapter 12.
Self-check questions¶
1. gpu-to-llvm ran without error on the original host function and left gpu.alloc and gpu.memcpy in place. What fixed it, and why was the failure silent?
Worked answer: gpu-async-region fixed it, by rewriting the ops into their async form with !gpu.async.tokens. The failure was silent because the lowering patterns only match async operations and a pattern that does not match simply leaves the op alone; the first visible error came later, when a conversion pass hit an unrealized_conversion_cast it could not legalize. The evidence in this chapter is the before/after listing: nine gpu.alloc/gpu.memcpy/gpu.dealloc ops survive without the pass and none with it. Chapter 12 later read the pattern's source and found the explicit "Can convert only async version" test.
2. The final mgpuLaunchKernel call ends in i64 23. What is that number, and what are the arguments before it?
Worked answer: it is the number of kernel parameters, 23, the same expanded count as the kernel's own signature (2 index arguments plus 3 memrefs of seven scalars each). The arguments in order are the kernel function handle; the grid x, y, z (2, 1, 1); the block x, y, z (2, 1, 1); the shared-memory size (0); the stream; a pointer to an array of 23 pointers, one per kernel parameter; and a null extra. The grid and block match Chapter 10's kernel exactly.
3. Where is gpu.launch_func actually turned into runtime calls in this toolchain, and how did the chapter find out?
Worked answer: at LLVM IR translation time, in mlir-translate-18, not in gpu-to-llvm. The evidence is a before/after observation: after the packaged pipeline, the MLIR still contains a gpu.launch_func (now with a stream operand and the 23 expanded operands) and no launch runtime call, while the translated host.ll contains mgpuModuleLoadJIT, mgpuLaunchKernel and mgpuModuleUnload. The chapter reached this from the outputs, not from documentation.
4. What does linking against the stub runtime establish, and what does it leave open?
Worked answer: it establishes that the host glue is right: three device allocations of 32 bytes, two copies in, a launch with grid (2,1,1) and block (2,1,1), the parameter array packed in the layout the stub reads, one copy out, and the memref descriptor returned correctly through Chapter 8's calling convention. It leaves open everything about the PTX: the stub's kernel is hand-written C mirroring the PTX's five loaded parameters and arithmetic, so a defect in the PTX, or a difference between how a real driver packs the parameter array and how the stub reads it, would be invisible.
5. Why was the negative control worth running, and what exactly did it show?
Worked answer: a check that cannot fail proves nothing. The same program was rebuilt with the stub's one arithmetic line removed. It printed zeros (the contents of a fresh malloc, which happened to be zero and are not guaranteed to be), not 6 8 / 10 12. So the correct answer in the real run does depend on the launch being packed and executed correctly, rather than being produced by some other path.