r
""
此模块公开了可调操作接口。
一些操作,例如 GEMMs,可以使用多个库来实现
或者使用多种技术。例如,GEMM 可以针对 CUDA 实现
ROCm 使用 BLAS 或 BLASLt 库。此外,ROCm 的 rocblas 和 hipblaslt 库允许用户查询所有可能的算法,然后选择一个。如何知道哪个实现是最快的,应该选择哪个?这正是 TunableOp 提供的作用。
hipblaslt 库允许用户查询所有可能的算法,然后选择一个。如何知道哪个实现是最快的,应该选择哪个?这正是 TunableOp 提供的作用。
如何知道哪个实现是最快的,应该选择哪个?这正是 TunableOp 提供的作用。
这正是 TunableOp 提供的作用。
启用可调 Op 并单独调整
========================================
可调 Op 功能是独立于启用调整阶段来启用的
。启用可调 Op 意味着 PyTorch 将用其可调实现替换任何标准
操作符。对 TunableOp 的任何调用首先
检查是否已经针对给定的操作符输入进行了调整。如果是的话,
它将立即调用调整后的操作;不会进行进一步调整
即使启用调谐设置。如果没有找到调谐结果,
启用和调整时,TunableOp 将基准测试每个已注册
给定输入集实现该运算符并选择最快的
最快的。
文件输入和输出
=====================
第一次调用任何可调运算符时,内部数据库将进行调优
操作将通过尝试从给定结果中读取来准备
文件。默认文件名为'tunableop_results.csv'。为了支持调整时
多个 GPU 在多个进程中使用,GPU 设备序号
自动插入到文件名中以避免多个进程覆盖
同一文件。
如果启用了调优并且在你使用过程中发现了新的调优
工作负载,它还会将所有调整写入相同的文件名
启动时读取的以及运行时发现的新内容。
可以用于通过多个工作负载构建调优文件。
通过重用相同的文件。当应用程序终止时,输出文件将自动创建。
此行为可以通过 C++和 Python API 进行控制,但不能通过环境变量。
控制此行为。
假设您指定了文件名,您将得到一个包含以下内容的 CSV 文件
看起来是这样的:
验证器,PT 版本,2.2.0
验证器,ROCM 版本,6.0.0.0-12969-1544e39
验证器,HIPBLASLT 版本,0.6.0-a9c5cc7
验证器,ROCBLAS 版本,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
注意“验证器”行。如果您更改库版本或 ROCm 版本,或 PyTorch 版本,
TunableOp 将检测到这一点并拒绝调整文件,因为之前的调整可能受到其他软件更改的影响。
剩余的行是遇到的每个 TunableOp 的调整解决方案。
这些是每个遇到的 TunableOp 的调整解决方案。
在您的执行过程中。每行包含 4 个由逗号分隔的字段:操作符
姓名、操作参数、解决方案名称和平均执行时间。
执行时间是可选字段。CSV 文件可以编辑,但需要
注意。例如,解决方案名称(字段 3)可以更改为“默认”
它将回退到原始的未调优的 PyTorch 实现。或者,在 ROCm 的 hipBLAS 或 hipBLASLt 库的情况下,如果您知道具体的解决方案索引,您可以替换 TunableOp 选择的解决方案。
如果您知道具体的解决方案索引,您可以替换 TunableOp 选择的解决方案。
通过替换值来覆盖。操作符名称和参数(字段 1 和 2)在内部被命名为。
操作符名称和参数(字段 1 和 2)在内部被命名为。
和不应修改。在 GemmTunableOp 的情况下,字段 1 表示
数据类型以及输入是否转置(T)或未转置(N)以及字段 2
指示 M、N、K 输入形状。
开启详细输出选项,但仅建议用于
用于调试目的。这将产生大量的诊断信息,但有助于查看是否使用了 TunableOp。否则,TunableOp 将完全静默,除非在使用过程中出现警告或错误
这可能有助于了解是否真的使用了 TunableOp。否则,TunableOp 除了文件输出外,将完全无声,除非在使用过程中出现警告或错误
除了文件输出外,除非在使用过程中出现警告或错误,否则 TunableOp 将完全静默
详细选项仅通过设置环境变量才能获得
变量 PYTORCH_TUNABLEOP_VEROBSE=1.
关于调整行为、预热和缓存效应的笔记
====================================================
调整操作符包括遍历列表或已注册
实现和配置每个实例。配置是通过运行
单次循环多次实现并取平均执行时间
时间。在调整之前还有一个可选的预热阶段,可以帮助
通过硬件达到稳定的功率状态。在调整工作负载时
各种硬件缓存在不调整时更可能产生命中。
有选项用于刷新指令缓存和旋转输入张量
可能有助于生成更忠实于调整操作符的配置文件,就像是在
操作员在更大的工作负载中运行,而不是在紧密、重复的循环中。
默认情况下,给定操作符的每个可能解决方案都将运行
100 次迭代或 30 毫秒内能运行的迭代次数,取其较小者,然后计算其平均执行时间。在所有成功进行性能分析的解决方案中,将选择最快的方案。如果给定的解决方案没有达到默认方案的相同精度,则性能分析可能会失败。
越小,其平均执行时间将越小。在所有成功进行性能分析的解决方案中,将选择最快的方案。如果给定的解决方案没有达到默认方案的相同精度,则性能分析可能会失败。
在所有成功进行性能分析的解决方案中,将选择最快的方案。如果给定的解决方案没有达到默认方案的相同精度,则性能分析可能会失败。
如果给定的解决方案没有达到默认方案的相同精度,则性能分析可能会失败。
实现或如果解决方案返回错误代码。
当前可调算子
=========================
ROCm 上的 TunableGemm
--------------------
目前仅实现了 ROCm 上的 TunableGemm。注意,CUDA 构建版本
PyTorch 在使用可调操作时将正常工作,但唯一可用的解决方案是 CUDA 构建的'Default'实现,即原始的 cuBLAS 默认版本,现在通过 TunableOp 调用。当启用时,对 at::cuda::blas::gemm()或::bgemm()的任何调用都将通过 TunableOp 路由。调用 gemm()进行
可用的解决方案是 CUDA 构建的'Default'实现,即原始的 cuBLAS 默认版本,现在通过 TunableOp 调用。任何调用 at::cuda::blas::gemm()或::bgemm()的
或::bgemm()将路由通过 TunableOp。调用 gemm()进行
时,将通过 TunableOp 路由。调用 gemm()进行
给定的输入参数集(transa, transb, m, n, k)将尝试使用
可用最快的实现,无论是 rocblas 还是 hipblaslt。
离线调优
==============
动机
----------
离线调优有几种用例。
其中一个用例涉及高内存利用率的工作负载,在这种情况下,常规调优可能会导致内存耗尽。
另一个用例是针对计算密集型工作负载。在这种情况下,收集
一次为工作负载设置 GEMM,然后使用不同的调整参数或库反复调整。
工作流程
--------
基本上有两个步骤:
1) 设置环境变量以收集未调整的 GEMM,这将生成`tunableop_untuned0.csv`文件:
.. 代码块 :: python
PYTORCH_TUNABLEOP_ENABLED=1
PYTORCH_TUNABLEOP_TUNING=0
PYTORCH_TUNABLEOP_RECORD_UNTUNED=1
...
2) 运行一个 Python 脚本,读取`tunableop_untuned0.csv`并生成`tunableop_results0.csv`,如下所示:
.. 代码块 :: python
import torch.cuda.tunable as tunable
导入 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")
也可能将多个未调优文件取出来,并将 GEMMs 分配给多个 GPU 进行调优
在单个节点内。首先,将 GEMM 收集起来,并消除重复的 GEMM。
然后,将 GEMM 分配到不同的 GPU 进行调优。所有 GEMM 调优完成后,
然后将所有 GPU 的结果收集到一个文件中,该文件的基文件名后附加 ``_full0``。
(例如 ``tunableop_results_full0.csv``)。最后,这个包含收集结果的新的文件,
每个 GPU 重复 N 次,方便用户,将使用调整过的配置运行工作负载。
在 N 个 GPU 上配置。
.. 代码块 :: python
if __name__ == "__main__":
num_gpus = 8 # 调优过程中使用的 GPU 数量
tunable.mgpu_tune_gemm_in_file("tunableop_untuned?.csv", num_gpus)
注意,`mgpu_tune_gemm_in_file` API 的使用与其单 GPU 版本不同
(`tune_gemm_in_file`)。调用 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 设置的任何设置。
环境变量优先于您使用 C++或 Python API 设置的任何设置。
环境变量接口
------------------------------
环境变量首次读取时会被缓存。您不能使用
环境变量接口以编程方式设置,因为设置变得固定。
使用 C++或 Python API 代替。
""
导入 concurrent.futures
导入
全局
导入
多进程
是 mp
导入
操作系统
导入 shutil
导入
警告
来自
打字
导入
可选
导入
火炬
全部 = [
启用,
是否启用,
调谐启用,
调谐是否启用,
记录未调谐启用,
记录未调谐是否启用,
设置最大调谐时长,
"获取最大调优时长",
"设置最大调优迭代次数",
"获取最大调优迭代次数",
"设置文件名",
"获取文件名",
"获取结果",
"获取验证器",
"退出时写入文件",
写入文件,
读取文件,
调整文件中的 gemm,
多 GPU 调整文件中的 gemm,
"设置旋转缓冲区大小",
"获取旋转缓冲区大小",
]
[docs]def 启用(val: bool = True) -> None:
r"""这是所有可调操作实现的大开关,用于开启/关闭功能。"""
torch._C._cuda_tunableop_enable(val) # 忽略未定义的属性
[文档]def is_enabled() -> bool:
r"""返回是否启用了 TunableOp 功能。"""
return torch._C._cuda_tunableop_is_enabled() # 忽略未定义的属性
[文档]def tuning_enable(val: bool = True) -> None:
r"""启用对可调操作的调整。
当启用时,如果找不到调整条目,则运行调整步骤并记录条目。
条目。
"""
"""
torch._C._cuda_tunableop_tuning_enable(val) # type: ignore[attr-defined]
[文档]def tuning_is_enabled() -> bool:
r"""返回是否可以调整 TunableOp 实现。"""
return torch._C._cuda_tunableop_tuning_is_enabled() # type: ignore[attr-defined]
[文档]def record_untuned_enable(val: bool = True) -> None:
r"""启用记录 TunableOp 操作的未调优信息以进行离线调优。
当启用时,如果找不到调优条目,则将其写入未调优文件。
"""
"""
torch._C._cuda_record_untuned_enable(val) # 忽略属性定义
[文档]def record_untuned_is_enabled() -> bool:
r"""返回是否记录了可调操作以进行离线调整。"""
return torch._C._cuda_record_untuned_is_enabled() # type: ignore[attr-defined]
[文档]def set_max_tuning_duration(duration: int) -> None:
r"""设置用于调整给定解决方案的最大时间(以毫秒为单位)。
如果同时设置了最大调整时间和迭代次数,则取两者中较小的一个
将获得荣誉。至少将运行 1 次调优迭代。
```python
# 假设输入文本为:
input_text = '"""'
# 翻译函数(此处仅为示例,实际翻译功能需要调用真实的翻译 API)
def translate_to_simplified_chinese(text):
# 这里应该调用真实的翻译 API 进行翻译
# 由于示例中不使用真实的 API,以下为模拟翻译结果
return text
# 输出翻译结果
translated_text = translate_to_simplified_chinese(input_text)
print(translated_text)
```
torch._C._cuda_tunableop_set_max_tuning_duration(duration) # 类型:忽略[属性定义]
[文档]def get_max_tuning_duration() -> int:
获取给定解决方案调优的最大时间。
return torch._C._cuda_tunableop_get_max_tuning_duration() # type: ignore[attr-defined]
[文档]def set_max_tuning_iterations(iterations: int) -> None:
设置给定解决方案调优的最大迭代次数。
如果同时设置了最大调优持续时间和迭代次数,则将尊重两者中较小的一个。至少将运行 1 次调优迭代。
将始终运行至少 1 次调优迭代。
"""
torch._C._cuda_tunableop_set_max_tuning_iterations(iterations) # type: ignore[attr-defined]
[文档]def get_max_tuning_iterations() -> int:
获取用于调整给定解决方案的最大迭代次数。
return torch._C._cuda_tunableop_get_max_tuning_iterations() # type: ignore[attr-defined]
[文档]def set_filename(filename: str, insert_device_ordinal: bool = False) -> None:
设置用于输入/输出调优结果的文件名。
如果 :attr:`insert_device_ordinal` 为 ``True``,则当前设备序号将自动添加到给定的文件名中。这可以在 1-process-per-gpu 场景中使用,以确保所有进程都写入单独的文件。
这可以用于确保在 1-process-per-gpu 场景中所有进程都写入不同的文件。
这可以用于确保在 1-process-per-gpu 场景中所有进程都写入不同的文件。
"""
"""
torch._C._cuda_tunableop_set_filename(filename, insert_device_ordinal) # type: ignore[attr-defined]
[文档]def get_filename() -> str:
r"""获取结果文件名。"""
return torch._C._cuda_tunableop_get_filename() # 忽略未定义的属性
[文档]def get_results() -> tuple[str, str, str, float]:
r"""返回所有可调操作的结果。"""
return torch._C._cuda_tunableop_get_results() # 忽略未定义的属性
[文档]def get_validators() -> tuple[str, str]:
r"""返回 TunableOp 验证器。"""
return torch._C._cuda_tunableop_get_validators() # type: ignore[attr-defined]
[文档]def write_file_on_exit(val: bool) -> None:
在调整上下文销毁期间,将文件写入磁盘。
这是一个在您的应用程序将结果最终写入磁盘时的有用操作
终止于正常操作或错误的结果。手动刷新
您的结果可以通过手动调用 `write_file()` 函数来实现。
torch._C._cuda_tunableop_write_file_on_exit(val) # 忽略未定义的属性
[文档]def write_file(filename: Optional[str] = None) -> bool:
r"""将结果写入 CSV 文件。
如果没有提供 :attr:`filename`,则调用 ``get_filename()``。
"""
如果文件名是 None:
文件名 = get_filename()
return torch._C._cuda_tunableop_write_file(filename) # type: ignore[attr-defined]
[文档]def read_file(filename: Optional[str] = None) -> bool:
r"""从可调操作 CSV 文件中读取结果。
如果没有提供 :attr:`filename`,则调用 ``get_filename()``。
"""
如果文件名是 None:
文件名 = get_filename()
return torch._C._cuda_tunableop_read_file(filename) # type: ignore[attr-defined]
[文档]def 设置旋转缓冲区大小(buffer_size: int) -> None:
将旋转缓冲区大小设置为 MB,如果缓冲区大小大于零。
如果小于零,查询 L2 缓存大小。如果等于零,表示停用旋转缓冲区。
```python
# 假设输入文本为:
input_text = '"""'
# 翻译函数(此处仅为示例,实际翻译功能需要调用真实的翻译 API)
def translate_to_simplified_chinese(text):
# 这里应该调用真实的翻译 API 进行翻译
# 由于示例中不使用真实的 API,以下为模拟翻译结果
return text
# 输出翻译结果
translated_text = translate_to_simplified_chinese(input_text)
print(translated_text)
```
返回 torch._C._cuda_tunableop_set_rotating_buffer_size(buffer_size) # type: ignore[attr-defined]
[文档]def get_rotating_buffer_size() -> int:
r"""获取旋转缓冲区大小(单位:千字节)。"""
return torch._C._cuda_tunableop_get_rotating_buffer_size() # type: ignore[attr-defined]
[文档]def tune_gemm_in_file(filename: str) -> None:
调整文件中的 GEMM。
断言已启用。
断言调整已启用。
deviceid = torch.cuda.current_device()
使用 open(filename) as file:
for line in file:
if line.startswith(("Gemm", "ScaledGemm")):
_process_single_offline_gemm(line, deviceid)
def 从文件中收集独特的未调优 gemm(
文件名模式: str) ->
设置[str
]
r处理多个未调优的结果文件,并返回一个去重后的集合。
独特的 gemm 条目 =
设置()
集合将避免重复
for 文件路径
在
球.
球(
文件名模式):
与
打开(
文件路径)
是
文件:
for 行
在
文件:
如果
行.
以...开头(("Gemm", "ScaledGemm")):
独特的 gemm 条目.
添加(
行)
返回
独特的 gemm 条目
def _收集可调操作结果() ->
无:
r从多个可调操作结果文件中收集结果并创建一个单独的文件。
gemm 行 =
设置()
验证器行 =
输入文本为空,请提供需要翻译的文本
需要允许结果文件名可能
使用 Python API 设置,而不是环境变量。
也可能结果文件名根本未设置。
有多个测试案例需要检查,但最终我们
需要一个可匹配的表达式
结果文件名 =
获取文件名()
# 空字符串可能返回此处
如果 (
结果文件名
是
不是
无
和
结果文件名 !=
请提供需要翻译的文本
): # 使用 Python API 设置文件名的情况
点位置 =
结果文件名.
找到(
“。”)
如果
点位置 != -1
和
点位置 > 0:
替换点左侧的字符
文件名模式 = (
结果文件名
[
点位置 - 1] +
"?" +
结果文件名[
点位置
]
)
否则:
文件名模式 =
请提供需要翻译的文本
# 需要使代码检查器满意
否则:
环境变量用于设置文件名的情况。
结果文件环境变量 = os.
获取环境变量("PYTORCH_TUNABLEOP_FILENAME")
如果
结果文件环境变量
是
无
或者
结果文件环境变量 ==
输入文本翻译为简体中文为:"":
文件名模式 =
"可调操作结果?.csv"
如果...否则 "%d"
在
结果文件环境变量名:
文件名模式 =
结果文件环境变量名.
替换("%d",
“?”)
否则:
文件名模式 =
结果文件环境变量名.
替换(
“。”, "?.")
断言
"?"
在
文件名模式
第一个文件 =
假
匹配的文件 =
球.
球(
文件名模式)
匹配文件数量 =
长度(
匹配文件)
for 文件路径
在
匹配文件:
与
打开(
文件路径)
是
文件:
for 行
在
文件:
如果
行.
以...开头(
验证器):
如果
不是 (
第一个文件):
# 只从第一个文件中读取验证器
验证行.append(
行)
否则:
GEMM 行.
添加(
行)
第一个文件 =
真实
输出文件 =
文件名模式.
替换(
“?”,
“全 0”)
与
打开(
输出文件,
w)
是
输出文件:
for 行
在
验证行:
输出文件.
写(
行)
for 行
在
GEMM 行:
输出文件.
写(
行)
创建 num_matching_copies 个结果文件副本
for i 在
范围(1,
匹配文件数量):
重复文件 =
输出文件.
替换("0", str(i))
shutil.复制(
输出文件,
重复文件)
def _处理单个离线 gemm(
未调优的 gemm 行: str, gpu_id: int) ->
无:
r处理单个未调优的 GEMM。
deviceid = "cuda:" + str(gpu_id)
dtype_dict = {
浮点数: torch.float32,
双精度浮点数: torch.float64,
"BFloat16": torch.bfloat16,
"Half": torch.半,
"c10::complex<double>": torch.complex128,
"c10::complex<float>": torch.complex64,
"Float8_e4m3fn": torch.float8_e4m3fn,
"Float8_e5m2": torch.float8_e5m2,
"Float8_e4m3fnuz": torch.float8_e4m3fnuz,
"Float8_e5m2fnuz": torch.float8_e5m2fnuz,
}
未调优的 gemm =
未调优的 gemm 行.strip().
分割(",")[:]
下划线数量 =
未调优的 gemm[0].
数量(
“_”)
# 初始化 dtype 以使代码检查器满意
dtype = 无
数据类型 A =
无
数据类型 B =
无
数据类型 C =
无
如果
下划线数量 == 2:
[操作符签名,
数据类型,
布局] =
未调优的 gemm[0].
分割(
“_”)
传输 A =
布局[0] == "T"
transB = 布局[1] == "T"
dtype = 数据类型字典.
获取(
数据类型)
否则:
# 缩放 GEMM
未调优的 gemm_temp =
未调优的 gemm[0].
分割(
“_”)
# dtypeC 可能不是 FP8 类型,请记录
# 下划线数量的
计算 =
未调优的 gemm_temp.
数量(
“_”)
操作符签名 =
未调优的 gemm_temp[0]
数据类型 A =
未调优的 gemm_temp[1] +
_ +
未调优的 gemm_temp[2]
数据类型 B =
未调优的 gemm_temp[3] +
_ +
未调优的 gemm_temp[4]
如果
计算 == 7:
数据类型 C =
未调优的 gemm_temp[5] +
_ +
未调优的 gemm_temp[6]
否则:
数据类型 C =
未调优的 gemm_temp[5]
传输 A =
未调优的 gemm_temp[
数量
]
[0] == "T"
transB = 未调优的 gemm_temp[
数量
]
[1] == "T"
数据类型 A =
数据类型字典.
获取(
数据类型 A)
数据类型 B =
数据类型字典.
获取(
数据类型 B)
数据类型 C =
数据类型字典.
获取(
数据类型 C)
未调优的 gemm_temp =
未调优的 gemm[1].
分割(
“_”)
[n, m, k] = [int(g) for g 在
未调优的 gemm_temp[1:4]]
如果
操作符签名 ==
"Gemm 可调操作":
矩阵 A = (
torch.随机(k, m,
数据类型=
数据类型,
设备=
设备 ID).t()
如果 transB
否则 torch.
随机(m, k,
数据类型=
数据类型,
设备=
设备 ID)
)
矩阵 B = (
torch.随机(n, k,
数据类型=
数据类型,
设备=
设备 ID).t()
如果
传输 A
否则 torch.
随机(k, n,
数据类型=
数据类型,
设备=
设备 ID)
)
torch.mm(材料 A,
材料 B)
如果...否则
操作符签名 ==
"可调 Gemm 分块批处理操作":
[b] = [int(g) for g 在
未调优的 gemm_temp[5:6]]
矩阵 A = (
torch.随机(b, k, m,
数据类型=
数据类型,
设备=
设备 ID)
如果 transB
否则 torch.
随机(b, m, k,
数据类型=
数据类型,
设备=
设备 ID)
)
矩阵 B = (
torch.随机(b, n, k,
数据类型=
数据类型,
设备=
设备 ID)
如果
传输 A
否则 torch.
随机(b, k, n,
数据类型=
数据类型,
设备=
设备 ID)
)
矩阵 A =
材料 A.
转置(1, 2)
如果 transB
否则
矩阵 A
矩阵 B =
材料 B.
转置(1, 2)
如果
传输 A
否则
矩阵 B
torch.bmm(材料 A,
材料 B)
如果...否则
操作符签名 ==
"缩放 Gemm 可调操作":
fillA = 0.25
fillB = 0.75
矩阵 A = (
torch.full((k, m), fillA, 数据类型=dtypeA,
设备=
设备 ID).t()
如果 transB
否则 torch.full((m, k), fillA,
数据类型=dtypeA,
设备=
设备 ID)
)
矩阵 B = (
torch.full((n, k), 填充 B,
数据类型=dtypeB,
设备=
设备 ID)
如果
传输 A
否则 torch.full((k, n),
填充 B,
数据类型=dtypeB,
设备=
设备 ID).t()
)
断言
未调优的 gemm_temp[8] ==
"读写"
如果
未调优的 gemm_temp[9] == "1":
行方向 =
真实
否则:
行方向 =
假
如果
行内:
标度 A = torch.
一((
材料 A.
形状[0
] 1),
设备=
设备 ID)
标度 B = torch.
一((1,
材料 B.
形状[0
)]
设备=
设备 ID)
否则:
标度 A = torch.
张量(0.8,
设备=
设备 ID)
标度 B = torch.
张量(0.9,
设备=
设备 ID)
断言
未调优的 gemm_temp[10] ==
偏差
如果
未调优的 gemm_temp[11] ==
无:
# 无偏置向量
torch.缩放毫米(
材料 A,
材料 B,
缩放_a=
规模 A, scale_b=
规模 B,
输出数据类型=
数据类型 C
)
否则:
# 偏置向量存在
fillbias = 0.10
偏置数据类型 =
数据类型字典.
获取(
未调优的 gemm_temp[11])
bias = (
torch.full((n,), 填充偏差,
数据类型=
偏差数据类型,
设备=
设备 ID)
如果
传输 A
否则 torch.full((m,),
填充偏差,
数据类型=
偏差数据类型,
设备=
设备 ID)
)
torch.缩放毫米(
材料 A,
材料 B,
缩放_a=
规模 A, scale_b=
规模 B,
输出数据类型=
数据类型 C,
偏差=bias
)
如果...否则
操作符签名 == "GemmAndBiasTunableOp":
# y = x*A^T + b
断言
传输 A != transB
X = (
torch.随机(k, m,
数据类型=
数据类型,
设备=
设备 ID).t()
如果 transB
否则 torch.
随机(m, k,
数据类型=
数据类型,
设备=
设备 ID)
)
矩阵 A = (
torch.随机(n, k,
数据类型=
数据类型,
设备=
设备 ID)
如果
传输 A
否则 torch.
随机(k, n,
数据类型=
数据类型,
设备=
设备 ID).t()
)
bias = (
torch.随机(n,
数据类型=
数据类型,
设备=
设备 ID)
如果
传输 A
否则 torch.
随机(m,
数据类型=
数据类型,
设备=
设备 ID)
)
torch.神经网络.
功能性.
线性(X,
材料 A,
偏差)
否则:
警告.
警告(f
错误:未知操作{
操作符签名}")
def _检查调整断言() ->
无:
r多 GPU 调整案例的辅助函数。需要检查 TunableOp 功能是否启用以及调整是否启用。
TunableOp 已被禁用。现在尝试启用。
"沉浸式翻译"
如果
启用()
是
错误:
警告.
警告(
"TunableOp 已被禁用。现在尝试启用。")
启用(
是)
断言
启用()
是
真实
断言
调试是否启用()
是
真实
断言
记录未调试是否启用()
是
假
[文档]def mgpu_tune_gemm_in_file(filename_pattern: str, num_gpus: int) -> None:
处理一个或多个文件,并将工作分配到一个或多个 GPU 上。
unique_gemm_entries = _gather_unique_untuned_gemm_from_files(filename_pattern)
total_gpus = torch.cuda.device_count()
assert 1 <= num_gpus <= total_gpus
mp_context = mp.get_context("spawn")
futures = [] # 空列表,用于存储 futures
flush_results = [] # 空列表,用于存储 flush_results
GEMM 以轮询方式分配给 GPU
h = 0
with concurrent.futures.ProcessPoolExecutor(
最大工作进程数 = num_gpus,
mp_context=mp_context,
initializer=_check_tuning_assertions,
) as executor:
# 工作进程是一个单独的进程。TunableOp 将被
# 在子进程中启用,如果 PYTORCH_TUNABLEOP_ENABLED=1
初始化器中,我们还尝试启用可调 OP
环境变量未设置。
for line in unique_gemm_entries:
future = executor.submit(_process_single_offline_gemm, line, h)
futures.append(future)
h = (h + 1) % num_gpus
for future in concurrent.futures.as_completed(futures):
future.result()
for g in range(num_gpus):
flush_result = executor.submit(write_file)
flush_results.append(flush_result)
for flush_result in concurrent.futures.as_completed(flush_results):
flush_result.result()
torch.cuda.synchronize()
_gather_tunableop_results()