# ViDiT-Q: Efficient and Accurate Quantization of Diffusion Transformers for
Image and Video Generation

###### Abstract

Diffusion transformers have exhibited remarkable performance in visual
generation tasks, such as generating realistic images or videos based on
textual instructions. However, larger model sizes and multi-frame processing
for video generation lead to increased computational and memory costs, posing
challenges for practical deployment on edge devices. Post-Training
Quantization (PTQ) is an effective method for reducing memory costs and
computational complexity. When quantizing diffusion transformers, we find that
applying existing diffusion quantization methods designed for U-Net faces
challenges in preserving quality. After analyzing the major challenges for
quantizing diffusion transformers, we design an improved quantization scheme:
ViDiT-Q (Video & Image Diffusion Transformer Quantization) to address these
issues. Furthermore, we identify highly sensitive layers and timesteps hinder
quantization for lower bit-widths. To tackle this, we improve ViDiT-Q with a
novel metric-decoupled mixed-precision quantization method (ViDiT-Q-MP). We
validate the effectiveness of ViDiT-Q across a variety of text-to-image and
video models. While baseline quantization methods fail at W8A8 and produce
unreadable content at W4A8, ViDiT-Q achieves lossless W8A8 quantization.
ViDiT-Q-MP achieves W4A8 with negligible visual quality degradation, resulting
in a 2.5x memory optimization and a 1.5x latency speedup. The code can be
found in: https://github.com/A-suozhang/ViDiT-Q.

Figure 1: We introduce ViDiT-Q, a quantization method specialized for
diffusion transformers used in image and video generation. ViDiT-Q achieves
lossless W8A8 quantization and minimal visual quality degradation at W4A8,
gaining 2.5x model size reduction and a 1.5x latency speedup.222All videos in
the Figures are provided in : https://a-suozhang.xyz/viditq.github.io/. Figure
2: The Challenges for existing diffusion quantization methods, and ViDiT-Q’s
solutions. ViDiT-Q introduced improved quantization scheme tailored for DiT to
achieve lossless W8A8, and metric decoupled mixed precision tailored for video
generation to mitigate degradation for W4A8.

##  1 Introduction

Diffusion Transformers [1, 44] and video generation tasks [53, 37, 22] have
garnered significant research interest since the impressive performance of
SORA [41]. However, the increasing model size poses challenges for application
and deployment on edge devices. In the realm of video generation, processing
multiple frames imposes a significant burden on both memory and computation.
For example, the OpenSORA [22] model consumes over 10 GB of GPU memory to
generate a single 512$\times$512 resolution video with only 16 frames, taking
about 50 seconds on an Nvidia A100 GPU.

Model quantization [25, 38] has proven to be an effective compression method.
By compressing high bit-width floating-point (FP) data into lower bit-width
integers, the computational and memory costs of the model can be effectively
reduced. While many prior studies [27, 16] have explored quantization for CNN
(U-Net)-based diffusion models [49], we empirically observe challenges when
quantizing Diffusion Transformers (DiTs) for image and video generation.
Specifically, as illustrated in Fig. 2, the W8A8 quantized open-sora (STDiT
[22]) model experiences notable quality degradation (the turtle has multiple
fins), and the W4A16 quantized model produces single-colored blank frames.

By observing the data distribution (See Fig. 4), we conclude the primary issue
with DiT quantization is adopting the same set of fixed quantization
parameters for highly variant values in 4 levels (See Sec. 3.1): (1) Input-
channel level, (2) Token level, (3) Timestep level, and (4) Classifier-free-
guidance level. Correspondingly, we design an improved quantization scheme:
“ViDiT-Q”, tailored for DiTs, to address this challenge. This scheme includes
timestep-aware channel balancing, token-wise and dynamic quantization. ViDiT-Q
enables W8A8 quantization without visual degradation.

Furthermore, in Fig. 2, when we continue to explore lower-bit quantization
with uniform bit-widths (W6A6 and W4A8), notable degradation is witnessed.
Through investigation, we discover that the quantization sensitivity varies
significantly across layers and timesteps. Certain layers exhibit high
sensitivity, acting as “bottlenecks” for quantization. To address this issue,
inspired by prior mixed precision research [42, 62, 10, 12], we adopt higher
bit-widths to “protect” these sensitive layers. When assigning proper bit-
width for each layer, we find that simple Mean Square Error with FP-generated
videos [68] cannot adequately distinguish quantization’s effects on various
aspects of the video, such as quality [59, 64], text-video alignment [17], and
temporal consistency [32]. To overcome this, we propose to “decouple” these
aspects by utilizing corresponding evaluation metrics, and introduce “metric-
decoupled” quantization sensitivity estimation method to assist in mixed
precision bit-width allocation. We summarize our contributions as follows:

  1. 1.

We conduct extensive analysis and identify challenges for DiT quantization and
design a specialized quantization scheme, “ViDiT-Q”, to address these issues.

  2. 2.

We identify the quantization “bottlenecked” in lower bit-widths and propose a
metric-decoupled mixed-precision method to mitigate this problem.

  3. 3.

We validate the effectiveness of “ViDiT-Q” on extensive DiT models for both
image and video generation. It achieves W8A8 with negligible metrics loss, and
“ViDiT-Q-MP” achieves W4A8 with minimal visual degradation, resulting in
2-2.5x memory saving and 1.5x latency speedup.

##  2 Related Works

###  2.1 Diffusion Transformers for Image and Video Generation

Diffusion Transformers (DiTs), which employ Transformers [60] to replace the
CNN-based diffusion backbones (U-Net [50]) in prior research [49, 45], have
achieved remarkable performance in visual generation. Image Generation: DiT
[44] and UViT [1] pioneer the use of transformers as diffusion backbones.
Subsequent research [14, 36, 71] makes further advancements in architecture
design and training methods. PixArt-$\alpha$ [6] explores text-to-image
generation with DiTs, PixArt-$\Sigma$ [5] extends generation resolution to 4K.
Video Generation: Early video generation models [21, 20, 13] mainly adopted
CNN backbones. Latte [37] pioneer the use of transformers for text-to-video
generation. The success of SORA [41] inspire the development of video
diffusion transformers such as GenTron [7], and open-sora STDiT [22]. Both
high-resolution image generation and multi-frame video generation add to
hardware costs, necessitating efficiency improvements.

###  2.2 Image and Video Generation Evaluation Metrics

Visual generation can be evaluated from multiple aspects, and many metrics are
introduced accordingly. Image Metrics: FID [18] and IS [51] are two commonly
adopted metrics for measuring the Inception network feature difference between
generated and reference images for quality and fidelity assessment. ClipScore
[17] evaluates how well the generated image follows the prompt instruction
(text-image alignment), while ImageReward [67], HPS [65] incorporates human
preference by collecting actual user data to train the reward model. Video
Metrics: FVD extends the feature-based metric FID to the video domain. CLIPSIM
[63] estimates the similarity between video and text instructions. CLIP-temp
[11] measures the semantic similarity between video frames. Flow-score is
proposed as part of the video evaluation benchmark EvalCrafter [33] to assess
motion quality. EvalCrafter also adopts DOVER [64] for video quality
assessment. These metrics from multiple aspects should be considered when
evaluating the effect of quantization.

###  2.3 Model Quantization

Post Training Quantization (PTQ) has proven to be an efficient and effective
model compression method [38]. It converts the floating-point data into low-
bit integers, the process could be represented as:
$x_{q}=\text{round}(\text{clamp}((x-z)/s,-2^{B-1},2^{B-1}-1))$. The $s$
(scale) and $z$ (zero point) are quantization parameters, which are determined
offline based on a set of calibration data with
$S=\text{max}(\text{abs}(x));z=(\text{max}+\text{min})/2$. Diffusion Model:
Prior research Q-Diffusion [27] and PTQD [16] apply post training quantization
for diffusion models. Other research [16, 68, 54, 56, 31, 4] continue to
improve diffusion quantization from the unique timestep dimension.
Transformer: Prior research has made advances in quantizing transformers for
both the ViT [35, 70, 23, 30] and language model quantization [69, 9, 66, 26].
To the best of our knowledge, no prior research has investigated the DiT
quantization. We identify the challenges of applying existing methods to DiT
quantization, and design a novel specialized technique as a solution.

Figure 3: The overall framework of ViDiT-Q. We design quantization scheme
tailored for DiT’s unique challenges, and introduce mixed precision
specialized for video generation.

##  3 Methods

The framework of the ViDiT-Q quantization method is illustrated in Fig. 3. We
start by analyzing the unique challenges for quantizing DiTs (Sec. 3.1). Then,
we introduce an improved quantization scheme, ViDiT-Q, specialized for DiTs to
address these challenges (Sec. 3.2). Further, we develop a metric-decoupled
mixed-precision method to improve quantization under lower bit-widths (Sec.
3.3).

Figure 4: The challenges of DiT quantization. The key unique challenge of DiT
quantization lies in the variance of data at multiple levels. The above
figures depicts the distribution of activation along different dimensions for
STDiT model’s layer “blocks.26.ffn.fc1”.

###  3.1 Challenges for DiT Quantization

As presented in Sec. 1 and Fig. 2, directly applying existing diffusion
quantization methods to DiTs leads to notable degradation for W8A8 and
produces blank images for W4A8. We delve into the reasons for these failures
by examining the data distributions. As illustrated in Fig. 4, we discover
that the unique challenge for DiT quantization lies in the data variance in
multiple levels. Existing methods use fixed and coarse-grained quantization
parameters, which struggle to handle highly variant data. We summarize the
data variance in four levels as follows:

  * •

Token-wise Variance: In DiTs, the activation consists of visual tokens, We
observe notable difference exists between these tokens.

  * •

CFG-wise Variance: For conditional generation, the classifier-free-guidance
[19] (CFG) conducts two separate forwards with and without the control signal
(often implemented with batch). We observe notable difference in activation
range between the conditional part (red square) and the unconditional part
(blue square).

  * •

Timestep-wise Variance: Diffusion method iterates the model for multiple
timesteps. We observe notable variance in activation for the same layer across
timesteps.

  * •

Input Channel-wise Variance: For both the weight and activation, we witness
significant difference across different input-channel.

###  3.2 ViDiT-Q: Improved DiT Quantization Scheme

####  3.2.1 Token-wise Quantization

The primary distinction between DiT and previous UNet-based diffusion models
lies in the feature extraction operator. While U-Net employs convolution
layers for local pixel feature aggregation, DiT utilizes linear and attention
layers for processing image tokens. In DiT quantization, we perform hardware
profiling on the GPU (Ref the appendix A.4 for more details) and observe that,
when combined with flash attention [8], linear layers account for over 77% of
the latency and nearly 100% of the memory cost. Therefore, following prior
literature on transformer-based Language Model (LM) quantization [66, 3, 26],
we focus on quantizing linear layers.

In convolutions, neighboring pixel values are summed together. During
quantization, these values need to share the same quantization parameters for
practical acceleration [38]. Therefore, current CNN-based diffusion
quantization methods primarily adopt tensor-wise quantization parameters.
However, as discussed earlier, token values exhibit notable variation in DiT
activation. Applying the same set of quantization parameters to them leads to
large quantization errors. In contrast to convolution layers, linear layers do
not require summation between tokens. Thus, token-wise quantization parameter
assignment is applicable, which can mitigate “token-wise variance” challenge.
We further analyze the extra cost of introducing token-wise quantization
parameters. The token-wise quantization parameter size is negligible compared
with the activation (around $1/1000$).

####  3.2.2 Dynamic Quantization

The “CFG-wise variance” and “Timestep-wise variance” present two unique
challenges for diffusion models. To the best of our knowledge, no prior
research addresses the CFG-wise variation. Regarding timestep-wise variance,
prior diffusion quantization methods [54, 15] tackle this issue by adopting
timestep-wise quantization parameters through calibration or gradient-based
optimization. However, such a process is costly and may encounter challenges
when generalizing to unseen conditions.

Inspired by recent advances in language model quantization [69, 66], we adopt
dynamic quantization that calculates the quantization parameters online for
activation quantization. We utilize the simple minmax quantization scheme, so
the cost of obtaining quantization parameters online is acceptable (obtaining
data’s max and min). Profiling results (refer to the appendix A.4 for details)
confirm that the additional latency is less than $1/100$ of the layer
inference. Moreover, the cost could be further mitigated with kernel fusion
techniques [61]. While dynamic quantization has been utilized for language
models, we emphasize its importance for diffusion models due to the unique
“CFG-wise variance” and “timestep-wise variance” challenges.

Figure 5: The illustration of timestep-wise channel balancing. Because the
activation quantization difficulty changes across timesteps, we assign
different $\alpha$ for each range to adjust the difficulty.

####  3.2.3 Timestep-aware Channel Balancing

The above challenges exist in activation quantization only. However, for both
weights and activations, there exists input channel-wise variance. A few
channels have significantly larger values ($>100\times$) than others, leading
to the majority of values being quantized to zero. Following prior channel
balancing quantization methods [66, 28], we introduce a per-channel balancing
mask $s\in\mathcal{R}^{C_{i}}$. By dividing the activation with $s$ and
multiplying $s$ with weights, it shifts the quantization difficulty from
activation to weights, and vice versa. The channel balancing mask could be
calculated as follows:

$Y=(X\text{diag}(s)^{-1})\cdot((\text{diag}(s)W))=\hat{X}\cdot\hat{W};\\\
\quad s_{i}=\text{max}(|X_{i}|)^{\alpha}/\text{max}(|W_{i}|)^{1-\alpha},$ (1)

where $X,Y,W$ represents the inputactivation , output activation, and weights.
The $s$ is the aforementioned channel-wise balancing mask, and $\alpha$ is a
hyperparameter. Channel balancing can effectively alleviate the input channel-
wise variance. However, we empirically discover that it is very sensitive to
$\alpha$ choices, and still yields unsatisfactory performance. We further
investigate such limitations and observe that the activation distribution is
notably different across timesteps. As shown in Fig. 5 for the
blocks.27.ffn.fc2 layer, the outlier values are more prominent for later
denoising steps (t=100). The “quantization difficulty” changes significantly
across timesteps. Therefore, employing the same alpha for earlier stages may
shift too much difficulty from activations to weights, harming the weight
quantization (especially for lower bit-width such as W4). For latter stages,
the $\alpha$ is not big enough, and the activation quantization is not
sufficiently compensated.

To address this issue, as presented in Fig. 5, we design a timestep-aware
channel balancing that assigns different $\alpha$ for different timesteps. We
empirically find that the distribution of most layers starts to become acute
after halfway. To minimize the cost of $\alpha$ tuning, we equally divide the
diffusion process into two halves. Specifically, it could be inferred from the
distribution that the former $\alpha_{1}$ should be smaller than globally
tuned $\alpha$, and vice versa for the latter $\alpha_{2}$. With the
assistance of such prior, we can efficiently acquire a well-performing
$\alpha$ within 20 PTQ processes ($<$ 1 GPU hour).

Figure 6: The quantization “bottleneck” phenonmenon and the motivation of
metric decoupled analysis. Some highly sensitive layers hinder the
quantization, and quantization’s effect on multiple aspects needs decoupling
for accurate sensitivity. The values in the lower-right heatmap are calculated
by applying z-score normalization [43] to the metric changes when certain
layer types are quantized.

###  3.3 Metric Decoupled Mixed Precision Design

####  3.3.1 Arch-level and Timestep-level Mixed Precision

As seen in Fig. 2, for lower bit-width settings (e.g., W6A6 and W4A8), quality
degradation is still observed. When investigating the reasons for such
failure, we find that some layers, despite having relatively low quantization
error, could significantly affect quantization. As presented in Fig. 6, W4A16
quantization produces blank images, but simply skipping quantizing the
“blocks.27.ffn.fc2” layer in the STDiT model results in output containing
readable content. This reflects that quantization is “bottlenecked” by some
highly sensitive layers. To address such “bottleneck” phenomenon, an intuitive
solution is to assign higher bit-widths to “protect” these sensitive layers.
We identify variance in quantization sensitivity across the layers and
timesteps, and propose applying mixed precision in these two levels.

####  3.3.2 Metric Decoupled Sensitivity Analysis

The main challenge with mixed precision allocation lies in identifying the
sensitive layers. Prior literature measures quantization sensitivity by
quantizing specific layers alone and then calculating the Mean Squared Error
(MSE) with the floating-point output to indicate their sensitivity. However,
video evaluation has multiple aspects, the MSE error alone may not accurately
reflect overall generation performance. As illustrated in the upper-right part
of Fig. 6, although the left video has a higher MSE error, it exhibits better
visual quality than the right video, which fails to correctly generate birds
but shows some white fractures (denoted with red squares). This underscores
the necessity to enhance sensitivity analysis metrics for video generation.

In exploring sensitivity analysis design, we find that quantization’s impact
on various aspects of generation quality is strongly correlated with layer
types. As observed in the lower-right part of Fig. 6, cross-attention layers
contribute significantly more than other layers to “content change”.
Similarly, “visual quality” is mainly affected by spatial attention and FFN
layers, while “temporal consistency” is primarily influenced by temporal
attention layers. Building upon this finding, we propose to “disentangle” the
mixed influences of quantization on multiple aspects of generation performance
and develop a metric-decoupled sensitivity analysis method. We categorize
layers into three groups based on their effects on different aspects and use
the corresponding metrics as sensitivity measures. Subsequently, we identify
layers with significantly higher relative sensitivity and adjust their bit-
width accordingly. For timestep mixed precision, we divide the entire process
into four equal ranges and quantize them separately to identify the most
sensitive part. Following these principles, we design mixed-precision versions
of ViDiT-Q for W6A6 and W4A8.

Table 1: Performance of ViDiT-Q text-to-video generation on VBench evaluation
benchmark suite. The bit-width “16” represents FP16 without quantization. The
“ViDiT-Q-MP” denotes quantization with mixed precision, which keeps a small
portion of layers as W8A8.

Model | Method | Bit-width | Imaging | Aesthetic | Motion | Dynamic | BG. | Subject | Scene | Overall |
---|---|---|---|---|---|---|---|---|---|---|---
(W/A) | Quality | Quality | Smooth. | Degree | Consist. | Consist. | Consist. | Consist. |
STDiT | - | 16/16 | 63.68 | 57.12 | 96.28 | 56.94 | 96.13 | 90.28 | 39.61 | 26.21 |
Naive PTQ | 8/8 | 56.52 | 53.64 | 95.85 | 69.44 | 93.54 | 85.30 | 29.28 | 25.42 |
(OpenSORA) | ViDiT-Q | 8/8 | 63.48 | 56.95 | 96.14 | 61.11 | 95.84 | 90.24 | 38.22 | 26.06 |
| ViDiT-Q-MP | 6/6 | 62.07 | 57.03 | 95.86 | 62.50 | 95.86 | 89.34 | 39.46 | 26.41 |
| ViDiT-Q-MP | 4/8 | 61.07 | 55.37 | 95.69 | 58.33 | 95.23 | 88.72 | 36.19 | 25.94 |

Table 2: Performance of text-to-video generation on UCF-101 Dataset.. The
description of metrics is provided in Sec. 4.1, unless specified with
$\downarrow$, higher metric values denote better performance.

Model | Method | Bit-width | FVD($\downarrow$) | FVD-FP16($\downarrow$) | CLIPSIM | CLIP-T | VQA- | VQA- | Flow | Temp.
---|---|---|---|---|---|---|---|---|---|---
(W/A) | Aesthetic | Technical | Score. | Flick.
STDiT | - | 16/16 | 136.87 | 0.00 | 0.1996 | 0.9978 | 41.63 | 56.64 | 2.24 | 97.53
Naive PTQ | 8/8 | 154.92 | 50.72 | 0.1993 | 0.9968 | 27.52 | 35.50 | 2.61 | 97.02
ViDiT-Q | 8/8 | 141.13 | 15.52 | 0.1995 | 0.9978 | 43.59 | 55.36 | 2.32 | 97.45
Naive PTQ | 4/8 | 544.34 | 637.02 | 0.1868 | 0.9982 | 0.16 | 0.13 | 1.61 | 99.90
ViDiT-Q-MP | 4/8 | 129.10 | 60.13 | 0.1995 | 0.9977 | 33.98 | 47.65 | 1.89 | 97.57
Latte | - | 16/16 | 99.90 | 0.00 | 0.1970 | 0.9963 | 36.33 | 91.23 | 3.37 | 96.22
Naive PTQ | 8/8 | 98.75 | 73.82 | 0.1981 | 0.9950 | 27.62 | 50.52 | 3.53 | 95.35
ViDiT-Q | 8/8 | 110.96 | 20.83 | 0.1959 | 0.9962 | 30.26 | 80.32 | 3.14 | 95.95
Naive PTQ | 4/8 | 183.52 | 239.08 | 0.1719 | 0.9929 | 5.62 | 0.41 | 66.06 | 65.14
ViDiT-Q | 4/8 | 95.04 | 79.11 | 0.1943 | 0.9971 | 21.76 | 32.17 | 2.84 | 95.57

##  4 Experiments

###  4.1 Implementation Details and Experimental Settings

We evaluate the performance of ViDiT-Q on a variety of models, bit-widths and
evaluation settings. For more detailed settings, please refer to Appendix B
and C.

Quantization Scheme: We adopt the simple minmax quantization scheme. The
quantization parameters for activation are dynamic and computed online with
negligible overhead. The channel balancing $\alpha$ and the mixed precision
plan are determined offline based on the calibration data.

Mixed Precision Strategy For W6A6, we discover the FFN layers for blocks 6 and
26 are significantly more sensitive, since they account for < 1% of the
overall latency, we maintain them as FP16, and keep the rest layers as W6A6.
For W4A8, we witness notable higher sensitivity for most FFNs, for simplicity,
we set all FFN layers (15$\%$ layers) as W8A8, and set the rest as W4A8. Also,
we discover the first quarter of timesteps is more sensitive, and set them to
W8A8.

Video Genration Evaluation Settings: We apply ViDiT-Q to Latte [37] and STDiT
(open-sora) [22] and evaluate them in two settings. (1) Benchmark suite: We
evaluate the quantized model on VBench [24] to provide comprehensive results .
Following prior research [48], we select 8 major dimensions from Vbench. (2)
Multi-aspects metrics: For more flexible evaluation, we select representative
metrics for each aspect, and measure them on UCF-101 [55] and open-sora [22]
prompt sets. Following EvalCrafter [33], we select CLIPSIM and CLIP-Temp to
measure the text-video alignment and temporal semantic consistency, and DOVER
[64]’s video quality assessment (VQA) metrics to evaluate the generation
quality from aesthetic (high-level) and technical (low-level) perspectives,
Flow-score and Temporal Flickering are used for evaluating the temporal
consistency. The commonly adopted FVD [59] is also provided. Specifically, due
to the lack of ground-truth videos for prompt-only datasets, inspired by [57],
we also report FVD-FP16 which chooses the FP16 generated video as ground-
truth. The above metrics are evaluated on 101 prompts (1 for each class) for
UCF-101, and 10 example prompts for open-sora. We adopt the class-conditioned
Latte model trained on UCF-101. We use the 20-steps DDIM solver with CFG scale
of 7.0 for Latte, and 100-steps DDIM with CFG scale of 4.0 for STDiT.

Image Evaluation Settings: We apply ViDiT-Q to PixArt-$\alpha$ and
PixArt-$\Sigma$ model on the the first 1024 prompts of COCO annotations [29].
We choose FID [18] for fidelity evaluation, Clipscore [17] for text-image
alignment, and ImageReward [67] for human preference in multiple aspects.
Following the original code implementation, we use the 20-steps DPM-solver
with CFG scale of 4.5.

Table 3: Comparison of quantization methods on open-sora prompt set. Due to
the lack of existing DiT quantization method, we adapt baseline quantization
methods to DiT.

Model | Method | Bit-width | CLIPSIM | CLIP-Temp | VQA- | VQA- | Flow
---|---|---|---|---|---|---|---
(W/A) | Aesthetic | Technical | Score.
STDiT | - | 16/16 | 0.1797 | 0.9988 | 63.4014 | 50.4597 | 0.9751
Naive PTQ | 8/8 | 0.1956 | 0.9988 | 48.2358 | 25.5961 | 0.6081
PTQ4DM [52] | 8/8 | 0.1812 | 0.9984 | 50.0674 | 25.1344 | 0.6335
Q-Diffusion [27] | 8/8 | 0.1781 | 0.9987 | 51.6834 | 38.2680 | 0.6473
ZeroQuant [69] | 8/8 | 0.1938 | 0.9988 | 58.3127 | 51.5312 | 0.6442
ViDiT-Q | 8/8 | 0.1950 | 0.9991 | 60.7025 | 54.6361 | 0.8865
Naive PTQ | 6/6 | 0.1912 | 0.9964 | 14.6374 | 10.6919 | 33.611
PTQ4DM [52] | 6/6 | 0.1804 | 0.9977 | 20.9610 | 8.6566 | 0.7600
ViDiT-Q-MP | 6/6 | 0.1794 | 0.9983 | 60.8124 | 53.4413 | 1.2266
Naive PTQ | 4/8 | 0.2010 | 0.9986 | 0.1765 | 0.0863 | 1.5722
PTQ4DM [52] | 4/8 | 0.1727 | 0.9981 | 0.4781 | 0.2672 | 1.5298
SmoothQuant [66] | 4/8 | 0.1910 | 0.9989 | 31.9626 | 22.8467 | 0.5594
ViDiT-Q-MP | 4/8 | 0.1809 | 0.9989 | 60.6158 | 49.3838 | 1.1278

Model | Method | Bit-width | FID($\downarrow$) | CLIP($\uparrow$) | IR($\uparrow$)
---|---|---|---|---|---
(W/A)
Pixart-$\alpha$ | - | 16/16 | 73.338 | 0.258 | 0.901
Naive | 8/8 | 115.14 | 0.226 | -0.953
4/8 | 108.40 | 0.234 | -0.725
ViDiT-Q | 8/8 | 75.613 | 0.259 | 0.917
4/8 | 74.326 | 0.257 | 0.887
Pixart-$\Sigma$ | - | 16/16 | 72.699 | 0.262 | 0.929
Naive | 8/8 | 302.12 | 0.163 | -2.236
4/8 | 287.80 | 0.165 | -2.219
ViDiT-Q | 8/8 | 72.845 | 0.263 | 0.926
4/8 | 71.936 | 0.264 | 0.944

Figure 7: Performance of ViDiT-Q text-to-image generation on COCO. Left: The
metric scores of PixArt-$\alpha$ and PixArt-$\Sigma$ quantization. Right:
Generated image comparison of W8A8 quantization.

###  4.2 Main Results

We briefly conclude key findings here, for more detailed analysis, please
refer to Appendix Sec. C.

Text-to-video generation on VBench and UCF-101: As presented in Tab. 1 and
Tab. 2, while naive W6A6 and W4A8 quantization generates pure noise or blank
images (Ref Fig. 2 and Appendix for visual example), ViDiT-Q consistently
preserves generation quality for all bit-widths. ViDiT-Q’s W4A8 even
outperforms baseline W8A8. It’s worth noting that some abnormal scores are due
to the baseline’s failure to generate meaningful contents. Besides, having a
flow-score that is either too high or too low than FP videos indicates
degradation. Also, we discover that the FVD fluctuates heavily with less data
(101 videos). We find that FVD-FP16 shows better correlation with visual
perception.

Comparison with other quantization methods: We compare ViDiT-Q with other
quantization methods in Tab. 3. Due to the lack of a baseline DiT quantization
method, we adapt prior diffusion and transformer quantization methods to DiTs.
For W8A8, ViDiT-Q outperforms other quantization methods and achieves similar
metric scores to FP16. For the challenging W4A8, only SmoothQuant [66] could
prevent generation failure. ViDiT-Q with mixed precision further improves
performance to be comparable with baseline W8A8 quantization.

Text-to-image generation on COCO: As presented in Fig. 7, ViDiT-Q
significantly improves generation quality from blurred and unreadable to
almost identical to the FP generation. The metric scores are similar and even
higher than FP16 baseline for W8A8 and W4A8. More analysis and generated
images are presented in Appendix Sec. C.3.

Ablation Studies: We conduct ablation studies by gradually incorporating
ViDiT-Q’s techniques to the challenging W4A8 quantization. As shown in Tab. 4,
the introduction of dynamic quantization marks the beginning of the generation
meaningful content. The subsequent integration of channel balancing, along
with its timestep-wise version, further improves generation quality.
Nevertheless, we still observe significant degradation, underscoring the
necessity for mixed precision strategies. Detailed analysis and video examples
are presented in Appendix Sec. C.4.

###  4.3 Hardware Resource Savings

Memory footprint reduction. Fig. 8 (a) shows the GPU memory usage of ViDiT-Q
and the FP16 baseline. ViDiT-Q can reduce the memory from two aspects: (1)
Weight quantization reduces the allocated memory for storing model weights.
(2) Activation quantization reduces allocated memory to store intermediate
activations. Combining the two benefits, ViDiT-Q can effectively reduce the
peak memory footprint by 2$\times$, 2.42$\times$ and 2.67$\times$ under the
W8A8, W4A8-MP and W6A6-MP.

Latency speedup. We present the latency speedup in Fig. 8-(b). Due to lack of
W6A6/W4A8 or dynamic quantization kernels, we adopt the INT8 GPU kernels to
estimate the latency speedup. Replacing FP16 layers with INT8 achieves $\sim
2\times$ acceleration. Considering the unquantizable layers (norms,
nonlinears, attention), and the FP to INT conversion cost, the overall speedup
is 1.47$\times$. Detailed description of hardware resource measurements are
provided in appendix A.4.

Table 4: Ablation studies of ViDiT-Q techniques.. By gradually incorporating
ViDiT-Q’s techniques, W4A8 quantization improves from failure to acceptable.

Methods | Bit-width | CLIPSIM | CLIP-Temp | VQA- | VQA- | Flow
---|---|---|---|---|---|---
Dynamic | Channel Balance | Timestep-aware CB | (W/A) | Aesthetic | Technical | Score.
- | - | - | 16/16 | 0.180 | 0.998 | 64.198 | 51.904 | 1.427
- | - | - | 4/8 | 0.201 | 0.997 | 0.178 | 0.086 | 1.572
✓ | - | - | 4/8 | 0.196 | 0.998 | 32.217 | 10.994 | 0.884
✓ | ✓ | - | 4/8 | 0.191 | 0.999 | 31.963 | 22.847 | 1.117
✓ | ✓ | ✓ | 4/8 | 0.199 | 0.999 | 51.823 | 23.215 | 1.189

Bit-width | BitOps | Memory | Latency
---|---|---|---
(W/A) | Opt. | Opt.
16/16 | 1.00$\times$ | 1.00$\times$ | 1.00$\times$
8/8 | 4.00$\times$ | 1.99$\times$ | 1.47$\times$
6/6-MP | 7.11$\times$ | 2.67$\times$ | -
4/8-MP | 5.05$\times$ | 2.42$\times$ | -

(a)

(b)

Figure 8: The illustration of ViDiT-Q’s hardware resource savings. Left: The
comparison of efficiency and performance under different ViDiT-Q quantization
configurations. The “6/6-MP” and “4/8-MP”’s latency could not be measured for
now due to lack of GPU kernels (Ref Sec. A.4 for details). Right: (a)
ViDiT-Q’s memory optimization, (b) ViDiT-Q’s latency breakdown under W8A8.

##  5 Conclusion and Limitations

We design ViDiT-Q, a quantization method tailored for DiTs. With the
assistance of metric decoupled mixed precision, ViDiT-Q achieves W4A8
quantization with minimum visual degradation. Despite achieving good
performance, it still faces the following limitations: the timestep-wise mixed
precision requires additional system-level optimization, the efficiency and
precision of quantization sensitivity analysis still needs improvement, lower
activation bit-width is essential for fully utilizing the acceleration
potential of 4-bit weight. We aim to address these issues and further improve
ViDiT-Q.

## Acknowledgements

This work was supported by National Natural Science Foundation of China (No.
62325405, 62104128, U19B2019, U21B2031, 61832007, 62204164), Tsinghua EE
Xilinx AI Research Fund, and Beijing National Research Center for Information
Science and Technology (BNRist). We thank for all the support from
Infinigence-AI.

## References

  * [1] Fan Bao, Shen Nie, Kaiwen Xue, Yue Cao, Chongxuan Li, Hang Su, and Jun Zhu.  All are worth words: A vit backbone for diffusion models.  In CVPR, 2023.
  * [2] Andreas Blattmann, Robin Rombach, Huan Ling, Tim Dockhorn, Seung Wook Kim, Sanja Fidler, and Karsten Kreis.  Align your latents: High-resolution video synthesis with latent diffusion models.  In Proceedings of the IEEE/CVF Conference on Computer Vision and Pattern Recognition, pages 22563–22575, 2023.
  * [3] Yaohui Cai, Zhewei Yao, Zhen Dong, Amir Gholami, Michael W Mahoney, and Kurt Keutzer.  Zeroq: A novel zero shot quantization framework.  In Proceedings of the IEEE/CVF Conference on Computer Vision and Pattern Recognition, pages 13169–13178, 2020.
  * [4] Hanwen Chang, Haihao Shen, Yiyang Cai, Xinyu Ye, Zhenzhong Xu, Wenhua Cheng, Kaokao Lv, Weiwei Zhang, Yintong Lu, and Heng Guo.  Effective quantization for diffusion models on cpus.  arXiv preprint arXiv:2311.16133, 2023.
  * [5] Junsong Chen, Chongjian Ge, Enze Xie, Yue Wu, Lewei Yao, Xiaozhe Ren, Zhongdao Wang, Ping Luo, Huchuan Lu, and Zhenguo Li.  Pixart-$\sigma$: Weak-to-strong training of diffusion transformer for 4k text-to-image generation, 2024.
  * [6] Junsong Chen, Jincheng Yu, Chongjian Ge, Lewei Yao, Enze Xie, Yue Wu, Zhongdao Wang, James Kwok, Ping Luo, Huchuan Lu, and Zhenguo Li.  Pixart-$\alpha$: Fast training of diffusion transformer for photorealistic text-to-image synthesis, 2023.
  * [7] Shoufa Chen, Mengmeng Xu, Jiawei Ren, Yuren Cong, Sen He, Yanping Xie, Animesh Sinha, Ping Luo, Tao Xiang, and Juan-Manuel Perez-Rua.  Gentron: Delving deep into diffusion transformers for image and video generation.  arXiv preprint arXiv:2312.04557, 2023.
  * [8] Tri Dao, Dan Fu, Stefano Ermon, Atri Rudra, and Christopher Ré.  Flashattention: Fast and memory-efficient exact attention with io-awareness.  Advances in Neural Information Processing Systems, 35:16344–16359, 2022.
  * [9] Tim Dettmers, Mike Lewis, Younes Belkada, and Luke Zettlemoyer.  Llm. int8 (): 8-bit matrix multiplication for transformers at scale, 2022.  CoRR abs/2208.07339.
  * [10] Zhen Dong, Zhewei Yao, Amir Gholami, Michael W Mahoney, and Kurt Keutzer.  Hawq: Hessian aware quantization of neural networks with mixed-precision.  In Proceedings of the IEEE/CVF International Conference on Computer Vision, pages 293–302, 2019.
  * [11] Patrick Esser, Johnathan Chiu, Parmida Atighehchian, Jonathan Granskog, and Anastasis Germanidis.  Structure and content-guided video synthesis with diffusion models.  In Proceedings of the IEEE/CVF International Conference on Computer Vision, pages 7346–7356, 2023.
  * [12] Zhao et. al.  Mixdq: Memory-efficient few-step text-to-image diffusion models with metric-decoupled mixed precision quantization.  https://a-suozhang.xyz/mixdq.github.io/, 2024.  Accessed: 2024-05-20.
  * [13] Yuwei Guo, Ceyuan Yang, Anyi Rao, Yaohui Wang, Yu Qiao, Dahua Lin, and Bo Dai.  Animatediff: Animate your personalized text-to-image diffusion models without specific tuning.  arXiv preprint arXiv:2307.04725, 2023.
  * [14] Ali Hatamizadeh, Jiaming Song, Guilin Liu, Jan Kautz, and Arash Vahdat.  Diffit: Diffusion vision transformers for image generation.  arXiv preprint arXiv:2312.02139, 2023.
  * [15] Yefei He, Jing Liu, Weijia Wu, Hong Zhou, and Bohan Zhuang.  Efficientdm: Efficient quantization-aware fine-tuning of low-bit diffusion models.  arXiv preprint arXiv:2310.03270, 2023.
  * [16] Yefei He, Luping Liu, Jing Liu, Weijia Wu, Hong Zhou, and Bohan Zhuang.  Ptqd: Accurate post-training quantization for diffusion models.  Advances in Neural Information Processing Systems, 36, 2024.
  * [17] Jack Hessel, Ari Holtzman, Maxwell Forbes, Ronan Bras, and Choi Yejin.  Clipscore: A reference-free evaluation metric for image captioning.  pages 7514–7528, 01 2021.
  * [18] Martin Heusel, Hubert Ramsauer, Thomas Unterthiner, Bernhard Nessler, and Sepp Hochreiter.  Gans trained by a two time-scale update rule converge to a local nash equilibrium.  In Proceedings of the 31st International Conference on Neural Information Processing Systems, NIPS’17, page 6629–6640, Red Hook, NY, USA, 2017. Curran Associates Inc.
  * [19] Jonathan Ho.  Classifier-free diffusion guidance.  ArXiv, abs/2207.12598, 2022.
  * [20] Jonathan Ho, William Chan, Chitwan Saharia, Jay Whang, Ruiqi Gao, Alexey Gritsenko, Diederik P Kingma, Ben Poole, Mohammad Norouzi, David J Fleet, et al.  Imagen video: High definition video generation with diffusion models.  arXiv preprint arXiv:2210.02303, 2022.
  * [21] Jonathan Ho, Tim Salimans, Alexey Gritsenko, William Chan, Mohammad Norouzi, and David J Fleet.  Video diffusion models.  Advances in Neural Information Processing Systems, 35:8633–8646, 2022.
  * [22] HPC-AI Advisory Council.  Open-Sora.  https://github.com/hpcaitech/Open-Sora, 2024.
  * [23] Xijie Huang, Zhiqiang Shen, and Kwang-Ting Cheng.  Variation-aware vision transformer quantization.  arXiv preprint arXiv:2307.00331, 2023.
  * [24] Ziqi Huang, Yinan He, Jiashuo Yu, Fan Zhang, Chenyang Si, Yuming Jiang, Yuanhan Zhang, Tianxing Wu, Qingyang Jin, Nattapol Chanpaisit, et al.  Vbench: Comprehensive benchmark suite for video generative models.  arXiv preprint arXiv:2311.17982, 2023.
  * [25] Benoit Jacob, Skirmantas Kligys, Bo Chen, Menglong Zhu, Matthew Tang, Andrew Howard, Hartwig Adam, and Dmitry Kalenichenko.  Quantization and training of neural networks for efficient integer-arithmetic-only inference.  In Proceedings of the IEEE conference on computer vision and pattern recognition, pages 2704–2713, 2018.
  * [26] Shiyao Li, Xuefei Ning, Luning Wang, Tengxuan Liu, Xiangsheng Shi, Shengen Yan, Guohao Dai, Huazhong Yang, and Yu Wang.  Evaluating quantized large language models, 2024.
  * [27] Xiuyu Li, Yijiang Liu, Long Lian, Huanrui Yang, Zhen Dong, Daniel Kang, Shanghang Zhang, and Kurt Keutzer.  Q-diffusion: Quantizing diffusion models.  In Proceedings of the IEEE/CVF International Conference on Computer Vision, pages 17535–17545, 2023.
  * [28] Ji Lin, Jiaming Tang, Haotian Tang, Shang Yang, Xingyu Dang, and Song Han.  Awq: Activation-aware weight quantization for llm compression and acceleration.  arXiv preprint arXiv:2306.00978, 2023.
  * [29] Tsung-Yi Lin, Michael Maire, Serge Belongie, Lubomir Bourdev, Ross Girshick, James Hays, Pietro Perona, Deva Ramanan, C. Lawrence Zitnick, and Piotr Dollár.  Microsoft coco: Common objects in context, 2015.
  * [30] Yang Lin, Tianyu Zhang, Peiqin Sun, Zheng Li, and Shuchang Zhou.  Fq-vit: Post-training quantization for fully quantized vision transformer.  arXiv preprint arXiv:2111.13824, 2021.
  * [31] Xuewen Liu, Zhikai Li, Junrui Xiao, and Qingyi Gu.  Enhanced distribution alignment for post-training quantization of diffusion models.  arXiv preprint arXiv:2401.04585, 2024.
  * [32] Yaofang Liu, Xiaodong Cun, Xuebo Liu, Xintao Wang, Yong Zhang, Haoxin Chen, Yang Liu, Tieyong Zeng, Raymond Chan, and Ying Shan.  Evalcrafter: Benchmarking and evaluating large video generation models.  arXiv preprint arXiv:2310.11440, 2023.
  * [33] Yaofang Liu, Xiaodong Cun, Xuebo Liu, Xintao Wang, Yong Zhang, Haoxin Chen, Yang Liu, Tieyong Zeng, Raymond Chan, and Ying Shan.  Evalcrafter: Benchmarking and evaluating large video generation models.  2023\.
  * [34] Zechun Liu, Zhiqiang Shen, Marios Savvides, and Kwang-Ting Cheng.  Reactnet: Towards precise binary neural network with generalized activation functions.  In Computer Vision–ECCV 2020: 16th European Conference, Glasgow, UK, August 23–28, 2020, Proceedings, Part XIV 16, pages 143–159. Springer, 2020.
  * [35] Zhenhua Liu, Yunhe Wang, Kai Han, Wei Zhang, Siwei Ma, and Wen Gao.  Post-training quantization for vision transformer.  Advances in Neural Information Processing Systems, 34:28092–28103, 2021.
  * [36] Zeyu Lu, Zidong Wang, Di Huang, Chengyue Wu, Xihui Liu, Wanli Ouyang, and Lei Bai.  Fit: Flexible vision transformer for diffusion model, 2024.
  * [37] Xin Ma, Yaohui Wang, Gengyun Jia, Xinyuan Chen, Ziwei Liu, Yuan-Fang Li, Cunjian Chen, and Yu Qiao.  Latte: Latent diffusion transformer for video generation.  arXiv preprint arXiv:2401.03048, 2024.
  * [38] Markus Nagel, Marios Fournarakis, Rana Ali Amjad, Yelysei Bondarenko, Mart Van Baalen, and Tijmen Blankevoort.  A white paper on neural network quantization.  arXiv preprint arXiv:2106.08295, 2021.
  * [39] NVIDIA.  Nsight Systems.  https://docs.nvidia.com/nsight-systems/index.html, 2024.
  * [40] Nvidia.  Nvidia cutlass release v3.4, 2024.
  * [41] OpenAI.  Video generation models as world simulators.  https://openai.com/index/video-generation-models-as-world-simulators/, Year of Access.
  * [42] Nilesh Prasad Pandey, Markus Nagel, Mart van Baalen, Yin Huang, Chirag Patel, and Tijmen Blankevoort.  A practical mixed precision algorithm for post-training quantization.  arXiv preprint arXiv:2302.05397, 2023.
  * [43] SGOPAL Patro and Kishore Kumar Sahu.  Normalization: A preprocessing stage.  arXiv preprint arXiv:1503.06462, 2015.
  * [44] William Peebles and Saining Xie.  Scalable diffusion models with transformers, 2023.
  * [45] Dustin Podell, Zion English, Kyle Lacey, A. Blattmann, Tim Dockhorn, Jonas Muller, Joe Penna, and Robin Rombach.  Sdxl: Improving latent diffusion models for high-resolution image synthesis.  ArXiv, abs/2307.01952, 2023.
  * [46] PyTorch.  PyTorch Memory Management, 2023.
  * [47] Alec Radford, Jong Wook Kim, Chris Hallacy, Aditya Ramesh, Gabriel Goh, Sandhini Agarwal, Girish Sastry, Amanda Askell, Pamela Mishkin, Jack Clark, et al.  Learning transferable visual models from natural language supervision.  In International conference on machine learning, pages 8748–8763. PMLR, 2021.
  * [48] Weiming Ren, Harry Yang, Ge Zhang, Cong Wei, Xinrun Du, Stephen Huang, and Wenhu Chen.  Consisti2v: Enhancing visual consistency for image-to-video generation.  arXiv preprint arXiv:2402.04324, 2024.
  * [49] Robin Rombach, Andreas Blattmann, Dominik Lorenz, Patrick Esser, and Björn Ommer.  High-resolution image synthesis with latent diffusion models.  In Proceedings of the IEEE/CVF conference on computer vision and pattern recognition, pages 10684–10695, 2022.
  * [50] Olaf Ronneberger, Philipp Fischer, and Thomas Brox.  U-net: Convolutional networks for biomedical image segmentation.  In Medical image computing and computer-assisted intervention–MICCAI 2015: 18th international conference, Munich, Germany, October 5-9, 2015, proceedings, part III 18, pages 234–241. Springer, 2015.
  * [51] Tim Salimans, Ian Goodfellow, Wojciech Zaremba, Vicki Cheung, Alec Radford, and Xi Chen.  Improved techniques for training gans.  NIPS’16, page 2234–2242, Red Hook, NY, USA, 2016. Curran Associates Inc.
  * [52] Yuzhang Shang, Zhihang Yuan, Bin Xie, Bingzhe Wu, and Yan Yan.  Post-training quantization on diffusion models.  In Proceedings of the IEEE/CVF Conference on Computer Vision and Pattern Recognition, pages 1972–1981, 2023.
  * [53] Uriel Singer, Adam Polyak, Thomas Hayes, Xiaoyue Yin, Jie An, Songyang Zhang, Qiyuan Hu, Harry Yang, Oron Ashual, Oran Gafni, Devi Parikh, Sonal Gupta, and Yaniv Taigman.  Make-a-video: Text-to-video generation without text-video data.  ArXiv, abs/2209.14792, 2022.
  * [54] Junhyuk So, Jungwon Lee, Daehyun Ahn, Hyungjun Kim, and Eunhyeok Park.  Temporal dynamic quantization for diffusion models.  Advances in Neural Information Processing Systems, 36, 2024.
  * [55] Khurram Soomro, Amir Roshan Zamir, and Mubarak Shah.  Ucf101: A dataset of 101 human actions classes from videos in the wild.  arXiv preprint arXiv:1212.0402, 2012.
  * [56] Haojun Sun, Chen Tang, Zhi Wang, Yuan Meng, Xinzhu Ma, Wenwu Zhu, et al.  Tmpq-dm: Joint timestep reduction and quantization precision selection for efficient diffusion models.  arXiv preprint arXiv:2404.09532, 2024.
  * [57] Siao Tang, Xin Wang, Hong Chen, Chaoyu Guan, Zewen Wu, Yansong Tang, and Wenwu Zhu.  Post-training quantization with progressive calibration and activation relaxing for text-to-image diffusion models.  arXiv preprint arXiv:2311.06322, 2023.
  * [58] Zachary Teed and Jia Deng.  Raft: Recurrent all-pairs field transforms for optical flow.  In Computer Vision–ECCV 2020: 16th European Conference, Glasgow, UK, August 23–28, 2020, Proceedings, Part II 16, pages 402–419. Springer, 2020.
  * [59] Thomas Unterthiner, Sjoerd van Steenkiste, Karol Kurach, Raphaël Marinier, Marcin Michalski, and Sylvain Gelly.  Fvd: A new metric for video generation.  2019\.
  * [60] Ashish Vaswani, Noam Shazeer, Niki Parmar, Jakob Uszkoreit, Llion Jones, Aidan N Gomez, Łukasz Kaiser, and Illia Polosukhin.  Attention is all you need.  Advances in neural information processing systems, 30, 2017.
  * [61] Guibin Wang, YiSong Lin, and Wei Yi.  Kernel fusion: An effective method for better power efficiency on multithreaded gpu.  In 2010 IEEE/ACM Int’l Conference on Green Computing and Communications & Int’l Conference on Cyber, Physical and Social Computing, pages 344–350. IEEE, 2010.
  * [62] Kuan Wang, Zhijian Liu, Yujun Lin, Ji Lin, and Song Han.  Haq: Hardware-aware automated quantization with mixed precision.  In Proceedings of the IEEE/CVF conference on computer vision and pattern recognition, pages 8612–8620, 2019.
  * [63] Chenfei Wu, Lun Huang, Qianxi Zhang, Binyang Li, Lei Ji, Fan Yang, Guillermo Sapiro, and Nan Duan.  Godiva: Generating open-domain videos from natural descriptions.  arXiv preprint arXiv:2104.14806, 2021.
  * [64] Haoning Wu, Erli Zhang, Liang Liao, Chaofeng Chen, Jingwen Hou, Annan Wang, Wenxiu Sun, Qiong Yan, and Weisi Lin.  Exploring video quality assessment on user generated contents from aesthetic and technical perspectives.  In Proceedings of the IEEE/CVF International Conference on Computer Vision, pages 20144–20154, 2023.
  * [65] Xiaoshi Wu, Yiming Hao, Keqiang Sun, Yixiong Chen, Feng Zhu, Rui Zhao, and Hongsheng Li.  Human preference score v2: A solid benchmark for evaluating human preferences of text-to-image synthesis.  arXiv preprint arXiv:2306.09341, 2023.
  * [66] Guangxuan Xiao, Ji Lin, Mickael Seznec, Hao Wu, Julien Demouth, and Song Han.  Smoothquant: Accurate and efficient post-training quantization for large language models, 2024.
  * [67] Jiazheng Xu, Xiao Liu, Yuchen Wu, Yuxuan Tong, Qinkai Li, Ming Ding, Jie Tang, and Yuxiao Dong.  Imagereward: Learning and evaluating human preferences for text-to-image generation, 2023.
  * [68] Yuewei Yang, Xiaoliang Dai, Jialiang Wang, Peizhao Zhang, and Hongbo Zhang.  Efficient quantization strategies for latent diffusion models.  arXiv preprint arXiv:2312.05431, 2023.
  * [69] Zhewei Yao, Reza Yazdani Aminabadi, Minjia Zhang, Xiaoxia Wu, Conglong Li, and Yuxiong He.  Zeroquant: Efficient and affordable post-training quantization for large-scale transformers.  Advances in Neural Information Processing Systems, 35:27168–27183, 2022.
  * [70] Zhihang Yuan, Chenhao Xue, Yiqi Chen, Qiang Wu, and Guangyu Sun.  Ptq4vit: Post-training quantization for vision transformers with twin uniform quantization.  In European conference on computer vision, pages 191–207. Springer, 2022.
  * [71] Hongkai Zheng, Weili Nie, Arash Vahdat, and Anima Anandkumar.  Fast training of diffusion models with masked transformers, 2024.

##  Appendix A Hardware Experiments Details

###  A.1 Profiling Settings

We evaluate the latency and memory usage of ViDiT-Q on the Nvidia RTX 4080 GPU
using CUDA 12.1. All profiling is conducted with a batch size of 1. Since
there are no open-source kernels supporting dynamic W8A8 quantization on GPU,
we demonstrate in Sec. A.3 that the additional cost of dynamic quantization
over static quantization is negligible. For our analysis, we utilize static
quantization GPU kernels implemented based on the Cutlass library [40] for
latency and memory measurement. Memory usage is estimated using PyTorch Memory
Management APIs [46], while inference latency is estimated with NVIDIA Nsight
tools [39].

###  A.2 Motivation for quantizing linear layers only

In Sec. 3, we mention that we focus on quantizing the linear layers and leave
the attention computation unquantized. We elaborate on the reason for this
focus here. In Fig. 9, we visualize the detailed latency breakdown for an
STDiT model block. The ’attention computation’ includes the matrix
multiplication for query and key embedding to generate the attention map, and
the multiplication of the attention map with the value embedding. The QKV
linear mapping and the projection after attention aggregation are not
included, as these are linear layers that can be quantized. As shown, when
utilizing FlashAttention, the latency cost of attention computation accounts
for only 14.3$\%$ of the overall latency. Additionally, FlashAttention
minimizes the activation memory usage for storing the attention map.
Therefore, we focus on the primary cost: the linear layers. We quantize all
linear layers except for the “t embedding”, “y embedding” and “final layer”,
they appear at the start or end of the model, and have smaller channel sizes.
They account for only negligible amount of computation ($<1/1000$ overall
latency), therefore we maintain them as FP16.

Figure 9: The latency comparison of linear layers and attention computation.
When utilizing FlashAttention, the attention computation only takes up a small
portion (14.3$\%$) of the latency.

###  A.3 Cost of ViDiT-Q quantization

Dynamic quantization’s extra cost As discussed in Sec. 3.2.2, dynamically
calculating the quantization parameters online is vital for preserving the
performance of diffusion transformers. ViDiT-Q adopts the easy-to-implement
minmax quantization scheme. In this case, the extra cost of determining the
quantization parameter online is obtaining the data’s min and max, which is
negligible compared with the linear layer computation. To verify this, we
conduct profiling with NVIDIA Nsight System. As seen from Fig. 10, the latency
of the linear layer takes $\sim 10$ ms, while the minmax operator
(“reduce_kernel”) only costs $\sim 100$us (around 1/100$\times$). Furthermore,
the extra cost of obtaining the minmax could be further reduced with the
kernel-fusion technique [61]. The minmax operator is an elment-wise reduced
operation, which is similar to the “normalization” and “nonlinear activation”
layers, which often occurs before the linear layer. By fusing the minmax
operator into the previous layer, the overhead could be significantly reduces.
To sum up, the extra cost of dynamic quantization over static quantization is
negligible.

Figure 10: The latency comparison of minmax operator and linear layer. The
minmax operator takes 1/100 of the linear layer’s latency.

Channel balancing’s cost: As discussed in Sec. 3.2.3, the channel balancing
methods introduce an additional channel-wise mask $s$ to balance the
’quantization difficulty’ between weight and activation. Prior literature [66,
28] has shown that the extra cost of multiplying $s$ with weights and dividing
the activation by $s$ can be eliminated by merging $s$ into the prior layer’s
computation. Most linear layers are preceded by normalization layers
(LayerNorm), which contain an affine transform. By dividing the weights in
this affine transform, we achieve the same effect as dividing the linear
layer’s activation by $s$. The multiplication of weights by $s$ can be
performed offline. This approach eliminates the extra cost of channel
balancing.

Post training quantization process’s cost ViDiT-Q adopts a simple min-max
quantization scheme, and the activation’s quantization parameters are computed
online. Timestep-wise channel balancing requires several model inferences for
$\alpha$ tuning. As discussed in Sec. 3.2.3, guided by distribution analysis,
proper $\alpha$ values can be achieved with 10-30 PTQ inferences, taking less
than 1 GPU hour. In contrast, some baseline quantization methods, such as
Q-Diffusion, require gradient-based tuning of quantization parameters, which
is costly and takes tens of GPU hours. For mixed precision, the major cost is
determining sensitivity. Current ViDiT-Q’s sensitivity analysis requires
iterative model testing, taking about 10 GPU hours. Enhancing the accuracy and
speed of sensitivity analysis remains a future direction.

###  A.4 Estimation of Hardware Resource Savings

In Sec. 4.3, we present the efficiency improvement of ViDiT-Q and provide a
more detailed analysis here. The efficiency improvement is shown in the table
within Fig. 8 in the main paper. Due to the lack of W6A6 and W4A8 GPU kernels,
we report the computational savings in BitOps, following prior literature
[34]. Their latency is represented by W8A8’s latency. With negligible visual
degradation under W6A6-MP and W4A8-MP, ViDiT-Q achieves $7.11\times$ and
$5.05\times$ computation resource savings, respectively.

Memory Optimization: GPU memory usage comprises two parts: allocated memory
for storing model weights and intermediate activations. With FlashAttention,
the need for storing the huge attention map is eliminated. We find that the
current peak memory for activation storage is the QKV embeddings. ViDiT-Q’s
weight quantization reduces the model size, and activation quantization
effectively reduces the size of QKV embeddings. By combining these benefits,
ViDiT-Q effectively reduces the peak memory of the model. Both the “W6A6-MP”
and “W4A8-MP” models adopt mixed precision and set some highly sensitive
layers with higher bit-widths. Therefore, the memory optimization achieved is
less than the theoretical statistics (W4A8-MP: $2.42\times$ < W4A8:
$2.67\times$).

Latency Optimization: As shown in Fig. 8-(b) in the main paper, we measure the
latency of the three parts of the quantized model: the unquantized layers, the
quantized layers, and the quantization overhead of FP16-to-INT8 conversion.
These are summed to produce the overall latency. ViDiT-Q achieves a
1.47$\times$ latency speedup compared to the FP16 baseline. Further
improvements in latency speedup with lower bit-width or system-level
optimizations (e.g., W4A8 GPU kernels) remain for future exploration.

##  Appendix B Detailed Description of Evaluation Metrics

###  B.1 Benchmark Suite

Following VBench [24], our benchmark suite encompasses three key dimensions.
(1) Frame-wise Quality assesses the quality of each individual frame without
taking temporal quality into concern.

  * •

Aesthetic Quality evaluates the artistic and beauty value perceived by humans
towards each video frame.

  * •

Imaging Quality assesses distortion (e.g., over-exposure, noise) presented in
the generated frames

(2) Temporal Quality assesses the cross-frame temporal consistency and
dynamics.

  * •

Subject Consistency assesses whether appearance of subjects in the video
remain consistent throughout the whole video.

  * •

Background Consistency evaluates the temporal consistency of the background
scenes.

  * •

Motion Smoothness evaluates whether the motion in the generated video is
smooth and follows the physical law of the real world.

  * •

Dynamic Degree evaluates the degree of dynamics by calculating average optical
flow on each video frame.

(3) Semantics evaluates the video’s adherence to the text prompt given by the
user. consistency.

  * •

Scene Consistency evaluates whether the video is consistent with the intended
scene described by the text prompt.

  * •

Overall Consistency reflects both semantics and style consistency of the
video.

We utilize three prompt sets provided by official github repository of VBench.
We generate one video for each prompt for evaluation.

  * •

subject_consistency.txt: include 72 prompts, used to evaluate subject
consistency, dynamic degree and motion smoothness.

  * •

overall_consistency.txt: include 93 prompts, used to evaluate overall
consistency, aesthetic quality and imaging quality.

  * •

scene.txt: include 86 prompts, used to evaluate scene and background
consistency.

###  B.2 Selected Metrics

FVD and FVD-FP16: FVD measures the similarity between the distributions of
features extracted from real and generated videos. We employ one randomly
selected video per label from the UCF-101 dataset (101 videos in total) as the
reference ground-truth videos for FVD evaluation. We follow [2] to use a
pretrained I3D model to extract features from the videos. Lower FVD scores
indicate higher quality and more realistic video generation. However, due to
relatively smaller video size (e.g. 101 videos in our case), employing FVD to
evaluate video generation models faces several limitations. Small sample size
cannot adequately represent either the diversity of the entire dataset or the
complexity and nuances of video generation, leading to inaccurate and unstable
results. To mitigate limitations above, we propose an enhanced metric, FVD-
FP16, for assessing the semantic loss in videos generated by quantized models
relative to those produced by pre-quantized models. Specifically, we utilize
101 videos generated by the FP16 model as ground-truth reference videos. The
FVD-FP16 has significantly higher correlation with human perception.

CLIPSIM and CLIP-temp: The CLIPSIM and CLIP-temp metrics are computed using
implementation from EvalCrafter [33]. For CLIPSIM, We use the CLIP-VIT-B/32
model [47] to compute the image-text CLIP similarity for all frames in the
generated videos and report the averaged results. The metric quantify the
discrepancy between input text prompts and generated videos. For CLIP-temp, we
use the same model to compute the CLIP similarity of each two consecutive
frames of the generated videos and then get the averages on each two frames.
The metric indicates semantics consistency of generated videos.

DOVER’s VQA: We employ the Dover [64] method to assess generated video quality
in terms of aesthetics and technicality. The technical rating(VQA-T) measures
common distortions like noise, blur and over-exposure. The aesthetic
rating(VQA-A) reflects aesthetic aspects such as the layout, the richness and
harmony of colors, the photo-realism, naturalness, and artistic quality of the
frames.

Flow Score: We employ flow score proposed by [33] to measure the general
motion information of the video. we use RAFT [58], to extract the dense flows
of the video in every two frames. Then, we calculate the average flow on these
frames to obtain the average flow score of each generated video.

Temporal Flickering: We utilize the temporal flickering score provided by
VBench [24] to measure temporal consistency at local and high-frequency
details of generated videos. We calculate the average MAE(mean absolute
difference) value between each frame.

##  Appendix C Detailed Analysis of Experimental Results

In this section, we present more detailed analysis of the experimental results
in Sec. 4.

Figure 11: The radar chart corresponding to the data presented in Table 1 from
Sec. 4.1. ViDiT-Q has a superior performance on VBench compared with the naive
PTQ.

###  C.1 Text-to-Video Performance on VBench

VBench is a comprehensive benchmark suite for video generation models,
covering a wide range of dimensions, such as motion smoothness and subject
consistency. The metric values of ViDiT-Q’s performance on VBench is presented
in Tab. 1 Sec. 4. We visualize the Radar plot of the VBench performance in
Fig. 11, the metric values are normalized by the maximum value in each
diemsnion. It’s clearly illustrated that ViDiT-Q achieves similar performance
with FP16 for all bit-widths (W8A8 , W6A6 mixed precision, W4A8 mixed
precision), outperforming the Naive PTQ W8A8. We further analyze the generated
video’s performance from three aspects as follows:

Dynamic Degree:  Dynamic degree indicates the range of motion in the video,
higher dynamic degree denotes more dynamic movement in the video. Lower
dynamic degree denotes that the video barely moves, resembling a static image.
Normally, higher dynamic degree is favored. However, in the quantization
scenario, we discover that quantization often causes the generated videos to
jitter and tremble. It is not favorable but results in notable dynamic degree
value increase. In our experimental setting, too high or too low dynamic
degree means degradation. Therefore, in the radar plot, using FP16 generated
videos as the ground-truth reference, we use the $(f_{Q}-f_{FP})/f_{FP}$ to
denote “relative dynamic degree changes from FP generated videos”, and use
$1-(f_{Q}-f_{FP})/f_{FP}$ as dynamic degree scoring in the radar plot. As
illustrated Fig. 11 dynamic degree dimension, Naive PTQ W8A8’s scoring
($<0.8$) is notably lower than ViDiT-Q results. The video examples in Fig. 12
supports this finding. In Fig. 12(c), the navive PTQ W8A8 generated buildings
have jittering and glitches, and changes significantly across frames (ref the
supplementary for the video). In contrast, both the FP16 and ViDiT-Q W8A8
generated buildings moves acutely.

Consistency:  The consistency denotes whether some object remains consistent
(does not disappear, change significantly) across frames. Vbench evaluates
consistency from the subject, scene, background, and overall level. From the
Radar plot, we witness ViDiT-Q also notably outperforms naive PTQ, especially
in the “scene consistency” dimension ($<0.8$). As seen in the aforementioned
video example in Fig. 12(c), the buildings (act as the “scene”) changes
significantly across frames. It violates the scene consistency and lead to
lower scoring. Also, as presented in Fig. 13(c), the generated bear’s ear does
not exist in earlier frames, and suddenly appears. This also reflects the
degradation of subject consistency.

Quality:  VBench evaluates the quality from both the aesthetic (composition
and color), and imaging quality (clarity, exposure) dimension. Fig. 14 shows
the example of Naive PTQ W8A8’s quality degradation. The color notably turns
blue, and the mountain on the left is blurred. Similar color shifting
degradation is also witnessed in Fig. 13(c).

(a) FP16

(b) ViDiT-Q: W8A8

(c) Naive PTQ: W8A8

Figure 12: The qualitative results on VBench about the ViDiT-Q’s ability to
maintain the dynamic degree.

(a) FP16

(b) ViDiT-Q: W8A8

(c) Naive PTQ: W8A8

Figure 13: The qualitative results on VBench about the ViDiT-Q’s ability to
maintain the consistency.

(a) FP16

(b) ViDiT-Q: W8A8

(c) Naive PTQ: W8A8

Figure 14: The qualitative results on VBench about the ViDiT-Q’s ability to
maintain the image quality.

###  C.2 Comparison with other quantization methods

In this section, we provide more detailed analysis of comparison with other
quantizatio methods in Table. 3. We also visualize the radar plot in Fig. 15
of the statistics within the table. Similar to the abovementioned “dynamic
degree” in VBench. The “flow Score” metric measures video dynamics. Therefore,
we adopt similar $1-(f_{Q}-f_{FP})/f_{FP}$ for the flow score metric. For all
bit-width. ViDiT-Q achieves similar performance with FP16, while baseline
methods notably degrades or even fails. We will further elaborate the
performance from each dimension with the assistance of example videos in Fig.
16.

(a) W8A8

(b) W6A6

(c) W4A8

Figure 15: The radar chart corresponding to the data presented in Table 3 from
Sec. 4.2. We compare ViDiT-Q with other quantization method designed for LLMs
and UNet based diffusion models.

CLIP-Temp and CLIPSIM: CLIP-temp measures the similarity between frames.
Despite baseline quantization methods faces subject and scene consistency
degradation between frames, the main content of the frame does not change
much. Therefore, we witness similar CLIPSIM score (close to 1.) for all
methods. CLIPSIM measures the text-video alignment, which represents how well
the generated video follows the text instruction. Except for some extreme
cases (Naive,PTQ4DM for W6A6,W8A8), which fail to generate reasonable
contents, the CLIPSIM score remain similar.

Flow Score: Flow Score measures the dynamic degree of the video. We witness
both higher and lower flow score for quantized model. Some videos witness
flickering and glitches similar to Fig. 13(c), and others have reduced motion
(Refer to the videos in the supplementary for visual effect). The W6A6 Naive
quantization generates pure noise, it significantly changes over frames and
results in abnormally large Flow Score values.

VQA:  The VQA metric we adopted evaluates video quality from two perspectives:
the aesthetic and the technical (similar to VBench). As could be seen from
Fig. 15, ViDiT-Q achieves VQA-A and VQA-T similar to FP16, while other
baseline methods notably degrades. It’s consistent with the visual perception
from examples in Fig. 16. Naive PTQ could only generate readable content in
W8A8, it generates pure noise for W6A6, and blank image for W4A8. PTQ4DM
generates turtle that are hardly recognizable for W8A8 and W6A6, and also
produce nearly blank images for W4A8. For improved quantization techniques
Q-Diffusion and SmoothQuant, although the main object (turtle) is
distinguishable, its movement is still unnatural and violates physics.

Figure 16: Example videos of different quantization methods for different bit-
widths.

###  C.3 Text-to-image Generation on COCO

We present more qualitative results of generated images by baseline
quantization and ViDiT-Q quantization in Fig. 17(a). As shown, the Naive PTQ’s
generated images are highly blurred. While the W8A8 images depict outlines of
objects, the W4A8 images generate nearly pure noises. In contrast, ViDiT-Q
generates images nearly identical to the FP16 ones, preserving both visual
quality and text-image alignment.

(a) Figure 17: Qulitative results of text-to-image generation

(b)

###  C.4 Ablation Studies

We present the generated videos for the ablation studies in Tab. 4. As seen in
Fig. 18, video quality improves from blank images to similar to the FP16
baseline. For the challenging W4A8 quantization, the baseline method generates
blank images. After adding dynamic quantization, some meaningful background
(deep ocean) appears, but the main object (turtle) is still missing. Channel
balancing reduces color deviation (from dark blue to green-blue), but the main
object remains unrecognizable and changes significantly across frames (please
refer to the supplementary materials for the video). Timestep-wise channel
balancing improves the consistency of the main object, but notable degradation
is still observed compared to the FP16 video. Finally, with mixed precision, a
similar generation quality to the FP16 baseline is achieved."

Figure 18: Generated videos of ablation studies.

We also present the comparison of quantization with naive channel balance and
timestep-aware channel balance. As seen in Fig. 19, The main object in text
instruction (“turtle” and “wispy grasses”) in the left images are hardly
recognizable. After introducing timestep-aware channel balancing, notable
improvements are witnessed.

Figure 19: Comparison of generated videos with naive and timestep-aware
channel balancing.


# GPU Offloading in ExaHyPE Through C++ Standard Algorithms

Uzmar Gomez  Department of Computer Science, Durham University, United Kingdom
[ Gonzalo Brito Gadeschi  NVIDIA, Munich, Germany  Tobias Weinzierl
Department of Computer Science, Durham University, United Kingdom

###### Abstract

The ISO C++17 standard introduces _parallel algorithms_ , a parallel
programming model promising portability across a wide variety of parallel
hardware including multi-core CPUs, GPUs, and FPGAs. Since 2019, the NVIDIA
HPC SDK compiler suite supports this programming model for multi-core CPUs and
GPUs. ExaHyPE is a solver engine for hyperbolic partial differential equations
for complex wave phenomena. It supports multiple numerical methods including
Finite Volumes and ADER-DG, and employs adaptive mesh refinement with dynamic
load balancing via space-filling curves as well as task-based parallelism and
offloading to GPUs. This study ports ExaHyPE’s tasks over blocks of Finite
Volumes to the ISO C++ parallel algorithms programming model, and compares its
performance and usability against an OpenMP implementation with offloading via
OpenMP target directives. It shows that ISO C++ is a feasible programming
model for non-trivial applications like our task-based AMR code. The
realisation is bare of vendor-specific or non-C++ extensions. It however is
slower than its OpenMP counterpart.

Uzmar<EMAIL_ADDRESS>
##  1 Introduction

Graphics Processing Units (GPUs) have become the main accelerator for data-
parallel processing tasks driving Exascale systems [1]. Yet, porting
applications to GPUs and designing GPU algorithms remains challenging. The
main HPC programming paradigms for GPU-accelerated systems are ISO standard
languages (C++ and Fortran) that run on GPUs “automatically”, annotation-based
programming models like OpenMP and OpenACC, and GPU programming languages or
extensions like CUDA or SYCL. The different approaches have pros and cons.

The NVIDIA HPC SDK supports automatic offloading of ISO C++ _parallel
algorithms_ to multi-core CPUs and GPUs [2]. Recent publications demonstrate
good performance for mini-applications [3] and NVIDIA’s documentation enlists
success stories for larger applications, such as the multi-physics
applications m-AIA [4] and STLBM [5].

The goal of this study is to examine the feasibility of parallelizing a large
application to portably target multi-core CPU and multi-GPU systems using
plain ISO C++. ExaHyPE is a solver engine for hyperpolic PDE equations [6] for
complex wave phenomena. It is a large and complex application with support for
adaptive mesh refinement with dynamic load balancing via space-filling curves.
On top of this MPI+X parallelism for the CPU, it can phrase the updates of
individual cells as tasks: We restrict ourselves to the computational kernels
of ExaHyPE’s Finite Volume solver for small Cartesian patches in a block-
structured AMR context [7]—these are embedded into the cells and hance form
tasks—for which pre-existing OpenMP offloading and SyCL versions are
available. More complex offloading strategies—notably porting execution logic
to a GPU—are beyond scope.

The main contributions of this work are discovering how to map individual
tasks to the new ISO C++ _parallel algorithms_ programming model, and
comparing the productivity, usability, and performance of this programming
model against OpenMP. We also report on the technical challenges we
encountered using this novel programming model for offloading ISO C++ to GPUs.

This study is organized as follows. Section 2 describes the problem domain,
the numerical scheme, and the ExaHyPE software. Section 3 discusses the core
compute kernel C++ implementation, GPU execution policies, and implementation
challenges. Section 4 shows our preliminary results from which we draw our
conclusions.

##  2 Block-structured Finite Volumes in ExaHyPE

ExaHyPE is an engine to write simulation codes for hyperbolic, first-order
equation systems over functions
$\vec{q}\left(\vec{x},t\right)=\left(q_{1},\dots,q_{s}\right)$

$\frac{\partial\vec{q}}{\partial t}+\sum_{j=1}^{d}\frac{\partial}{\partial
x_{j}}\vec{f^{j}}(\vec{q})=0,$ (1)

where $\vec{x}\in\mathbb{R}^{d}$ and the directional terms $\vec{f^{j}}$ form
the flux function. In ExaHyPE, users adopt such generic PDE formulations by
injecting function implementations evaluating $f^{j}(\vec{q})$. How these
function evaluations are used within a numerical scheme, on which core, and
when, is hidden from the users. This is the Hollywood principle from [8]. For
the present experiments, we benchmark the Euler equations. In conservative
form, these equations read as

$\frac{\partial}{\partial t}\begin{pmatrix}\rho\\\ \vec{j}\\\
E_{t}\end{pmatrix}+\nabla\cdot\begin{pmatrix}\vec{j}\\\
\frac{1}{\rho}\vec{j}\otimes\vec{j}+pI\\\
\left(E_{t}+p\right)\frac{1}{\rho}\end{pmatrix}=\vec{0},$ (2)

where $\vec{j}=\rho\vec{u}$ is the momentum density, $\rho$ is the fluid mass
density, $p$ is the pressure, and $E_{t}$ is the total energy density.

ExaHyPE’s spatial discretisation is fixed: The code employs a block-structured
mesh [9], resulting from spacetrees [8]: A cubic computational domain is split
into three chunks along each coordinate axis and this process then is repeated
locally and recursively. We end up with a dynamically adaptive grid of cubes.
On top of this spatial discretisation, ExaHyPE offers various numerical
schemes. We stick to block-structured Finite Volumes here: Each cube of the
grid hosts a small Cartesian patch of $p\times p$ or $p\times p\times p$
volumes. Each patch is supplemented with a halo layer.

It is ExaHyPE’s responsibility to maintain, store and traverse the block-
structured mesh. Per Cartesian patch, the code invokes an update kernel to
advance a patch in time. This kernel is tailored towards the PDE of interest
by the user-provided flux functions. After that, it is ExaHyPE’s
responsibility to keep all (halo) data consistent.

### Parallelisation and GPU offloading

ExaHyPE employs a traditional non-overlapping domain decomposition realised
through MPI. Per MPI rank, it then applies a non-overlapping domain
decomposition again. The resulting subdomains are distributed per rank among
the threads.

Each thread runs through its subdomain per time step and advances the
numerical solution. Throughout the traversal, the code prioritises patches
along AMR boundaries or the subdomain boundaries: These patches are updated
immediately while we run throught the domain, as their output is urgent as it
feeds into domain boundary exchange or is involved in potentially expensive
AMR routines.

The remaining patches are held back and put into a thread-local buffer.
Whenever this buffer exceeds a given threshold $N$, we take the $N$ patches
and deploy them as a whole to an accelerator. All remaining patches are
deployed to normal tasks at the end of a traversal to exploit the host cores.
The classification of patches realises the principle of enclave tasking [10],
while the buffering uses a concept from [7].

### A generic Finite Volume kernel

Per volume $\mathcal{C}_{i}$ within each $p\times p\times p$ patch, we use a
Finite Difference approximation of the time derivative in (1) subject to a
Finite Volume scheme for the spatial discretisation:

$\frac{Q_{i}^{n+1}-Q_{i}^{n}}{\Delta
t}+\frac{F_{i+1/2}^{n}-F_{i-1/2}^{n}}{\Delta x}=0,$ (3)

where $F_{i-1/2}^{n}$ is the flux over the faces of $\mathcal{C}_{i}$.
Subsequently, we use Rusanov’s approximation to the flux from (3):

$F_{i-1/2}^{n}=\frac{1}{2}\left[f\left(Q_{i-1}^{n}\right)+f\left(Q_{i}^{n}\right)-a_{i-1/2}\left(Q_{i}^{n}-Q_{i-1}^{n}\right)\right]$
(4)

where

$\displaystyle a_{i-1/2}$ $\displaystyle=$
$\displaystyle\max\left(\left|f^{\prime}(q)\right|\right)\ \text{over all
}q\text{ between }Q_{i-1}^{n}\text{ and }Q_{i}^{n}$ (5) $\displaystyle=$
$\displaystyle\max\left(\left|\lambda_{k}\right|\right).$

$\lambda_{k}$ are the eigenvalues of the PDE’s Jacobian matrix.

With this generic numerical scheme, we can implement a generic patch update
kernel which is tailored towards (2) by setting the number of quantities of
interest to $d+2$ and by injecting two routines: one evaluates Euler’s flux,
the other one returns the maximum eigenvalue.

##  3 A C++ GPU computer kernel

Algorithm 1 Compute steps for one patch. A kernel for multiple patches in one
rush adds an additional outer loop.

1:$Q^{\text{new}}=Q^{\text{old}}$

2:Compute $a_{i-1/2}Q^{\text{old}}$ where $a_{i-1/2}$ is given by eq. 5.

3:$a=\max\left(\lambda_{\max}Q^{\text{old}}\right)$ for face in boundary of
$Q^{\text{old}}$.

4:Update the volume by $a\left(Q^{\text{old}}-Q^{\text{old-1}}\right)$.

5:Compute the flux given by eq. eq. 4.

6:Compute the average of the fluxes between $Q^{\text{old}}$ and the adjacent
volumes.

7:Update the volume with the fluxes from their 2d adjacent faces subject to a
scaling with $\frac{\Delta t}{\Delta x}$.

8:Compute the Finite Volume’s non-conservative product.

9:Update the value of $Q^{\text{new}}$.

10:Compute the patch global maximum eigenvalue.

ExaHyPE offers a generic compute kernel implementation based upon functors,
but we also have a templated version handling $N$ patches in one batch. It
consists of a sequence of nested loops (Algorithm 1). We run over all patches.
Per patch, we execute different steps, whereas each step consists of loops
over cells or faces, respectively, as well as the $s=5$ unknowns of the PDE.

⬇

// Detecting the cartesian product type automatically

using cart_prod_3d = decltype(tl::views::cartesian_product(

std::views::iota(0,1), std::views::iota(0,1),

std::views::iota(0,1)

));

// Allocating in heap

auto* range2d_cell = new cart_prod_3d;

// Defining components

*range2d_cell = tl::views::cartesian_product(

std::views::iota(0,numberOfVolumesPerAxisInPatch),

std::views::iota(0,numberOfVolumesPerAxisInPatch),

std::views::iota(0,hostPatchData.numberOfCells)

);

...

delete[] range2d_cell;

Code 1: Allocation of a Cartesian product (iteration space) on the heap. This
way, it can be used by the GPU.

### SIMD loop traversals and functors

⬇

std::for_each(

std::execution::par_unseq,

range2d_cell->begin(), range2d_cell->end(),

[

=,

QIn=hostPatchData.QIn,

cellCentre=hostPatchData.cellCentre,

cellSize=hostPatchData.cellSize,

t=hostPatchData.t, dt=hostPatchData.dt,

QOut=hostPatchData.QOut

](auto ids){

auto [x,y,patchIndex] = ids;

// function call has to refer to a template, i.e.

// implementation has to be in header. Otherwise,

// compiler does not know that it has to create a

// GPU variant of function.

// step (1) from algorithm blueprint

internal::copySolution_LoopBody<Solver>(

QIn[patchIndex], QInEnumerator,

cellCentre[patchIndex], cellSize[patchIndex],

patchIndex, volumeIndex2d(x,y),

t[patchIndex], dt[patchIndex],

QOut[patchIndex], QOutEnumerator

);

[...]

// steps (2--10) from algorithm blueprint

}

);

Code 2: Standard std::for_each algorithm offloaded to GPU.

The original kernel implementation used virtual-function dynamic dispatch and
separately compiled implementations - i.e. in separate .cpp files - to reduce
compilation times and improve developer productivity. However, the NVIDIA
implementation of ISO C++ does not support virtual-function dispatch in device
code and separate compilation. We hence replace virtual-function dispatch in
ExaHyPE with template functions defined in header files and static dispatch
(Code LABEL:foreachcode). In the case of the maximum in (Algorithm 5), we
replace the standard C++ function invocation by a manual comparison.

While the NVIDIA implementation supports separate compilation via non-standard
CUDA-like __host__ __device__ annotations, these require the programmer to
choose between using a non-standard extension or moving implementation details
to header files, which significantly increases compile-times. A manual
implementation of maximum is not “genuine C++”.

### Multi-dimensional iteration spaces with Cartesian Product

Efficient GPU kernels require large iteration spaces of collapsed loops to
harvest all parallelism available in modern accelerators. Hence, a single C++
parallel algorithm that collapses several for loops is built in our Code
LABEL:allocatingcode. It combines multiple std::views::iota structures using
C++23’s std::views::cartesian_product. This implementation takes pointers to
the host-thread stack. However, according to the HPC SDK documentation [11],
only the host heap can be accessed from GPU kernels on the hardware platform
used in this study. As a consequence, our code LABEL:allocatingcode had to
allocate the underlying Cartesian product on the heap.

The absence of an implementation of Cartesian product in the C++ standard
library used in this study requires us to include a preliminary implementation
that is not tuned for accelerators. The heap placement is solely due to
technical constraints.

### Data transfer

The ISO C++ programming model relies on unified memory to enable the
accelerator devices to access all program memory. The NVIDIA implementation is
non-conforming on non-coherent hardware platforms like the one used in this
study: devices are only allowed to access host heap memory. While ExaHyPE’s
data is mainly stored on the heap, pointers to this data are conveniently
aggregated in structs like hostPatchData, which are often stored on the host’s
stack. Device code is not allowed to access these values to load pointers to
heap data. ExaHyPE hence captures these pointers explicitly in the lambdas
passed to the parallel algorithms. Lambda captures are copied to the GPU on
kernel launch, while the heap data these pointers refer to is migrated on-
demand as the GPU threads access it. This data stays on the GPU until the CPU
access it again, in which case it is automatically migrated back. There is no
portable C++ API for prefetching heap data between the host and the device.

### Software stack

Our work uses the NVIDIA HPC SDK 22.3 which leverages the system’s C++
standard library (libstdc++) but uses its own _C++ parallel algorithms_
implementation based on the Thrust library. Many HPC systems do not default to
a sufficiently new version of libstdc++, which requires us to manually pick a
different one via compiler flags. We reported some genuine nvc++ bugs - e.g.
[12] \- which will be solved in future releases.

##  4 Results

Figure 1:  Time per Finite Volume update for different implementation variants
and choices of the number of patches $N$ handled per kernel invocation.

Fig. 1 shows the benchmark results of our compute kernels for $p=17$ in a 2d
setup on a NVIDIA 2080TI for different sizes of patch sets $N$. All timings
include all data transfers between host and device. The kernels over $N$
patches are available as patch-wise implementation where the outermost
Cartesian loop runs over the patches, and a batched version, where we run
through the individual steps of Algorithm 1 and invoke a loop over all
patches, all volumes, and all unknowns per step.

Both kernel variants exploit the symmetries from (4): The face flux enters the
left and right adjacent volumes’ solutions. We compute $f$ (as well as the
eigenvalues) per volume and determine the linear combination of these values
from (4) redundantly for the left and right adjacent cell. The intermediate
$f$ and $\lambda$ values can be held as Array of Structures (AoS), Structure
of Arrays (SoA) or Array of Structures of Arrays (AoSoA).

The batched vs. patch-wise flavours result from a permutation of the loops
within the Cartesian product (Code LABEL:allocatingcode). However, our AMR
code employs local and adaptive timestepping and hence requires the maximum
eigenvalue per patch as additional output per kernel invocation. We need a
per-patch reduction. This reduction is performed by using a sequential
std::for_each loop to avoid data races, but it could be parallelized by using
another standard algorithm or atomic memory operations.

Our results show that the batched variant is always faster than the patch-wise
variant for the C++ implementations. It is advantageous to split an algorithm
into its steps and to use a large iteration space per step. As the time per
Finite Volume update decreases with the number of patches offloaded, we may
assume that the GPU kernel does a good job in keeping the data in-between two
steps on the GPU. Data data layout of temporary data within the computation
has limited performance impact.

For OpenMP, we see that it can be advantageous to run over the patches one-by-
one. We assume that different patches in this case end up on different SMs on
the GPU, and run more or less independent of each other. Overall, the OpenMP
port is significantly faster—notably for small $N$-counts—than the C++ port.
NSight studies suggest that this is due to a bespoke OpenMP memory management,
where we have explicit control via device pointers when data are offloaded.
The C++ version in contrast has the on-demand memory transfers baked into the
implementation.

Though our software is ready to allow multiple tasks to offload to the GPU
simultaneously, neither C++ nor OpenMP did support this feature. Both ran into
segmentation faults.

##  5 Outlook and conclusion

GPU offloading via C++ execution policies within C++ standard algorithms is,
in principle, feasible for non-trivial test cases as our task-based AMR code.
The resulting realisation is bare of specific (OpenMP) extensions and clean;
though this is a subjective assessment. Performance-wise, the C++ port is
slower than its OpenMP counterpart. It is hard to pinpoint a reason for this
behaviour. The two-dimensional Euler equations have a very low arithmetic
intensity. The performance discrepancies might disappear for more complex
calculations such as the ADER-DG ExaHyPE solvers; a phenomenon to be studied
in the future.

It is early days to comment on the pros and cons of C++ GPU offloading.
Besides obvious teething problems, our work highlights particular challenge
for developers: While C++ may be syntactically and semantically correct, its
GPU transfer might fail (as data is not allocated “artifically” on the heap,
e.g.), or conform by running the code sequentially, violating the intention of
the programmer to offload computations to an accelerator. On top of this, the
C++ community needs a portable alternative to NVIDIA’s vendor-specific
__host__ __device__ annotations for separate compilation. In the ideal case,
multi-pass compilation or link-time compilation for GPUs would make these
annotations unnecessary.

As alternative to this automatisation, the C++ language might supplement
existing execution policies with declarative GPU offloading policies. Using
such policies would render otherwise correct loop bodies invalid if they do
not fit to the GPU. It is not clear how this is properly reflected within the
language, or exclusively should be handled by the compiler throughout
translation.

## References

  * [1] Volodymyr V. Kindratenko, Jeremy J. Enos, Guochun Shi, Michael T. Showerman, Galen W. Arnold, John E. Stone, James C. Phillips, and Wen-mei Hwu.  Gpu clusters for high-performance computing.  In 2009 IEEE International Conference on Cluster Computing and Workshops, pages 1–8, 2009.
  * [2] David Olsen.  C++17 parallel algorithms on nvidia gpus with pgi c++.  In GPU Technology Conference, 2019.
  * [3] Wei-Chen Lin, Tom Deakin, and Simon McIntosh-Smith.  Evaluating iso c++ parallel algorithms on heterogeneous hpc systems.  In 13th IEEE International Workshop on Performance Modeling, Benchmarking and Simulation of High Performance Computer Systems, 2022.
  * [4] European Exascale Computing Center of Excellence.  Research on AI- and Simulation-Based Engineering at Exascale.  https://www.hpccoe.eu/2021/06/04/m-aia/, 2021.  Last updated: Sep 1, 2022.
  * [5] Jonas Latt, Christophe Coreixas, and Joël Beny.  Cross-platform programming model for many-core lattice boltzmann simulations.  PLOS ONE, 16(4):e0250306, apr 2021.
  * [6] Anne Reinarz, Dominic E. Charrier, Michael Bader, Luke Bovard, Michael Dumbser, Kenneth Duru, Francesco Fambri, Alice-Agnes Gabriel, Jean-Matthieu Gallard, Sven Köppel, Lukas Krenz, Leonhard Rannabauer, Luciano Rezzolla, Philipp Samfass, Maurizio Tavelli, and Tobias Weinzierl.  Exahype: An engine for parallel dynamically adaptive simulations of wave problems.  Computer Physics Communications, 254:107251, 2020.
  * [7] Han Zhang, Tobias Weinzierl, Holger Schulz, and Baojiu Li.  Spherical accretion of collisional gas in modified gravity I: self-similar solutions and a new cosmological hydrodynamical code.  Monthly Notices of the Royal Astronomical Society, 515(2):2464–2482, 07 2022.
  * [8] T. Weinzierl.  The peano software—parallel, automaton-based, dynamically adaptive grid traversals.  ACM Transactions on Mathematical Software, 45(2):14, 2019.
  * [9] A. Dubey, A. S. Almgren, J. B. Bell, M. Berzins, S. R. Brandt, G. Bryan, P. Colella, D. T. Graves, M. Lijewski, F. Löffler, B. O’Shea, E. Schnetter, B. van Straalen, and K. Weide.  A survey of high level frameworks in block-structured adaptive mesh refinement packages.  Journal of Parallel and Distributed Computing, 74(12):3217–3227, 2016.
  * [10] Dominic Etienne Charrier, Benjamin Hazelwood, and Tobias Weinzierl.  Enclave tasking for dg methods on dynamically adaptive meshes.  SIAM Journal on Scientific Computing, 42(3):C69–C96, 2020.
  * [11] NVIDIA HPC Compilers. C++ Parallel Algorithms.  https://docs.nvidia.com/hpc-sdk/compilers/c++-parallel-algorithms/index.html, 2022\.
  * [12] NVIDIA Bug report.  https://nvbugs/3706250, 2022.

##  Appendix A OpenMP reference implementation

⬇

const int numberOfCells = hostPatchData.numberOfCells;

double* rawPointerToCellCentres = hostPatchData.cellCentre[0].data();

double* rawPointerToCellSizes = hostPatchData.cellSize[0].data();

double* t = hostPatchData.t;

double* dt = hostPatchData.dt;

double* maxEigenvalue = hostPatchData.maxEigenvalue;

#pragma omp target enter data
map(to:rawPointerToCellCentres[0:hostPatchData.numberOfCells*Dimensions])
device(targetDevice)

#pragma omp target enter data
map(to:rawPointerToCellSizes[0:hostPatchData.numberOfCells*Dimensions])
device(targetDevice)

#pragma omp target enter data
map(to:mappedPointersToQIn[0:hostPatchData.numberOfCells])
device(targetDevice)

#pragma omp target enter data
map(to:mappedPointersToQOut[0:hostPatchData.numberOfCells])
device(targetDevice)

#pragma omp target enter data map(to:t[0:hostPatchData.numberOfCells])
device(targetDevice)

#pragma omp target enter data map(to:dt[0:hostPatchData.numberOfCells])
device(targetDevice)

#pragma omp target enter data
map(alloc:maxEigenvalue[0:hostPatchData.numberOfCells]) device(targetDevice)

[...]

#pragma omp target exit data
map(delete:rawPointerToCellCentres[0:hostPatchData.numberOfCells*Dimensions])
device(targetDevice)

#pragma omp target exit data
map(delete:rawPointerToCellSizes[0:hostPatchData.numberOfCells*Dimensions])
device(targetDevice)

#pragma omp target exit data
map(delete:mappedPointersToQIn[0:hostPatchData.numberOfCells])
device(targetDevice)

#pragma omp target exit data
map(delete:mappedPointersToQOut[0:hostPatchData.numberOfCells])
device(targetDevice)

#pragma omp target exit data map(delete:t[0:hostPatchData.numberOfCells])
device(targetDevice)

#pragma omp target exit data map(delete:dt[0:hostPatchData.numberOfCells])
device(targetDevice)

#pragma omp target exit data
map(from:maxEigenvalue[0:hostPatchData.numberOfCells]) device(targetDevice)

Code 3: Allocation and deallocation of variables on the device for OpenMP.

Our OpenMP vanilla version hosts the data of $N$ (number of cells) patches
within a struct on the host. Data here comprises the actual Finite Volume data
as well as meta data such as cell (patch) sizes and positions.

The OpenMP variant maps these data explicitly onto the GPU prior to the kernel
invocation (Code LABEL:allocatingcodeopenmp). An analogous loop over the
individual patches transfers the Finite Volume data per patch (now shown).

⬇

#pragma omp target teams distribute parallel for simd collapse(3)
device(targetDevice)

for (int patchIndex=0; patchIndex<numberOfCells; patchIndex++)

for (int y = 0; y < numberOfVolumesPerAxisInPatch; y++)

for (int x = 0; x < numberOfVolumesPerAxisInPatch; x++) {

GPUCellData patchData;

patchData.numberOfCells = numberOfCells;

patchData.cellCentre =
(tarch::la::Vector<Dimensions,double>*)(rawPointerToCellCentres);

patchData.cellSize =
(tarch::la::Vector<Dimensions,double>*)(rawPointerToCellSizes);

patchData.QIn = mappedPointersToQIn;

patchData.QOut = mappedPointersToQOut;

patchData.t = t;

patchData.dt = dt;

patchData.maxEigenvalue = maxEigenvalue;

[...]

copySolution_LoopBody<Solver>(

patchData.QIn[patchIndex],

QInEnumerator,

patchData.cellCentre[patchIndex],

patchData.cellSize[patchIndex],

patchIndex,

volumeIndex2d(x,y),

patchData.t[patchIndex],

patchData.dt[patchIndex],

patchData.QOut[patchIndex],

QOutEnumerator

);

[...]

}

Code 4: Main loop body for the OpenMP code. Equivalent to Code
LABEL:foreachcode.

Our OpenMP reference implementation realises the nested Cartesian iteration
space via pure nested for loops (Code LABEL:forloopopenmp). They are collapsed
via a pragma. In this sense, the realisation is closer to C than its C++ GPU
offloading counterpart.

Within the loop body, the code reconstructs the host struct manually using
device pointers, i.e. it creates a struct counterpart on the GPU and befills
its translated GPU pointers manually. Alternatively, a realisation could have
supplemented the struct hosting the $N$ patch data with an OpenMP mapping.

The function copySolutionAndAddSourceTerm_LoopBody and other compute step
realisations from Algorithm 1 are templates and therefore found in headers.
Consequently, the compiler can directly inline them and deduce from the
context that they have to be translated for the GPU.

##  Appendix B Test environment

We provide a docker image to reproduce the results obtained in this paper. The
custom image is based on the latest HPC SDK compiler, and it includes some
Python libraries and other necessary software to run ExaHyPE and its
underlying Peano AMR framework. If the software is going to be run on a GPU,
use the following command to pull the image and start a bash terminal

⬇

docker run \--gpus all -it ghcr.io/uzmargomez/peano:hpcsdk22.9 bash

Without access to a GPU, the \--gpus all flag has to be omitted. In the
container’s terminal, we next run the commands

⬇

cd /

git clone https://gitlab.lrz.de/hpcsoftware/Peano.git && cd Peano && git
checkout gpus && git checkout c57ef3b5

export PYTHONPATH=/Peano/python

libtoolize && aclocal && autoconf && autoheader && cp src/config.h.in . &&
automake \--add-missing

./configure \--enable-exahype \--enable-blockstructured \--enable-
loadbalancing \--with-gpu=cpp \--with-multithreading=cpp CXX=$NVCPP CC=$NVCPP
CXXFLAGS="-O4 -std=c++20 -stdpar=gpu -gpu=cc75" LDFLAGS="-stdpar=gpu
-gpu=cc75"

make -j4

cd benchmarks/exahype2/euler/fv-kernel-benchmarks

python3 euler-fv-rusanov-kernel-benchmarks.py -a 1e-4

make -j8

./fv-kernel-benchmarks

If a GPU is not available, the HPC SDK compiler can use a multicore CPU to
parallelize standard algorithms. In this case, the image’s application
benchmarks solely the CPU kernel performance. In order to do this, we change
the flags -stdpar=gpu to -stdpar=multicore.