可调操作 ¶

概述 ¶

此模块公开了可调操作接口。

一些操作，如 GEMM，可以使用多个库或多种技术来实现。例如，GEMM 可以针对 CUDA 或 ROCm 使用 BLAS 或 BLASLt 库来实现。此外，ROCm 的 rocblas 和 hipblaslt 库允许用户查询所有可能的算法，然后选择一个。如何知道哪个实现是最快的，应该选择哪个？这正是 TunableOp 提供的作用。

启用 TunableOp 和单独调整

TunableOp 功能是独立于启用调整阶段本身来启用的。启用 TunableOp 意味着 PyTorch 将用其可调整实现替换任何标准操作。任何对 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 字段）可以更改为“默认”，它将回退到原始 PyTorch 未调整实现。或者，在 ROCm 的 hipBLAS 或 hipBLASLt 库的情况下，如果您知道特定的解决方案索引，可以通过替换值来覆盖 TunableOp 选择的解决方案。操作符名称和参数（第 1 和第 2 字段）是内部命名的，不应修改。在 GemmTunableOp 的情况下，第 1 字段表示数据类型以及输入是否转置（T）或未转置（N），第 2 字段表示 M、N、K 输入形状。

有一个选项可以启用详细输出，但仅推荐用于调试目的。这将产生大量诊断信息，但可能有助于查看是否使用了 TunableOp。否则，TunableOp 在文件输出之外完全静默，除非在使用过程中出现警告或错误。详细输出选项只能通过设置环境变量 PYTORCH_TUNABLEOP_VEROBSE=1 来获得。

关于调整行为、预热和缓存效果的说明

调整操作符包括遍历注册的实现列表并对每个实现进行性能分析。性能分析是通过多次运行单个实现并取平均执行时间来建立的。在调整之前还有一个可选的预热阶段，可以帮助硬件达到稳定的功耗状态。在调整工作负载时，各种硬件缓存更有可能产生命中，而不是在不调整时。还有选项可以刷新指令缓存和旋转输入张量，这可能有助于产生更真实的调整操作符性能分析，就像操作符在更大的工作负载中而不是在紧密的重复循环中运行一样。

默认情况下，对于给定的运算符的每个可能的解决方案，将运行 100 次迭代或 30 毫秒内可以运行的迭代次数，取较小者，并计算其平均执行时间。将选择所有成功分析中速度最快的解决方案。如果给定的解决方案没有达到默认实现的相同精度或返回错误代码，则分析可能会失败。

当前可调运算符

TunableGemm for ROCm

目前仅实现了 ROCm 的 TunableGemm。请注意，当使用 TunableOp 时，CUDA 构建的 PyTorch 将正常工作，但 CUDA 构建可用的唯一解决方案是“默认”实现，即现在通过 TunableOp 调用的原始 cuBLAS 默认实现。启用 TunableOp 时，对 at::cuda::blas::gemm()或::bgemm()的任何调用都将通过 TunableOp 路由。对于给定的一组输入参数（transa、transb、m、n、k），gemm()将尝试使用 rocblas 和 hipblaslt 中可用的最快实现。

离线调优

动机 ¶

离线调优有几种用例。

其中一个用例涉及高内存利用率的工作负载，在这种情况下，常规调优可能会导致内存耗尽。

另一个用例是针对计算密集型工作负载。在这种情况下，一次性收集工作负载的 GEMMs，然后使用不同的调优参数或库反复调优，会更加资源高效。

工作流程

基本上分为两个步骤：1) 设置环境变量以收集未调优的 GEMM，这将生成 tunableop_untuned0.csv ：

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

2. 运行一个 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 次，每次复制一次，以便用户可以在 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)

注意， mgpu_tune_gemm_in_file API 的使用与其单 GPU 对应版本（ tune_gemm_in_file ）不同。由于使用了 concurrent futures 模块，调用 API 的 Python 脚本主体必须像下面这样使用 main() 进行包装。 mgpu_tune_gemm_in_file 的参数必须包含通配符表达式（ ? 或 * ），以生成包含 GEMMs 的待处理未调优文件列表。 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)[source][source]

这是所有可调操作实现的开关。

torch.cuda.tunable.is_enabled()[source][source]

返回是否启用了 TunableOp 功能。

返回类型:

布尔型

torch.cuda.tunable.tuning_enable(val=True)[source][source]

启用对 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()[source][source]

获取用于调整给定解决方案的最大时间。

返回类型:

int

torch.cuda.tunable.set_max_tuning_iterations(iterations)

设置用于调整给定解决方案的最大迭代次数。

如果同时设置了最大调整持续时间和迭代次数，则将尊重两者中较小的一个。至少将始终运行 1 次调整迭代。

torch.cuda.tunable.get_max_tuning_iterations()

获取用于调整给定解决方案的最大迭代次数。

返回类型:

int

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

设置用于调整结果输入/输出的文件名。

如果 insert_device_ordinal 是 True ，则当前设备序号将自动添加到给定的文件名中。这可以在每个进程对应一个 GPU 的场景中使用，以确保所有进程都写入不同的文件。

torch.cuda.tunable.get_filename()

获取结果文件名。

返回类型:

str

torch.cuda.tunable.get_results()

返回所有可调操作结果。

返回类型:

tuple[str, str, str, float]

torch.cuda.tunable.get_validators()

返回可调操作的验证器。

返回类型:

tuple[str, str]

torch.cuda.tunable.write_file_on_exit(val)

在调整上下文销毁期间，将文件写入磁盘。

如果您的应用程序因正常操作或错误而终止，这很有用，可以作为最终将结果刷新到磁盘的操作。可以通过手动调用 write_file() 来手动刷新您的结果。

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

将结果写入 CSV 文件。

如果未提供 filename ，则调用 get_filename() 。

返回类型:

布尔型

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

从 TunableOp CSV 文件中读取结果。

如果未提供 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)

将旋转缓冲区大小设置为 MB，如果缓冲区大小大于零。

如果小于零，则查询 L2 缓存大小。如果等于零，则表示禁用旋转缓冲区。

torch.cuda.tunable.get_rotating_buffer_size()

获取旋转缓冲区大小（以千字节为单位）。

返回类型:

int