JOURNAL OF LATEX CLASS FILES, VOL. 14, NO. 8, AUGUST 2021
1
DPIFrame: A Dual-Level Parallelism Acceleration Framework for CTR Model Inference
arXiv:2606.21101v1 [cs.DC] 19 Jun 2026
Dezhi Yi, Huifeng Guo, Kunpeng Xie, Zhaolong Jian, Haochi Yu, Wenxuan He, Zhenhua Dong, Ruiming Tang, Ye Lu
Abstract—Deep learning technology has enhanced the ability of Click-through rate (CTR) prediction models to learn features and improve prediction accuracy. However, it is challenging to deploy CTR models on GPU smoothly and perform inference efficiently, because there is a huge mismatch between the serial computational pattern and the parallel model structure. In this paper, we propose DPIFrame, the first dual parallelizable framework to accelerate CTR model inference. In DPIFrame, a) a dual parallelizable architecture is proposed to perform parallel CTR model inference in both intra-module and inter-module; b) an efficient multi-table lookup algorithm is presented for embedding operations through anticipating the whole workload in advance; c) a breadth-first stream scheduling strategy is designed for finegrained management of parallel computation on GPU to further supporting the dual parallel execution. Extensive experiments are conducted on two real-world datasets, and the results highlight that DPIFrame can reduce the embedding latency efficiently by 23.0× compared to PyTorch. Compared with PyTorch, TorchRec, HugeCTR, and OneFlow, DPIFrame can achieve state-of-the-art inference performance on GPU with speedups of 5.83×, 4.29×, 2.15×, and 2.0×, respectively. Index Terms—Dual-Level Parallelism, Parallel Lookup, MultiStream Scheduling for GPU, CTR Prediction.
I. INTRODUCTION Industrial recommender systems usually involve millions of advertisers, billions of users, and even nearly a trillion dollars in revenues [1]. The most important part in those recommender systems is the Click-through rate (CTR) prediction model, since it can improve the user experience effectively [2] and increase advertising incomes greatly [3]. Modern typical CTR model structure generally includes two key parts, the This work was supported by the National Natural Science Foundation of China (No. 62372253) and the Fundamental Research Funds for the Central Universities (No. 079-63263260). (Corresponding author: Ye Lu.) Dezhi Yi, Kunpeng Xie, Zhaolong Jian, Haochi Yu, and Wenxuan He are with College of Computer Science, Nankai University, Tianjin 300350, China, also with Tianjin Key Laboratory of Network and Data Science Technology, Tianjin 300350, China, and also with Key Laboratory of Data and Intelligent System Security, Ministry of Education, Tianjin 300350, China (e-mail: [email protected]; [email protected]; [email protected]; [email protected]; [email protected]). Huifeng Guo and Zhenhua Dong are with Huawei Technologies Co., Ltd., Shenzhen 518129, China (e-mail: [email protected]; [email protected]). Ruiming Tang is with Kuaishou Technology Co., Ltd., Beijing 100085, China (e-mail: [email protected]). This work was done while Ruiming Tang was with Huawei Technologies Co., Ltd. Ye Lu is with College of Cryptology and Cyber Science, Nankai University, Tianjin 300350, China, also with Tianjin Key Laboratory of Network and Data Science Technology, Tianjin 300350, China, and also with Key Laboratory of Data and Intelligent System Security, Ministry of Education, Tianjin 300350, China (e-mail: [email protected]).
embedding module, which consists of independent sparse operations [4, 5], and the neural network module, which consists of explicit interaction layers and implicit interaction layers [6]. The frequently used operators, such as vector operations in both types of layers, can help calculate features, enhance representation features, and establish correlation relationships effectively. By integrating the embedding module with the neural network module, the CTR model can leverage these benefits to achieve higher inference accuracy [7]. However, as the number of operators and layers inside both the neural network module and embedding module increases significantly, the deep and complex structures of CTR models often result in parameter data explosion by hundreds or even thousands of times. To handle these large amounts of data, the acceleration hardware GPU is widely deployed [8, 9], which can leverage CUDA kernels to enhance the processing of model computation tasks [10]. Nevertheless, the CTR model inference on the GPU needs to satisfy several stringent performance requirements. The first is the stringent accuracy demand [11]. Even a 0.1% error on the CTR model inference accuracy can lead to hundreds of millions of dollars in economic losses in advertising recommender applications [12, 13, 14]. The second is the strict response latency limitation [15]. To process the escalated data daily, recommender systems usually require fast model inference speed, which is typically in the millisecond range [16, 17]. The third is the high hardware computing resource utilization requirement [18]. The operations in the aforementioned two modules are made up of massive short tasks. When CTR model inference, they critically require a reasonable parallel mechanism to improve GPU hardware utilization to avoid wasted investment [19], but even though the most popular PyTorch [20] framework can only achieve under 40% single GPU utilization, as our previous actual profiling. Thus, it is so challenging to take into account the three above aspects at the same time. Existing approaches struggle to balance those three goals. Compression techniques such as quantization [11, 21], pruning [22, 23], and hashing [24, 25] can effectively reduce the computational demands of the CTR model. However, those aggressive compression methods will significantly degrade the accuracy of the model, thus making the model unusable for industrial recommendation systems [26]. Additionally, data parallelism [27, 28, 29], model parallelism [30, 31], and hybrid parallelism [32, 33] techniques are deployed across multiple GPUs to improve the computational efficiency of the CTR model. Their computation scheduling strategies are typically applied at the level of computation graph, module, or layer,
JOURNAL OF LATEX CLASS FILES, VOL. 14, NO. 8, AUGUST 2021
rather than the finer-grained operator level. This limitation, coupled with the communication overhead, frequently leads to insufficient GPU utilization. So, optimizing the efficiency of models at the algorithm level or increasing the number of devices simply cannot achieve the balance goals among the three aspects. To this end, we should turn the perspective on the underlying architecture for the CTR models in the recommender system and overcome the following threefold challenges to meet the aforementioned performance requirements. Firstly, the sparse and irregular access pattern of the embedding module conflicts with the execution model of GPU. Recommendation models exhibit fundamental architectural differences from traditional machine learning models, with over 95% of their parameters concentrated in the embedding module. It requires independent sparse lookup operations across tens to hundreds of feature fields, which are inherently mismatched with the dense, compute-intensive execution model of GPU. As industrial-scale data continues to grow exponentially, these discrete and sparse embedding lookup operations have emerged as the primary performance bottleneck. Secondly, the frequent and fine-grained vector operations in CTR models incur considerable launch overhead and fail to exploit GPU parallelism fully. The explicit interaction module in recommendation models involves largely frequent vector operations to capture complex feature relationships. This computation pattern introduces substantial operator launch overhead and redundant memory accesses, thus limiting the potential for parallel execution with operators in the implicit interaction module. For instance, in our previous profiling, the actual execution time for a single lookup operator is often about 4.75µs, but the overhead of launching the operator is about 21.92µs. So many small computational workloads are too brief and insufficient for GPU to accelerate computation by parallelism. Thirdly, the coarse-grained scheduling mechanisms of existing frameworks prevent effective operator-level parallelism in CTR models. Both the explicit and implicit interaction modules in recommendation models exhibit a parallel characteristic from an architectural perspective, theoretically enabling parallel execution via multiple CUDA streams. However, due to the relatively low computation of individual operators in recommendation models, the existing deep learning frameworks which primarily employ coarse-grained control mechanisms, struggle to effectively parallelize those operations. For example, although the PyTorch framework can provide the advanced APIs to create distinct CUDA streams [34], this kind of scheduling orients the stream-wise rather than the operatorwise. It makes CUDA stream parallel execution still cannot be controlled according to higher-level logic, due to considerations of computational stability and result determinacy by PyTorch [35]. Besides, the scheduling strategy PyTorch exploited follows a depth-first principle by default [36], which can cause the different operators in various streams to miss earlier startup opportunities, thus resulting in the CTR model’s low performance. A reasonable choice is the breadth-first scheduling strategy, which can alternate the launching of operators across different streams, but it is too difficult for
2
PyTorch to precisely schedule operators. To overcome the above difficulties, we propose DPIFrame, which is the first dual-level parallelism acceleration framework for CTR model inference in industrial recommender systems. Aiming to boost inference performance, we redesign and optimize the underlying system architecture and mechanism by addressing the root causes of inefficiency in GPU-based CTR inference. Our contributions are summarized as follows: • A parallel execution architecture that operates both intraand inter-module is proposed to close the gap between serial computational patterns and parallel model structures in CTR inference on GPU. The dual parallelizable architecture innovatively improves CTR inference performance from an underlying system-level perspective. • An efficient and novel multi-table lookup algorithm is presented for embedding operations by anticipating their whole workload in advance. Based on our designed output-first workload allocation, we can ensure address continuity during data read/write. Embedding latency compared with PyTorch can be reduced greatly by an average of 12.9× and a maximum of 23.0×. • An operator-wise stream scheduling strategy at system level with breadth-first principle is designed to further supporting intra- and inter- module parallelism, thus deeply addressing the mismatch between hardware execution and model structure. The alternately overlapping scheduling strategy help DPIFrame achieve operatorlevel parallel execution, the maximum GPU utilization improvement can be of 36.4% and 60.6% on Avazu and Criteo datasets, respectively. • DPIFrame is implemented as a building block and is embedded seamlessly into PyTorch. Compared to original PyTorch, TorchRec, HugeCTR, and OneFlow, DPIFrame can achieve end2end inference speedup by 5.83×, 4.29×, 2.15×, and 2.0×, respectively. II. BACKGROUND In this section, we first introduce the fundamental concepts related to recommender systems and CTR model. For clarity, the typical DCN [13] is shown in Figure 1 as an example to illustrate the preliminaries and CTR model characteristics. We then introduce and analyze CUDA stream mechanism. A. Recommender System Recommender systems have become foundational in modern digital platforms, including e-commerce, online advertising, content streaming, and social networks. Their primary goal is to predict user preferences and provide personalized content, thereby improving user experience and increasing business revenue. A core task in these systems is CTR prediction, which estimates the probability that a user will interact with a recommended item, such as clicking on a product or watching a video. Accurate CTR prediction is crucial, especially in scenarios where only a limited number of items can be recommended within strict time constraints. To meet the demands of large-scale and dynamic recommendation environments, Deep Learning Recommendation Models
JOURNAL OF LATEX CLASS FILES, VOL. 14, NO. 8, AUGUST 2021
Input
Feature 1 encode
Embedding module
…
Feature 2 encode
[0,1,…,0]
[0,0,…,1]
Lookup
Lookup c
[0.5,2.4,…,1.8] [3.6,1.2,…,1.6] … [4.7,5.1,…,0.8]
[5.4,6.7,…,3.5] [9.2,4.6,…,8.3] … [8.1,4.5,…,4.8]
[3.6,1.2,…,1.6]
[8.1,4.5,…,4.8]
…
… …
3
Feature n encode [1,0,…,0]
One-hot vector
Lookup
Serial execution
[7.6,5.4,…,1.7] [3.8,9.6,…,3.5] … [4.6,6.2,…,7.3]
Embedding tables
[7.6,5.4,…,1.7]
Real-valued Vector
Concatenate [3.6,1.2,…,1.6, 8.1,4.5,…,4.8, …,7.6,5.4,…,1.7]
Neural network module
GEMM
ReLU
Deep network module
GEMM
Ele-Mul
Ele-Add
Cross network module
Serial execution
GEMM & Sigmoid
Output
probability of clicking
Fig. 1. Take DCN [13] as an example of CTR model inference. General Matrix Multiplication is denoted as GEMM.
(DLRMs) have become the industry standard. As shown in Figure 1, a typical DLRM consists of two main components: the embedding module and the neural network module. The embedding module handles high-dimensional categorical features by projecting them into dense representations using large embedding tables. These embeddings are concatenated and fed into the neural network module for feature interaction and final CTR prediction. The inference phase of recommender systems must prioritize low latency, as even a few milliseconds of delay can significantly impact user engagement and system throughput. Optimizing inference is therefore critical to enable ranking more candidate items within a fixed time window, leading to better recommendations and higher click-through rates. Beyond latency, the efficient utilization of computational resources is equally important. Embedding module often involves frequent and irregular memory accesses, which place substantial pressure on memory bandwidth. Simultaneously, the neural network module is computationally intensive and demands high parallel efficiency from modern hardware such as the GPU. Inefficient resource usage can lead to bottlenecks in production systems, increase infrastructure costs, and limit the scalability of the recommender service. Overall, a highquality recommender system must achieve strong performance in three key aspects: prediction accuracy, response latency, and efficient use of computational resources. B. CTR Model Inference For the sake of consistency, three terms are first clarified: model, module, and operator. The important modules, key processes, and structure are also illustrated. • Model The complete model consisting of several modules, e.g., DCN [13], DCNv2 [14]. • Module The component including a certain number of operators in the model, e.g., Embedding, Cross Network.
The module is the traditional smallest unit that can be executed in parallel by PyTorch on GPU. • Operator The function that performs computational tasks on the device, e.g., embedding lookup(), element-wise add(). The small computational operators can be fused into several single operators to be executed. With these concepts, the typical inference process of the CTR model can be described in three main steps, as shown in Figure 1. First, the original sparse data is input into the embedding module for multiple table lookup operations, resulting in a dense vector of real values. Next, the concatenated dense vector is input to the neural network module, which performs more complex computations, such as matrix multiplication. Finally, the computed results are fed into a prediction layer to generate predictions. The roles of the two core modules are as follows: Embedding Module transforms sparse high-dimensional one-hot vectors of categorical features into dense lowdimensional real-valued vectors. Each feature field has a unique embedding table E ∈ Rn×d , where n is the number of features and d is the embedding dimension. After performing a table lookup operation for each feature field separately, their vector representations are obtained. Then these embedding vectors of different feature fields are concatenated as input into neural network module. Neural Network Module typically consists of two parallel modules: explicit interaction module (e.g., cross network in Figure 1) and implicit interaction module (e.g., deep network in Figure 1). The two modules share the same input and have no data dependencies on each other and are computationally independent of each other. The cross network module learns explicit interaction features, and the deep neural network learns implicit interaction features. Finally, the outputs of these two modules are concatenated into a prediction layer to make predictions. By analyzing the inference process, we can observe that the CTR model inherently possesses the parallel structure. Firstly, all lookup operators in embedding module are independent of each other. They are each associated with their respective feature fields and do not query the same embedding table. Secondly, the computations in the explicit interaction module and the implicit interaction module are mutually independent. They are executed without inter-module dependencies and can be performed in parallel without requiring intermediate results from each other. This independence implies that these operations can be executed concurrently, providing the foundation for our parallel speedup. C. CUDA Stream Scheduling The stream is a core abstraction in the NVIDIA CUDA programming model, used to achieve asynchronous concurrent execution and task scheduling. A stream refers to a sequence of operations delivered to a device for execution [37]. These operations include data transmission and computational tasks, also referred to as the aforementioned operators. The operators are fed into the stream. When resources are available, the GPU executes them using the CUDA kernel from the stream. As
JOURNAL OF LATEX CLASS FILES, VOL. 14, NO. 8, AUGUST 2021
Single Stream
CPU
Stream 0 (Default Stream)
Kernel 1
Stream 1
Kernel 1
Multiple Streams
Stream 2
Kernel 2
Kernel 2
timeline
saved timeline
Fig. 2. An example of CUDA streams.
shown in Figure 2, a single CUDA stream can contain multiple operations, which must be launched strictly in the order that they are added to the stream. Once sufficient resources are available, operations from multiple streams can be scheduled for concurrent execution [37], with the default scheduling following the depth-first strategy. Existing high-level frameworks, such as PyTorch, typically use only the default stream (“stream 0”) for CTR model inference, meaning the operators are executed on the GPU in serial. Even if developers explicitly specify a stream in PyTorch, the execution may still not follow the developers’ instructions precisely. The limitation in CUDA stream scheduling, especially the reliance on a single default stream and the coarse-grained scheduling control in high-level frameworks, results in suboptimal parallelism and inefficiencies in CTR model inference on GPU. III. MOTIVATION Modern CTR prediction models have evolved toward increasingly complex architectures, yet their inference efficiency on GPU remains unsatisfactory. Through detailed analysis, we find that the root cause lies in the mismatch between the intrinsic characteristics of CTR models and the execution paradigm of GPU systems. In the following, we identify four key challenges and corresponding opportunities that motivate the design of DPIFrame. Challenge 1: Structural parallelism vs. serialized execution. CTR models inherently exhibit a high degree of parallelism at both the operator and module levels. For example, lookup operations in the embedding module are independent across feature fields, and the explicit and implicit interaction modules in the neural network component can be executed without interdependencies. However, existing deep learning frameworks typically rely on coarse-grained execution models and default CUDA stream, resulting in largely serialized operator execution. This prevents the model from leveraging its intrinsic parallel structure and leads to low GPU utilization. Opportunity 1: This observation suggests that the inherent independence across operators and modules can be explicitly exploited to enable concurrent execution. By exposing and coordinating both intra-module and inter-module parallelism, it is possible to align the execution model with the parallel nature of CTR architectures and significantly improve hardware efficiency. Challenge 2: Irregular memory access vs. hardware efficiency. The embedding module introduces sparse and irregular memory access patterns due to large-scale table lookups. These accesses lack spatial and temporal locality, which conflicts with the design of GPU memory hierarchies and leads to
4
Depth-first Scheduling
CPU Stream 1
Kernel 1
Stream 2
Kernel 2 Kernel 3
Kernel 4
CPU
Breadth-first Scheduling Stream 1 (Ours) Stream 2
Kernel 1
Kernel 2
Kernel 3
Kernel 4
timeline
saved
timeline
Fig. 3. Depth-first and breadth-first scheduling. Depth-first scheduling prioritizes launching all operators within a CUDA stream before initiating operators in another CUDA stream. Breadth-first scheduling entails interleaving the launch of operators from different CUDA streams.
inefficient bandwidth utilization. Consequently, embedding operations often become the dominant performance bottleneck in CTR inference. Opportunity 2: Despite the irregularity of individual accesses, the overall embedding workload is deterministic and known in advance during inference. This enables a global view of memory access patterns and provides an opportunity to reorganize embedding operations into a structured computation process. By transforming scattered lookups into coordinated, sequentialized memory accesses, we can significantly improve memory efficiency and alleviate bandwidth bottlenecks. Challenge 3: Fine-grained operators vs. GPU execution efficiency. CTR model inference is dominated by numerous lightweight operators such as element-wise vector computations. These operators typically have very short execution times, while the overhead of launching CUDA kernels is comparatively high. As a result, a significant portion of execution time is spent on kernel invocation rather than actual computation, leading to poor overall efficiency. Moreover, the low computational load of these operators makes it difficult to directly parallelize them to effectively overlap execution latency. Opportunity 3: To address this inefficiency, it is critical to increase computational granularity. By fusing multiple small operators and reorganizing workloads into larger computational units, we can reduce kernel launch overhead, improve data locality, and better match GPU execution characteristics. Such transformations enable more efficient utilization of GPU compute resources and create sufficiently large workloads that are amenable to parallel execution and latency hiding. Challenge 4: Coarse-grained scheduling vs. fine-grained parallelism. Existing frameworks typically employ coarsegrained scheduling strategies at the module level, often following depth-first execution orders. Such strategies fail to fully exploit the fine-grained parallelism present in CTR models and delay the execution of independent operators that could otherwise run concurrently. As a result, many GPU resources remain underutilized due to insufficient overlap between operations. Opportunity 4: This limitation highlights the need for a more fine-grained scheduling mechanism. By introducing operator-level scheduling across multiple CUDA streams and carefully coordinating their execution order, it is possible to maximize concurrency and overlap computation effectively. In particular, as shown in Figure 3, adopting a breadth-first scheduling strategy allows independent operators
JOURNAL OF LATEX CLASS FILES, VOL. 14, NO. 8, AUGUST 2021
Multi-table Lookup Module
Lookup Algorithm
Intra-Module Parallel
5
Workload Allocation
Output-first strategy
Parallel
b
Input
k
....
Table 1
? ? ?
Outputs Table N
? ? ? ?
Neural Network ×3 GE MM
ReLU
Operators
GE MM×3 ReLU×3
GE MM
GE MM×2
nonGE MM
FusedOP
Stream 1 Stream 2
GEMM
Output Kernel 1 Kernel 2 Kernel 3 Kernel 4
Sigmoid
d
b k×d
timeline
GE MM
Prediction Sigmoid
concatenate
Breadth-first
FusedOP
nonGE MM
b
Stream Scheduler
Task CUDA
n
Tables
Vectors
Inter-Module Parallel
×2
d
look up
Other Modules CTR Model Embedding ... Lookup Lookup
b: samples number k: feature fields number n: feature number d: feature dim
GPU Block
... ...
Fig. 5. An example of embedding table lookups. Where the number of samples is 2, the number of feature fields is 4, and the number of features per feature field is 3.
Fig. 4. The overview of DPIFrame.
B. Multi-table Lookup Algorithm from different modules to be launched earlier, thereby improving overall execution efficiency and GPU utilization. In summary, the inefficiency of CTR model inference arises from the mismatch between fine-grained, parallel model structures and coarse-grained execution mechanisms. These challenges motivate a systematic redesign that jointly considers parallelism extraction, workload restructuring, and execution scheduling. Based on these insights, we propose DPIFrame, which leverages dual-level parallelism and optimized scheduling to bridge this gap and achieve efficient CTR model inference on GPU. IV. DPIFRAME DESIGN As shown in Figure 4, DPIFrame, as the first dual parallel execution inference system for CTR prediction, supports both intra-module and inter-module parallel execution. In this section, we first present the overall architecture of DPIFrame. We then outline the core parallel design following the CTR model inference process. In the embedding stage, we design the multi-table lookup algorithm and focus on workload allocation. In the neural network stage, we describe the fusion of non-GEMM operators and the corresponding stream scheduling strategy. A. Dual Parallel Architecture As aforementioned, we identify the critical limitation of the CTR model inference inefficiency as the mismatch between the serial computational pattern and the parallel model structure. Furthermore, we capture that the computational load of the operators in the CTR model is short tasks, making it difficult to directly parallelize them to overlap the latency. Therefore, we propose a dual parallelizable architecture of the CTR model inference system that performs parallel CTR model inference both intra-module and inter-module. We also process the nonGEMM operators in a module by merging them into a few ones, then their workloads are intensively allocated to threads, thus enabling our multi-table parallel lookup algorithm and achieving intra-module parallelism. Besides, the operators of parallel modules are alternately scheduled to different CUDA streams to maximize overlap, thereby boosting the intermodule parallelism.
In the embedding stage, we design the multi-table lookup algorithm that leverages the independence of lookup operations to unleash the parallel potential in the model structure. Suppose the input X of the CTR model is of size b × k, where b represents the number of samples and k represents the number of feature fields. Each feature field has a separate embedding vector table Ei ∈ Rni ×d , where i represents i-th feature field, ni represents the number of features contained in the i-th feature field. For example, the ”gender” feature field contains male and female features, so the n of this feature field is 2. Some feature fields have a huge number of features, which can reach millions or even tens of millions, so their corresponding embedding vector table is also very large. d represents the dimension of the embedding table, that’s the length of the vector used to represent each feature. For the input samples, the k feature fields are first encoded with onehot encoding, respectively. To reduce sparsity and memory consumption, they are often further represented as numeric IDs. An example of embedding operations is shown in Figure 5, which looks up its corresponding embedding table for each of these k feature fields. After a total of k lookup operations to obtain their vector representations. Then, these vectors are concatenated into a matrix Xembed ∈ Rb×(k×d) as the input to the next layer of the model. All the embedding lookups are as a whole, so it is entirely possible to replace the original k serial lookups with a single parallel operation. We can compute the size of the embedded output based on the embedding configuration parameters as well as the size of the input. Therefore, we only need to initialize the output matrix and populate it element by element, without any additional intermediate storage vectors. To match each position of the matrix with the corresponding position of its embedding table, we design Algorithm 1 to accomplish the global embedding table lookup. Firstly, we compute the size of the output matrix as well as the number of its columns based on the size of the input and the embedding configuration parameters (line 1-2). Then start traversing the output matrix. Next, iterate through the matrix row by row (line 3), based on the step, derive the position of the current element in terms of the row and column to which it belongs in the matrix (line 4-5). Further, we can compute to know which
JOURNAL OF LATEX CLASS FILES, VOL. 14, NO. 8, AUGUST 2021
6
Algorithm 1: Multi-table Lookup Input: IDs ∈ Nb×k : Input feature IDs Embs: Embedding tables, where the i-th table is of size ni × d Output: EmbedOut ∈ Rb×(k×d) : Lookup results total elements ← b × k × d row width ← k × d 3 for idx ← 0 to total elements − 1 do idx 4 row ← row width 5 col ← idx mod row width col 6 table id ← d 7 emb row ← IDs[row × k + table id] 8 emb col ← col mod d 9 EmbedOut[idx] ← Embs[table id][emb row × d + emb col] 10 end 11 return EmbedOut
Threads
Table
Output
Input
Threads
Table
Output
Input-first
1
2
Input
Output-first (Ours)
Fig. 6. Output-first workload allocation’s advantages on fused embedding lookup operation.
addressed. This ensures optimal utilization of the hardware memory bandwidth and brings performance improvements. D. Fusion Mechanism for non-GEMM Operators
embedding table to look up (line 6). Finally, leveraging the gathered information along with the input IDs, we determine the specific row and column in the embedding table (lines 78), allowing us to find the value to be filled in the currently traversed output matrix element (line 9). The optimization design consolidates multiple table lookups across different feature fields into a single operation. The algorithm only requires very few computations to derive the indices. To face multi-hot encoded features, it can also support sequential features by marking offset information. C. Workload Allocation Once the total workload of fused operators is determined, the focus shifts to allocating the workload across threads, aiming at ensuring that it is GPU-friendly and further enhancing intra-module parallel execution. We design the output-first workload allocation methods, and take the fused embedding lookup operator as an example to explain its advantages, as shown in Figure 6. As a comparison, we also introduce the drawbacks of the input-first approach. The input-first scheme involves maximizing the allocation of different input samples among threads, while our output-first scheme involves coordinating threads to process the same sample together. According to the NVIDIA GPU programming model [38], the smallest hardware execution unit on a GPU is a warp, which typically consists of 32 threads. When the 32 threads execute the same instruction, they operate in a Single Instruction, Multiple Threads (SIMT) fashion, allowing simultaneous reads and writes. When adopting an input-first scheme where different threads prioritize different samples, even though the input reads have a sequentially addressed pattern, the table lookups and output writes will involve random reads and writes instead of sequentially addressed ones. This is hardware-unfriendly because the memory bandwidth is not fully utilized, resulting in wasted resources and low performance. Our output-first approach enables different threads to share the same sample and can perform embedding table lookups feature by feature. As a result, both input reads and table lookups, as well as output writes, remain sequentially
Different from embedding, neural network includes more varied operators. There are two typical modules in neural network: an implicit feature interaction module and an explicit feature interaction module. The implicit interaction module remains relatively consistent and typically consists of several layers of multi-layer perceptrons (MLPs), which contain GEMM and activation operators. But due to different design objectives, the explicit interaction modules of different deep CTR models often differ. For example, the explicit interaction module of DCNv2 [14] is a cross network, which contains element-wise addition and element-wise multiplication operators besides GEMM. The explicit interaction module of DeepFM [12] consists of a logistic regression and a factorization machine, which contain various operators such as table lookup, element-wise square and reduce-sum. To address the variety of neural network among CTR models, we propose a uniform fusion mechanism: we combine all neighbouring non-GEMM operators into a single operator. DPIFrame represents the model forward propagation by constructing a directed acyclic graph, in which nodes are operators and edges are tensors. Starting from the root node, we traverse the graph to mark all non-GEMM nodes connected by edges. We define the set of these nodes and the edges connecting them as a subgraph. Within a subgraph, we fuse the operators into a new operator. Moreover, the workload allocation follows a similar principle as described in Section IV-C. We traverse the output matrix element by element, assigning the filling tasks for each matrix element to individual threads, ensuring that data reads and writes are contiguous in memory. E. Stream Scheduling Strategy Beyond intra-module parallelism, when deploying models on the GPU, the opportunity for parallel execution exists for operators that are placed into different CUDA streams. However, existing deep learning frameworks typically use only one default stream to schedule operators. Operators are executed sequentially according to the scheduling order, even for models with parallel structures. To achieve inter-module parallel execution and further address the mismatch between hardware execution and model
JOURNAL OF LATEX CLASS FILES, VOL. 14, NO. 8, AUGUST 2021
Algorithm 2: Breadth-first Stream Scheduling Input: Explicit: Explicit Interaction Module Implicit: Implicit Interaction Module Output: Sexplicit : CUDA stream for Explicit Simplicit : CUDA stream for Implicit Q: Breadth-first scheduling queue Opsexplicit ← Explicit.operators() Opsimplicit ← Implicit.operators() 3 nexplicit ← |Opsexplicit | 4 nimplicit ← |Opsimplicit | 5 Sexplicit ← cudaStreamCreate() 6 Simplicit ← cudaStreamCreate() 7 Sexplicit .Add(Opsexplicit ) 8 Simplicit .Add(Opsimplicit )
1 2
if nimplicit > nexplicit then for i ← 0 to nexplicit − 1 do 11 Q.Push(Opsimplicit [i]) 12 Q.Push(Opsexplicit [i]) 13 end 14 for j ← nexplicit to nimplicit − 1 do 15 Q.Push(Opsimplicit [j]) 16 end 17 end 18 else 19 for i ← 0 to nimplicit − 1 do 20 Q.Push(Opsexplicit [i]) 21 Q.Push(Opsimplicit [i]) 22 end 23 for j ← nimplicit to nexplicit − 1 do 24 Q.Push(Opsexplicit [j]) 25 end 26 end 9
10
27
return Sexplicit , Simplicit , Q
structure, we design an operator-level stream scheduling strategy to fine-tune the parallel computation of the key modules in CTR models. The core idea of the stream scheduling strategy is to assign the operators of key modules to separate CUDA streams using a breadth-first scheduling approach. The operators from different modules are added to different CUDA streams. The operators are then pushed into the launch queue in an alternating pattern by module wise. This can ensure operators are executed on GPU as early as possible to avoid blocking or waiting. The module that has more operators launches first in our DPIFrame, it can help hide the startup costs of many small computation tasks. The specific design details are described in Algorithm 2. V. EVALUATION In this section, we evaluate DPIFrame through a set of comprehensive experiments. We introduce the experimental setup (V-A), to verify that DPIFrame meets the requirements for accuracy (V-B) and efficiency (V-C), analyze each key design component of DPIFrame improvement(V-D), evaluate the hardware utilization improvements(V-E), We also analyze the embedding sensitivity, the advantages of the designed steam scheduling strategy, and point out an interesting phenomenon observed during multi-stream scheduling, aiming to inspire further research. Since embedding is an essential module of all CTR models, we focus on embedding sensitivity studies to ensure that DPIFrame delivers consistent performance gains under various complex configurations (V-F). We then evaluate the advantages of the designed stream scheduling strategy
7
compared to single-stream and general multi-stream scheduling (V-G). Finally, we test different startup sequences for the configuration, aiming to achieve a comprehensive assessment (V-H). A. Experimental Setup Hardware. We deploy DPIFrame on a server equipped with an Intel Core i9-10980XE CPU (3.0 GHz, 18 Cores) and 256GB DRAM in total, and a single NVIDIA RTX 3090 GPU with 24GB DRAM. Software. The original implementation of CTR models is based on an open-source project, FuxiCTR 1 . We implemented DPIFrame as a building block and embedded it into PyTorch (v1.13.0) using pybind11. The environment we used includes CUDA 11.6, and the version of gcc is 9.4.0. Datasets. To thoroughly evaluate the accuracy and efficiency of DPIFrame, while ensuring its effectiveness in realworld scenarios, we conducted experiments on two real-world benchmark datasets released by internationally renowned advertising companies: Avazu 2 and Criteo 3 . They contain 24 and 39 feature fields, respectively, with significant variations in the number of features within each field, reflecting real-world scenarios. Models. We validate DPIFrame using four industryrecognized CTR models: DCN [13], DCNv2 [14], Wide&Deep [39], and DeepFM [12]. These models have been validated in real application scenarios, and their diverse structures enable a comprehensive evaluation of DPIFrame designs. Configuration parameter. To validate that DPIFrame provides effective acceleration in different scenarios, our experiments used various model configuration parameters. Embedding dimension: 16 and 32. Hidden size units of deep network module: [256, 256, 256], [512, 512, 512], and [1024, 1024, 1024]. Batch size is typically set to 2048. For DCN [13] and DCNv2 [14], the number of cross layers is 3. In summary, each model is validated under 6 sets of configuration parameters. Metrics. We evaluate DPIFrame from the following metrics: (1) AUC and LogLoss. These are two classic metrics for evaluating the accuracy of CTR models [12]. The premise of DPIFrame acceleration is to rigorously ensure the accuracy of the model. (2) Speedup and latency. The main goal of DPIFrame is to accelerate the inference of CTR models. (3) Hardware utilization. DPIFrame pursues higher hardware utilization to achieve fast model inference. Baselines. We compare DPIFrame with four strong frameworks: PyTorch [20], TorchRec [40], HugeCTR [41], and OneFlow [42]. TorchRec is a PyTorch-based domain library for recommendation systems. It enables researchers to build stateof-the-art personalized models and deploy them in production environments. TorchRec includes optimized recommender system kernels driven by FBGEMM [43], which enables highspeed GPU inference. Currently, TorchRec is used in Meta’s 1 https://github.com/reczoo/FuxiCTR 2 https://www.kaggle.com/c/avazu-ctr-prediction 3 https://www.kaggle.com/c/criteo-display-ad-challenge
JOURNAL OF LATEX CLASS FILES, VOL. 14, NO. 8, AUGUST 2021
6
PyTorch
Speedup (x)
5
TorchRec
HugeCTR
OneFlow
8
DPIFrame
4 3 2 1 0
16_256
16_512 16_1024 32_256
32_512 32_1024 16_256
16_512 16_1024 32_256
DCN
32_512 32_1024 16_256
16_512 16_1024 32_256
DCNv2
32_512 32_1024 16_256
Wide&Deep
16_512 16_1024 32_256
32_512 32_1024
DeepFM
(a) Speedup on the Avazu Dataset 7
PyTorch
Speedup (x)
6
TorchRec
HugeCTR
OneFlow
DPIFrame
5 4 3 2 1 0
16_256
16_512 16_1024 32_256
32_512 32_1024 16_256
16_512 16_1024 32_256
DCN
32_512 32_1024 16_256
16_512 16_1024 32_256
DCNv2
32_512 32_1024 16_256
Wide&Deep
16_512 16_1024 32_256
32_512 32_1024
DeepFM
(b) Speedup on the Criteo Dataset Fig. 7. The overall speedup comparison and DPIFrame on four models with different computational scales.
TABLE I P REDICTION ACCURACY OF DPIF RAME . DCN DCNv2 Wide&Deep DeepFM
AUC 0.7922 0.7935 0.7925 0.7928
Avazu LogLoss 0.3725 0.3718 0.3723 0.3722
AUC 0.8094 0.8113 0.8089 0.8061
Criteo LogLoss 0.4422 0.4405 0.4432 0.4461
recommendation system platform. HugeCTR is an opensource GPU-accelerated integration framework developed by NVIDIA. It is specifically designed to optimize both the training and inference processes for CTR models. The OneFlow team has introduced an efficient, scalable, and highly flexible recommendation system component called OneEmbedding 4 . Its performance exceeds that of typical frameworks and even outperforms the dedicated recommendation framework developed by NVIDIA, HugeCTR. B. Prediction Accuracy We train CTR models on two datasets using PyTorch and evaluate the accuracy of these models on test sets when performing inference using both PyTorch and DPIFrame. The purpose of this evaluation is to verify that DPIFramestrictly maintains model accuracy. The specific performance is shown in Table I. The experimental results show that for both the AUC and LogLoss metrics, DPIFrame and PyTorch achieve identical values to the fourth decimal place (in fact, they are identical to the sixth decimal place). This shows that the speedup achieved by DPIFrame can be trusted. C. End-to-End Speedup We evaluate the end-to-end speedup of DPIFrame with PyTorch, TorchRec, HugeCTR, and OneFlow on four CTR models with six sets of configuration parameters on two datasets, 4 https://docs.oneflow.org/en/master/cookies/one embedding.html
TABLE II S PEEDUP OF DPIF RAME COMPARED TO BASELINES ON FOUR MODELS . DCN DCNv2 Wide&Deep DeepFM
PyTorch Avg. Max. 3.14× 4.08× 2.65× 3.43× 4.32× 5.83× 4.08× 5.82×
DPIFrame VS. TorchRec HugeCTR Avg. Max. Avg. Max. 2.41× 2.77× 1.62× 1.88× 1.94× 2.41× 1.73× 1.95× 3.03× 3.94× 1.72× 2.12× 3.13× 4.29× 1.71× 2.15×
OneFlow Avg. Max. 1.51× 1.71× 1.46× 1.71× 1.60× 2.00× 1.54× 1.96×
as shown in Figure 7. X Y means embedding dimension and hidden layer size of MLP separately. The experimental results show that DPIFrame comprehensively outperforms the comparative baselines. Framework-wise comparison. The results show that DPIFrame achieves average speedups of 3.55×, 2.63×, 1.69×, and 1.53× compared to PyTorch, TorchRec, HugeCTR, and OneFlow, with maximum speedups of 5.83×, 4.29×, 2.15×, and 2.0×, respectively. DPIFrame can significantly improve the performance of inference with the help of intra-module and inter-module parallel execution. Dataset-wise comparison. As shown in Figure 7(a), on the Avazu dataset, DPIFrame achieves average speedups of 3.32×, 2.59×, 1.61×, and 1.60× compared to PyTorch, TorchRec, HugeCTR, and OneFlow, with maximum speedups of 4.95×, 3.89×, 1.88×, and 2.0×, respectively. Figure 7(b) shows on the Criteo dataset, DPIFrame achieves average speedups of 3.77×, 2.67×, 1.78×, and 1.46× compared to PyTorch, TorchRec, HugeCTR, and OneFlow, with maximum speedups of 5.83×, 4.29×, 2.15×, and 1.77×, respectively. The acceleration performance achieved on these two datasets is similar, with slightly better performance on the Criteo dataset compared to the Avazu dataset. This observation underscores DPIFrame’s ability to deliver excellent performance even in large data scenarios. Model-wise comparison. The results of the speedup for the different models are shown in Table II. Overall, the acceleration achieved for Wide&Deep and DeepFM is slightly
Speedup (x)
JOURNAL OF LATEX CLASS FILES, VOL. 14, NO. 8, AUGUST 2021
7 6 5 4 3 2 1 0
PyTorch-A
16_256
16_512 16_1024 32_256 DCN
PyTorch-B
DPIFrame-A
32_512 32_1024 16_256
DPIFrame-B
9
DPIFrame-C
16_512 16_1024 32_256
32_512 32_1024 16_256
16_512 16_1024 32_256
DCNv2
32_512 32_1024 16_256
16_512 16_1024 32_256
Wide&Deep
32_512 32_1024
DeepFM