Module Rig_cuda

A device of this library is one GPU, opened on its primary context with two CUDA streams of its own, its queues "COMPUTE:0" and "COPY:0". "COMPUTE:0" runs launches, copies and fills, and "COPY:0" runs copies and fills. The device's timeline word is the word its opener's make is given, in rig's array of words, and the streams write it themselves. The first device that opens registers that array for every context (cuMemHostRegister, portable and mapped) in its GPU's primary context, which the process keeps, so the registration lasts for the process. A GPU has one device at a time.

A kernel reaches the GPU in a cubin, a fatbin or PTX text. Rig.Image.load loads it as a CUDA module and places every function's code on the GPU. Each Rig.submit enqueues its parts on the streams and returns without waiting for them, a launch with cuLaunchKernel and a copy with cuMemcpyAsync. CUDA events order a value's work after the previous value's work, on both streams. To commit a value, the device enqueues a write of the value into the word (cuStreamWriteValue64), which runs once the work before it has completed. It commits when a wait asks for a value, and on its own (Work). Work that waits for another device's point waits on the GPU, as a stream wait on that device's word, where the facts allow it (Facts).

Given a device g that Rig.open_ (Rig_cuda.gpu 0) opened, a buffer x on it that holds n floats, where n is a multiple of 256, and a cubin bin compiled from

extern "C" __global__ void twice(float *x) {
  x[blockIdx.x * blockDim.x + threadIdx.x] *= 2;
}

this doubles x in place:

let image = Rig.Image.load g bin |> Result.get_ok in
let twice = Rig.Image.resolve image "twice" |> Option.get in
let twice =
  Rig.Submission.(launch ~refs:[| write 0 ~at:0 |] twice ~params:8)
in
let s = Rig.Submission.make g [| twice |] in
let run = Rig.Submission.take s in
let b = Rig.Submission.block s 0 in
Rig.Submission.Run.groups run b (n / 256) 1 1;
Rig.Submission.Run.threads run b 256 1 1;
Rig.Point.wait (Rig.submit run ~buffers:[| x |])

The kernel's parameters are its one pointer, and the ref puts the address of x's first byte in those 8 bytes.

A device lives until Rig.close or a loss, and a fault of its work ends the GPU's use for the rest of the process (Failure and stopping). Other CUDA libraries in the process share the primary context. Every call that needs a current context makes the device's current on the calling thread, and restores the thread's previous context before it returns. Any domain may call any function of this module, at the same time as other domains. A function that waits for CUDA lets other domains run while it waits.

count, an opener's make and the graph function of capability, called from a callback inside a lock (Rig's domains section), raise Invalid_argument, having changed nothing. rig calls the functions of The driver, but locate, peer and signaled, holding no lock (Drivers); one called from such a callback raises Invalid_argument, having changed nothing, where it would take a lock that the interrupted call holds.

Rig_edge states how rig drives a device through its driver, and The C functions how it hands the driver work.

Opening

A device needs:

• the CUDA library of CUDA 12.4 or later (driver 550);
• a GPU with 64-bit stream memory operations, with which its streams write and wait on timeline words;
• unified addressing, under which the GPU's memory and the host's share one address space.

The CUDA library (libcuda) is loaded at the first call of count, buses or an opener's make, by its standard name from the platform's library search path: libcuda.so.1 (or libcuda.so) on Linux, nvcuda.dll on Windows, libcuda.dylib on macOS. A library of that name earlier on the path is loaded in its place, for the whole process. A load that fails is tried again at the next call. Building and linking this library needs no CUDA installation.

The driver

rig calls these functions. On CUDA, alloc, free, locate, map_host, peer, map_peer and dma behave as in Memory; image, entry and unload as in Code; sleep and stop as in Failure and stopping; and the C functions as in Work.

Facts

arch is "sm_89" for compute capability 8.9. "COMPUTE:0" runs Fill, Copy and Launch, and "COPY:0" runs Fill and Copy. The streams write the word, and the host does not address Device memory. The word holds an unsigned 64-bit integer in the host's byte order, and the device's values start after the value it holds when make returns. alignment is 256 since the CUDA C++ Programming Guide (Device Memory Accesses) states that every address an allocation routine of the driver or runtime API returns "is always aligned to at least 256 bytes". On a GPU that registers no host memory, or when the word make was given lies outside rig's array, the word is eight bytes of the device's own page-locked memory, which starts at that value and has no host address in its location: rig reads it through signaled. That word, one 8-byte cuMemHostAlloc allocation, goes back when rig frees it (Rig_edge.facts). Its free, as every free of host memory, may wait for the GPU's work.

The streams wait on any aligned 64-bit word the device's work addresses, whoever writes it. A submission's waits go to its stream in batches, as many as it carries. hang_ms is None where a display's watchdog bounds the GPU's kernels (CU_DEVICE_ATTRIBUTE_KERNEL_EXEC_TIMEOUT), and Some 30_000, 30 s, where nothing does (Failure and stopping).

Work

The device's struct rig_driver holds rig_cuda_room, rig_cuda_submit, rig_cuda_commit and rig_cuda_settle, which rig_cuda.h declares. A part on queue 0 runs on "COMPUTE:0", one on queue 1 on "COPY:0".

The device commits a value by enqueueing a write of it into the word (cuStreamWriteValue64) on the stream the last value ended on. The write follows a system-wide memory fence (CU_STREAM_WRITE_VALUE_DEFAULT), so the work's writes are visible once the word shows the value. A commit (The C functions) writes nothing if a commit already wrote that value or a later one. rig_cuda_submit also commits, and returns RIG_COMMITTED, for a value 64 values past the last commit, so the device commits each value at most 64 values late.

A fill f is called as f stream arg v, with the device's context current on the calling thread (Rig_cuda_abi). A copy part moves its rows between the handles of two regions of the device, at their offsets, in one 2D copy, or a copy per row where a pitch exceeds the GPU's CU_DEVICE_ATTRIBUTE_MAX_PITCH or CUDA refuses the pitches, and a copy part that fills sets its bytes with one memset. A part must not write a timeline word, the device's or another's. Only the devices' streams write their words, and a word that work wrote could move backwards or show work that has not completed.

A wait holds the work back until the aligned 64-bit word at at holds at least value, compared circularly, in the sense that x is at least w if x - w, as a signed 64-bit integer, is not negative.

rig_cuda_room never returns RIG_LATER, since a stream takes any amount of work. rig_cuda_submit and rig_cuda_commit instead block while a stream is full, until the device's earlier work frees it. A timed submit that measures blocks for its round trips, tens of microseconds, and one timed 1,024 values after an unconverted value waits for that value. While a profile is taken, rig_cuda_submit records a timing event on each stream a value's work enters, after the value's waits, and one after its work on the stream it ends on, into one of 1,024 stamps, whose events it makes at the stamp's first use. It waits for the value 1,024 values before it, whose events it reuses, if that value was timed. rig_cuda_settle, or a later timed rig_cuda_submit, converts the earliest start and the end to the host's clock, in value order, as The C functions requires of settle. A failed conversion loses the device.

The device measures its clock against the host's when it opens, and again in a timed rig_cuda_submit that finds the last measure 8 s old. A measure records a reference event on a stream of the device's own between two readings of the host's clock, asking CUDA until the event has completed, eight times, and keeps the closest; while that is over 10 µs it tries again, up to sixteen times. An event's time is its elapsed time since the reference, a float of milliseconds, plus the host's time at the reference. A time is thus within half the closest round trip of the host's clock, 5 µs once a round trip within 10 µs is seen, plus the GPU timer's rate error since the measure, plus half a float step at its distance from the reference, half a microsecond within 8 s.

Launches

A launch runs its function with cuLaunchKernel over the groups, threads and dynamic shared memory of its block, and passes its parameters as one buffer (CU_LAUNCH_PARAM_BUFFER_POINTER). A launch's bytes past those its function reads (Rig.Image.params) are not passed.

rig_cuda_room refuses a launch if:

• it is on queue 1;
• its block has a size of 0 along an axis;
• it has more groups or threads along an axis than the GPU allows (CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_X to _Z, CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_X to _Z);
• it has more threads per group, or more dynamic shared memory, than its function's limits.

Memory

A region is an allocation of the device, host memory it maps, or another device's memory it maps. alloc returns:

• for Device, GPU memory, which the host does not address;
• for Pinned, page-locked host memory that the host and every CUDA device address (cuMemHostAlloc, portable and mapped). It is not write-combined, so the host reads it through its caches;
• for Mapped, None, since CUDA maps no GPU memory for the host.

An allocation is None if CUDA refuses it, such as for lack of memory. alloc and map_host raise Invalid_argument for fewer than one byte.

A region's locate address is the address of its first byte in the process's unified address space, which every CUDA device's work uses. Its handle is the address by which CUDA names its allocation: the device address for GPU memory, the host address for host memory. Its host is that host address for host memory, and None for GPU memory.

free frees an allocation, and ends a region that map_peer or map_host gave.

peer g g' is true if g and g' are devices of one GPU, and otherwise iff CUDA gave g's GPU access to the GPU memory of g''s when the later of the two GPUs first opened. That open enables the access both ways where CUDA gives it, for the rest of the process. map_peer maps host memory always, and GPU memory where peer is true, and is Never otherwise. Freeing a map_peer region ends only that region. Peer access between two GPUs is untested.

map_host g a n page-locks the n bytes at a for every CUDA device and maps them, at the address CUDA gives for the memory (cuMemHostGetDevicePointer). Page-locking belongs to the process, and the region's pages stay locked until every region map_host gave over them, on any device, is freed. The last free unlocks them. One made while a map_host or an open of a CUDA device runs, on another thread or in the call its callback interrupted, may instead leave them locked until no such call runs. The result depends on what is already locked:

• A range inside one that map_host locked is mapped and counts against it.
• A range inside one allocation that CUDA locked for another owner, such as an allocation of alloc or of another library, is mapped as it is, uncounted. That memory must stay locked until the region is freed.
• Any other range is locked. Where CUDA refuses, it is No_room if CUDA had no memory to lock or map the range, and Never otherwise, as for read-only memory, a shared mapping of a file and a range that overlaps locked memory without lying inside one of its ranges.

It is also Never where g does not map host memory (maps_host).

On Linux, CUDA locks host pages for writing, which Linux refuses for a shared mapping of a file. A CUDA device therefore cannot borrow a region that an io device holds for writing through a shared mapping of a file (Rig_edge.Io.pages), and copies to and from that region go through staging memory.

dma g r ~fn exports the whole allocation r lies in as a dma-buf (cuMemGetHandleForAddressRange, with no flags), at r's distance from the allocation's start. CUDA exports Device memory only where its kernel driver can, as CU_DEVICE_ATTRIBUTE_DMA_BUF_SUPPORTED states: on Linux, with NVIDIA's open kernel modules. It exports Pinned memory only where CU_DEVICE_ATTRIBUTE_HOST_ALLOC_DMA_BUF_SUPPORTED holds, an attribute from CUDA 13 on. The kernel maps the memory for fn where fn registers the dma-buf, and refuses there a topology it does not support. With no flags it maps a coherent GPU's memory as the host's pages, over the GPU's link to the host, so a NIC behind the GPU's own PCIe switch reaches it through the host's memory system, at that link's bandwidth, where a BAR1 mapping would cross only the switch. fn reads what the device's work wrote once that work's value is reached, since the word's write (cuStreamWriteValue64_v2) fences the stream's writes for every reader first. It is Error, naming the cause: for an fn that is no bus address or no function under /sys/bus/pci/devices; where CUDA exports no dma-buf, naming the attribute that decides it; and for an allocation whose start or size is off the host's page, as CUDA's export requires and its allocations, aligned to 256 bytes, need not be.

Code

CUDA may wait for all of the GPU's work before free frees an allocation, and before image or unload loads or unloads a module. Launches in the GPU's context, from any thread, wait for that wait. The calling thread blocks while CUDA waits: an open sets its primary context's scheduling to CU_CTX_SCHED_BLOCKING_SYNC where the context's flags hold CU_CTX_SCHED_AUTO, with which CUDA spins unless the process has more contexts than processors. The setting holds for every wait CUDA makes in that context, those of other libraries in the process included, such as cuStreamSynchronize: each returns after its thread wakes, later than a spin would. A scheduling another library set before the open stays, but one that set CU_CTX_SCHED_AUTO cannot be told from the default and is changed.

image g bin is Ok (Loaded m) with m the CUDA module of bin: a cubin, a fatbin, or PTX text, which CUDA compiles for the GPU. It places every function's code before it returns, whatever CUDA_MODULE_LOADING asks, so neither entry nor a launch loads code later. Each function may take as much dynamic shared memory as a block of the GPU can have, less its static shared memory. The result is Error with CUDA's error if CUDA refuses bin, for instance a cubin for another GPU, or lacks the memory for its code.

entry m f is the entry of kernel f of m. Its code is the kernel's CUfunction, which compiled code passes to cuLaunchKernel. Its launch is C memory that holds the function and its limits: the most threads a group of it has (CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK), and the most dynamic shared memory a group of it takes. Both are valid until m is unloaded, and a second entry for f returns the same. Its params end at the kernel's last parameter (cuFuncGetParamInfo), and its threads are the most threads above. A launch hands CUDA those params bytes of its parameters, since CUDA refuses a buffer of another size. entry raises Invalid_argument for a kernel of more parameter bytes than CUDA launches on the GPU: 32,764 on compute capability 7.0 or later, and 4,096 below.

After unload, no new graph may record the image's functions. Its CUDA module stays loaded until every graph that recorded one of its functions is destroyed (Graphs).

Failure and stopping

rig loses a device when CUDA reports a fault of its work, such as an illegal address or a kernel that a display's watchdog ended, at a later call. A fault leaves the process in an inconsistent state: every later CUDA call returns the fault's error, and to use CUDA again the process must be terminated and relaunched. So the GPU cannot be opened again, and every other CUDA device of the process is lost at its first call that CUDA answers with the fault's error (Loss).

On a GPU whose kernels no display's watchdog bounds, rig also loses the device once its work has kept the word still for 30 s, even work that would have ended later.

Before a fault, alloc, map_host, image and entry return CUDA's refusals as their result, None, Never or No_room, Error or Absent, and unload raises Fault if CUDA refuses to unload the module. After a fault, the calls of alloc, map_host, image, entry, unload and sleep that call CUDA raise Fault with its error, and the others answer as before the fault, such as an entry for a function already found or a sleep once the word has moved. rig_cuda_submit and rig_cuda_commit then return RIG_FAILED once they call CUDA, and the graph function of capability returns Error. free and stop ignore CUDA's errors and raise nothing.

A rig_cuda_submit or rig_cuda_commit that fails returns RIG_FAILED with what it was doing and the error of the first CUDA call that failed, including one a fill made, as in "running a fill: CUDA_ERROR_ILLEGAL_ADDRESS: an illegal memory access was encountered". The parts enqueued before the failure may run, and the others never do.

A failed device stays failed, since CUDA may keep the context's error for the process. Every later rig_cuda_submit enqueues none of its parts, and every later rig_cuda_submit and rig_cuda_commit returns the same failure. The word still reaches every value, so that work waiting on it on other devices runs on. A failed rig_cuda_submit, and every later one, writes its value after its waits, every earlier value and the work it enqueued. It writes nothing once the context has failed, whose work no longer runs, or if CUDA refuses a call that orders the write.

sleep asks CUDA each millisecond whether g's streams have met an error (cuStreamQuery), and raises Fault with CUDA's error once one has.

stop ignores fault. If g's work has ended (every value's work has completed, its streams are idle, or a fault has ended the context's work), it writes the last value. Otherwise it commits that value, and the word reaches it once the work ends, unless the work waits on a word of another device that never reaches its value, and it keeps an event behind the work on each stream. Either way it destroys g's streams and its other events, which CUDA frees once their work ends. The GPU can be opened again once that work ends, by completing or by a fault. CUDA signals no fault, so after one the GPU's next open raises the word to the last value, and until then the device does not count as stopped; that open's make answers the fault's error. The stop frees g's C state but for what reading the word needs, the 64-byte edge, which stays for the process. Once the word shows the last value and the stop has returned, no work of g writes it: a stream's write of the last value is its last command, and the stop and the next open raise the word only once both streams are idle or one has met an error, after which CUDA runs none of the context's queued work, as the 615 driver does (CUDA's reference states only that such an error is returned to every later call). Compiled code may still call the symbol and graph functions of capability after stop, and rig's free and unload may come after the GPU has opened again, since the device's modules lie in the GPU's primary context, which the process keeps.

References

• CUDA Driver API: Primary Context Management, Stream Management, Event Management, Stream Memory Operations (cuStreamWaitValue64, cuStreamWriteValue64, cuStreamBatchMemOp, CUstreamWaitValue_flags), Memory Management (cuMemHostAlloc, cuMemHostRegister), Unified Addressing, Peer Context Memory Access, Module Management, Graph Management (cuGraphAddKernelNode, cuGraphInstantiateWithFlags), Error Handling (CUDA_ERROR_ILLEGAL_ADDRESS), Version Management (cuDriverGetVersion).
• CUDA 12.1 Supports Large Kernel Parameters: a kernel's parameters take up to 32,764 bytes on Volta and later under a CUDA 12.1 driver or newer, and 4,096 below Volta.