Saba Shahrukh September 13, 2026 0 If you want to keep track of your post-reading status, please register on the site.

Hardware Compilers, CUDA Runtimes, and OCI Hooks

Deploying modern artificial intelligence models at scale requires a clear understanding of the full software and hardware execution pipeline. While deep learning frameworks present high-level mathematical abstractions like tensors and computational graphs, the physical execution of these operations relies on a complex chain of hardware compilers, low-level execution runtimes, host kernel drivers, and container virtualization hooks.

To optimize AI workloads for maximum throughput and minimum latency, performance engineers must understand how high-level framework code is translated into low-level machine instructions, how execution queues interact with GPU hardware, and how containerized environments safely expose bare-metal accelerators without added virtualization overhead.

1. Hardware Compilers & Graph Lowering: From High-Level IR to SASS

High-level deep learning frameworks (such as PyTorch or JAX) express compute jobs as Directed Acyclic Graphs (DAGs) of operations (e.g., MatMul, Add, Relu). Executing these operations as individual, discrete kernels introduces significant overhead due to memory bandwidth constraints.

+--------------------------------------------------------------------------+
| HIGH-LEVEL GRAPH (Framework Layer: PyTorch / JAX / TorchDynamo)          |
|  - Abstract Directed Acyclic Graph (DAG)                                 |
|  - High-level operators: Conv2D -> Add -> ReLU -> GEMM                   |
+--------------------------------------------------------------------------+
                                     |
                                     v Lowering / Optimization
+--------------------------------------------------------------------------+
| GRAPH COMPILER & DIALECT IR (MLIR / TorchInductor / XLA / Triton)       |
|  - Dead code elimination & layout transformations                        |
|  - Operator Fusion (combining elementwise ops into single memory pass)   |
|  - Tiling & loop unrolling strategies                                    |
+--------------------------------------------------------------------------+
                                     |
                                     v Code Generation (LLVM Backend)
+--------------------------------------------------------------------------+
| VIRTUAL ISA / INTERMEDIATE ASSEMBLY (NVVM IR -> PTX)                     |
|  - Parallel Thread Execution (PTX) virtual instruction set               |
|  - Architecture-agnostic register allocation & thread block mapping      |
+--------------------------------------------------------------------------+
                                     |
                                     v Hardware Assembler (ptxas)
+--------------------------------------------------------------------------+
| HARDWARE NATIVE ISA (SASS: Streaming Assembler Code)                     |
|  - Target-specific machine code (e.g., SM_80 / SM_90 execution units)    |
|  - Exact register assignment, instruction scheduling, warp allocation    |
+--------------------------------------------------------------------------+

Intermediate Representations (IR)

To convert a framework compute graph into machine code, hardware compiler toolchains (such as MLIR, TorchInductor, XLA, and Triton) lower the graph through several stages of Intermediate Representation:

  1. High-Level Graph IR: Represents mathematical operations independently of the targeted hardware. At this layer, compilers perform layout transformations (e.g., NCHW to NHWC) and algebraic simplifications.
  2. Dialect IR (e.g., MLIR / LLVM-NVVM): Breaks down abstract operations into structured loops, memory allocations, and parallel execution blocks.
  3. Virtual ISA (PTX): Parallel Thread Execution (PTX) is a low-level virtual machine and instruction set architecture (ISA) for NVIDIA GPUs. PTX defines operations for a generic parallel processor, abstracting away specific GPU microarchitectures.
  4. Hardware Native ISA (SASS): The ptxas assembler compiles PTX into SASS (Streaming Assembler), which represents the binary machine instructions executed directly by the GPU’s Streaming Multiprocessors (SMs). SASS is specific to individual GPU architectures (e.g., Ampere SM_80, Hopper SM_90).

Memory Bandwidth Reduction via Operator Fusion

The primary objective of hardware compilers is to maximize Arithmetic Intensity—the ratio of floating-point operations (FLOPs) to memory transactions (Bytes transferred).

Consider an unfused sequence of an elementwise matrix addition followed by a ReLU activation on an $M \times N$ matrix:

$$\text{Memory Traffic}_{\text{unfused}} = \underbrace{2 \cdot M \cdot N \cdot b}_{\text{Read } A, B} + \underbrace{M \cdot N \cdot b}_{\text{Write } A+B} + \underbrace{M \cdot N \cdot b}_{\text{Read } A+B} + \underbrace{M \cdot N \cdot b}_{\text{Write } \text{ReLU}(A+B)} = 5 \cdot M \cdot N \cdot b \text{ bytes}$$

(where $b$ is the byte size per element)

A graph compiler fuses these two operations into a single GPU kernel loop. The intermediate sum $A+B$ is stored inside fast, on-chip warp registers rather than round-tripping through High Bandwidth Memory (HBM):

$$\text{Memory Traffic}_{\text{fused}} = \underbrace{2 \cdot M \cdot N \cdot b}_{\text{Read } A, B} + \underbrace{M \cdot N \cdot b}_{\text{Write } \text{ReLU}(A+B)} = 3 \cdot M \cdot N \cdot b \text{ bytes}$$

By fusing the operations, the compiler cuts HBM traffic by 40%, preventing the GPU compute engines from stalling while waiting for memory transfers.

2. The CUDA Execution Boundary: Driver API vs. Runtime API

Once compilers generate device-specific kernels, application code uses the CUDA user-space stack to schedule and execute those kernels on the accelerator.

+-------------------------------------------------------------------------+
| APPLICATION LAYER (Python / C++ Application Code)                      |
+-------------------------------------------------------------------------+
                                     |
                                     v
+-------------------------------------------------------------------------+
| CUDA RUNTIME API (libcudart.so)                                         |
|  - Implicit Context & Device Initialization                             |
|  - High-Level Abstractions: cudaMalloc(), cudaMemcpy(), cudaLaunchKernel()|
|  - Managed memory allocation & stream management                        |
+-------------------------------------------------------------------------+
                                     |
                                     v Wraps & Delegates
+-------------------------------------------------------------------------+
| CUDA DRIVER API (libcuda.so)                                            |
|  - Explicit Context Management (CUcontext, CUmodule, CUfunction)        |
|  - Direct Hardware Control: cuMemAlloc(), cuLaunchKernel()              |
|  - Directly interfaces with Linux Kernel via ioctl system calls         |
+-------------------------------------------------------------------------+
                                     |
                                     v ioctl() system calls
+-------------------------------------------------------------------------+
| HOST KERNEL MODULE (nvidia.ko)                                          |
|  - Manages GPU Virtual Memory page tables and DMA allocation            |
|  - Pushes command buffers to GPU Hardware Work Queues (Ring Buffers)    |
+-------------------------------------------------------------------------+

Runtime API (libcudart.so) vs. Driver API (libcuda.so)

The CUDA software layer provides two interfaces for managing hardware:

  • CUDA Runtime API (libcudart.so): A high-level C/C++ interface built on top of the Driver API. It abstracts device management by automatically initializing CUDA contexts, tracking active devices, and exposing simplified function calls like cudaMalloc() and cudaLaunchKernel().
  • CUDA Driver API (libcuda.so): A low-level, object-oriented C API supplied directly by the NVIDIA driver package. It requires explicit management of CUDA Contexts (CUcontext), JIT-compiled Modules (CUmodule), and Execution Functions (CUfunction). Framework runtimes (such as PyTorch and TensorRT) use the Driver API directly to retain fine-grained control over execution pipelines.

Command Queuing and Hardware Ring Buffers

When a application issues a kernel launch request (via cudaLaunchKernel or cuLaunchKernel), the host CPU does not wait for the GPU to execute the work:

  1. Work Submission: The CUDA driver constructs a hardware-formatted command buffer containing the kernel’s execution parameters (grid dimensions, block sizes, shared memory allocations, and function pointers).
  2. Ring Buffer Insertion: The driver writes these commands into host-mapped shared memory buffers known as Hardware Ring Buffers (or Work Queues).
  3. MMIO Doorbell Notification: The host writes to a specific Memory-Mapped I/O (MMIO) register on the GPU (a “doorbell switch”).
  4. Asynchronous Execution: The GPU’s hardware scheduler (GigaThread engine) reads the commands directly from the ring buffer via Direct Memory Access (DMA) and assigns thread blocks to available Streaming Multiprocessors (SMs) asynchronously.

3. OCI Runtime Hooks & Container Hardware Virtualization

Running GPU-accelerated workloads inside isolated container environments (such as Docker or Kubernetes) requires passing hardware access through the container boundary without introducing performance penalties.

+-------------------------------------------------------------------------+
| CONTAINER PROCESS (Isolated Namespaces: mnt, pid, net, ipc)             |
|                                                                         |
|  +-------------------------------------------------------------------+  |
|  | PyTorch Process / CUDA Runtime Application                       |  |
|  +-------------------------------------------------------------------+  |
|                                    |                                    |
|                                    v Calls libcuda.so                   |
|  +-------------------------------------------------------------------+  |
|  | Injected User-Space Driver Library (libcuda.so bind-mounted)      |  |
|  +-------------------------------------------------------------------+  |
+-------------------------------------------------------------------------+
                                     |
                                     v Accesses character devices
+-------------------------------------------------------------------------+
| OCI CONTAINER ISOLATION BOUNDARY                                        |
|                                                                         |
|  +-------------------------------------------------------------------+  |
|  | Injected Device Nodes: /dev/nvidia0, /dev/nvidiactl, /dev/nvidia-uvm|  |
|  +-------------------------------------------------------------------+  |
|  | Linux cgroups (devices subsystem): Whitelists major/minor numbers  |  |
|  +-------------------------------------------------------------------+  |
+-------------------------------------------------------------------------+
                                     |
                                     v ioctl system calls
+-------------------------------------------------------------------------+
| HOST KERNEL & PHYSICAL ACCELERATOR                                      |
|  - Host Kernel Module (nvidia.ko) handles interrupts & DMA              |
|  - Physical GPU Hardware Execution (PCIe Bus Communication)             |
+-------------------------------------------------------------------------+

OCI Hook Architecture

Containers use Linux namespaces to isolate processes, filesystem mounts, and network interfaces. However, standard Linux namespaces cannot manage GPU devices. Instead of fully virtualizing hardware, the NVIDIA Container Toolkit uses the Open Container Initiative (OCI) runtime specification hooks to expose physical host GPUs to isolated containers.

When launching a GPU-enabled container, the container runtime (e.g., containerd or runc) follows a three-step setup process:

  1. Prestart Interception: The container engine pauses startup after creating the container’s isolated filesystem namespace, but before executing the entrypoint application. It then executes the nvidia-container-runtime-hook.
  2. Cgroup Whitelisting: The hook updates the container’s Linux devices Control Group (cgroup). It grants read, write, and character device creation permissions (c *:* rwm) specifically for NVIDIA device major numbers (e.g., major character device 195).
  3. Node Injection and Bind-Mounting: The hook populates the container’s private /dev directory with physical character device nodes (/dev/nvidia0, /dev/nvidia-ctl, /dev/nvidia-uvm) using the mknod system call. It then bind-mounts the host’s user-space driver libraries (libcuda.so, libnvidia-ml.so) directly into the container’s dynamic library paths.

Because the containerized application uses injected device nodes to make direct ioctl calls to the host kernel module (nvidia.ko), hardware execution incurs zero virtualization overhead.

4. Unified Execution Pipeline Trace

Tracing an execution call across the entire hardware and software stack reveals how each system component processes a forward pass operation:

LayerSystem ComponentLow-Level SubsystemArtifact / Action State
1. ApplicationFramework APIPython Runtimetorch.compile() builds execution graph for processing.
2. Graph CompilerTorchInductor / TritonHigh-Level IR / MLIRFuses operations and emits optimized PTX assembly code.
3. Hardware Assemblerptxas / CUDA DriverLLVM BackendCompiles PTX assembly down to native architecture SASS machine instructions.
4. CUDA Runtimelibcudart.soUser-Space APIAllocates host/device buffers and schedules work onto CUDA streams.
5. CUDA Driverlibcuda.so (OCI Injected)User-Space Driver APIFormulates command buffers and issues ioctl calls to device files.
6. Isolation BoundaryOCI Hook / cgroupsLinux Device SubsystemValidates process access to /dev/nvidia0 via cgroup rules.
7. Kernel Drivernvidia.koLinux Kernel SpaceUpdates hardware ring buffers and triggers MMIO doorbell ring on PCIe.
8. Physical SiliconStreaming MultiprocessorHardware Execution UnitsFetches SASS instructions, loads data from HBM, and runs matrix compute engines.

5. End-to-End Inspection & Compilation Profiling Script

The following executable Python script inspects the lower levels of the CUDA execution stack. It uses ctypes to interface directly with the low-level CUDA Driver API (libcuda.so), inspects local Linux cgroup device permissions to verify container hardware passthrough, inline-compiles a custom CUDA C++ kernel into PTX assembly code, and benchmarks its execution against an unfused PyTorch implementation.

Python

import os
import sys
import time
import ctypes
import torch
from torch.utils.cpp_extension import load_inline

def inspect_cuda_driver_api():
    """Queries low-level driver status directly via libcuda.so system library."""
    print("=== 1. Low-Level CUDA Driver API Telemetry ===")
    
    # Load user-space driver library injected via OCI hook / system path
    try:
        cuda_driver = ctypes.CDLL("libcuda.so")
    except OSError:
        try:
            cuda_driver = ctypes.CDLL("libcuda.so.1")
        except OSError:
            print("[ERROR] Could not load user-space libcuda.so driver library!")
            return

    # Initialize CUDA Driver API (cuInit)
    status = cuda_driver.cuInit(0)
    if status != 0:
        print(f"[ERROR] cuInit failed with status code: {status}")
        return

    # Query Driver Version
    version = ctypes.c_int()
    cuda_driver.cuDriverGetVersion(ctypes.byref(version))
    driver_version = version.value
    major = driver_version // 1000
    minor = (driver_version % 1000) // 10

    print(f"CUDA Driver API Status   : Successfully Initialized (cuInit = 0)")
    print(f"Driver API Version       : {major}.{minor} (Raw Code: {driver_version})")

    # Query Active Device Count via Driver API
    count = ctypes.c_int()
    cuda_driver.cuDeviceGetCount(ctypes.byref(count))
    print(f"Driver-Reported Devices  : {count.value}\n")

def verify_oci_cgroup_isolation():
    """Inspects Linux cgroup device permissions to verify OCI container setup."""
    print("=== 2. OCI Container Hardware Passthrough Telemetry ===")
    
    # Check for presence of character device nodes
    device_nodes = ["/dev/nvidia0", "/dev/nvidiactl", "/dev/nvidia-uvm"]
    found_nodes = [node for node in device_nodes if os.path.exists(node)]
    
    print(f"Injected Device Nodes    : {found_nodes}")

    # Inspect Cgroup Controller Path (cgroup v1 or v2)
    cgroup_path = "/sys/fs/cgroup/devices/devices.list" # cgroup v1
    if not os.path.exists(cgroup_path):
        cgroup_path = "/sys/fs/cgroup/cgroup.controllers" # cgroup v2
        
    print(f"Detected Cgroup Path     : {cgroup_path}")
    print(f"Container Mode           : Isolated hardware passthrough active\n")

def compile_custom_cuda_kernel():
    """
    Compiles a custom C++/CUDA kernel inline, demonstrating PTX code generation
    and low-level runtime kernel execution vs unfused framework operations.
    """
    print("=== 3. Kernel Compilation & Execution Comparison ===")
    
    if not torch.cuda.is_available():
        print("[ERROR] Active CUDA accelerator required for compilation test.")
        return

    # Define low-level CUDA C++ source code for Fused Add + ReLU
    cuda_source = """
    #include <cuda_runtime.h>

    __global__ void fused_add_relu_kernel(const float* x, const float* y, float* out, int size) {
        int idx = blockIdx.x * blockDim.x + threadIdx.x;
        if (idx < size) {
            float sum = x[idx] + y[idx];
            out[idx] = sum > 0.0f ? sum : 0.0f; // Fused elementwise operation
        }
    }

    torch::Tensor fused_add_relu_cuda(torch::Tensor x, torch::Tensor y) {
        auto size = x.numel();
        auto out = torch::empty_like(x);

        const int block_size = 256;
        const int num_blocks = (size + block_size - 1) / block_size;

        fused_add_relu_kernel<<<num_blocks, block_size>>>(
            x.data_ptr<float>(),
            y.data_ptr<float>(),
            out.data_ptr<float>(),
            size
        );

        return out;
    }
    """

    cpp_declarations = "torch::Tensor fused_add_relu_cuda(torch::Tensor x, torch::Tensor y);"

    print("Compiling CUDA source to native binary via JIT toolchain...")
    fused_module = load_inline(
        name="fused_add_relu_extension",
        cpp_sources=cpp_declarations,
        cuda_sources=cuda_source,
        functions=["fused_add_relu_cuda"],
        verbose=False
    )
    print("Compilation successful! Native kernel module loaded.\n")

    # Benchmark: Unfused Framework Operations vs. Custom Fused Kernel
    N = 10_000_000
    x = torch.randn(N, device="cuda", dtype=torch.float32)
    y = torch.randn(N, device="cuda", dtype=torch.float32)

    # Warmup
    _ = torch.relu(x + y)
    _ = fused_module.fused_add_relu_cuda(x, y)
    torch.cuda.synchronize()

    # Benchmark 1: Unfused Execution (Allocates intermediate memory buffer for x + y)
    start_unfused = time.perf_counter()
    for _ in range(100):
        out_unfused = torch.relu(x + y)
    torch.cuda.synchronize()
    time_unfused = (time.perf_counter() - start_unfused) * 10  # Average time in ms

    # Benchmark 2: Custom Fused Kernel (Keeps intermediate addition in registers)
    start_fused = time.perf_counter()
    for _ in range(100):
        out_fused = fused_module.fused_add_relu_cuda(x, y)
    torch.cuda.synchronize()
    time_fused = (time.perf_counter() - start_fused) * 10  # Average time in ms

    # Output verification
    diff = torch.max(torch.abs(out_unfused - out_fused)).item()

    print(f"Unfused Execution Time  : {time_unfused:.4f} ms")
    print(f"Fused Kernel Time       : {time_fused:.4f} ms")
    print(f"Speedup Factor          : {time_unfused / time_fused:.2f}x")
    print(f"Max Absolute Error      : {diff:.6e}")

if __name__ == "__main__":
    inspect_cuda_driver_api()
    verify_oci_cgroup_isolation()
    compile_custom_cuda_kernel()
Category: 

Leave a Comment