CUDA 11功能清單
基於NVIDIA Ampere GPU架構的新型NVIDIA A100 GPU在加速計算方面實現了最大的飛躍。A100 GPU具有革命性的硬件功能，CUDA 11與A100一起發布。
CUDA 11能夠利用新的硬件功能來加速HPC，基因組學，5G，渲染，深度學習，數據分析，數據科學，機器人技術以及更多不同的工作負載。
CUDA 11包含了所有功能-從平台系統軟件到入門和開發GPU加速的應用程序所需的一切。本文概述了此版本中的主要軟件功能：
- 支持NVIDIA Ampere GPU架構，包括新的NVIDIA A100 GPU，用於加速擴展和擴展AI和HPC數據中心；具有NVSwitch結構的多GPU系統，例如DGX A100和HGX A100。
- 多實例GPU（MIG）分區功能特別有利於雲服務提供商（CSP）改善GPU利用率。
- 新的第三代Tensor內核可加快對不同數據類型（包括TF32和Bfloat16）的混合精度矩陣操作。
- 任務圖，異步數據移動，細粒度同步和L2緩存駐留控制的編程和API。
- CUDA庫中針對線性代數，FFT和矩陣乘法的性能優化。
- Nsight產品工具系列的更新，用於跟蹤，分析和調試CUDA應用程序。
- 全面支持所有主要的CPU體系結構，包括x86_64，Arm64服務器和POWER體系結構。
本文不能完全代表CUDA 11中可用的所有功能。在本文的末尾，有指向GTC Digital會議的鏈接，這些鏈接提供了對CUDA新功能的更深入研究。
CUDA和NVIDIA Ampere微架構GPU
NVIDIA Ampere GPU微架構采用台積電7nm N7制造工藝制造，包括更多的流多處理器（SM），更大，更快的內存，以及與第三代NVLink互連的帶寬，以提供巨大的計算吞吐量。
A100的40 GB（5站點）高速HBM2內存的帶寬為1.6 TB /秒，比V100快1.7倍以上。A100上的40 MB L2緩存幾乎是Tesla V100的7倍，並提供了2倍的L2緩存讀取帶寬。CUDA 11在A100上提供了新的專用L2緩存管理和駐留控制API。A100中的SM包含更大，更快的L1高速緩存和共享內存單元（每個SM 192 KB）的組合，提供的容量是Volta V100 GPU的1.5倍。
A100配備了專用的硬件單元，包括第三代Tensor內核，更多的視頻解碼器（NVDEC）單元，JPEG解碼器和光流加速器。所有這些都由各種CUDA庫使用，以加速HPC和AI應用程序。
接下來將討論NVIDIA A100中引入的主要創新，以及CUDA 11如何充分利用這些功能。無論是管理集群的平台DevOps工程師，還是編寫GPU加速應用程序的軟件開發人員，CUDA 11都能為每個人提供一些功能。有關NVIDIA Ampere GPU微體系結構的更多信息，請參閱深度NVIDIA Ampere Architecture。
多實例GPU
MIG功能可以將單個A100 GPU物理上划分為多個GPU。使多個客戶端（例如VM，容器或進程）能夠同時運行，同時在這些程序之間提供錯誤隔離和高級服務質量（QoS）。
Figure 1. New MIG feature in A100.
A100 is the first GPU that can either scale up to a full GPU with NVLink or scale out with MIG for many users by lowering the per-GPU instance cost. MIG enables several use cases to improve GPU utilization. This could be for CSPs to rent separate GPU instances, running multiple inference workloads on the GPU, hosting multiple Jupyter notebook sessions for model exploration, or resource sharing of the GPU among multiple internal users in an organization (single-tenant, multi-user).
MIG is transparent to CUDA and existing CUDA programs can run under MIG unchanged to minimize programming effort. CUDA 11 enables configuration and management of MIG instances on Linux operating systems using the NVIDIA Management Library (NVML) or its command-line interface nvidia-smi (nvidia-smi mig subcommands).
Using the NVIDIA Container Toolkit and A100 with MIG enabled, you can also run GPU containers with Docker (using the --gpus option starting with Docker 19.03) or scale out with the Kubernetes container platform using the NVIDIA device plugin.
The following command shows MIG management using nvidia-smi:
A100是第一款可以通過NVLink擴展到完整GPU或通過MIG擴展許多用戶的GPU，降低了每個GPU實例的成本。MIG支持多種用例，以提高GPU利用率。這可能是CSP可以租用單獨的GPU實例，在GPU上運行多個推理工作負載，托管多個Jupyter筆記notebook會話以進行模型探索，或者在組織中的多個內部用戶（單租戶，多用戶）之間共享GPU的資源。 。
MIG對CUDA是透明的，並且現有的CUDA程序可以在MIG下運行，而無需更改，以最大程度地減少編程工作。CUDA 11使用NVIDIA管理庫（NVML）或其命令行界面nvidia-smi（nvidia-smimig子命令）在Linux操作系統上配置和管理MIG實例。
使用啟用了MIG的NVIDIA Container Toolkit和A100，可以在Docker上運行GPU容器（使用--gpus從Docker 19.03開始的選項），或者使用NVIDIA設備插件通過Kubernetes容器平台進行橫向擴展。
以下命令顯示使用以下命令進行MIG管理nvidia-smi：
| # List gpu instance profiles: # nvidia-smi mig -i 0 –lgip |
系統軟件平台支持
為了在企業數據中心中使用，NVIDIA A100引入了新的內存錯誤恢復功能，這些功能可以提高彈性並避免影響正在運行的CUDA應用程序。先前架構上無法糾正的ECC錯誤會影響GPU上所有正在運行的工作負載，從而需要重置GPU。
在A100上，影響僅限於遇到錯誤並已終止的應用程序，而其它正在運行的CUDA工作負載則不受影響。GPU不再需要重置即可恢復。NVIDIA驅動程序執行動態頁面黑名單，以將頁面標記為不可用，以便當前和新應用程序都不會訪問受影響的內存區域。
重置GPU后，作為常規GPU / VM服務窗口的一部分，A100配備了一種稱為行重新映射的新硬件機制，該機制用備用單元替換內存中退化的單元，並避免在物理內存地址空間中造成任何漏洞。
帶有CUDA 11的NVIDIA驅動程序現在報告與帶內（使用NVML / nvidia-smi）和帶外（使用系統BMC）行重新映射有關的各種度量。A100包括新的帶外功能，包括更多可用的GPU和NVSwitch遙測，控制以及改進的GPU和BMC之間的總線傳輸數據速率。
為了提高DGX A100和HGX A100等多GPU系統的彈性和高可用性，該系統軟件支持禁用發生故障的GPU或NVSwitch節點的能力，而不是像上一代系統那樣禁用整個baseboard。
CUDA 11是第一個添加對Arm服務器的生產支持的版本。通過將Arm的高能效CPU架構與CUDA相結合，Arm生態系統將從GPU加速的計算中受益，適用於各種用例：從邊緣，雲和游戲到為超級計算機提供動力。CUDA 11支持Marvell的高性能基於ThunderX2的服務器，並與Arm和生態系統中的其他硬件和軟件合作伙伴緊密合作，以快速實現對GPU的支持。
第三代多精度Tensor核心
NVIDIA A100中每個SM的四個大型Tensor核心（總共432個Tensor核心）為所有數據類型提供了更快的矩陣乘法累加（MMA）操作：Binary，INT4，INT8，FP16，Bfloat16，TF32和FP64。
可以通過不同的深度學習框架，CUTLASS提供的CUDA C ++模板抽象或CUDA庫（例如cuBLAS，cuSOLVER，cuTENSOR或TensorRT）訪問Tensor Core 。
CUDA C ++使用warp級矩陣（WMMA）API使Tensor Core可用。這種可移植的API抽象提供了專門的矩陣加載，矩陣乘法和累加以及矩陣存儲操作，以有效地使用CUDA C ++程序中的Tensor Core。nvcuda::wmma名稱空間中提供了WMMA的所有功能和數據類型。可以使用mma_syncPTX指令直接訪問Tensor Cores for A100（即具有計算能力compute_80及更高版本的設備）。
CUDA 11增加了對新輸入數據類型格式的支持：Bfloat16，TF32和FP64。Bfloat16是另一種FP16格式，但與FP32數值范圍匹配的精度降低。使用導致較低的帶寬和存儲要求，從而提高吞吐量。Bfloat16__nv_bfloat16通過WMMA在cuda_bf16.h中作為新的CUDA C ++數據類型公開，並受各種CUDA數學庫支持。
TF32是一種特殊的浮點格式，旨在與Tensor Core一起使用。TF32包含一個8位指數（與FP32相同），10位尾數（與FP16相同的精度）和一個符號位。這是默認的數學模式，可在進行DL訓練時獲得FP32之上的加速，而無需對模型進行任何更改。最后，A100為MMA操作帶來了雙精度（FP64）支持，WMMA接口也支持該功能。
Figure 2. Table of supported data types, configurations, and performance for matrix operations.
編程NVIDIA Ampere架構GPU
為了改善GPU的可編程性並利用NVIDIA A100 GPU的硬件計算功能，CUDA 11包括用於內存管理，任務圖加速，新指令和線程通信構造的新API操作。這里介紹了其中的一些新操作以及如何能夠利用A100和NVIDIA Ampere微體系結構。
內存管理
最大化GPU內核性能的優化策略之一是最小化數據傳輸。如果內存駐留在全局內存中，則將數據讀取到L2高速緩存或共享內存中的等待時間可能需要數百個處理器周期。
例如，在GV100上，共享內存提供的帶寬比全局內存快17倍，比L2快3倍。因此，某些具有生產者－消費者范例的算法可能會觀察到在內核之間將數據，持久存儲在L2中時的性能優勢，從而獲得更高的帶寬和性能。
在A100上，CUDA 11提供API操作以預留一部分40 MB的L2緩存，持久保存對全局內存的數據訪問。持久性訪問優先使用了L2緩存的此預留部分，而對全局存儲器的常規訪問或流式訪問，僅在L2的這一部分通過持久訪問而未使用時才可以使用。
可以將L2持久性設置為在CUDA流或CUDA圖形內核節點中使用。當預留L2緩存區域時，需要考慮一些注意事項。例如，在具有不同訪問策略窗口的同時，在不同流中同時執行的多個CUDA內核共享L2預留緩存。下面的代碼示例顯示為持久性預留L2緩存比率。
cudaGetDeviceProperties( &prop, device_id);
// Set aside 50% of L2 cache for persisting accesses
size_t size = min( int(prop.l2CacheSize * 0.50) , prop.persistingL2CacheMaxSize );
cudaDeviceSetLimit( cudaLimitPersistingL2CacheSize, size);
// Stream level attributes data structure
cudaStreamAttrValue attr ;
attr.accessPolicyWindow.base_ptr = /* beginning of range in global memory */ ;
attr.accessPolicyWindow.num_bytes = /* number of bytes in range */ ;
// hitRatio causes the hardware to select the memory window to designate as persistent in the area set-aside in L2
attr.accessPolicyWindow.hitRatio = /* Hint for cache hit ratio */
// Type of access property on cache hit
attr.accessPolicyWindow.hitProp = cudaAccessPropertyPersisting;
// Type of access property on cache miss
attr.accessPolicyWindow.missProp = cudaAccessPropertyStreaming;
cudaStreamSetAttribute(stream,cudaStreamAttributeAccessPolicyWindow,&attr);
將虛擬內存管理API業務已經擴展到支持固定GPU內存壓縮到L2減少DRAM帶寬。這對於深度學習訓練和推理用例可能很重要。使用創建共享內存句柄cuMemCreate時，會向API操作提供分配提示。
算法（例如3D模板或卷積）的有效實現涉及內存復制和計算控制流模式，其中數據從全局內存傳輸到線程塊的共享內存中，然后使用該共享內存進行計算。全局到共享內存的副本被擴展為從全局內存到寄存器的讀取，然后是對共享內存的寫操作。
CUDA 11可以利用新的異步副本（async-copy）范例。實際上與通過計算將數據從全局復制到共享內存重疊，並且避免了使用中間寄存器或L1高速緩存。異步復制有很多好處：控制流不再遍歷內存管道兩次，並且不使用中間寄存器可以減少寄存器壓力，從而增加內核占用率。在A100上，異步復制操作是硬件加速的。
以下代碼示例顯示了使用異步復制的簡單示例。生成的代碼雖然性能更高，但是可以通過流水線處理多批異步復制操作來進一步優化。這種額外的流水線操作可以消除代碼中的同步點之一。
異步復制作為CUDA 11中的一項實驗功能提供，並使用協作組集合公開。在CUDA C ++編程指南包括使用異步拷貝用多級流水線和在A100硬件加速屏障操作的更先進的實施例。
//Without async-copy
using namespace nvcuda::experimental;
__shared__ extern int smem[];
// algorithm loop iteration
while ( ... ) {
__syncthreads();
// load element into shared mem
for ( i = ... ) {
// uses intermediate register
// {int tmp=g[i]; smem[i]=tmp;}
smem[i] = gldata[i];
}
//With async-copy
using namespace nvcuda::experimental;
__shared__ extern int smem[];
pipeline pipe;
// algorithm loop iteration
while ( ... ) {
__syncthreads();
// load element into shared mem
for ( i = ... ) {
// initiate async memory copy
memcpy_async(smem[i],
gldata[i],
pipe);
}
// wait for async-copy to complete
pipe.commit_and_wait();
__syncthreads();
/* compute on smem[] */
}
任務圖加速
CUDA 10中引入的CUDA圖形代表了一種使用CUDA提交工作的新模型。圖由一系列操作組成，例如內存副本和內核啟動，通過依賴關系進行連接並與其執行分開定義。
圖啟用了一次定義一次運行重復執行流程。可以減少累積的啟動開銷並提高應用程序的整體性能。對於深度學習應用程序尤其如此，因為深度學習應用程序可能會啟動多個具有減小的任務大小和運行時間的內核，或者可能具有任務之間的復雜依賴性。
從A100開始，GPU提供了任務圖硬件加速功能，以預取網格啟動描述符，指令和常量。與以前的GPU（例如V100）相比，可以在A100上使用CUDA圖形來改善內核啟動延遲。
CUDA Graph API操作現在具有輕量級機制，可以支持對實例化圖進行就地更新，而無需重建圖。在圖的重復實例化期間，節點參數（例如內核參數）通常會發生變化，而圖拓撲保持不變。圖形API操作提供了一種用於更新整個圖形的機制，其中可以為拓撲相同的cudaGraph_t對象提供更新的節點參數，或者為各個節點進行顯式更新。
此外，CUDA圖現在支持協作內核啟動（cuLaunchCooperativeKernel），包括與CUDA流奇偶校驗的流捕獲。
線程集合
以下是CUDA 9中引入的CUDA 11對協作組的增強。協作組是一種集體編程模式，旨在能夠明確表示線程可以進行通信的粒度。這將在CUDA中啟用協作並行性的新模式。
在CUDA 11中，合作組集體展示了新的A100硬件功能，並添加了一些API增強功能。有關更改的完整列表的更多信息，請參見《CUDA C ++編程指南》。
A100引入了一條新的reduce指令，該指令對每個線程提供的數據進行操作。這是使用協作組的新集合，提供了可移植的抽象，也可以在較舊的體系結構上使用。reduce操作支持算術（例如加）和邏輯（例如AND）操作。以下代碼示例顯示了reduce集合。
// Simple Reduction Sum
#include <cooperative_groups/reduce.h>
...
const int threadId = cta.thread_rank();
int val = A[threadId];
// reduce across tiled partition
reduceArr[threadId] = cg::reduce(tile, val, cg::plus<int>());
// synchronize partition
cg::sync(cta);
// accumulate sum using a leader and return sum
合作組提供labeled_partition了將父組划分為一維子組（在其中合並線程）的集體操作（）。這對於試圖通過條件語句的基本塊跟蹤活動線程的控制流特別有用。
例如，可以使用labeled_partition原子添加操作並在其中使用一個warp級別組（不限於2的冪）來形成多個分區。labeled_partitionAPI操作評估條件，標簽，並分配具有用於標簽到同一組的值相同的線程。
以下代碼示例顯示了自定義線程分區：
// Get current active threads (that is, coalesced_threads())
cg::coalesced_group active = cg::coalesced_threads();
// Match threads with the same label using match_any()
int bucket = active.match_any(value);
cg::coalesced_group subgroup = cg::labeled_partition(active, bucket);
// Choose a leader for each partition (for example, thread_rank = 0)
//
if (subgroup.thread_rank() == 0) {
threadId = atomicAdd(&addr[bucket], subgroup.size());
}
// Now use shfl to transfer the result back to all threads in partition
return (subgroup.shfl(threadId, 0));
CUDA C ++語言和編譯器改進
CUDA 11還是第一個正式將CUB納入CUDA Toolkit的版本。CUB是受支持的CUDA C ++核心庫之一。
nvcc for CUDA 11的主要功能之一是對鏈接時間優化（LTO）的支持，以提高單獨編譯的性能。LTO使用--dlink-time-opt或-dlto選項在編譯期間存儲中間代碼，然后在鏈接時執行更高級別的優化，例如跨文件內聯代碼。
CUDA 11中的nvcc增加了對ISO C ++ 17的支持，並支持了跨PGI，gcc，clang，Arm和Microsoft Visual Studio的新主機編譯器。如果要試驗尚不支持的主機編譯器，則nvcc--allow-unsupported-compiler在編譯構建工作流程中會支持一個新標志。nvcc添加了其他新功能，其中包括：
- 改進的lambda支持
- 依賴文件生成增強（
-MD，-MMD選項） - 傳遞給主機編譯器的選項
Figure 4. Platform support in CUDA 11.
CUDA庫
CUDA 11中的庫通過在線性代數，信號處理，基本數學運算和圖像處理中使用最新最先進的A100硬件功能超越熟悉的嵌入式API，從而不斷突破性能和開發人員生產力的界限。
在線性代數庫中，將看到Tensor Core加速度，了解A100上所有可用的精度，包括FP16，Bfloat16，TF32和FP64。這包括cuBLAS中的BLAS3操作，cuSOLVER中的分解和密集線性求解器，以及cuTENSOR中的張量收縮。
除了提高精度范圍外，還消除了對矩陣尺寸和Tensor Core加速度對齊的限制。為了獲得適當的精度，加速度現在是自動的，不需要用戶選擇。當在A100上帶有MIG的GPU實例上運行時，cuBLAS的啟發式方法會自動適應資源。
Figure 6. Mixed-precision matrix multiply on A100 with cuBLAS.
CUTLASS是高性能GEMM的CUDA C ++模板抽象，支持A100提供的所有各種精度模式。借助CUDA 11，CUTLASS現在與cuBLAS達到了95％以上的性能均等性。可以編寫自己的自定義CUDA內核，以便在NVIDIA GPU中對Tensor Core進行編程。
cuFFT利用了A100中較大的共享內存大小，從而在批量較大時為單精度FFT提供了更好的性能。最后，在多GPU A100系統上，與V100相比，cuFFT可以擴展並提供每個GPU 2倍的性能。
nvJPEG是用於JPEG解碼的GPU加速庫。結合數據增強和圖像加載庫NVIDIA DALI，可以加速對圖像分類模型（尤其是計算機視覺）的深度學習訓練。這些庫加快了深度學習工作流程的圖像解碼和數據增強階段。
A100包含5核硬件JPEG解碼引擎，nvJPEG利用硬件后端對JPEG圖像進行批處理。專用硬件模塊進行JPEG加速可緩解CPU瓶頸，並提高GPU利用率。
nvjpegDecode對於給定的圖像，硬件自動選擇硬件解碼器，或者使用nvjpegCreateExinit函數明確選擇硬件后端。nvJPEG提供基線JPEG解碼和各種顏色轉換格式（例如YUV 420、422和444）的加速。
圖8顯示，與僅使用CPU的處理相比，這將使圖像解碼速度提高18倍。如果使用DALI，則可以直接從此硬件加速中受益，因為nvJPEG是抽象的。
Figure 9. nvJPEG Speedup vs. CPU.
(Batch 128 with Intel Platinum 8168 @2GHz 3.7GHz Turbo HT on; with TurboJPEG)
CUDA數學庫中的功能比單個帖子中涵蓋的功能還要多。
開發工具
CUDA 11繼續在現有的開發人員工具組合中添加豐富的功能。包括用於Visual Studio的熟悉的插件，以及用於Visual Studio的NVIDIA Nsight集成，以及具有Nsight Eclipse插件版的Eclipse。它還包括獨立的工具，例如用於內核概要分析的Nsight Compute和用於系統范圍的性能分析的Nsight系統。CUDA支持的所有三種CPU體系結構現在都支持Nsight Compute和Nsight系統：x86，POWER和Arm64。
One of the key features of Nsight Compute for CUDA 11 is the ability to generate the Roofline model of the application. A Roofline model is a visually intuitive method for you to understand kernel characteristics by combining floating-point performance, arithmetic intensity, and memory bandwidth into a two-dimensional plot.
Nsight Compute for CUDA 11的主要功能之一是能夠生成應用程序的Roofline模型。Roofline模型是一種視覺直觀的方法，通過將浮點性能，算術強度和內存帶寬組合成二維圖，可以了解內核特性。
通過查看Roofline模型，可以快速確定內核是受計算限制還是受內存限制。還可以了解進一步優化的潛在方向，例如，靠近車頂線的內核可以最佳地利用計算資源。
有關更多信息，請參見Roofline性能模型。
Figure 11. A Roofline model in Nsight Compute.
CUDA 11包括Compute Sanitizer，這是下一代的功能正確性檢查工具，可對越界內存訪問和競爭條件提供運行時檢查。Compute Sanitizer旨在替代該cuda-memcheck工具。
以下代碼示例顯示了Compute Sanitizer檢查內存訪問的示例。
//Out-of-bounds Array Access
__global__ void oobAccess(int* in, int* out)
{
int bid = blockIdx.x;
int tid = threadIdx.x;
if (bid == 4)
{
out[tid] = in[dMem[tid]];
}
}
int main()
{
...
// Array of 8 elements, where element 4 causes the OOB
std::array<int, Size> hMem = {0, 1, 2, 10, 4, 5, 6, 7};
cudaMemcpy(d_mem, hMem.data(), size, cudaMemcpyHostToDevice);
oobAccess<<<10, Size>>>(d_in, d_out);
cudaDeviceSynchronize();
...
$ /usr/local/cuda-11.0/Sanitizer/compute-sanitizer --destroy-on-device-error kernel --show-backtrace no basic
========= COMPUTE-SANITIZER
Device: Tesla T4
========= Invalid __global__ read of size 4 bytes
========= at 0x480 in /tmp/CUDA11.0/ComputeSanitizer/Tests/Memcheck/basic/basic.cu:40:oobAccess(int*,int*)
========= by thread (3,0,0) in block (4,0,0)
========= Address 0x7f551f200028 is out of bounds
以下代碼示例顯示了用於競爭條件檢查的Compute Sanitizer示例。
//Contrived Race Condition Example
__global__ void Basic()
{
__shared__ volatile int i;
i = threadIdx.x;
}
int main()
{
Basic<<<1,2>>>();
cudaDeviceSynchronize();
...
$ /usr/local/cuda-11.0/Sanitizer/compute-sanitizer --destroy-on-device-error kernel --show-backtrace no --tool racecheck --racecheck-report hazard raceBasic
========= COMPUTE-SANITIZER
========= ERROR: Potential WAW hazard detected at __shared__ 0x0 in block (0,0,0) :
========= Write Thread (0,0,0) at 0x100 in /tmp/CUDA11.0/ComputeSanitizer/Tests/Racecheck/raceBasic/raceBasic.cu:11:Basic(void)
========= Write Thread (1,0,0) at 0x100 in /tmp/CUDA11.0/ComputeSanitizer/Tests/Racecheck/raceBasic/raceBasic.cu:11:Basic(void)
========= Current Value : 0, Incoming Value : 1
=========
========= RACECHECK SUMMARY: 1 hazard displayed (1 error, 0 warnings)
最后，即使CUDA 11不再支持在macOS上運行應用程序，仍將開發工具提供給macOS主機上的用戶使用：
- 使用遠程目標調試 cuda-gdb
- NVIDIA Visual Profiler
- Nsight Eclipse插件
- Nsight系列工具，用於遠程分析或跟蹤
Summary
CUDA 11提供了一個基礎開發環境，用於為NVIDIA Ampere GPU架構和基於NVIDIA A100構建的強大服務器平台構建應用程序，以用於本地（DGX A100）和雲（HGX A100）的AI，數據分析和HPC工作負載部署。
Figure 12. Different ways to get CUDA 11.
CUDA 11現在可用。與往常一樣，可以通過多種方式獲得CUDA 11：下載本地安裝程序軟件包，使用軟件包管理器進行安裝或從各個注冊表中獲取容器。對於企業部署，CUDA 11還包括使用模塊化流的RHEL 8驅動程序包裝改進，以提高穩定性並減少安裝時間。