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,1219,1.262
GemmTunableOp_float_NT,nt_4096_4096_64,1216,0.033
請注意“驗證器”行。如果您更改了函式庫版本、ROCm 版本或 PyTorch 版本,TunableOp 將會偵測到這一點並拒絕調整檔案,因為先前的調整可能會受到其他軟體變更的影響。
其餘的行是您執行期間遇到的每個 TunableOp 的調整解決方案。每一行都包含 4 個逗號分隔的欄位:運算子名稱、運算子參數、解決方案名稱和平均執行時間。執行時間是一個可選欄位。可以編輯 CSV 檔案,但要小心。例如,解決方案名稱(欄位 3)可以更改為“預設”,並且它將回到原始的 PyTorch 未調整實現。或者,在 ROCm 的 hipBLAS 或 hipBLASLt 函式庫的情況下,如果您知道特定的解決方案索引,則可以透過替換值來覆寫 TunableOp 選擇的解決方案。運算子名稱和參數(欄位 1 和 2)是內部命名的,不應修改。在 GemmTunableOp 的情況下,欄位 1 表示資料類型以及輸入是轉置 (T) 還是不轉置 (N),欄位 2 表示 M、N、K 輸入形狀。
有一個選項可以啟用詳細輸出,但建議僅用於除錯目的。這將產生大量診斷訊息,但可能有助於查看是否正在使用 TunableOp。否則,除了檔案輸出之外,TunableOp 是完全靜默的,除非在使用過程中出現警告或錯誤。詳細選項只能透過設定環境變數 PYTORCH_TUNABLEOP_VEROBSE=1 來啟用。
關於調整行為的注意事項¶
調整運算子包括迭代已註冊實現的清單或設定檔,並對每個實現進行性能分析。設定檔是透過在迴圈中多次執行單個實現並採用平均執行時間來建立的。
根據預設,給定運算子的每個可能解決方案將執行 100 次迭代或在 30 毫秒內可以執行的盡可能多的迭代(以較小者為準),並計算其平均執行時間。將選擇已成功設定檔的所有解決方案中最快的解決方案。如果給定的解決方案沒有達到與預設實現相同的精度,或者如果解決方案返回錯誤代碼,則設定檔可能會失敗。
目前的 Tunable 運算子¶
適用於 ROCm 的 TunableGemm¶
目前僅實作了 ROCm 的 TunableGemm。請注意,使用 TunableOp 時,CUDA 建置的 PyTorch 將正常運作,但 CUDA 建置唯一可用的解決方案是「預設」實作,即原始的 cuBLAS 預設值,現在透過 TunableOp 呼叫。啟用後,任何對 at::cuda::blas::gemm() 或 ::bgemm() 的呼叫都將透過 TunableOp 進行路由。針對一組給定的輸入參數(transa、transb、m、n、k)呼叫 gemm() 將嘗試使用 rocblas 和 hipblaslt 中最快的可用實作。
調整上下文¶
TunableOp 的行為目前透過環境變數、at::cuda::tunable::getTuningContext() 的 C++ 介面或包裝 C++ TuningContext 的 torch.cuda.tunable Python 介面進行操作。環境變數優先於您使用 C++ 或 Python API 操作的任何設定。
API 參考¶
- torch.cuda.tunable.set_max_tuning_duration(duration)[原始碼]¶
設定調整給定解決方案所花費的最大時間(以毫秒為單位)。
如果同時設定了最大調整持續時間和迭代次數,則將採用兩者中較小者。至少會始終執行 1 次調整迭代。
- torch.cuda.tunable.set_max_tuning_iterations(iterations)[原始碼]¶
設定調整給定解決方案所花費的最大迭代次數。
如果同時設定了最大調整持續時間和迭代次數,則將採用兩者中較小者。至少會始終執行 1 次調整迭代。
- torch.cuda.tunable.set_filename(filename, insert_device_ordinal=False)[原始碼]¶
設定用於輸入/輸出調整結果的檔案名稱。
如果
insert_device_ordinal為True,則目前的裝置序號將自動新增到給定的檔案名稱中。這可以用於每個 GPU 一個程序的 cenario 中,以確保所有程序都寫入到單獨的檔案中。
- torch.cuda.tunable.write_file_on_exit(val)[原始碼]¶
在調整上下文銷毀期間,將檔案寫入磁碟。
如果您的應用程式因正常操作或錯誤而終止,這對於將結果最終刷新到磁碟非常有用。您可以透過手動呼叫
write_file()來手動刷新結果。