Asynchronous Execution and Scheduling

Modern ML workloads, especially those deployed on heterogeneous hardware, often involve complex sequences of operations including data movement, pre/post-processing on the CPU, and intensive computations on accelerators like GPUs or specialized NPUs. Executing these operations sequentially can lead to significant underutilization of hardware resources, as one component often waits idly while another completes its task. Asynchronous execution and sophisticated scheduling are fundamental runtime capabilities required to mitigate these inefficiencies and maximize throughput.

The core idea is to represent the ML model inference or training step as a directed acyclic graph (DAG) of tasks, where nodes represent operations (e.g., kernel launch, memory copy, synchronization event) and edges represent dependencies. The runtime's responsibility is then to execute this graph as efficiently as possible, respecting dependencies while exploiting opportunities for parallelism and overlap.

Representing Tasks and Dependencies

Each unit of work managed by the runtime is encapsulated as a task. A task might correspond to:

• A GPU kernel launch (e.g., cudaLaunchKernel, SYCL kernel submission).
• A memory transfer operation (e.g., cudaMemcpyAsync, sycl::queue::memcpy).
• A CPU-bound computation.
• Synchronization primitives (e.g., recording or waiting on an event).

Dependencies dictate the execution order. For example, a GPU kernel cannot execute until its input data has been copied to the GPU memory, and a host-to-device copy cannot begin until the source data is ready on the host. These dependencies form the structure of the task DAG.

A simplified task graph showing dependencies. Kernels 1 and 2 can potentially run concurrently on different streams after their respective data transfers complete. Kernel 3 depends on the completion of both Kernel 1 and Kernel 2, potentially synchronized via events.

Execution Primitives: Streams and Events

Hardware acceleration APIs provide primitives for asynchronous execution. In CUDA, these are streams, while in SYCL/OpenCL, they are queues. Operations enqueued onto the same stream/queue are typically executed sequentially by the device, but operations on different streams/queues can execute concurrently, subject to hardware resource availability and explicit dependencies.

• Streams/Queues: Represent independent sequences of operations submitted to a device. Submitting tasks to different streams allows the hardware scheduler to potentially overlap their execution.
• Events: Act as synchronization points within or between streams. An event can be recorded after a task completes on one stream, and another stream can be instructed to wait for that event before executing subsequent tasks. This is important for managing cross-stream dependencies, as shown in the graph above where Kernel 3 waits for events signaling the completion of Kernels 1 and 2.

Scheduling Algorithms for Task Graphs

Once the task graph is defined, the runtime scheduler determines the actual execution order and timing. Common strategies include:

1. Topological Sort: The fundamental approach for executing DAGs. Tasks are executed in an order that respects all dependencies. A simple implementation maintains a count of pending dependencies for each task and adds tasks to a "ready queue" when their count reaches zero.
2. Ready Queue Management: Schedulers often maintain a queue of tasks whose dependencies have been met. The scheduler picks tasks from this queue to dispatch to available execution resources (e.g., GPU streams, CPU threads). The policy for picking tasks can vary (FIFO, LIFO, priority-based).
3. Priority-Based Scheduling: Tasks can be assigned priorities. For instance, tasks on the critical path of the graph might receive higher priority to minimize overall latency. Memory transfers might be prioritized to keep compute units fed.
4. Data Locality Aware Scheduling: In systems with complex memory hierarchies (e.g., NUMA on CPUs, HBM vs. DRAM on GPUs), schedulers might prioritize tasks that operate on data already resident in a favorable memory location.

The scheduler's goal is often to maximize resource utilization by keeping compute units busy and overlapping data movement with computation whenever possible.

Achieving Compute and Communication Overlap

Asynchronous operations via streams/queues are the primary enabler for overlap. Consider a typical pattern: copy input data (Host-to-Device, H2D), compute on the device, copy results back (Device-to-Host, D2H).

• Synchronous Execution: Each step waits for the previous one to complete. The CPU might be idle during GPU compute, and the GPU might be idle during data transfers.
• Asynchronous Execution: Using separate streams for computation and memory copies allows the runtime to issue these operations without waiting. For example:
  1. Issue H2D copy of chunk 1 (Stream M).
  2. Issue Compute on chunk 0 (Stream C), dependent on H2D chunk 0 completion.
  3. Issue D2H copy of chunk -1 results (Stream M), dependent on Compute chunk -1 completion.

By pipelining the processing of data chunks using multiple streams and appropriate event synchronization, the runtime can significantly hide the latency of data transfers behind computation.

Comparison of execution timelines. In synchronous execution, operations happen sequentially. In asynchronous execution, data transfers (H2D, D2H) for the next/previous chunk can overlap with the current chunk's GPU compute, reducing total execution time.

Synchronization Approaches

While asynchronicity is powerful, explicit synchronization is sometimes necessary. For example, the CPU might need to wait for a GPU result before proceeding with subsequent logic, or results from multiple concurrent streams must be available before a final reduction step.

Runtime APIs provide synchronization primitives:

• cudaStreamSynchronize(stream) / queue.wait(): Blocks the calling CPU thread until all previously submitted operations on the specified stream/queue are complete.
• cudaEventSynchronize(event) / event.wait(): Blocks the calling CPU thread until the specified event is recorded (i.e., the associated task is complete).
• cudaStreamWaitEvent(stream, event) / queue.ext_oneapi_submit_barrier(event_list): Enqueues a wait operation on a stream/queue; subsequent operations on that stream/queue will not begin until the event is recorded. This does not block the host thread.

Excessive or improperly placed synchronization points can negate the benefits of asynchronous execution, re-introducing serialization points and reducing parallelism. Runtime designers must carefully manage synchronization to ensure correctness without sacrificing performance. Minimizing host-device synchronization points is often a significant optimization goal.

Challenges in Asynchronous Scheduling

Designing effective asynchronous schedulers involves several challenges:

• Overhead: Managing task graphs, dependencies, events, and streams incurs runtime overhead. This overhead must be minimal compared to the execution time of the tasks themselves.
• Resource Contention: Even with overlap, operations might compete for limited hardware resources like memory bandwidth or specific execution units on the accelerator. The scheduler might need heuristics to mitigate this contention.
• Dynamicism: If the task graph changes dynamically (e.g., due to dynamic shapes or control flow), the scheduler must adapt quickly and efficiently.
• Complexity: Implementing and debugging multi-stream, event-driven asynchronous execution logic is inherently complex.

Mastering asynchronous execution and scheduling is essential for building high-performance ML runtimes capable of fully exploiting modern hardware capabilities. It requires a deep understanding of the target hardware's concurrency mechanisms, careful management of dependencies, and intelligent scheduling strategies to overlap operations effectively.