TunableOp
Создано: 03 июня 2024 г. | Последнее обновление: 29 июня 2026 г.
Обзор
Этот модуль предоставляет интерфейс TunableOp.
Некоторые операции, например GEMM, можно реализовать с помощью нескольких библиотек или разных методов. Например, GEMM для CUDA или ROCm можно реализовать с использованием библиотек blas или blasLt. Кроме того, библиотеки rocblas и hipblaslt для ROCm позволяют пользователю запросить все возможные алгоритмы и выбрать один из них. Как определить, какая реализация работает быстрее и должна быть выбрана? Для этого и предназначен TunableOp.
Раздельное включение TunableOp и настройки
Функция TunableOp включается отдельно от самого этапа настройки. Включение TunableOp означает, что PyTorch заменит все стандартные операторы их настраиваемыми реализациями. При каждом вызове TunableOp сначала проверяется, выполнялась ли настройка для заданных входных данных оператора. Если да, сразу вызывается настроенная операция; дальнейшая настройка не выполняется, даже если параметр настройки включен. Если же результат настройки не найден и настройка включена, TunableOp проверит производительность всех зарегистрированных реализаций этого оператора для заданного набора входных данных и выберет самую быструю.
Ввод и вывод файлов
При первом вызове любого TunableOp будет подготовлена внутренняя база данных настроенных операций: система попытается прочитать результаты из указанного файла. Имя файла по умолчанию — ‘tunableop_results.csv’. Чтобы поддержать настройку при использовании нескольких GPU в нескольких процессах, порядковый номер устройства GPU автоматически добавляется к имени файла, чтобы процессы не перезаписывали один и тот же файл.
Если настройка включена и во время выполнения рабочей нагрузки обнаружены новые настройки, они также будут записаны в этот же файл вместе со всеми настройками: как прочитанными при запуске, так и найденными во время выполнения. Например, так можно постепенно собрать файл настроек для нескольких рабочих нагрузок, повторно используя один и тот же файл. Выходной файл создается автоматически при завершении приложения. Этим поведением можно управлять с помощью API C++ и Python, но не с помощью переменных среды.
Если вы указали имя файла, в итоге будет создан 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. Для библиотек hipBLAS или hipBLASLt в ROCm, если вам известен индекс конкретного решения, можно переопределить выбранное TunableOp решение, заменив соответствующее значение. Имена оператора и его параметры (поля 1 и 2) задаются внутри системы и не должны изменяться. В случае GemmTunableOp поле 1 указывает тип данных и то, транспонированы ли входные данные (T) или нет (N), а поле 2 указывает размеры входных данных M, N, K.
Можно включить подробный вывод, но он рекомендуется только для отладки. При этом будет выведено множество диагностических сообщений, которые могут помочь понять, используется ли вообще TunableOp. В остальных случаях TunableOp работает полностью без вывода, за исключением записи в файл, если при его использовании не возникает предупреждения или ошибки. Подробный вывод можно включить только с помощью переменной среды PYTORCH_TUNABLEOP_VERBOSE=1.
Примечание о поведении настройки, прогреве и влиянии кэша
Настройка оператора заключается в переборе списка зарегистрированных реализаций и профилировании каждой из них. Профиль формируется путем многократного выполнения одной реализации в цикле и вычисления среднего времени выполнения. Перед настройкой также можно выполнить необязательный этап прогрева, который помогает оборудованию достичь стабильных режимов энергопотребления. Во время настройки рабочей нагрузки попадания в различные аппаратные кэши будут происходить чаще, чем при отсутствии настройки. Предусмотрены параметры для очистки кэша инструкций и циклической смены входных тензоров; это может помочь получить более точный профиль настроенного оператора, соответствующий его работе в составе большой рабочей нагрузки, а не в тесном повторяющемся цикле.
По умолчанию каждое возможное решение для заданного оператора выполняется либо 100 итераций, либо столько итераций, сколько можно выполнить за 30 мс, в зависимости от того, какое значение меньше; затем вычисляется среднее время выполнения. Выбирается самое быстрое из всех решений, профилирование которых завершилось успешно. Профилирование может завершиться неудачей, если заданное решение не обеспечивает точность, равную точности реализации по умолчанию, или возвращает код ошибки.
Поддержка CUDA cuBLASLt использует кэш результатов TunableOp и механизм профилирования для замера времени настраиваемого числа эвристических кандидатов cuBLASLt.
Текущие настраиваемые операторы
TunableGemm для ROCm
При включенном TunableOp любой вызов at::cuda::blas::gemm() или ::bgemm() будет перенаправлен через него. Вызов gemm() для заданного набора входных аргументов (transa, transb, m, n, k) в ROCm попытается использовать самую быструю доступную реализацию из rocblas и hipblaslt. В CUDA TunableGemm регистрирует эвристические кандидаты cuBLASLt для путей GEMM, которые уже используют cuBLASLt.
Эвристическая настройка cuBLASLt для CUDA
Количество эвристических кандидатов cuBLASLt задается с помощью set_cublaslt_requested_algo_count() или PYTORCH_TUNABLEOP_CUBLASLT_REQUESTED_ALGO_COUNT; значение по умолчанию — 8. Если это значение равно 1, доступен только самый предпочтительный эвристический кандидат cuBLASLt.
Автономная настройка
Причины использования
Существует несколько сценариев использования автономной настройки.
Один из сценариев — рабочая нагрузка с высоким потреблением памяти, при которой обычная настройка может привести к нехватке памяти.
Другой сценарий — вычислительно интенсивные рабочие нагрузки. В таких случаях эффективнее с точки зрения использования ресурсов один раз собрать GEMM для рабочей нагрузки, а затем повторять настройку с различными параметрами настройки или библиотеками.
Рабочий процесс
Он состоит из двух основных шагов: 1) задайте переменные среды для сбора ненастроенных GEMM — в результате будет создан tunableop_untuned0.csv:
export PYTORCH_TUNABLEOP_ENABLED=1 export PYTORCH_TUNABLEOP_TUNING=0 export PYTORCH_TUNABLEOP_RECORD_UNTUNED=1 ...
- Запустите скрипт Python, который считывает
tunableop_untuned0.csvи создаетtunableop_results0.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")
Также можно взять несколько ненастроенных файлов и распределить GEMM для настройки между несколькими GPU в одном узле. На первом шаге GEMM сначала собираются, а повторяющиеся GEMM удаляются. Затем GEMM распределяются между разными GPU для настройки. После настройки всех GEMM результаты со всех GPU собираются в один файл, к базовому имени которого добавляется _full0 (например, tunableop_results_full0.csv). Наконец, этот новый файл со всеми собранными результатами дублируется N раз — по одному экземпляру для каждого GPU. Это сделано для удобства пользователя, который будет запускать рабочую нагрузку с настроенной конфигурацией на N GPU.
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)
Обратите внимание, что использование API mgpu_tune_gemm_in_file отличается от использования его варианта для одного GPU (tune_gemm_in_file). Тело скрипта Python, вызывающего API, необходимо, как показано, обернуть в main() из-за использования модуля concurrent futures. Аргумент для mgpu_tune_gemm_in_file должен содержать выражение с подстановочным знаком (? или *) для создания списка ненастроенных файлов с GEMM, которые требуется обработать. Значение num_gpus должно быть от 1 до общего количества доступных GPU.
Контекст настройки
В настоящее время поведением TunableOp можно управлять с помощью переменных среды, интерфейса C++ at::cuda::tunable::getTuningContext() или интерфейсов Python torch.cuda.tunable. Переменные среды имеют приоритет над любыми параметрами, заданными с помощью API C++ или Python.
Интерфейс переменных среды
Переменные среды кэшируются при первом чтении. Программно использовать интерфейс переменных среды нельзя, поскольку заданные параметры становятся неизменными. Вместо этого используйте API C++ или Python.
Справочник API
-
torch.cuda.tunable.enable(val=True)[исходный код] -
Это главный переключатель включения и отключения всех реализаций TunableOp.
-
torch.cuda.tunable.is_enabled()[исходный код] -
Возвращает, включена ли функция TunableOp.
- Тип возвращаемого значения:
-
torch.cuda.tunable.tuning_enable(val=True)[исходный код] -
Включает настройку реализаций TunableOp.
Если настройка включена и настроенная запись не найдена, выполняется этап настройки и запись сохраняется.
-
torch.cuda.tunable.tuning_is_enabled()[исходный код] -
Возвращает, можно ли настраивать реализации TunableOp.
- Тип возвращаемого значения:
-
torch.cuda.tunable.record_untuned_enable(val=True)[исходный код] -
Включает запись ненастроенных операций TunableOp для автономной настройки.
Если настройка включена и настроенная запись не найдена, операция записывается в файл ненастроенных операций.
-
torch.cuda.tunable.record_untuned_is_enabled()[исходный код] -
Возвращает, записываются ли операции TunableOp для автономной настройки.
- Тип возвращаемого значения:
-
torch.cuda.tunable.set_max_tuning_duration(duration)[исходный код] -
Задает максимальное время в миллисекундах, отводимое на настройку заданного решения.
Если заданы и максимальная продолжительность настройки, и количество итераций, будет использоваться меньшее из двух значений. В любом случае будет выполнена как минимум 1 итерация настройки.
-
torch.cuda.tunable.get_max_tuning_duration()[исходный код] -
Возвращает максимальное время, отводимое на настройку заданного решения.
- Тип возвращаемого значения:
-
torch.cuda.tunable.set_max_tuning_iterations(iterations)[исходный код] -
Задает максимальное количество итераций, отводимых на настройку заданного решения.
Если заданы и максимальная продолжительность настройки, и количество итераций, будет использоваться меньшее из двух значений. В любом случае будет выполнена как минимум 1 итерация настройки.
-
torch.cuda.tunable.get_max_tuning_iterations()[исходный код] -
Возвращает максимальное количество итераций, отводимых на настройку заданного решения.
- Тип возвращаемого значения:
-
torch.cuda.tunable.set_cublaslt_requested_algo_count(count)[исходный код] -
Задает количество эвристических алгоритмов cuBLASLt, запрашиваемых в CUDA.
Значения меньше 1 ограничиваются значением 1.
-
torch.cuda.tunable.get_cublaslt_requested_algo_count()[исходный код] -
Возвращает количество эвристических алгоритмов cuBLASLt, запрошенных в CUDA.
- Тип возвращаемого значения:
-
torch.cuda.tunable.set_filename(filename, insert_device_ordinal=False)[исходный код] -
Задает имя файла для ввода и вывода результатов настройки.
Если
insert_device_ordinalравноTrue, порядковый номер текущего устройства автоматически добавляется к заданному имени файла. Это можно использовать в сценарии «один процесс на GPU», чтобы процессы записывали данные в отдельные файлы.
-
torch.cuda.tunable.get_filename()[исходный код] -
Возвращает имя файла с результатами.
- Тип возвращаемого значения:
-
torch.cuda.tunable.get_results()[исходный код] -
Возвращает все результаты TunableOp.
-
torch.cuda.tunable.get_validators()[исходный код] -
Возвращает валидаторы TunableOp.
-
torch.cuda.tunable.read_file(filename=None)[исходный код] -
Читает результаты из CSV-файла TunableOp.
Если
filenameне задан, вызываетсяget_filename().- Тип возвращаемого значения:
-
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)[исходный код] -
Задает размер циклического буфера в МБ, если размер буфера больше нуля.
Если значение меньше нуля, запрашивается размер кэша L2. Значение, равное нулю, отключает циклический буфер.
-
torch.cuda.tunable.get_rotating_buffer_size()[исходный код] -
Возвращает размер циклического буфера в килобайтах.
- Тип возвращаемого значения:
-
torch.cuda.tunable.set_numerical_check_tolerances(enable, atol=1e-05, rtol=1e-05)[исходный код] -
Задает значения atol и rtol для проверки численной точности.
© 2026, PyTorch Contributors
PyTorch has a BSD-style license, as found in the LICENSE file.
https://docs.pytorch.org/docs/2.14/cuda.tunable.html