环境变量¶
NCCL 拥有大量的环境变量,可以针对特定用途进行调整。
环境变量也可以静态地设置在 /etc/nccl.conf 中(供管理员设置系统范围的值)或 ${NCCL_CONF_FILE} 中(自 2.23 起;见下文)。例如,这些文件可以包含
NCCL_DEBUG=WARN
NCCL_SOCKET_IFNAME==ens1f0
环境变量分为两类。有些是使 NCCL 遵循系统特定配置所必需的,可以保存在脚本和系统配置中。“调试”部分列出的其他参数不应在生产环境中使用,也不应保留在脚本中,或仅作为权宜之计,并在问题解决后立即删除。保持设置它们可能会导致次优行为、崩溃或挂起。
系统配置¶
NCCL_SOCKET_IFNAME¶
NCCL_SOCKET_IFNAME
变量指定用于通信的 IP 接口。
接受的值¶
定义为前缀列表,以过滤 NCCL 要使用的接口。
可以提供多个前缀,用 ,
符号分隔。
使用 ^
符号,NCCL 将排除以该列表中任何前缀开头的接口。
要匹配(或不匹配)确切的接口名称,请以 =
字符开始前缀字符串。
示例
eth
: 使用所有以 eth
开头的接口,例如 eth0
, eth1
, …
=eth0
: 仅使用接口 eth0
=eth0,eth1
: 仅使用接口 eth0
和 eth1
^docker
: 不要使用任何以 docker
开头的接口
^=docker0
: 不要使用接口 docker0
。
注意:默认情况下,环回接口 (lo
) 和 docker 接口 (docker*
) 将不会被选择,除非没有其他可用接口。如果您更喜欢使用 lo
或 docker*
而不是其他接口,您需要使用 NCCL_SOCKET_IFNAME
显式选择它们。默认算法也会优先选择以 ib
开头的接口,而不是其他接口。设置 NCCL_SOCKET_IFNAME
将绕过自动接口选择算法,并可能使用所有匹配手动选择的接口。
NCCL_SOCKET_FAMILY¶
NCCL_SOCKET_FAMILY
变量允许用户强制 NCCL 仅使用 IPv4 或 IPv6 接口。
接受的值¶
设置为 AF_INET
以强制使用 IPv4,或 AF_INET6
以强制使用 IPv6。
NCCL_SOCKET_RETRY_CNT¶
(自 2.24 起)
NCCL_SOCKET_RETRY_CNT
变量指定在 ETIMEDOUT
、ECONNREFUSED
或 EHOSTUNREACH
错误后,NCCL 重试建立套接字连接的次数。
接受的值¶
默认值为 34,任何正值均有效。
NCCL_SOCKET_RETRY_SLEEP_MSEC¶
(自 2.24 起)
NCCL_SOCKET_RETRY_SLEEP_MSEC
变量指定在第一次 ETIMEDOUT
、ECONNREFUSED
或 EHOSTUNREACH
错误后,NCCL 等待多少毫秒才重试建立套接字连接。对于后续错误,等待时间与错误计数线性缩放。因此,总时间将为 (N+1) * N/2 * NCCL_SOCKET_RETRY_SLEEP_MSEC
,其中 N 由 NCCL_SOCKET_RETRY_CNT
给出。使用默认值 NCCL_SOCKET_RETRY_CNT
和 NCCL_SOCKET_RETRY_SLEEP_MSEC
,总重试时间约为 60 秒。
接受的值¶
默认值为 100 毫秒,任何正值均有效。
NCCL_SOCKET_NTHREADS¶
(自 2.4.8 起)
NCCL_SOCKET_NTHREADS
变量指定每个网络连接用于套接字传输的 CPU 辅助线程数。增加此值可能会提高套接字传输性能,但会增加 CPU 使用率。
接受的值¶
1 到 16。在 AWS 上,默认值为 2;在具有 gVNIC 网络接口的 Google Cloud 实例上,默认值为 4(自 2.5.6 起);在其他情况下,默认值为 1。
对于通用的 100G 网络,此值可以手动设置为 4。但是,NCCL_SOCKET_NTHREADS
和 NCCL_NSOCKS_PERTHREAD
的乘积不能超过 64。另请参阅 NCCL_NSOCKS_PERTHREAD
。
NCCL_NSOCKS_PERTHREAD¶
(自 2.4.8 起)
NCCL_NSOCKS_PERTHREAD
变量指定套接字传输的每个辅助线程打开的套接字数。在每个套接字速度受限的环境中,将此变量设置为大于 1 可能会提高网络性能。
接受的值¶
在 AWS 上,默认值为 8;在其他情况下,默认值为 1。
对于通用的 100G 网络,此值可以手动设置为 4。但是,NCCL_SOCKET_NTHREADS
和 NCCL_NSOCKS_PERTHREAD
的乘积不能超过 64。另请参阅 NCCL_SOCKET_NTHREADS
。
NCCL_CROSS_NIC¶
NCCL_CROSS_NIC
变量控制 NCCL 是否应允许环/树使用不同的 NIC,从而导致节点间通信在不同的节点上使用不同的 NIC。
为了在使用多个 NIC 时最大化节点间通信性能,NCCL 尝试在节点之间通信时使用相同的 NIC,以允许网络设计中节点上的每个 NIC 连接到不同的网络交换机(网络轨道),并避免任何流量干扰的风险。NCCL_CROSS_NIC
设置因此取决于网络拓扑,尤其取决于网络结构是否针对轨道进行了优化。
这对于只有一个 NIC 的系统没有影响。
接受的值¶
0:始终为同一环/树使用相同的 NIC,以避免跨越网络轨道。适用于具有每个 NIC 交换机(轨道)的网络,且轨道间连接速度较慢。请注意,如果通信器在每个节点上不包含相同的 GPU,NCCL 可能仍然需要跨 NIC 进行通信。
1:允许为同一环/树使用不同的 NIC。这适用于节点上的所有 NIC 都连接到同一交换机的网络,因此尝试跨相同的 NIC 进行通信并不能帮助避免流量冲突。
2:(默认)尝试为同一环/树使用相同的 NIC,但如果这样可以获得更好的性能,仍然允许使用不同的 NIC。
NCCL_IB_HCA¶
NCCL_IB_HCA
变量指定用于通信的 RDMA 接口。
接受的值¶
定义为过滤 NCCL 要使用的 IB Verbs 接口。列表以逗号分隔;端口号可以使用 :
符号指定。可选前缀 ^
表示该列表是排除列表。第二个可选前缀 =
表示标记是确切的名称,否则默认情况下 NCCL 会将每个标记视为前缀。
示例
mlx5
: 使用所有以 mlx5
开头的卡的所有端口
=mlx5_0:1,mlx5_1:1
: 使用卡 mlx5_0
和 mlx5_1
的端口 1。
^=mlx5_1,mlx5_4
: 不要使用卡 mlx5_1
和 mlx5_4
。
注意:使用不带前导 =
的 mlx5_1
将选择 mlx5_1
以及 mlx5_10
到 mlx5_19
(如果它们存在)。因此,始终建议添加 =
前缀以确保完全匹配。
NCCL_IB_TIMEOUT¶
NCCL_IB_TIMEOUT
变量控制 InfiniBand Verbs 超时。
超时时间计算为 4.096 µs * 2 ^ timeout,正确的值取决于网络的大小。增加该值可能有助于非常大的网络,例如,如果 NCCL 在调用 ibv_poll_cq 时因错误 12 而失败。
有关更多信息,请参阅 InfiniBand 规范卷 1(本地确认超时)的第 12.7.34 节。
NCCL_IB_RETRY_CNT¶
(自 2.1.15 起)
NCCL_IB_RETRY_CNT
变量控制 InfiniBand 重试计数。
有关更多信息,请参阅 InfiniBand 规范卷 1 的第 12.7.38 节。
接受的值¶
默认值为 7。
NCCL_IB_GID_INDEX¶
(自 2.1.4 起)
NCCL_IB_GID_INDEX
变量定义 RoCE 模式下使用的全局 ID 索引。请参阅 InfiniBand show_gids 命令以设置此值。
有关更多信息,请参阅 InfiniBand 规范卷 1 或供应商文档。
接受的值¶
默认值为 -1。
NCCL_IB_ADDR_FAMILY¶
(自 2.21 起)
当 NCCL_IB_GID_INDEX
保持未设置状态时,NCCL_IB_ADDR_FAMILY
变量定义与 NCCL 动态选择的 infiniband GID 关联的 IP 地址族。
接受的值¶
默认值为 “AF_INET”。
NCCL_IB_ADDR_RANGE¶
(自 2.21 起)
当 NCCL_IB_GID_INDEX
保持未设置状态时,NCCL_IB_ADDR_RANGE
变量定义 NCCL 动态选择的有效 GID 范围。
NCCL_IB_ROCE_VERSION_NUM¶
(自 2.21 起)
当 NCCL_IB_GID_INDEX
保持未设置状态时,NCCL_IB_ROCE_VERSION_NUM
变量定义与 NCCL 动态选择的 infiniband GID 关联的 RoCE 版本。
接受的值¶
默认值为 2。
NCCL_IB_FIFO_TC¶
(自 2.22.3 起)
定义控制消息的 InfiniBand 流量类别。控制消息是短 RDMA 写入操作,用于控制信用返回,这与其他传输大数据段的 RDMA 操作相反。此设置允许这些消息使用高优先级、低延迟的流量类别,并避免被其余流量延迟。
接受的值¶
默认值是由 NCCL_IB_TC 设置的流量类别,如果未设置,则默认为 0。
NCCL_IB_RETURN_ASYNC_EVENTS¶
(自 2.23 起)
IB 事件作为警告报告给用户。如果启用,NCCL 还将在发生致命 IB 异步事件时停止 IB 通信。
接受的值¶
默认值为 1,设置为 0 以禁用
NCCL_OOB_NET_ENABLE¶
(自 2.23 起)变量 NCCL_OOB_NET_ENABLE
启用将 NCCL net 用于带外通信。启用 NCCL net 的使用将更改通信器初始化期间执行的 allgather 的实现。
接受的值¶
将变量设置为 0 以禁用,设置为 1 以启用。
NCCL_OOB_NET_IFNAME¶
(自 2.23 起)如果为带外通信启用了 NCCL net(请参阅 NCCL_OOB_NET_ENABLE
),则 NCCL_OOB_NET_IFNAME
变量指定要使用的网络接口。
接受的值¶
定义为过滤 NCCL 用于带外通信的接口。接受的接口列表取决于 NCCL 使用的网络。列表以逗号分隔;端口号可以使用 :
符号指定。可选前缀 ^
表示该列表是排除列表。第二个可选前缀 =
表示标记是确切的名称,否则默认情况下 NCCL 会将每个标记视为前缀。如果指定了多个设备,NCCL 将选择列表中第一个匹配的设备。
示例
NCCL_NET="IB" NCCL_OOB_NET_ENABLE=1 NCCL_OOB_NET_IFNAME="=mlx5_1"
将使用 Infiniband NET,接口为 mlx5_1
NCCL_NET="IB" NCCL_OOB_NET_ENABLE=1 NCCL_OOB_NET_IFNAME="mlx5_1"
将使用 Infiniband NET,接口为在 mlx5_1
、mlx5_10
、mlx5_11
等列表中找到的第一个接口。
NCCL_NET="Socket" NCCL_OOB_NET_ENABLE=1 NCCL_OOB_NET_IFNAME="ens1"
将使用套接字 NET,接口为在 ens1f0
、ens1f1
等列表中找到的第一个接口。
NCCL_UID_STAGGER_THRESHOLD¶
(自 2.23 起)NCCL_UID_STAGGER_THRESHOLD
变量用于触发 NCCL 秩和 ncclUniqueId 之间的通信交错,以避免 ncclUniqueId 溢出。如果通信的 NCCL 秩数超过指定阈值,则使用秩值交错通信(请参阅下面的 NCCL_UID_STAGGER_RATE)。如果每个 ncclUniqueId 的 NCCL 秩数小于或等于阈值,则不执行交错。
例如,如果我们有 128 个 NCCL 秩、1 个 ncclUniqueId 和 64 的阈值,则执行交错。但是,如果 2 个 ncclUniqueIds 与 128 个 NCCL 秩一起使用,并且阈值为 64,则不执行交错。
接受的值¶
NCCL_UID_STAGGER_THRESHOLD
的值必须是严格的正整数。如果未指定,则默认值为 256。
NCCL_UID_STAGGER_RATE¶
(自 2.23 起)
NCCL_UID_STAGGER_RATE
变量用于定义在交错 NCCL 秩和 ncclUniqueId 之间的通信时,作为目标的邮件速率。如果使用交错(请参阅上面的 NCCL_UID_STAGGER_THRESHOLD),则邮件速率用于计算给定 NCCL 秩必须等待的时间。
接受的值¶
NCCL_UID_STAGGER_RATE
的值必须是严格的正整数,以消息/秒表示。如果未指定,则默认值为 7000。
NCCL_NET¶
(自 2.10 起)
强制 NCCL 使用特定网络,例如确保 NCCL 使用外部插件,并且不会自动回退到内部 IB 或套接字实现。设置此环境变量将覆盖所有通信器中的 netName
配置(请参阅 ncclConfig_t);如果未设置(未定义),则网络模块将由配置确定;如果未传递配置,NCCL 将自动选择最佳网络模块。
接受的值¶
NCCL_NET 的值必须与使用的 NCCL 网络的名称完全匹配(不区分大小写)。内部网络名称为 “IB”(通用 IB verbs)和 “Socket”(TCP/IP 套接字)。外部网络插件定义自己的名称。默认值为未定义。
NCCL_NET_PLUGIN¶
(自 2.11 起)
- 将其设置为后缀字符串或库名称,以在多个 NCCL net 插件中进行选择。此设置将导致 NCCL 使用以下策略查找 net 插件库
- 如果设置了 NCCL_NET_PLUGIN,请尝试加载名称由 NCCL_NET_PLUGIN 指定的库;
- 如果设置了 NCCL_NET_PLUGIN 且先前失败,请尝试加载 libnccl-net-<NCCL_NET_PLUGIN>.so;
- 如果未设置 NCCL_NET_PLUGIN,请尝试加载 libnccl-net.so;
- 如果未找到任何插件(既不是用户定义的也不是默认的),请使用内部网络插件。
例如,设置 NCCL_NET_PLUGIN=foo
将导致 NCCL 尝试加载 foo
,如果找不到 foo
,则尝试加载 libnccl-net-foo.so
(前提是它在系统上存在)。
接受的值¶
插件后缀、插件文件名或 “none”。
NCCL_TUNER_PLUGIN¶
- 将其设置为后缀字符串或库名称,以在多个 NCCL tuner 插件中进行选择。此设置将导致 NCCL 使用以下策略查找 tuner 插件库
- 如果设置了 NCCL_TUNER_PLUGIN,则尝试加载由 NCCL_TUNER_PLUGIN 指定名称的库;
- 如果设置了 NCCL_TUNER_PLUGIN 且之前失败,则尝试加载 libnccl-net-<NCCL_TUNER_PLUGIN>.so;
- 如果未设置 NCCL_TUNER_PLUGIN,则尝试加载 libnccl-tuner.so;
- 如果未找到任何插件,请在 net 插件中查找 tuner 符号(请参阅
NCCL_NET_PLUGIN
); - 如果未找到任何插件(既不是通过 NCCL_TUNER_PLUGIN 也不是通过 NCCL_NET_PLUGIN),则使用内部 tuner 插件。
例如,设置 NCCL_TUNER_PLUGIN=foo
将导致 NCCL 尝试加载 foo
,如果找不到 foo
,则尝试加载 libnccl-tuner-foo.so
(前提是系统上存在该文件)。
接受的值¶
插件后缀、插件文件名或 “none”。
NCCL_PROFILER_PLUGIN¶
- 将其设置为后缀字符串或库名称,以在多个 NCCL profiler 插件中进行选择。此设置将导致 NCCL 使用以下策略查找 profiler 插件库
- 如果设置了 NCCL_PROFILER_PLUGIN,则尝试加载由 NCCL_PROFILER_PLUGIN 指定名称的库;
- 如果设置了 NCCL_PROFILER_PLUGIN 且之前失败,则尝试加载 libnccl-profiler-<NCCL_PROFILER_PLUGIN>.so;
- 如果未设置 NCCL_PROFILER_PLUGIN,则尝试加载 libnccl-profiler.so;
- 如果未找到任何插件(既不是用户定义的也不是默认的),则不启用 profiling。
- 如果 NCCL_PROFILER_PLUGIN 设置为
STATIC_PLUGIN
,则将在程序二进制文件中搜索插件符号。
例如,设置 NCCL_PROFILER_PLUGIN=foo
将导致 NCCL 尝试加载 foo
,如果找不到 foo
,则尝试加载 libnccl-profiler-foo.so
(前提是系统上存在该文件)。
接受的值¶
插件后缀、插件文件名或 “none”。
NCCL_IGNORE_CPU_AFFINITY¶
(自 2.4.6 版本起)
NCCL_IGNORE_CPU_AFFINITY
变量可用于使 NCCL 忽略作业提供的 CPU affinity,而仅使用 GPU affinity。
接受的值¶
默认值为 0,设置为 1 将使 NCCL 忽略作业提供的 CPU affinity。
NCCL_CONF_FILE¶
(自 2.23 起)
NCCL_CONF_FILE
变量允许用户指定包含静态配置的文件。这不接受路径中包含 ~
字符;请先转换为相对路径或绝对路径。
接受的值¶
如果未设置或版本早于 2.23,NCCL 将使用主目录中的 .nccl.conf(如果可用)。
NCCL_DEBUG¶
NCCL_DEBUG
变量控制从 NCCL 显示的调试信息。此变量通常用于调试。
接受的值¶
VERSION - 在程序启动时打印 NCCL 版本。
WARN - 每当任何 NCCL 调用出错时,打印显式错误消息。
INFO - 打印调试信息
TRACE - 在每次调用时打印可重放的跟踪信息。
NCCL_DEBUG_FILE¶
(自 2.2.12 版本起)
NCCL_DEBUG_FILE
变量将 NCCL 调试日志输出定向到一个文件。文件名格式可以设置为 filename.%h.%p,其中 %h 替换为主机名,%p 替换为进程 PID。这不接受路径中包含 ~
字符,请先转换为相对路径或绝对路径。
接受的值¶
除非设置了此环境变量,否则默认输出文件为 stdout。
设置 NCCL_DEBUG_FILE
将导致 NCCL 创建并覆盖任何同名的先前文件。
注意:如果文件名在所有作业进程中不是唯一的,则输出可能会丢失或损坏。
NCCL_DEBUG_SUBSYS¶
(自 2.3.4 版本起)
NCCL_DEBUG_SUBSYS
变量允许用户根据子系统过滤 NCCL_DEBUG=INFO
输出。该值应是以逗号分隔的子系统列表,以包含在 NCCL 调试日志跟踪中。
在子系统名称前加上 ‘^’ 将禁用该子系统的日志记录。
接受的值¶
默认值为 INIT,BOOTSTRAP,ENV。
支持的子系统名称包括 INIT(代表初始化)、COLL(代表 collectives,集合通信)、P2P(代表点对点)、SHM(代表共享内存)、NET(代表网络)、GRAPH(代表拓扑检测和图搜索)、TUNING(代表算法/协议调优)、ENV(代表环境设置)、ALLOC(代表内存分配)、CALL(标准函数调用)、PROXY(代表代理线程操作)、NVLS(标准 NVLink SHARP)、BOOTSTRAP(代表早期初始化)、REG(代表内存注册)、PROFILE(代表粗粒度初始化 profiling)、RAS(代表可靠性、可用性和可维护性子系统)和 ALL(包括每个子系统)。
调试¶
这些环境变量应谨慎使用。新版本的 NCCL 可能以不同的方式工作,强制它们使用特定值将阻止 NCCL 自动选择最佳设置。因此,从长远来看,它们可能会导致性能问题,甚至破坏某些功能。
它们可以用于实验或调试问题,但通常不应为生产代码设置。
NCCL_P2P_DISABLE¶
NCCL_P2P_DISABLE
变量禁用点对点 (P2P) 传输,该传输使用 CUDA 直接访问 GPU 之间,使用 NVLink 或 PCI。
接受的值¶
定义并设置为 1 以禁用直接 GPU 到 GPU (P2P) 通信。
NCCL_P2P_LEVEL¶
(自 2.3.4 版本起)
NCCL_P2P_LEVEL
变量允许用户精细控制何时在 GPU 之间使用点对点 (P2P) 传输。该级别定义了 NCCL 将使用 P2P 传输的 GPU 之间的最大距离。应使用表示路径类型的短字符串来指定使用 P2P 传输的拓扑截止值。
如果未指定此项,NCCL 将尝试根据其运行的架构和环境最佳地选择一个值。
接受的值¶
- LOC : 永不使用 P2P(始终禁用)
- NVL : 当 GPU 通过 NVLink 连接时使用 P2P
- PIX : 当 GPU 位于同一 PCI 交换机上时使用 P2P。
- PXB : 当 GPU 通过 PCI 交换机连接时使用 P2P(可能存在多个跃点)。
- PHB : 当 GPU 位于同一 NUMA 节点上时使用 P2P。流量将通过 CPU。
- SYS : 在 NUMA 节点之间使用 P2P,可能跨越 SMP 互连(例如 QPI/UPI)。
整数值(旧版)¶
还可以选择将 NCCL_P2P_LEVEL
声明为与路径类型对应的整数。这些数值是为了向后兼容性而保留的,适用于那些在允许使用字符串之前使用数值的人。
由于路径类型中的重大更改,不建议使用整数值 - 字面值可能会随时间变化。为避免调试配置时的麻烦,请使用字符串标识符。
- LOC : 0
- PIX : 1
- PXB : 2
- PHB : 3
- SYS : 4
大于 4 的值将被解释为 SYS。旧版整数值不支持 NVL。
NCCL_P2P_DIRECT_DISABLE¶
NCCL_P2P_DIRECT_DISABLE
变量禁止 NCCL 通过同一进程的 GPU 之间的 P2P 直接访问用户缓冲区。当用户缓冲区使用不自动使其可供同一进程管理的其他 GPU 访问且具有 P2P 访问权限的 API 分配时,这很有用。
接受的值¶
定义并设置为 1 以禁用跨 GPU 的直接用户缓冲区访问。
NCCL_SHM_DISABLE¶
NCCL_SHM_DISABLE
变量禁用共享内存 (SHM) 传输。当点对点无法发生时,在设备之间使用 SHM,因此使用主机内存。当禁用 SHM 时,NCCL 将使用网络(即 InfiniBand 或 IP 套接字)在 CPU 插槽之间进行通信。
接受的值¶
定义并设置为 1 以禁用通过共享内存 (SHM) 进行通信。
NCCL_BUFFSIZE¶
NCCL_BUFFSIZE
变量控制 NCCL 在 GPU 对之间通信数据时使用的缓冲区大小。
如果您在使用 NCCL 时遇到内存约束问题,或者您认为不同的缓冲区大小可以提高性能,请使用此变量。
NCCL_NTHREADS¶
NCCL_NTHREADS
变量设置每个 CUDA 线程块的 CUDA 线程数。NCCL 将为每个通信通道启动一个 CUDA 线程块。
如果您认为您的 GPU 时钟频率较低并且想要增加线程数,请使用此变量。
您也可以使用此变量来减少线程数,以降低 GPU 工作负载。
NCCL_MAX_NCHANNELS¶
(自 2.0.5 版本起为 NCCL_MAX_NRINGS,自 2.5.0 版本起为 NCCL_MAX_NCHANNELS)
NCCL_MAX_NCHANNELS
变量限制了 NCCL 可以使用的通道数。减少通道数也会减少用于通信的 CUDA 线程块的数量,从而减少对 GPU 计算资源的影响。
旧的 NCCL_MAX_NRINGS
变量(在 2.4 版本之前使用)在新版本中仍然作为别名工作,但如果设置了 NCCL_MAX_NCHANNELS
,则会被忽略。
此环境变量已被 NCCL_MAX_CTAS
取代,后者也可以使用 ncclCommInitRankConfig 以编程方式设置。
接受的值¶
任何大于或等于 1 的值。
NCCL_MIN_NCHANNELS¶
(自 2.2.0 版本起为 NCCL_MIN_NRINGS,自 2.5.0 版本起为 NCCL_MIN_NCHANNELS)
NCCL_MIN_NCHANNELS
变量控制您希望 NCCL 使用的最小通道数。增加通道数也会增加 NCCL 使用的 CUDA 线程块的数量,这可能有助于提高性能;但是,它会使用更多的 CUDA 计算资源。
当在通常 NCCL 只创建一个通道的平台上使用聚合 collectives 时,这尤其有用。
旧的 NCCL_MIN_NRINGS
变量(在 2.4 版本之前使用)在新版本中仍然作为别名工作,但如果设置了 NCCL_MIN_NCHANNELS
,则会被忽略。
此环境变量已被 NCCL_MIN_CTAS
取代,后者也可以使用 ncclCommInitRankConfig 以编程方式设置。
接受的值¶
默认值取决于平台。设置为整数值,最高为 12(最高到 2.2 版本)、16(2.3 和 2.4 版本)或 32(2.5 及更高版本)。
NCCL_CHECKS_DISABLE¶
(自 2.0.5 版本起,在 2.2.12 版本中已弃用)
NCCL_CHECKS_DISABLE
变量可用于禁用每次 collective 调用时的参数检查。检查在开发期间很有用,但会增加延迟。可以禁用它们以提高生产环境中的性能。
接受的值¶
默认值为 0,设置为 1 以禁用检查。
NCCL_CHECK_POINTERS¶
(自 2.2.12 版本起)
NCCL_CHECK_POINTERS
变量启用每次 collective 调用时对 CUDA 内存指针的检查。检查在开发期间很有用,但会增加延迟。
NCCL_IB_DISABLE¶
NCCL_IB_DISABLE
变量阻止 NCCL 使用 IB/RoCE 传输。相反,NCCL 将回退到使用 IP 套接字。
接受的值¶
定义并设置为 1 以禁用 InfiniBand Verbs 用于通信(并强制使用另一种方法,例如 IP 套接字)。
NCCL_IB_QPS_PER_CONNECTION¶
(自 2.10 起)
每个等级之间连接使用的 IB 队列对的数量。这在多级 fabric 上可能很有用,这些 fabric 需要多个队列对才能具有良好的路由熵。有关在多个 QP 上拆分数据的不同方法,请参阅 NCCL_IB_SPLIT_DATA_ON_QPS
,因为它会影响性能。
接受的值¶
介于 1 和 128 之间的数字,默认值为 1。
NCCL_IB_SPLIT_DATA_ON_QPS¶
(自 2.18 版本起)
此参数控制当我们创建多个队列对时如何使用队列对。设置为 1(拆分模式),每条消息将在每个队列对上均匀拆分。如果使用许多 QP,这可能会导致明显的延迟下降。设置为 0(轮询模式),队列对将在轮询模式下用于我们发送的每条消息。不发送多条消息的操作将不会使用所有 QP。
接受的值¶
0 或 1。默认值为 0(自 NCCL 2.20 版本起)。将其设置为 1 将启用拆分模式(2.18 和 2.19 版本中的默认值)。
NCCL_IB_CUDA_SUPPORT¶
(在 2.4.0 版本中已删除,请参阅 NCCL_NET_GDR_LEVEL)
NCCL_IB_CUDA_SUPPORT
变量用于强制或禁用 GPU Direct RDMA 的使用。默认情况下,如果拓扑允许,NCCL 会启用 GPU Direct RDMA。此变量可以禁用此行为或在所有情况下强制使用 GPU Direct RDMA。
NCCL_IB_PCI_RELAXED_ORDERING¶
(自 2.12 版本起)
为 IB Verbs 传输启用 Relaxed Ordering 的使用。Relaxed Ordering 可以极大地帮助虚拟化环境中 InfiniBand 网络的性能。
接受的值¶
设置为 2 以在可用时自动使用 Relaxed Ordering。设置为 1 以强制使用 Relaxed Ordering,如果不可用则失败。设置为 0 以禁用 Relaxed Ordering 的使用。默认值为 2。
NCCL_IB_ADAPTIVE_ROUTING¶
(自 2.16 版本起)
为 IB Verbs 传输启用支持自适应路由的数据传输的使用。自适应路由可以提高大规模通信的性能。必须相应地选择系统定义的启用自适应路由的 SL(参见 NCCL_IB_SL
)。
接受的值¶
在 IB 网络上默认启用 (1)。在 RoCE 网络上默认禁用 (0)。设置为 1 以强制使用支持自适应路由的数据传输。
NCCL_IB_ECE_ENABLE¶
(自 2.23 起)
在 IB Verbs 网络上启用增强连接建立 (ECE) 的使用。
接受的值¶
默认启用 (1)。设置为 0 以禁用 ECE 网络功能的使用。
NCCL_MEM_SYNC_DOMAIN¶
(自 2.16 版本起)
为 NCCL 内核设置默认内存同步域(CUDA 12.0 和 sm90 及更高版本)。当 NCCL 内核和应用程序计算内核使用不同的域时,内存同步域可以帮助消除它们之间的干扰。
接受的值¶
默认值为 cudaLaunchMemSyncDomainRemote
(1)。当前支持的值为 0 和 1。
NCCL_CUMEM_ENABLE¶
(自 2.18 版本起)
在 NCCL 中使用 CUDA cuMem* 函数分配内存。
接受的值¶
0 或 1。在 2.18 版本中默认值为 0(禁用);自 2.19 版本起,如果系统支持此功能,则默认情况下会自动启用此功能(NCCL_CUMEM_ENABLE 仍然可以用于覆盖自动检测)。
NCCL_CUMEM_HOST_ENABLE¶
(自 2.23 起)
在 NCCL 中使用 CUDA cuMem* 函数分配主机内存。
接受的值¶
0 或 1。在 2.23 版本中默认值为 0;自 2.24 版本起,如果 CUDA 驱动程序 >= 12.6 且 CUDA 运行时 >= 12.2,则默认值为 1
NCCL_NET_GDR_LEVEL(以前为 NCCL_IB_GDR_LEVEL)¶
(自 2.3.4 版本起。在 2.4.0 版本中,NCCL_IB_GDR_LEVEL 重命名为 NCCL_NET_GDR_LEVEL)
NCCL_NET_GDR_LEVEL
变量允许用户精细控制何时在 NIC 和 GPU 之间使用 GPU Direct RDMA。该级别定义了 NIC 和 GPU 之间的最大距离。应使用表示路径类型的字符串来指定 GpuDirect 的拓扑截止值。
如果未指定此项,NCCL 将尝试根据其运行的架构和环境最佳地选择一个值。
接受的值¶
- LOC : 永不使用 GPU Direct RDMA(始终禁用)。
- PIX : 当 GPU 和 NIC 位于同一 PCI 交换机上时使用 GPU Direct RDMA。
- PXB : 当 GPU 和 NIC 通过 PCI 交换机连接时使用 GPU Direct RDMA(可能存在多个跃点)。
- PHB : 当 GPU 和 NIC 位于同一 NUMA 节点上时使用 GPU Direct RDMA。流量将通过 CPU。
- SYS : 即使跨越 NUMA 节点之间的 SMP 互连(例如 QPI/UPI)也使用 GPU Direct RDMA(始终启用)。
整数值(旧版)¶
还可以选择将 NCCL_NET_GDR_LEVEL
声明为与路径类型对应的整数。这些数值是为了向后兼容性而保留的,适用于那些在允许使用字符串之前使用数值的人。
由于路径类型中的重大更改,不建议使用整数值 - 字面值可能会随时间变化。为避免调试配置时的麻烦,请使用字符串标识符。
- LOC : 0
- PIX : 1
- PXB : 2
- PHB : 3
- SYS : 4
大于 4 的值将被解释为 SYS。
NCCL_NET_GDR_READ¶
NCCL_NET_GDR_READ
变量启用 GPU Direct RDMA,以便在发送数据时使用,只要 GPU-NIC 距离在 NCCL_NET_GDR_LEVEL
指定的距离内即可。在 2.4.2 版本之前,GDR 读取默认禁用,即发送数据时,数据首先存储在 CPU 内存中,然后传输到 InfiniBand 卡。自 2.4.2 版本起,对于基于 NVLink 的平台,默认启用 GDR 读取。
注意:在某些平台(例如 PCI-E)上,当发送数据时,直接从 GPU 内存读取数据已知比从 CPU 内存读取数据稍慢。
接受的值¶
0 或 1。定义并设置为 1 以使用 GPU Direct RDMA 直接将数据发送到 NIC(绕过 CPU)。
在 2.4.2 版本之前,所有平台的默认值为 0。自 2.4.2 版本起,对于基于 NVLink 的平台,默认值为 1,否则为 0。
NCCL_SINGLE_RING_THRESHOLD¶
(自 2.1 版本起,在 2.3 版本中已删除)
NCCL_SINGLE_RING_THRESHOLD
变量设置了 NCCL 将仅使用一个环的限制。这将限制带宽,但提高延迟。
NCCL_ALGO¶
(自 2.5 版本起)
NCCL_ALGO
变量定义了 NCCL 将使用的算法。
接受的值¶
(自 2.5 版本起)
算法的逗号分隔列表(不区分大小写),包括
版本 | 算法 |
---|---|
2.5+ | 环 (Ring) |
2.5+ | 树 (Tree) |
2.5 到 2.13 | Collnet |
2.14+ | CollnetChain |
2.14+ | CollnetDirect |
2.17+ | NVLS |
2.18+ | NVLSTree |
2.23+ | PAT |
NVLS 和 NVLSTree 启用 NVLink SHARP 卸载。
要指定要排除的算法(而不是包含),请以 ^
开头列表。
(自 2.24 起)
接受的值已扩展以允许更大的灵活性,如果找到意外的标记,解析将发出警告并失败。此外,如果 ring
未指定为有效算法,则在没有其他有效算法用于该函数的情况下,它不会隐式回退到 ring
。相反,它将失败。
格式现在是以分号分隔的函数名称和算法列表对,其中第一个条目的函数名称是可选的。如果未提供函数名称,则该设置应用于所有后续未列出的函数。冒号分隔函数(如果存在)和逗号分隔的算法列表。此外,如果逗号分隔的算法列表的第一个字符是插入符号 (^
),则所有选择都会被反转。
例如,NCCL_ALGO="ring,collnetdirect;allreduce:tree,collnetdirect;broadcast:ring"
将为所有函数启用 ring 和 collnetdirect 算法,然后为 allreduce 启用 tree 和 collnetdirect 算法,并为 broadcast 启用 ring 算法。
并且,NCCL_ALGO=allreduce:^tree
将允许所有函数使用默认算法(所有可用算法),但 allreduce 函数除外,allreduce 函数将使用除 tree 之外的所有可用算法。
默认值是未设置,这将导致 NCCL 根据节点拓扑和架构自动选择可用的算法。
NCCL_PROTO¶
(自 2.5 版本起)
NCCL_PROTO
变量定义了 NCCL 将允许使用的协议。
不建议用户设置此变量,除非怀疑 NCCL 中存在错误而需要禁用特定协议。 特别是,在不支持 LL128 的平台上启用 LL128 可能会导致数据损坏。
接受的值¶
(自 2.5 版本起)逗号分隔的协议列表(不区分大小写),包括:LL
、LL128
和 Simple
。 要指定要排除的协议(而不是要包含的协议),请以 ^
开头列表。
默认行为是启用所有支持的算法:在支持 LL128 的平台上等同于 LL,LL128,Simple
,否则等同于 LL,Simple
。
(自 2.24 版本起)接受的值已扩展以允许更大的灵活性,就像上面针对 NCCL_ALGO
描述的那样,允许用户为每个函数指定协议。
NCCL_P2P_PXN_LEVEL¶
(自 2.12 版本起)
控制在哪些情况下将 PXN 用于发送/接收操作。
接受的值¶
值为 0 将禁用 PXN 用于发送/接收。值为 1 将在目标首选的 NIC 无法通过 PCI 交换机访问时启用 PXN。值为 2(默认值)将导致始终使用 PXN,即使 NIC 通过 PCI 交换机连接,也会将来自节点内所有 GPU 的数据存储在中间 GPU 上以最大化聚合。
NCCL_RUNTIME_CONNECT¶
(自 2.22 版本起)
在运行时(例如,调用 ncclAllreduce())而不是初始化阶段动态连接对等方。
接受的值¶
默认为 1,设置为 0 以在初始化阶段连接对等方。
NCCL_GRAPH_REGISTER¶
(自 2.11 起)
当 CUDA Graphs 捕获 NCCL 调用时,启用用户缓冲区注册。
仅在以下情况下有效:(i)正在使用 CollNet 算法;(ii)节点内的所有 GPU 彼此具有 P2P 访问权限;(iii)每个进程最多有一个 GPU。
用户缓冲区注册可以减少用户缓冲区和 NCCL 内部缓冲区之间的数据复制次数。 当 CUDA Graphs 被销毁时,用户缓冲区将自动注销。
接受的值¶
0 或 1。默认值为 1(启用)。
NCCL_LEGACY_CUDA_REGISTER¶
(自 2.24 起)
通过 cudaMalloc(和相关的内存分配器)分配的 Cuda 缓冲区是旧式缓冲区。 注册旧式缓冲区可能会导致隐式同步,这不安全,并可能导致 NCCL 挂起。 默认情况下,NCCL 禁用旧式缓冲区注册,用户应迁移到基于 cuMem 的内存分配器以进行缓冲区注册。
接受的值¶
0 或 1。默认值为 0(禁用)。
NCCL_SET_STACK_SIZE¶
(自 2.9 版本起)
将 CUDA 内核堆栈大小设置为所有 NCCL 内核中的最大堆栈大小。
它可以避免加载时 CUDA 内存重新配置。 如果您遇到因 CUDA 内存重新配置而导致的挂起,请设置为 1。
接受的值¶
0 或 1。默认值为 0(禁用)。
NCCL_GRAPH_MIXING_SUPPORT¶
(自 2.13 版本起)
启用/禁用对来自并行 CUDA 图形或 CUDA 图形和非捕获 NCCL 调用的多个未完成 NCCL 调用的支持。 NCCL 调用被认为是未完成的,从它们的主机端启动开始(例如,对于非捕获调用调用 ncclAllreduce() 或对于捕获调用调用 cudaGraphLaunch()),到设备内核执行完成时结束。 如果禁用图形混合支持,则不支持以下用例
- 从并行图形启动中使用 NCCL 通信器(或拆分共享通信器),其中并行是指在没有依赖关系的、会序列化其执行的不同流上。
- 在正在进行的图形启动期间启动非捕获 NCCL 集合通信,该图形启动使用相同的通信器(或拆分共享通信器),而与流排序无关。
禁用支持的能力是由于观察到当启用支持并且多个 rank 从同一线程通过 cudaGraphLaunch 启动工作时,CUDA 启动中出现挂起。
接受的值¶
0 或 1。默认为 1(启用)。
NCCL_DMABUF_ENABLE¶
(自 2.13 版本起)
使用 Linux dma-buf 子系统启用 GPU Direct RDMA 缓冲区注册。
Linux dma-buf 子系统允许支持 GPU Direct RDMA 的 NIC 直接读取和写入 CUDA 缓冲区,而无需 CPU 参与。
接受的值¶
0 或 1。默认值为 1(启用),但如果 Linux 内核或 CUDA/NIC 驱动程序不支持该功能,则会自动禁用该功能。
NCCL_P2P_LL_THRESHOLD¶
(自 2.14 版本起)
NCCL_P2P_LL_THRESHOLD
是 NCCL 将对 P2P 操作使用 LL 协议的最大消息大小。
接受的值¶
十进制数字。 默认值为 16384。
NCCL_ALLOC_P2P_NET_LL_BUFFERS¶
(自 2.14 版本起)
NCCL_ALLOC_P2P_NET_LL_BUFFERS
指示通信器为所有 P2P 网络连接分配专用 LL 缓冲区。 这使所有 rank 都可以对低于 NCCL_P2P_LL_THRESHOLD
大小的延迟受限的发送和接收操作使用 LL 协议。 节点内 P2P 传输始终分配专用 LL 缓冲区。 如果运行具有大量 rank 的 all-to-all 工作负载,这将导致较高的扩展内存开销。
接受的值¶
0 或 1。默认值为 0(禁用)。
NCCL_COMM_BLOCKING¶
(自 2.14 版本起)
NCCL_COMM_BLOCKING
变量控制是否允许 NCCL 调用阻塞。 这包括对 NCCL 的所有调用,包括 init/finalize 函数,以及由于发送/接收调用的连接延迟初始化而可能阻塞的通信函数。 设置此环境变量将覆盖所有通信器中的 blocking
配置(请参阅 ncclConfig_t); 如果未设置(未定义),则通信器行为将由配置决定; 如果未传递配置,则通信器是阻塞的。
接受的值¶
0 或 1。 1 表示阻塞通信器,0 表示非阻塞通信器。 默认值是未定义的。
NCCL_CGA_CLUSTER_SIZE¶
(自 2.16 版本起)
设置 CUDA Cooperative Group Array (CGA) 集群大小。 在 sm90 及更高版本中,我们有一个额外的层次结构,我们可以在 Grid 中将多个块组合在一起,称为线程块集群。 将此设置为非零值将导致 NCCL 启动通信内核,并相应地设置 Cluster Dimension 属性。 设置此环境变量将覆盖所有通信器中的 cgaClusterSize
配置(请参阅 ncclConfig_t); 如果未设置(未定义),则 CGA 集群大小将由配置决定; 如果未传递配置,则 NCCL 将自动选择最佳值。
接受的值¶
0 到 8。 默认值是未定义的。
NCCL_MAX_CTAS¶
(自 2.17 版本起)
设置 NCCL 应使用的最大 CTA 数量。 设置此环境变量将覆盖所有通信器中的 maxCTAs
配置(请参阅 ncclConfig_t); 如果未设置(未定义),则最大 CTA 将由配置决定; 如果未传递配置,则 NCCL 将自动选择最佳值。
接受的值¶
设置为高达 32 的正整数值。 默认值是未定义的。
NCCL_MIN_CTAS¶
(自 2.17 版本起)
设置 NCCL 应使用的最小 CTA 数量。 设置此环境变量将覆盖所有通信器中的 minCTAs
配置(请参阅 ncclConfig_t); 如果未设置(未定义),则最小 CTA 将由配置决定; 如果未传递配置,则 NCCL 将自动选择最佳值。
接受的值¶
设置为高达 32 的正整数值。 默认值是未定义的。
NCCL_NVLS_ENABLE¶
(自 2.17 版本起)
启用 NVLink SHARP (NVLS) 的使用。 NVLink SHARP 在第三代 NVSwitch 系统 (NVLink4) 中可用,具有 Hopper 及更高版本的 GPU 架构,允许将 ncclAllReduce
等集合通信卸载到 NVSwitch 域。 在不支持该功能的系统上,NVLS 将自动禁用。
接受的值¶
默认为自动检测,定义并设置为 0 以禁用 NVLink SHARP 的使用。
NCCL_IB_MERGE_NICS¶
(自 2.20 版本起)
使 NCCL 能够将双端口 IB NIC 合并为单个逻辑网络设备。 这使 NCCL 能够更轻松地聚合双端口 NIC 带宽。
接受的值¶
默认为 1(启用),定义并设置为 0 以禁用 NIC 合并
NCCL_MNNVL_ENABLE¶
(自 2.21 起)
启用 NCCL 在可用时使用多节点 NVLink (MNNVL)。 如果系统或驱动程序不支持多节点 NVLink,则 MNNVL 将自动禁用。 此功能还需要启用 NCCL CUMEM 支持 (NCCL_CUMEM_ENABLE
)。 MNNVL 需要为构成 NVLink 域的所有节点配置完整且可操作的 IMEX 域。 有关 IMEX 域的更多详细信息,请参阅 CUDA 文档。
接受的值¶
默认为自动检测,定义并设置为 0 以禁用 MNNVL 支持。
NCCL_RAS_ENABLE¶
(自 2.24 起)
启用 NCCL 的可靠性、可用性和可维护性 (RAS) 子系统,该子系统可用于在执行期间查询 NCCL 作业的健康状况(请参阅 RAS)。
接受的值¶
默认为 1(启用); 定义并设置为 0 以禁用 RAS。
NCCL_RAS_ADDR¶
(自 2.24 起)
指定 RAS 子系统将侦听客户端连接的套接字的 IP 地址和端口号。 RAS 可以在多个进程之间共享此套接字,但如果多个独立的 NCCL 作业共享单个节点(并且这些作业属于不同的用户,则操作系统将不允许共享套接字),则不希望这样做。 在这种情况下,每个作业都应使用不同的值启动(例如,localhost:12345
、localhost:12346
等)。 由于通常使用 localhost
,因此只有有权访问运行作业的节点的人员才能连接到套接字。 如果需要,可以改为指定外部可访问的网络接口的地址,这将使 RAS 可以从其他节点(例如集群的头节点)访问,但这具有应考虑的安全隐患。
接受的值¶
默认值为 localhost:28028
。 主机名或 IP 地址都可以用作第一部分; IPv6 地址需要用方括号括起来(例如,[::1]
)。