TunableOp

概述

此模組公開了一個 TunableOp 介面。

某些操作（例如 GEMM）可以使用不止一個庫或不止一種技術來實現。例如，可以使用 blas 或 blasLt 庫為 CUDA 或 ROCm 實現 GEMM。此外，ROCm 的 rocblas 和 hipblaslt 庫允許使用者查詢所有可能的演算法，然後選擇其中一個。如何知道哪種實現最快以及應該選擇哪種？這就是 TunableOp 所提供的。

啟用 TunableOp 和獨立調優

TunableOp 功能的啟用與調優階段本身的啟用是分開的。啟用 TunableOp 意味著 PyTorch 會將其任何標準運算子替換為 Tunable 實現。對 TunableOp 的任何呼叫首先會檢查是否已針對給定的運算子輸入進行了調優。如果是，它會立即呼叫已調優的操作；即使調優設定已啟用，也不會進行進一步的調優。如果未找到調優結果且調優已啟用，則 TunableOp 會針對給定輸入集對該運算子的每個註冊實現進行基準測試，並選擇最快的一個。

檔案輸入和輸出

首次呼叫任何 TunableOp 時，將透過嘗試從給定檔案中讀取結果來準備已調優操作的內部資料庫。預設檔名為“tunableop_results.csv”。為了支援在使用多個程序跨多個 GPU 進行調優時，GPU 裝置序號會自動插入到檔名中，以避免多個程序覆蓋同一個檔案。

如果啟用了調優，並且在工作負載執行過程中發現了新的調優結果，它也會將所有調優結果寫入同一個檔案，包括啟動時讀取的以及執行時發現的新結果。例如，這可以用於透過重複使用同一個檔案來在多個工作負載中構建一個調優檔案。輸出檔案在應用程式終止時自動建立。此行為可透過 C++ 和 Python API 控制，但不能透過環境變數控制。

假設您指定了一個檔名，您將得到一個內容如下的 CSV 檔案

Validator,PT_VERSION,2.2.0
Validator,ROCM_VERSION,6.0.0.0-12969-1544e39
Validator,HIPBLASLT_VERSION,0.6.0-a9c5cc7
Validator,ROCBLAS_VERSION,4.0.0-72e57364-dirty
GemmTunableOp_float_NT,nt_25088_4096_64,Gemm_Hipblaslt_1219,1.262
GemmTunableOp_float_NT,nt_4096_4096_64,Gemm_Rocblas_1216,0.033

請注意“Validator”行。如果您更改了庫版本、ROCm 版本或 PyTorch 版本，TunableOp 會檢測到這一點並拒絕調優檔案，因為之前的調優結果可能受到其他軟體更改的影響。

其餘行是您執行過程中遇到的每個 TunableOp 的已調優解決方案。每行包含 4 個逗號分隔的欄位：運算子名稱、運算子引數、解決方案名稱和平均執行時間。執行時間是可選欄位。CSV 檔案可以編輯，但要謹慎。例如，可以將解決方案名稱（欄位 3）更改為“Default”，它將回退到原始的 PyTorch 未調優實現。或者，在 ROCm 的 hipBLAS 或 hipBLASLt 庫的情況下，如果您知道特定的解決方案索引，可以透過替換該值來覆蓋 TunableOp 選擇的解決方案。運算子名稱和引數（欄位 1 和 2）是內部命名，不應修改。在 GemmTunableOp 的情況下，欄位 1 表示資料型別以及輸入是否轉置 (T) 或不轉置 (N)，欄位 2 表示 M、N、K 輸入形狀。

有一個選項可以啟用詳細輸出，但這僅建議用於除錯目的。這將產生許多診斷訊息，但可能有助於檢視 TunableOp 是否正在使用。否則，除了檔案輸出外，TunableOp 完全靜默，除非在使用過程中出現警告或錯誤。詳細輸出選項只能透過設定環境變數 PYTORCH_TUNABLEOP_VEROBSE=1 啟用。

關於調優行為、預熱和快取效應的注意事項

調優運算子包括遍歷註冊實現的列表並對每個實現進行效能分析。透過在一個迴圈中多次執行單個實現並計算平均執行時間來建立效能分析。調優之前還有一個可選的預熱階段，有助於硬體達到穩定的功耗狀態。在工作負載調優期間，各種硬體快取更有可能命中，而不是未調優時。有一些選項可以重新整理指令快取並旋轉輸入張量，這可能有助於更準確地反映已調優運算子的效能特徵，就像運算子在更大的工作負載中執行，而不是在緊密、重複的迴圈中執行一樣。

預設情況下，給定運算子的每個可能解決方案將執行 100 次迭代或在 30 毫秒內可以執行的最大迭代次數（以較小者為準），並計算其平均執行時間。將選擇所有成功分析的解決方案中最快的一個。如果給定解決方案未達到與預設實現相同的精度，或者如果解決方案返回錯誤程式碼，則效能分析可能會失敗。

當前可調優運算子

用於 ROCm 的 TunableGemm

目前僅實現了用於 ROCm 的 TunableGemm。請注意，在使用 TunableOp 時，PyTorch 的 CUDA 構建將正常工作，但 CUDA 構建唯一可用的解決方案是“Default”實現，即原始的 cuBLAS 預設實現，現在透過 TunableOp 呼叫。啟用後，對 at::cuda::blas::gemm() 或 ::bgemm() 的任何呼叫都將透過 TunableOp 路由。為給定輸入引數集 (transa, transb, m, n, k) 呼叫 gemm() 時，將嘗試使用 rocblas 和 hipblaslt 中最快的可用實現。

離線調優

動機

離線調優有幾種用例。

一種用例涉及記憶體利用率高的工作負載，其中常規調優可能會導致記憶體不足。

另一種用例是計算密集型工作負載。

在這種情況下，更具資源效率的做法是先收集工作負載的 GEMM 一次，然後使用不同的調優引數或庫重複進行調優。

工作流程

PYTORCH_TUNABLEOP_ENABLED=1
PYTORCH_TUNABLEOP_TUNING=0
PYTORCH_TUNABLEOP_RECORD_UNTUNED=1
...

2. 基本上有兩步：1) 設定環境變數以收集未調優的 GEMM，這將生成 tunableop_untuned0.csv

import torch.cuda.tunable as tunable
import os

os.putenv('PYTORCH_TUNABLEOP_ENABLED', '1')
os.putenv('PYTORCH_TUNABLEOP_TUNING', '1')
os.putenv('PYTORCH_TUNABLEOP_RECORD_UNTUNED', '0')
tunable.tune_gemm_in_file("tunableop_untuned0.csv")

執行一個 Python 指令碼，該指令碼讀取 tunableop_untuned0.csv 並生成 tunableop_results0.csv，如下所示

if __name__ == "__main__":
    num_gpus = 8 # number of GPUs that will be used during the tuning process
    tunable.mgpu_tune_gemm_in_file("tunableop_untuned?.csv", num_gpus)

還可以獲取多個未調優檔案，並將 GEMM 分佈到單個節點內的多個 GPU 進行調優。第一步，先收集 GEMM 並消除重複的 GEMM。接下來，將 GEMM 分配到不同的 GPU 進行調優。所有 GEMM 調優完成後，來自所有 GPU 的結果將被收集到一個檔案中，該檔案的基本檔名附加了 _full0（例如 tunableop_results_full0.csv）。最後，這個包含收集結果的新檔案將被複制 N 次，每個 GPU 一次，以便使用者在 N 個 GPU 上使用已調優的配置執行工作負載。

請注意，mgpu_tune_gemm_in_file API 的用法與其單 GPU 對應 API (tune_gemm_in_file) 不同。由於使用了 concurrent futures 模組，呼叫該 API 的 Python 指令碼的主體必須按所示包裝在 main() 中。傳遞給 mgpu_tune_gemm_in_file 的引數必須包含萬用字元表示式 (? 或 *)，以生成包含待處理 GEMM 的未調優檔案列表。num_gpus 必須介於 1 和可用 GPU 總數之間。

調優上下文

TunableOp 的行為目前透過環境變數、at::cuda::tunable::getTuningContext() 的 C++ 介面或 torch.cuda.tunable python 介面進行控制。環境變數優先於您使用 C++ 或 Python API 操作的任何設定。

環境變數介面

環境變數在首次讀取時會被快取。

您不能以程式設計方式使用環境變數介面，因為設定已固定。請改用 C++ 或 Python API。

API 參考

torch.cuda.tunable.enable(val=True)

這是所有 TunableOp 實現的總開關。

torch.cuda.tunable.is_enabled()

返回 TunableOp 功能是否已啟用。

返回型別

bool

torch.cuda.tunable.tuning_enable(val=True)

啟用 TunableOp 實現的調優。

啟用後，如果未找到已調優條目，則執行調優步驟並記錄該條目。

torch.cuda.tunable.is_enabled()

返回 TunableOp 功能是否已啟用。

torch.cuda.tunable.tuning_is_enabled()

返回 TunableOp 實現是否可以調優。

torch.cuda.tunable.record_untuned_enable(val=True)

啟用記錄未調優的 TunableOp 操作以進行離線調優。

啟用後，如果未找到已調優條目，則將其寫入未調優檔案。

torch.cuda.tunable.is_enabled()

返回 TunableOp 功能是否已啟用。

torch.cuda.tunable.record_untuned_is_enabled()

返回是否記錄 TunableOp 操作以進行離線調優。

torch.cuda.tunable.set_max_tuning_duration(duration)

設定調優給定解決方案的最大時間（毫秒）。

如果同時設定了最大調優持續時間和迭代次數，則以兩者中較小者為準。

torch.cuda.tunable.is_enabled()

至少會執行 1 次調優迭代。

torch.cuda.tunable.get_max_tuning_duration()

獲取調優給定解決方案的最大時間。

torch.cuda.tunable.set_max_tuning_duration(duration)

int

torch.cuda.tunable.set_max_tuning_iterations(iterations)

torch.cuda.tunable.is_enabled()

至少會執行 1 次調優迭代。

設定調優給定解決方案的最大迭代次數。

torch.cuda.tunable.get_max_tuning_iterations()

獲取調優給定解決方案的最大迭代次數。

torch.cuda.tunable.set_filename(filename, insert_device_ordinal=False)

設定用於調優結果輸入/輸出的檔名。

torch.cuda.tunable.is_enabled()

如果 insert_device_ordinal 為 True，則當前裝置序號將自動新增到給定檔名中。這可以在 1 個程序對應 1 個 GPU 的場景中使用，以確保所有程序寫入不同的檔案。

torch.cuda.tunable.get_filename()

獲取結果檔名。

torch.cuda.tunable.is_enabled()

str

torch.cuda.tunable.get_results()

返回所有 TunableOp 結果。

torch.cuda.tunable.is_enabled()

tuple[str, str, str, float]

torch.cuda.tunable.get_validators()

返回 TunableOp 驗證器。

tuple[str, str]

torch.cuda.tunable.write_file_on_exit(val)

在銷燬調優上下文期間，將檔案寫入磁碟。

如果您的應用程式因正常執行或錯誤而終止，這有助於將最終結果重新整理到磁碟。可以透過手動呼叫 write_file() 來手動重新整理結果。

torch.cuda.tunable.is_enabled()

返回 TunableOp 功能是否已啟用。

Manual flushing of your results can be achieved by manually calling write_file().

torch.cuda.tunable.write_file(filename=None)

如果您的應用程式因正常執行或錯誤而終止，這有助於將最終結果重新整理到磁碟。可以透過手動呼叫 write_file() 來手動重新整理結果。

torch.cuda.tunable.is_enabled()

返回 TunableOp 功能是否已啟用。

torch.cuda.tunable.tune_gemm_in_file(filename)

在檔案中調優 GEMM。

torch.cuda.tunable.mgpu_tune_gemm_in_file(filename_pattern, num_gpus)

處理一個或多個檔案，並將工作分佈到一個或多個 GPU 上。

torch.cuda.tunable.set_rotating_buffer_size(buffer_size)

將旋轉緩衝區大小設定為此 MB 值，前提是緩衝區大小大於零。

如果小於零，則查詢 L2 快取大小。如果等於零，則表示停用旋轉緩衝區。

torch.cuda.tunable.get_rotating_buffer_size()

獲取旋轉緩衝區大小（單位：千位元組）。

torch.cuda.tunable.is_enabled()

至少會執行 1 次調優迭代。