arXiv:2602.20826v1 [cs.OS] 24 Feb 2026
Exploiting Dependency and Parallelism: Real-Time Scheduling and Analysis for GPU Tasks Yuanhai Zhang
Songyang He
Ruizhe Gou
Sun Yat-sen University Guang Zhou, China
Sun Yat-sen University Guang Zhou, China
Hunan University Chang Sha, China
Mingyue Cui
Boyang Li
Shuai Zhao∗
Sun Yat-sen University Guangzhou, China
Sun Yat-sen University Guangzhou, China
Sun Yat-sen University Guangzhou, China
Kai Huang Sun Yat-sen University Guangzhou, China
Abstract With the rapid advancement of Artificial Intelligence, the Graphics Processing Unit (GPU) has become increasingly essential across a growing number of safety-critical application domains. Applying a GPU is indispensable for parallel computing; however, the complex data dependencies and resource contention across kernels within a GPU task may unpredictably delay its execution time. To address these problems, this paper presents a scheduling and analysis method for Directed Acyclic Graph (DAG)-structured GPU tasks. Given a DAG representation, the proposed scheduling scales the kernel-level parallelism and establishes inter-kernel dependencies to provide a reduced and predictable DAG response time. The corresponding timing analysis yields a safe yet nonpessimistic makespan bound without any assumption on kernel priorities. The proposed method is implemented using the standard CUDA API, requiring no additional software or hardware support. Experimental results under synthetic and real-world benchmarks demonstrate that the proposed approach effectively reduces the worst-case makespan and measured task execution time compared to the existing methods up to 32.8% and 21.3%, respectively.
1
Introduction
The rapid advancement of Artificial Intelligence (AI) has made Graphics Processing Units (GPUs) indispensable in many safetycritical domains such as autonomous driving, avionics, and industrial control, where both computational efficiency and timing predictability are crucial. GPU platforms provide massive parallel execution capabilities and are widely adopted in such systems to accelerate computation. A GPU task typically consists of multiple computation stages, i.e., kernels in the CUDA framework, which are often represented as a Directed Acyclic Graph (DAG) by modern AI compilers [8]. Although CUDA supports kernel-level parallel execution, the complex data dependencies and heterogeneous workload distributions within GPU tasks often delay the response time, particularly for tasks composed of numerous lightweight kernels with small computation loads [10]. Moreover, the black-box nature of GPU hardware execution further increases the pessimism of timing analysis and limits the flexibility of scheduling. ∗ Corresponding author. Email: [email protected]
Improving kernel-level concurrency is an effective way to reduce the overall response time of a GPU task, i.e., its makespan [9]. Several studies have enhanced concurrency using standard CUDA APIs such as streams and events [9, 13], while others have proposed resource allocators to accelerate kernel execution [14, 21]. However, these approaches do not consider the various computation load and resource requirements across kernels, resulting in unpredictable prolonged kernel execution time [20]. Existing response time analysis methods for GPU tasks are also limited in scope. Most analysis approaches focus on the GPU task with single or sequential kernels [11, 14, 18, 24], while others fail to consider the dependency delay between kernels [19]. To the best of our knowledge, the timing analysis for DAG-structured GPU tasks remains unseen. Although CUDA streams enforce a deterministic kernel submission order, the GPU’s hardware scheduler introduces nondeterminism in the actual execution order of concurrently launched kernels. Moreover, kernel-level preemption through prioritized CUDA streams is unreliable [1], meaning that the DAG scheduling and analysis approaches relying on node-level priorities cannot be applied [7, 22, 23]. Some studies instead schedule GPU tasks via priority mechanisms at the CPU process level [6, 17]; however, such approaches do not suit tasks that contain multiple kernels with complex dependencies. Therefore, developing a predictable scheduling and timing analysis framework for GPU tasks that fully exploits kernel dependency and parallelisms remains an open problem. Contribution. To address the aforementioned problems, this paper proposes a scheduling and timing analysis framework for DAG-structured GPU tasks, providing a reduced and predictable makespan. Given a DAG representation of a GPU task, the proposed method decomposes the task into a sequence of balanced groups, each containing kernels that can execute concurrently. We propose a parallelism scaling mechanism for kernels within a group by adjusting the computing resource requirement according to each node’s computation load, achieving balanced execution times and avoiding resource contention. To further reduce the makespan, parallel nodes of a balanced group are opportunistically launched. For large kernels that exceed the spare GPU capacity, a node segmentation mechanism is constructed to divide them into smaller sequential segments. The sequential execution order of each balanced group is ensured by the constructed extra dependency on the
DAC ’26, July 26-29, 2026, Long Beach, CA, USA
Yuanhai Zhang, Songyang He, Ruizhe Gou, Mingyue Cui, Boyang Li, Shuai Zhao, and Kai Huang
2
System Model
This section introduces the system model, which describes a GPU platform executing a task in a DAG-structured intermediate representation (IR). Figure 1 illustrates a typical GPU task, which begins with data transfer from the CPU to the GPU (memHtoD), proceeds through a computation phase modeled as a DAG, and concludes with data transfer back to the CPU (memDtoH ). The GPU platform contains 𝑀 homogeneous Streaming Multiprocessors (SMs), e.g. 𝑀 = 80 for NVIDIA V100. Although NVIDIA GPUs are assumed, the proposed method applies to any platform following the Single Instruction Multiple Thread (SIMT) execution model.
2.1
GPU Task Model
The proposed method focuses on a single periodic DAG-structured GPU task. The computation stage on the GPU is periodically released by a CPU process in the system. A periodic DAG task is defined as 𝜏 = {𝐺,𝑇 , 𝐷}, where 𝐺 = (𝑉 , 𝐸) is a directed acyclic graph, 𝑇 is the period, and 𝐷 = 𝑇 is the implicit deadline. Each node 𝑣𝑖 ∈ 𝑉 represents a kernel, and each edge 𝑒𝑖,𝑗 = (𝑣𝑖 , 𝑣 𝑗 ) ∈ 𝐸 captures a data dependency (e.g., write-after-read), ensuring that the successor 𝑣 𝑗 starts only after the predecessor 𝑣𝑖 completes [13]. A node is said to be released when all its predecessors have finished. Nodes in a set 𝑆 without predecessors or successors form the source and sink node sets, denoted by 𝑠𝑟𝑐 (𝑆) and 𝑠𝑖𝑛𝑘 (𝑆), respectively. Without loss of generosity, we assume a single source and a single sink node for 𝐺. If 𝑣𝑖 is a transitive predecessor of 𝑣 𝑗 , then 𝑣𝑖 is an ancestor (i.e., 𝑣𝑖 ∈ 𝑎𝑛𝑐𝑒 (𝑣 𝑗 )). Nodes that share no transitive dependencies with 𝑣𝑖 form its concurrent set, denoted as 𝑐𝑜𝑛(𝑣𝑖 ). A path 𝜆 = ⟨𝑣 1, 𝑣 2, . . . , 𝑣𝑘 ⟩ is defined as a sequence of nodes satisfying (𝑣𝑖 , 𝑣𝑖+1 ) ∈ 𝐸 for all 𝑣𝑖 ∈ 𝜆 \ 𝑣𝑘 . The length of a path is the sum of the computation loads of its nodes.
2.2
Execution Model
In the SIMT execution model, parallel execution for each kernel is configured through two launch parameters, 𝐷𝑔 and 𝐷𝑏 , which represent the thread grid and block dimensions, respectively. Kernel execution: In the DAG model, each node 𝑣𝑖 = {𝐶ˆ𝑖 , 𝐶𝑖 , 𝑚𝑖 } represents a kernel execute on GPU, where 𝐶ˆ𝑖 denotes the computation load, 𝐶𝑖 is the kernel execution time, and 𝑚𝑖 represents the number of SMs required for execution. The parallelism 𝑚𝑖 can be
CUDA Driver Backend
• A scheduling scheme integrating parallelism scaling, node segmentation , and extra dependency mechanism for DAGstructured GPU tasks to reduce the makespan. • A timing analysis of the worst-case makespan of the GPU tasks without the assumption of node priorities. • Experimental results show that the proposed method outperforms the existing methods by up to 32.8% in worst-case makespan and 21.3% in measured task execution time.
(a) Intermediate Representation (IR)
User scheduling
DAG, resulting in a predictable overall makespan. The proposed method can be entirely implemented within the scope of the standard CUDA API, without requiring any additional hardware or software support. A large-scale synthesized experiment and a case study demonstrate the effectiveness of the proposed scheduling and timing analysis framework. In summary, the contributions of this paper are as follows:
mem HtoD
1
4 2
𝑣
2
Node
𝑣
Data dependency
5
𝑣
3
𝑣
1
𝑣
mem DtoH
𝑣
Kernels on Streams Standard CUDA API Black-box to user
𝑣
(b) CUDA Structure Stream1 Stream2
cudaStreamWaitEv ent(event,S)
𝑣 𝑣
Stream3
𝑣
𝑣
𝑣
𝑣
𝑣
(c) Hardware work queue
buffer
buffer
⋯⋯
buffer
kernel <<< D , D , S>>>(*args)
Launch Setup & Dependency
Block Scheduling
Figure 1: The execution of a GPU task in DAG representation
adjusted by varying 𝐷𝑔 and 𝐷𝑏 in the kernel launch configuration. Each SM has a finite resource capacity, limited by factors such as register usage and shared memory. When the number of active threads exceeds the capacity of a single SM, more SMs are required to accommodate the workload. For example, if a kernel is launched with 𝐷𝑔 = 1 and 𝐷𝑏 = 256, fully occupying one SM. Increasing 𝐷𝑔 in this case scales up 𝑚𝑖 , thereby increasing the degree of parallelism. The execution time of a kernel follows Gustafson’s law, where the time per thread decreases as the number of threads increases [24]. The computation load 𝐶ˆ𝑖 is defined as the maximum kernel execution time when 𝑣𝑖 occupies one SM. The execution time under parallelism 𝑚𝑖 is modeled as l 𝑚 m 𝐶ˆ 𝑖 𝑖 𝐶𝑖 = max 𝑡 min, , (1) 𝑀 𝑚𝑖 where 𝑡 min denotes the theoretical lower bound constrained by memory throughput, according to the Roofline performance model [16]. The maximum useful parallelism for accelerating 𝑣𝑖 is 𝑚𝑖max = 𝐶ˆ𝑖 /𝑡 min . For simplicity, we set 𝑡 min = 1 as the basic time unit, which does not affect the generality of the proposed method. Note that increasing 𝑚𝑖 does not always minimize 𝐶𝑖 . If some SMs are already occupied, the node cannot achieve the desired parallelism and may incur unpredictable delay. Kernel concurrency: The standard CUDA API provides kernel concurrency and dependency control through streams and events. Each kernel in the DAG is launched on a stream. Kernels within the same stream execute sequentially, while kernels on different streams may run concurrently if resources allow. Inter-stream dependencies are enforced using CUDA events. Kernel dependencies can also be controlled using CUDA Graphs API, which additionally reduces kernel launch overhead, an important advantage for tasks composed of many small kernels [5]. When a GPU task is launched, the CUDA driver generates command buffers encoding kernel configurations and dependencies, and places them into the hardware
Exploiting Dependency and Parallelism: Real-Time Scheduling and Analysis for GPU Tasks
work queue. The block scheduler dispatches kernels whose dependencies are satisfied to available SMs, while others remain pending [12]. Due to the black-box nature of the hardware scheduler, the execution order of concurrent kernels and the resulting contention delays are unpredictable to users. Example: Figure 1 illustrates the execution of a DAG-structured GPU task with 7 nodes, 𝑉 = {𝑣 1, . . . , 𝑣 7 }, where 𝑣 1 and 𝑣 7 are the source and sink nodes of 𝐺, respectively. The number inside each node indicates its computation load, e.g., 𝐶ˆ3 = 3 time units. The concurrent node set of 𝑣 2 is 𝑐𝑜𝑛(𝑣 2 ) = {𝑣 3, 𝑣 4 }. The nodes in the same color are launched on the same streams as: 𝑆 1 = {𝑣 2 }, 𝑆 2 = {𝑣 1, 𝑣 3, 𝑣 5, 𝑣 7 }, and 𝑆 3 = {𝑣 4, 𝑣 6 }. Inter-kernel dependencies like {𝑣 1, 𝑣 4 } and {𝑣 2, 𝑣 7 } are implemented via CUDA events. These kernels are placed into the hardware work queue as command buffers and dispatched to SMs by the block scheduler. The actual execution order between concurrent nodes such as 𝑣 2 and 𝑣 3 , however, remains unpredictable. Problem formulation: Given a DAG-structured GPU task, the objective is to minimize the task execution time while guaranteeing a predictable DAG makespan. This is achieved by appropriately scaling each node’s parallelism 𝑚𝑖 and adjusting the DAG structure to balance concurrent execution under dependency constraints.
3
Sub-graph Division
To minimize the delay caused by complex data dependencies in the DAG, concurrent nodes without dependency constraints are extracted. This section presents a sub-graph division method that decomposes the DAG into a sequence of disjoint balanced groups, forming the basis for the proposed scheduling and timing analysis. First, the nodes that may experience dependency delay are identified, and their ancestors are organized into disjoint blocks. Then, within each block, concurrent nodes are further arranged into a sequence of balanced groups, the basic sub-graph for scheduling and analysis.
3.1
Block Formulation
A node with multiple predecessors is defined as a join node, and it cannot start execution until all predecessors finish. A join node may experience a dependency delay, which source from the unbalanced finish time of its predecessors. Join nodes are first identified and sorted according to their Cumulative ancestor workload. The notation 𝑊 𝑎𝑛𝑐 (𝑣𝑖 ) de notes the cumulative computation load cost when 𝑣𝑖 finished as: ∑︁ 𝑊 𝑎𝑛𝑐 (𝑣𝑖 ) = 𝐶ˆ 𝑗 , ∀𝑣𝑖 ∈ 𝐺 . (2)
DAC ’26, July 26-29, 2026, Long Beach, CA, USA
After all the 𝑁 blocks are formed, the remaining nodes are collected into a residual block 𝐵 𝑁 +1 . Notice that the blocks in 𝐵 also follow the order of Θ. Since all the 𝜃 𝑘 do not belong to their ancestor block 𝐵𝑘 , all the nodes in 𝐵𝑘 contain at most one predecessor in the same block. Any 𝑣𝑖 ∈ 𝐵𝑘 is a local source if all its predecessors lie outside 𝐵𝑘 , i.e. 𝑣𝑖 ∈ 𝑠𝑟𝑐 (𝐵𝑘 ) , and a local sink if all successors lie outside 𝐵𝑘 , i.e. 𝑣𝑖 ∈ 𝑠𝑖𝑛𝑘 (𝐵𝑘 ). A local complete path 𝜆𝑘𝑗 connects a local source and a local sink: 𝜆𝑘𝑗 = ⟨𝑣𝑠 , . . . , 𝑣𝑒 ⟩,
(4)
Each node in 𝐵𝑘 belongs to at least one local complete path, and each path ends with a predecessor of the join node 𝜃 𝑘 . The notation Λ𝑘 denotes the set of local complete paths of 𝐵𝑘 .
3.2
Balanced Groups Construction
To minimize the release time of 𝜃 𝑘 , the proposed scheduling balances the finish time of each local complete path in 𝐵𝑘 . The concurrent nodes in different local complete paths that can occupy all the 𝑀 SMs are extracted and organized into a balanced group. Following the topology order, each group can safely run sequentially and form the ordered balanced group list Π. The formulation of a balanced group 𝜋 𝑗 ∈ Π satisfies the following rules: Rule 1: |𝜋 𝑗 | ≤ 𝑀, ∀𝜋 𝑗 ∈ Π. Each kernel requires at least one SM; therefore, the number of concurrent nodes cannot exceed 𝑀. Otherwise, the resource contention among 𝜋 𝑗 can not be avoided. Rule 2: If |𝜋 𝑗 | > 1, then 𝑚𝑚𝑎𝑥 < 𝑀, ∀𝑣𝑖 ∈ 𝜋 𝑗 . For a large 𝑖 kernel with enough computation load to fully occupy all the SMs (𝑚𝑚𝑎𝑥 ≥ 𝑀), the kernel-level parallelism is unnecessary. 𝑖 Algorithm 1 outlines the process of complete sub-graph division in Section 3, which takes 𝐺 as the input and outputs the ordered balanced group list Π. First, the join nodes are detected and ordered by Cumulative ancestor workload (Lines 1). Then the blocks are constructed (Lines 2-6). For each block, after the local complete path set is identified (Lines 9), balanced groups are formed by iteratively extracting the head nodes of each local complete path until all nodes are grouped (Lines 10–17), following Rule 1 (Line 11) and Rule 2 (Lines 12-14). The time complexity of Algorithm 1 is O (|𝑉 |). Example: The DAG in Figure 1 contains 2 join nodes: 𝜃 1 = 𝑣 5 and 𝜃 2 = 𝑣 7 . The DAG is decomposed into 3 disjoint blocks: 𝐵 1 = {𝑣 1, 𝑣 3, 𝑣 4 }, 𝐵 2 = {𝑣 2, 𝑣 5, 𝑣 6 }, and 𝐵 3 = {𝑣 7 }. The corresponding local complete path sets are Λ1 = {⟨𝑣 1, 𝑣 3 ⟩, ⟨𝑣 1, 𝑣 4 ⟩} and Λ2 = {⟨𝑣 2 ⟩, ⟨𝑣 5 ⟩, ⟨𝑣 6 ⟩} . The balanced group list of this DAG is Π = ⟨{𝑣 1 }, {𝑣 3, 𝑣 4 }, {𝑣 2, 𝑣 5, 𝑣 6 }, {𝑣 7 }⟩.
4
∀𝑣 𝑗 ∈𝑎𝑛𝑐𝑒 (𝑣𝑖 )∪{𝑣𝑖 }
𝑣𝑠 ∈ 𝑠𝑜𝑢𝑟𝑐𝑒 (𝐵𝑘 ), 𝑣𝑒 ∈ 𝑠𝑖𝑛𝑘 (𝐵𝑘 ).
Scheduling and Analysis
The ordered list Θ = ⟨𝜃 1, . . . , 𝜃 𝑁 ⟩ contains all the 𝑁 join nodes sorted by descending cumulative ancestor workload, which also preserves the DAG’s topological order. Next, the DAG is partitioned into a sequence of disjoint node blocks 𝐵 = [𝐵 1, 𝐵 2, . . . , 𝐵 𝑁 +1 ] (not confused with the thread block in CUDA), where each block 𝐵𝑘 includes all ancestors of the join node 𝜃 𝑘 that are not assigned to previous blocks:
This section presents a scheduling and timing analysis framework for DAG-structured GPU tasks on the basis of the proposed subgraph division. The scheduling method determines the execution order and the required parallelism of each node. Building on the schedule scheme, the timing analysis derives the overall DAG makespan.
𝑘Ø −1
The proposed scheduling framework operates by scaling the required parallelism of each node and restructuring the DAG through node segmentation and extra dependency constraints. Importantly,
𝐵𝑘 = 𝑎𝑛𝑐𝑒 (𝜃 𝑘 ) \
𝑗=1
𝐵𝑗,
∀𝜃 𝑘 ∈ Θ
(3)
4.1
Scheduling
DAC ’26, July 26-29, 2026, Long Beach, CA, USA
Yuanhai Zhang, Songyang He, Ruizhe Gou, Mingyue Cui, Boyang Li, Shuai Zhao, and Kai Huang
Algorithm 1: Sub-graph_Division(𝐺) Input: DAG 𝐺 = (𝑉 , 𝐸) Output: Balanced group set Π 𝑎𝑛𝑐 ; 1 Θ ← join nodes sorted by 𝑊 2 Initialize 𝐵 ← [ ]; 3 for 𝜃 𝑘 ∈ Θ do 4 𝐵𝑘 ← Eq. (3); Append 𝐵𝑘 to 𝐵; 5 end 6 Append residual block to 𝐵; 7 Initialize Π ← [ ]; 8 for each 𝐵𝑘 ∈ 𝐵 do 9 Identify Λ𝑘 via Eq. (4); 10 while Λ𝑘 ≠ ∅ do 11 𝜋 ← 𝑇𝑜𝑝 𝑀 {head(𝜆) | 𝜆 ∈ Λ𝑘 } by 𝑊 𝑎𝑛𝑐 ; 12 if 𝑚𝑚𝑎𝑥 > 𝑀, ∀𝑣𝑖 ∈ 𝜋 then 𝑖 13 𝜋 ← node with max 𝑚𝑚𝑎𝑥 in 𝜋; 𝑖 14 end 15 Remove nodes in 𝜋 from Λ𝑘 ; 16 Append 𝜋 to Π; 17 end 18 end
it does not rely on GPU priority mechanisms and can be fully implemented using the standard CUDA API. The parallelism scaling scheme equalizes the execution times of concurrently executed nodes according to their computation loads, thereby minimizing the completion time of each balanced group. The node segmentation mechanism and the extra dependency further improve hardware utilization while ensuring a predictable overall DAG makespan. Parallelism scaling: For lightweight kernels within a balanced group, the scheduler aims to equalize their execution times. For any balanced group 𝜋 𝑗 ∈ Π, the parallelism 𝑚𝑖 of each node 𝑣𝑖 ∈ 𝜋 𝑗 is determined in proportion to its computation load: 𝐶ˆ𝑖 𝑚𝑖 = min 𝑚𝑚𝑎𝑥 , round · 𝑀 , ∀𝑣𝑖 ∈ 𝜋 𝑗 (5) 𝑖 𝑊 (𝜋 𝑗 ) Í where 𝑊 (𝜋 𝑗 ) = 𝑣𝑖 ∈𝜋 𝑗 𝐶ˆ𝑖 is the total computation load of the group. This configuration ensures that small kernels share GPU resources proportionally, while yielding balanced finishing times. This parallelism setup ensures the computing resource requirement Í of a group would not exceed the GPU capacity (i.e. 𝑚𝑖 ≤ 𝑀, ∀𝑣𝑖 ∈ 𝜋 𝑗 ), which guarantees each node can achieve its required number of SM, thereby achieving the execution time in Eq. (5). Launching parallel nodes: If a balanced group 𝜋 𝑗 does not fully occupy the GPU, parallel nodes may be launched on the spare SMs. The GPU spare capacity when 𝜋 𝑗 executing is: 𝑊 𝑆 𝑗 = max(𝐶𝑖 ) × 𝑆 𝑗 , ∀𝑣𝑖 ∈ 𝜋 𝑗 (6) Í where 𝑆 𝑗 = 𝑀 − 𝑚𝑖 denote the number of remaining SMs. The set of parallel nodes of a group 𝜋 𝑗 is defined as: ©Ø ª 𝑝𝑎𝑟𝑎(𝜋 𝑗 ) = 𝑠𝑟𝑐 𝑐𝑜𝑛(𝑣𝑖 ) \ 𝜋 𝑗 ® ∩ released_node 𝑣𝑖 ∈𝜋 𝑗 « ¬
(7)
where 𝑠𝑟𝑐 (·) returns the source nodes of the given set. Only nodes that are released at the moment when 𝜋 𝑗 is released are put into the parallel set, and whether a node is released is known during the process of scheduling. After the parallel set is formed, nodes in 𝑝𝑎𝑟𝑎(𝜋 𝑗 ) are kept chosen to launch in descending order of cumulative ancestor workload until the spare capacity is exhausted or no parallel nodes remain. For each chosen node 𝑣𝑐 , its parallelism is assigned greedily as 𝑚𝑐 = min(𝑚𝑚𝑎𝑥 , 𝑆 𝑗 ) to fully utilize the spare 𝑐 SM. Node segmentation: In the SIMT execution model, a node can be divided into multiple segments, each with its own computation load and parallelism configuration, but executing the same function as the original node. During the execution of 𝜋 𝑗 , if a selected node 𝑣𝑐 ∈ 𝑝𝑎𝑟𝑎(𝜋 𝑗 ) has a computation load exceeding the available capacity (i.e., 𝐶ˆ𝑐 > 𝑊 𝑆 𝑗 ), it is split into two segments 𝑣𝑐,𝑝 and 𝑣𝑐,𝑟 with 𝐶ˆ𝑐,𝑝 = 𝑊 𝑆 𝑗 and 𝐶ˆ𝑐,𝑟 = 𝐶ˆ𝑐 − 𝐶ˆ𝑐,𝑝 . The parallel segment 𝑣𝑐,𝑝 is scaled as the chosen node with 𝑚𝑐,𝑝 = min(𝑚𝑚𝑎𝑥 𝑐,𝑝 , 𝑆 𝑗 ) and finishes together with 𝜋 𝑗 , while the residual segment 𝑣𝑐,𝑟 inherits the forward dependencies of 𝑣𝑐 and execute after 𝜋 𝑗 finished. Since at most one segmentation occurs per group, the total number of segmentations in a DAG is bounded by the number of balanced groups |Π|, which is smaller than |𝑉 |. Extra dependency: To preserve the sequential execution order of each balanced group, extra dependencies are introduced. Let 𝑣 𝑅 = arg max𝑣𝑖 ∈𝜋 𝑗 𝐶𝑖 denote the node in 𝜋 𝑗 with the maximum execution time. For each unlaunched node 𝑣𝑛 ∈ 𝑝𝑎𝑟𝑎(𝜋 𝑗 ) (including unexecuted segments), an edge (𝑣 𝑅 , 𝑣𝑛 ) is added, forcing these nodes to start only after the current group finishes. For each launched node or segment 𝑣𝑛 ∈ 𝑝𝑎𝑟𝑎(𝜋 𝑗 ), extra dependencies (𝑣𝑛 , 𝑠𝑢𝑐𝑐 (𝑣 𝑅 )) are added to ensure they finish no later than the current group.
Algorithm 2: Scheduling(𝐺, Π) Input: DAG 𝐺 = (𝑉 , 𝐸), balanced group list Π Output: Parallelism 𝑚𝑖 , segmentation 𝑉 , extra deps 𝐸 1 Initialize 𝑉 ← ∅, 𝐸 ← ∅; 2 Released_nodes ← {𝑠𝑟𝑐 (𝐺)} ; 3 foreach 𝜋 𝑗 ∈ Π do 4 𝑚𝑖 ← Eq. (5), ∀𝑣𝑖 ∈ 𝜋 𝑗 ; 5 𝑊 𝑆 𝑗 ← Eq. (6) ; 6 𝑝𝑎𝑟𝑎(𝜋 𝑗 ) ← Eq. (7); 7 while 𝑊 𝑆 𝑗 > 0 and 𝑝𝑎𝑟𝑎(𝜋 𝑗 ) ≠ ∅ do 8 𝑣𝑐 ← Top(𝑝𝑎𝑟𝑎(𝜋 𝑗 )) by 𝑊 𝑎𝑛𝑐 ; 9 if 𝐶ˆ𝑐 > 𝑊 𝑆 𝑗 then 10 Add {𝑣𝑐 : (𝑣𝑐,𝑝 , 𝑣𝑐,𝑟 )} to 𝑉 ; 11 𝐶ˆ𝑐 ← 𝐶ˆ𝑐,𝑝 ; 12 end 13 𝑚𝑐 ← min{𝑚𝑚𝑎𝑥 , 𝑆 𝑗 }; 𝑐 14 Add extra dependencies to 𝐸; 15 𝑊 𝑆 𝑗 ← 𝑊 𝑆 𝑗 − 𝐶ˆ𝑐 ; 16 Remove 𝑣𝑐 from 𝑝𝑎𝑟𝑎(𝜋 𝑗 ); 17 end 18 Update Released_nodes ; 19 end
Exploiting Dependency and Parallelism: Real-Time Scheduling and Analysis for GPU Tasks
Algorithm 2 outlines the complete scheduling process of Section 4.1, which takes 𝐺 and Π as the input and outputs the scheduling scheme including parallelism setup 𝑚𝑖 for each node, node segmentation scheme 𝑉 , and the extra dependency set 𝐸. For each balanced group 𝜋 ∈ Π, the scheduler first assigns parallelism values for each node in 𝜋 (Line 4). Then, the spare capacity 𝑊 𝑆 𝑗 and parallel nodes 𝑝𝑎𝑟𝑎(𝜋 𝑗 ) are initialized (Line 5-6). The parallel nodes are launched following the descending cumulative ancestor workload order until spare capacity is filled or all the parallel nodes are launched (Line 7-17). Node segmentation is applied when the chosen node 𝑣𝑐 is too big to fit in 𝑊 𝑆 𝑗 (Line 9-12). The chosen node (segment) is always set with the maximum parallelism (Line 13). After that, the extra dependencies are inserted to ensure the execution order (Line 14), and the status is updated for the next parallel node chosen (Line 15-16). The release nodes are updated when the scheduling of a group finishes (Line 18). The time complexity of Algorithm 2 is O (|𝜋 |) ≺ O (|𝑉 |).
1 π 𝑣 π
2 𝑣 ,
2 𝑣 ,
2
5 𝑣
2
𝑣 𝑣
π
3
Kernels on SM
SM SM SM SM SM SM SM SM
Unpredictable delay
SM SM SM SM SM SM t
(a) Analysis for Greedy
𝑅(𝜋 ) 𝑅(𝜋 )
𝑅(𝜋 )
𝑅(𝜋 )
t
(b) Proposed Analysis
Figure 3: Example of timing analysis Theorem 4.2. The makespan of the DAG is upper-bounded as: ∑︁ 𝑅(𝐺) ≤ 𝑅(𝜋 𝑗 ). (9) 𝜋 𝑗 ∈Π
Proof. Since the proposed extra dependencies enforce a strictly sequential execution order over the groups in Π, the total makespan is bounded by the sum of all the relative response times. □
Node segmentation Original dependency Extra dependency
π
1 𝑣
Balanced group
𝑣
Figure 2: Example of proposed schedule Example. Figure 2 illustrates the proposed scheduling scheme for the example DAG in Figure 1 on a GPU with 𝑀 = 6 SMs. The balanced group 𝜋 2 = {𝑣 3, 𝑣 4 }, configured with 𝑚 3 = 𝑚 4 = 2, does not fully occupy the GPU. Thus, the only node in 𝑝𝑎𝑟𝑎(𝜋2 ) = {𝑣 2 } is selected to launch. Since 𝐶ˆ2 = 4 exceeds the available capacity (𝑊 𝑆 2 = 2), node 𝑣 2 is split into a parallel segment 𝑣 2,1 and a residual segment 𝑣 2,2 , each with load 𝑐ˆ2,1 = 𝑐ˆ2,2 = 2. The segment 𝑣 2,1 executes in parallel with 𝜋2 under 𝑚 2,1 = 2, while the residual segment 𝑣 2,2 is scheduled with its own group 𝜋3 . The extra dependencies {𝑣 2,1, 𝑣 5 } and {𝑣 3, 𝑣 2,2 } enforce the SM occupancy of each group and the sequential execution between 𝜋2 and 𝜋3 .
4.2
DAC ’26, July 26-29, 2026, Long Beach, CA, USA
Response Time Analysis
The extra dependencies introduced by scheduling enforce a determined group execution order, enabling a predictable upper bound on the DAG response time. Lemma 4.1. For any balanced group 𝜋 𝑗 ∈ Π, its relative response time 𝑅(𝜋 𝑗 ), defined as the interval from the completion of 𝜋 𝑗 −1 to the completion of 𝜋 𝑗 , is upper-bounded as: 𝑅(𝜋 𝑗 ) ≤ max 𝐶𝑖 .
(8)
𝑣𝑖 ∈𝜋 𝑗
Proof. The extra dependencies ensure that the nodes in 𝜋 𝑗 execute concurrently. The parallelism scaling and node segmentation mechanisms avoid unpredictable delay. Thus, 𝑅(𝜋 𝑗 ) is bounded by its maximum node execution time. □
Sustainability: If the actual execution time of any node 𝑣𝑖 is less than its analytical bound 𝐶𝑖 in Eq. (1), the corresponding group’s response time will not exceed the worst-case bound in Eq. (8). Such early completion also does not violate the sequential execution order enforced by the extra dependencies. Therefore, the timing analysis in Eq. (9) remains sound under early completion. Example. Figure 3 illustrates the timing analysis, where the size of each kernel on SM represents its computation load. Figure 3(a) shows the analysis under the Greedy schedule, where each 𝑣𝑖 is assigned 𝑚𝑖max . The concurrent node sets {𝑣 2, 𝑣 3, 𝑣 4 } and {𝑣 5, 𝑣 6 } require more SMs than the GPU can provide, leading to resource contention and unpredictable kernel execution delays. As a comparison, Figure 3(b) shows the analysis under the proposed scheduling. The proposed parallelism scaling and node segmentation ensure that the SM demands of 𝜋 2 and 𝜋3 never exceed the GPU capacity, eliminating contention-induced delays. The extra dependency enforces sequential execution between 𝜋2 and 𝜋3 , yielding a predictable makespan bounded by 𝑅(𝐺) ≤ 5. Notably, this bound does not make any assumptions about the hardware scheduling.
5 Evaluation 5.1 Experimental Setup To evaluate the worst-case makespan, a large set of synthesized DAGs was generated using the tool in [3]. Unless otherwise stated, the maximum DAG depth (i.e., number of layers) is randomly selected between 5 and 8. DAG construction starts from a single source node and proceeds layer by layer, with the number of nodes in each layer uniformly sampled between 2 and the maximum parallelism parameter 𝑃. For each setup, 1, 000 DAGs are randomly generated. Within each DAG, the node’s computation load is assigned uniformly, satisfying a given average load 𝐶ˆavg . To evaluate the GPU task execution time, a case study on realworld benchmarks is performed. The proposed scheduling scheme and the competitive method are implemented using the CUDA Graph API, reducing the overhead of repeated kernel launches [5]. The required parallelism 𝑚𝑖 for each node is realized by configuring the kernel launch parameters 𝐷𝑔 and 𝐷𝑏 on each target platform. To
emulate different computation loads across nodes, a single kernel function is used with various iteration counts.
5.2
Large-scale Experiment
Normalized Makespan
In this experiment, we compare the derived worst-case makespan bound in Eq. (9) (Proposed) against 3 representative methods: Greedy. Each kernel is assigned the maximum parallelism as 𝑚𝑖 = min{𝑚𝑚𝑎𝑥 , 𝑀 } for each 𝑣𝑖 . As the kernel execution order 𝑖 and delay caused by resource competition are unpredictable, a sequential node execution is assumed in the analysis. Greedy + unaware (Reference). The parallelism assignment ignore the platform structure as 𝑚𝑖 = 𝑚𝑚𝑎𝑥 for each 𝑣𝑖 , with a 𝑖 sequential node execution order. The resulting makespan is used as a reference for normalization in analysis. Graham_para. A parallel adaptation of Graham’s bound [4], which analyzes the DAG’s makespan without assuming node-level priorities. Each node is replaced with a set of parallel unit nodes with basic computation load (𝐶ˆ = 1), ensuring each unit occupies exactly one SM, and the Graham bound is then applied to the transformed DAG.
Normalized Makespan
Yuanhai Zhang, Songyang He, Ruizhe Gou, Mingyue Cui, Boyang Li, Shuai Zhao, and Kai Huang
1.0 0.8 0.6 0.4
Greedy
4
Graham_para
5
Proposed
7 8 Max Parallelism
9
10
1.0 0.8 0.6 0.4
Greedy
10
20
1.0 0.8
6
Figure 5: Makespan with various 𝑃 when 𝑀 = 32, 𝐶ˆavg = 20
Normalized Makespan
DAC ’26, July 26-29, 2026, Long Beach, CA, USA
Graham_para
Proposed
30 40 50 Number of Vertex
60
70
Figure 6: Makespan with various |𝑉 | when 𝑀 = 32, 𝐶ˆavg = 20
0.6 0.4 Greedy
0.2 4
Graham_para
8
16
Proposed
32 64 Number of SM
128
256
Figure 4: Makespan with various 𝑀 when 𝐶ˆavg = 20 Figure 4 reports the normalized makespan under various numbers of SMs, with 𝑀 ranging logarithmically from 4 to 256. The performance gap between the proposed method and the others generally widens as 𝑀 increases. When 𝑀 = 128, the proposed method outperforms Graham_para by 28.7% on average due to finer-grained exploitation of kernel concurrency; this gap narrows to 21.2% when 𝑀 = 256, where both methods benefit from abundant SM resources and reach near-maximal parallelism. For 𝑀 ≤ 8, the performance gap between Greedy and the proposed method remains below 3%, as kernels frequently saturate the GPU and execute mostly sequentially. As 𝑀 grows beyond 32, Greedy fails to capitalize on the additional SMs, leading to noticeable degradation. Greedy and Greedy_unaware converge when 𝑀 ≥ 64, where all kernels can achieve their maximum parallelism. Figure 5 reports the normalized makespan under different 𝑃. The proposed method consistently achieves the lowest makespan across all configurations, outperforming Graham_para and Greedy by 18.9% and 27.2% on average, demonstrating its robustness on various DAG structures. The performance gap between the proposed method and Greedy decreases from 23.9% when 𝑃 = 4 to 13.7% when 𝑃 = 10. This is because with a fixed 𝑀, increasing 𝑃 makes the DAG more likely to fully occupy the GPU and thereby reduces the advantage of the proposed parallelism scaling. Figure 6 illustrates the normalized makespan under various |𝑉 |. In this setup, |𝑉 | is scaled by increasing the DAG depth while fixing
𝑃 = 32. Across all configurations, the proposed method consistently yields the lowest makespan, showing robustness with the DAG size expansion. The advantage of the proposed method increases as |𝑉 | increases, up to 32.8% and 18.2% than Greedy and Graham_para, respectively. This is because a deeper DAG depth results in more balanced groups, amplifying the benefits of the proposed method.
5.3
Case Study
A case study was conducted on 3 DAG-structured benchmark tasks: Laplace, Gaussian elimination, and Stencil [2, 15], on two GPU platforms: an NVIDIA RTX 3060 (𝑀 = 30) and a Jetson Orin Nano (𝑀 = 8). Each benchmark was executed 100 times, with the measured task execution time summarized in Table 1 and 2. The Greedy approach assigns maximum parallelism to each node without enforcing sequential constraints among nodes. In most setups, the proposed method achieves more stable task execution times (lower standard deviation) due to its enforced execution order. It also achieves a lower execution time in most cases, with improvements up to an average 21.3% when 𝑀 = 8, 𝐶ˆavg = 4 and 15.0% when 𝑀 = 30, 𝐶ˆavg = 20 for Gaussian elimination. In these system setups, most kernels can neither saturate the GPU nor get its maximum parallelism, amplifying the benefit of the proposed parallelism scaling mechanism. When the average load is too small or too large for 𝑀, kernels either mostly reach maximum parallelism or fully occupy the GPU, reducing the proposed benefit, which is consistent with trends against Graham_para in Figure 4.
6
Conclusion
This paper presented a scheduling and timing analysis framework for DAG-structured GPU tasks. By decomposing the DAG into balanced groups, the method enables fine-grained scheduling and analysis, providing a reduced and predictable makespan. The scheduling is fully realizable within the scope of the standard CUDA API,
Exploiting Dependency and Parallelism: Real-Time Scheduling and Analysis for GPU Tasks
Table 1: Task execution time of benchmarks when 𝐶ˆ𝑎𝑣𝑔 = 4 Method Execution time(ms) Gaussian M=8 Laplace Stencil Gaussian M=30 Laplace Stencil
Proposed max. avg. 74.25 51.23 100.83 83.30 114.12 86.07 10.22 7.32 11.32 11.05 11.98 11.37
std. 7.47 7.20 9.13 0.33 0.14 0.15
Greedy max. avg. 85.61 65.07 121.72 104.39 122.76 106.12 12.25 7.38 11.35 11.08 11.51 11.09
std. 9.12 9.02 8.17 0.60 0.15 0.17
Table 2: Task execution time of benchmarks when 𝐶ˆ𝑎𝑣𝑔 = 20 Method Execution time(ms) Gaussian M=8 Laplace Stencil Gaussian M=30 Laplace Stencil
Proposed max. avg. 193.11 165.24 309.29 280.31 333.62 297.03 29.05 23.96 46.81 40.74 49.44 43.39
std. 11.36 12.17 12.12 0.99 1.43 3.11
Greedy max. avg. 194.50 167.90 331.42 306.65 353.54 324.89 31.71 28.20 51.30 44.82 54.00 48.45
std. 12.34 10.36 12.87 1.05 1.43 3.39
without additional hardware or software support. The corresponding timing analysis does not make any assumption about kernellevel priority. Extensive experiments on large-scale synthetic DAGs and a case study prove the effectiveness of the proposed approach. Future work will extend the framework to conditional DAGs on heterogeneous platforms.
References [1] Joshua Bakita and James H. Anderson. 2024. Demystifying NVIDIA GPU Internals to Enable Reliable GPU Management. In 2024 IEEE 30th Real-Time and Embedded Technology and Applications Symposium (RTAS). 294–305. [2] Anne Benoit, Mourad Hakem, and Yves Robert. 2009. Contention awareness and fault-tolerant scheduling for precedence constrained tasks in heterogeneous systems. Parallel Comput. 35, 2 (2009), 83–108. [3] Xiaotian Dai. 2022. dag-gen-rnd: A randomized multi-DAG task generator for scheduling and allocation research. https://doi.org/10.5281/zenodo.6334205 [4] R. L. Graham. 1969. Bounds on Multiprocessing Timing Anomalies. SIAM J. Appl. Math. 17, 2 (mar 1969), 416–429. [5] Alan Gray. 2020. Getting Started with CUDA Graphs. https://developer.nvidia. com/blog/cuda-graphs/. NVIDIA Developer Blog. [6] Mingcong Han, Rong Chen, Weihang Shen, Hanze Zhang, Jinrong Yang, and Haibo Chen. 2025. Real-time, Work-conserving GPU Scheduling for Concurrent DNN Inference. ACM Transactions on Computer Systems (2025). [7] Qingqiang He, Xu Jiang, Nan Guan, and Zhishan Guo. 2019. Intra-Task Priority Assignment in Real-Time Scheduling of DAG Tasks on Multi-Cores. IEEE Transactions on Parallel and Distributed Systems 30, 10 (2019), 2283–2295. [8] Mingzhen Li, Yi Liu, Xiaoyan Liu, Qingxiao Sun, Xin You, Hailong Yang, Zhongzhi Luan, Lin Gan, Guangwen Yang, and Depei Qian. 2020. The deep learning compiler: A comprehensive survey. IEEE Transactions on Parallel and Distributed Systems 32, 3 (2020), 708–727. [9] Zejia Lin, Zewei Mo, Xuanteng Huang, Xianwei Zhang, and Yutong Lu. 2023. KeSCo: Compiler-based Kernel Scheduling for Multi-task GPU Applications. In 2023 IEEE 41st International Conference on Computer Design (ICCD). IEEE, 247–254. [10] Lingxiao Ma, Zhiqiang Xie, Zhi Yang, Jilong Xue, Youshan Miao, Wei Cui, Wenxiang Hu, Fan Yang, Lintao Zhang, and Lidong Zhou. 2020. Rammer: Enabling holistic deep learning compiler optimizations with { rTasks } . In 14th USENIX Symposium on Operating Systems Design and Implementation (OSDI 20). 881–897. [11] Yinchen Ni, Tianrui Ma, Jintao Chen, Chongye Yang, Siwei Ye, Yuankai Xu, Yier Jin, and An Zou. 2025. HARD: Hardening Real-Time Scheduling and Analysis for Accelerator Enabled Computing. In 2025 IEEE 31st Real-Time and Embedded Technology and Applications Symposium (RTAS). 389–401. [12] Nathan Otterness, Ming Yang, Tanya Amert, James Anderson, and F Donelson Smith. 2017. Inferring the scheduling policies of an embedded CUDA GPU. OSPERT’17 (2017).
DAC ’26, July 26-29, 2026, Long Beach, CA, USA
[13] Alberto Parravicini, Arnaud Delamare, Marco Arnaboldi, and Marco D Santambrogio. 2021. Dag-based scheduling with resource sharing for multi-task applications in a polyglot gpu runtime. In 2021 IEEE International Parallel and Distributed Processing Symposium (IPDPS). IEEE, 111–120. [14] Sujan Kumar Saha, Yecheng Xiang, and Hyoseung Kim. 2019. STGM: SpatioTemporal GPU Management for Real-Time Tasks. In 2019 IEEE 25th International Conference on Embedded and Real-Time Computing Systems and Applications (RTCSA). 1–6. [15] H. Topcuoglu, S. Hariri, and Min-You Wu. 2002. Performance-effective and lowcomplexity task scheduling for heterogeneous computing. IEEE Transactions on Parallel and Distributed Systems 13, 3 (2002), 260–274. [16] Samuel Williams, Andrew Waterman, and David Patterson. 2009. Roofline: an insightful visual performance model for multicore architectures. Commun. ACM 52, 4 (2009), 65–76. [17] Yecheng Xiang and Hyoseung Kim. 2019. Pipelined Data-Parallel CPU/GPU Scheduling for Multi-DNN Real-Time Inference. In 2019 IEEE Real-Time Systems Symposium (RTSS). 392–405. [18] Yuankai Xu, Yinchen Ni, Tiancheng He, Ruiqi Sun, Yier Jin, and An Zou. 2025. Real-Time Scheduling and Analysis of Fixed-Priority Tasks on a Basic Heterogeneous Architecture With Multiple CPUs and Many PEs. IEEE Trans. Comput. 74, 8 (2025), 2785–2798. [19] Ming Yang and James H Anderson. 2017. Response-Time Bounds for Concurrent GPU Scheduling. Edited by Patrick Meumeu Yomsi (2017), 13. [20] Deze Zeng, Andong Zhu, Lin Gu, Peng Li, Quan Chen, and Minyi Guo. 2023. Enabling Efficient Spatio-Temporal GPU Sharing for Network Function Virtualization. IEEE Trans. Comput. 72, 10 (2023), 2963–2977. [21] Shulai Zhang, Quan Chen, Weihao Cui, Han Zhao, Chunyu Xue, Zhen Zheng, Wei Lin, and Minyi Guo. 2025. Improving GPU Sharing Performance through Adaptive Bubbleless Spatial-Temporal Sharing. In Proceedings of the Twentieth European Conference on Computer Systems. 573–588. [22] Shuai Zhao, Xiaotian Dai, and Iain Bate. 2022. DAG Scheduling and Analysis on Multi-Core Systems by Modelling Parallelism and Dependency. IEEE Transactions on Parallel and Distributed Systems 33, 12 (2022), 4019–4038. [23] Shuai Zhao, Xiaotian Dai, Iain Bate, Alan Burns, and Wanli Chang. 2020. DAG Scheduling and Analysis on Multiprocessor Systems: Exploitation of Parallelism and Dependency. In 2020 IEEE Real-Time Systems Symposium (RTSS). 128–140. [24] An Zou, Jing Li, Christopher D Gill, and Xuan Zhang. 2023. RTGPU: Real-time GPU scheduling of hard deadline parallel tasks with fine-grain utilization. IEEE Transactions on Parallel and Distributed Systems 34, 5 (2023), 1450–1465.