-
Notifications
You must be signed in to change notification settings - Fork 2.9k
Closed
Description
Checklist
- 1. I have searched related issues but cannot get the expected help.
- 2. The bug has not been fixed in the latest version.
- 3. Please note that if the bug-related issue you submitted lacks corresponding environment info and a minimal reproducible demo, it will be challenging for us to reproduce and resolve the issue, reducing the likelihood of receiving feedback.
- 4. If the issue you raised is not a bug but a question, please raise a discussion at https://github.com/sgl-project/sglang/discussions/new/choose Otherwise, it will be closed.
- 5. Please use English, otherwise it will be closed.
Describe the bug
We observe successful single-node (8×H20) deployment of DeepSeek-R1 with SGLang+DeepEP (deep-ep-mode=auto), but encounter failures in 2-node configuration with the following error:
[2025-04-09 10:35:31 TP2] Scheduler hit an exception: Traceback (most recent call last):
File "/usr/local/lib/python3.11/site-packages/sglang/srt/managers/scheduler.py", line 1999, in run_scheduler_process
scheduler = Scheduler(server_args, port_args, gpu_id, tp_rank, dp_rank)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/usr/local/lib/python3.11/site-packages/sglang/srt/managers/scheduler.py", line 249, in __init__
self.tp_worker = TpWorkerClass(
^^^^^^^^^^^^^^
File "/usr/local/lib/python3.11/site-packages/sglang/srt/managers/tp_worker_overlap_thread.py", line 63, in __init__
self.worker = TpModelWorker(server_args, gpu_id, tp_rank, dp_rank, nccl_port)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/usr/local/lib/python3.11/site-packages/sglang/srt/managers/tp_worker.py", line 74, in __init__
self.model_runner = ModelRunner(
^^^^^^^^^^^^
File "/usr/local/lib/python3.11/site-packages/sglang/srt/model_executor/model_runner.py", line 178, in __init__
self.initialize(min_per_gpu_memory)
File "/usr/local/lib/python3.11/site-packages/sglang/srt/model_executor/model_runner.py", line 188, in initialize
self.load_model()
File "/usr/local/lib/python3.11/site-packages/sglang/srt/model_executor/model_runner.py", line 400, in load_model
self.model = get_model(
^^^^^^^^^^
File "/usr/local/lib/python3.11/site-packages/sglang/srt/model_loader/__init__.py", line 22, in get_model
return loader.load_model(
^^^^^^^^^^^^^^^^^^
File "/usr/local/lib/python3.11/site-packages/sglang/srt/model_loader/loader.py", line 365, in load_model
model = _initialize_model(
^^^^^^^^^^^^^^^^^^
File "/usr/local/lib/python3.11/site-packages/sglang/srt/model_loader/loader.py", line 146, in _initialize_model
return model_class(
^^^^^^^^^^^^
File "/usr/local/lib/python3.11/site-packages/sglang/srt/models/deepseek_v2.py", line 1352, in __init__
self.model = DeepseekV2Model(
^^^^^^^^^^^^^^^^
File "/usr/local/lib/python3.11/site-packages/sglang/srt/models/deepseek_v2.py", line 1278, in __init__
[
File "/usr/local/lib/python3.11/site-packages/sglang/srt/models/deepseek_v2.py", line 1279, in <listcomp>
DeepseekV2DecoderLayer(
File "/usr/local/lib/python3.11/site-packages/sglang/srt/models/deepseek_v2.py", line 1077, in __init__
self.mlp = DeepseekV2MoE(
^^^^^^^^^^^^^^
File "/usr/local/lib/python3.11/site-packages/sglang/srt/models/deepseek_v2.py", line 259, in __init__
self.deepep_dispatcher = DeepEPDispatcher(
^^^^^^^^^^^^^^^^^
File "/usr/local/lib/python3.11/site-packages/sglang/srt/layers/moe/ep_moe/token_dispatcher.py", line 537, in __init__
self._low_latency_dispatcher = _DeepEPDispatcherImplLowLatency(
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/usr/local/lib/python3.11/site-packages/sglang/srt/layers/moe/ep_moe/token_dispatcher.py", line 350, in __init__
self.buffer_low_latency = _get_buffer_low_latency(
^^^^^^^^^^^^^^^^^^^^^^^^
File "/usr/local/lib/python3.11/site-packages/sglang/srt/layers/moe/ep_moe/token_dispatcher.py", line 79, in _get_buffer_low_latency
_buffer_low_latency = Buffer(
^^^^^^^
File "/usr/local/lib64/python3.11/site-packages/deep_ep/buffer.py", line 89, in __init__
self.runtime.sync(device_ids, ipc_handles, root_unique_id)
RuntimeError: Failed: Assertion error /home/moyun.zty/DeepEP/csrc/kernels/runtime.cu:56 'nvshmem_team_split_strided(NVSHMEM_TEAM_WORLD, rank % NUM_MAX_NVL_PEERS, NUM_MAX_NVL_PEERS, num_ranks / NUM_MAX_NVL_PEERS, &cpu_rdma_team_config, 0, &cpu_rdma_team) == 0'
Followed WARN
logs:
WARN: Invalid <start, stride, size>: child <6, 8, 2>, parent <0, 1, 2>
WARN: Invalid <start, stride, size>: child <4, 8, 2>, parent <0, 1, 2>
WARN: Invalid <start, stride, size>: child <5, 8, 2>, parent <0, 1, 2>
WARN: Invalid <start, stride, size>: child <7, 8, 2>, parent <0, 1, 2>
WARN: Starting PE (1) or ending PE (9) is invalid
WARN: Invalid <start, stride, size>: child <3, 8, 2>, parent <0, 1, 2>
WARN: Starting PE (0) or ending PE (8) is invalid
WARN: Invalid <start, stride, size>: child <2, 8, 2>, parent <0, 1, 2>
Reproduction
# node 0
NCCL_IB_GID_INDEX=3 SUPPORT_CUTLASS_BLOCK_FP8=1 python -m sglang.launch_server \
--model-path /path/to/deepseek-ai__DeepSeek-R1 \
--tp-size 16 \
--dist-init-addr ${master_ip}:20001 \
--nnodes 2 \
--node-rank 0 \
--enable-ep-moe \
--enable-deepep-moe \
--deepep-mode auto \
--disable-cuda-graph \
--trust-remote-code \
--quantization fp8 \
--log-level info \
--chunked-prefill-size -1 \
--disable-radix-cache \
--context-length 65535 \
--max-running-requests 128 \
--stream-output \
> sglang-deep-ep-auto.log 2>&1 &
# node 1
NCCL_IB_GID_INDEX=3 SUPPORT_CUTLASS_BLOCK_FP8=1 python -m sglang.launch_server \
--model-path /path/to/deepseek-ai__DeepSeek-R1 \
--tp-size 16 \
--dist-init-addr ${master_ip}:20001 \
--nnodes 2 \
--node-rank 1 \
--enable-ep-moe \
--enable-deepep-moe \
--deepep-mode auto \
--disable-cuda-graph \
--trust-remote-code \
--quantization fp8 \
--log-level info \
--chunked-prefill-size -1 \
--disable-radix-cache \
--context-length 65535 \
--max-running-requests 128 \
--stream-output \
> sglang-deep-ep-auto.log 2>&1 &
Environment
$ python3.11 -m sglang.check_env
Python: 3.11.11 (main, Dec 16 2024, 17:23:09) [GCC 10.2.1 20200825 (Alibaba 10.2.1-3.8 2.32)]
CUDA available: True
GPU 0,1,2,3,4,5,6,7: NVIDIA H20
GPU 0,1,2,3,4,5,6,7 Compute Capability: 9.0
CUDA_HOME: /usr/local/cuda
NVCC: Cuda compilation tools, release 12.4, V12.4.99
CUDA Driver Version: 550.127.08
PyTorch: 2.6.0+cu124
sglang: 0.4.5
sgl_kernel: 0.0.6
flashinfer: Module Not Found
triton: 3.2.0
transformers: 4.50.3
torchao: 0.9.0
numpy: 1.26.4
aiohttp: 3.11.16
fastapi: 0.115.12
hf_transfer: Module Not Found
huggingface_hub: 0.30.1
interegular: 0.3.3
modelscope: Module Not Found
orjson: 3.10.16
outlines: 0.1.11
packaging: 24.2
psutil: 7.0.0
pydantic: 2.11.1
multipart: Module Not Found
zmq: 0.0.0
uvicorn: 0.34.0
uvloop: 0.21.0
vllm: 0.8.2
xgrammar: 0.1.16
openai: 1.70.0
tiktoken: 0.9.0
anthropic: Module Not Found
litellm: Module Not Found
decord: 0.6.0
NVIDIA Topology:
GPU0 GPU1 GPU2 GPU3 GPU4 GPU5 GPU6 GPU7 NIC0 NIC1 NIC2 NIC3 CPU Affinity NUMA Affinity GPU NUMA ID
GPU0 X NV18 NV18 NV18 NV18 NV18 NV18 NV18 NODE NODE SYS SYS 0-47,96-143 0 N/A
GPU1 NV18 X NV18 NV18 NV18 NV18 NV18 NV18 PIX NODE SYS SYS 0-47,96-143 0 N/A
GPU2 NV18 NV18 X NV18 NV18 NV18 NV18 NV18 NODE NODE SYS SYS 0-47,96-143 0 N/A
GPU3 NV18 NV18 NV18 X NV18 NV18 NV18 NV18 NODE PIX SYS SYS 0-47,96-143 0 N/A
GPU4 NV18 NV18 NV18 NV18 X NV18 NV18 NV18 SYS SYS PIX NODE 48-95,144-191 1 N/A
GPU5 NV18 NV18 NV18 NV18 NV18 X NV18 NV18 SYS SYS NODE NODE 48-95,144-191 1 N/A
GPU6 NV18 NV18 NV18 NV18 NV18 NV18 X NV18 SYS SYS NODE PIX 48-95,144-191 1 N/A
GPU7 NV18 NV18 NV18 NV18 NV18 NV18 NV18 X SYS SYS NODE NODE 48-95,144-191 1 N/A
NIC0 NODE PIX NODE NODE SYS SYS SYS SYS X NODE SYS SYS
NIC1 NODE NODE NODE PIX SYS SYS SYS SYS NODE X SYS SYS
NIC2 SYS SYS SYS SYS PIX NODE NODE NODE SYS SYS X NODE
NIC3 SYS SYS SYS SYS NODE NODE PIX NODE SYS SYS NODE X
Legend:
X = Self
SYS = Connection traversing PCIe as well as the SMP interconnect between NUMA nodes (e.g., QPI/UPI)
NODE = Connection traversing PCIe as well as the interconnect between PCIe Host Bridges within a NUMA node
PHB = Connection traversing PCIe as well as a PCIe Host Bridge (typically the CPU)
PXB = Connection traversing multiple PCIe bridges (without traversing the PCIe Host Bridge)
PIX = Connection traversing at most a single PCIe bridge
NV# = Connection traversing a bonded set of # NVLinks
NIC Legend:
NIC0: mlx5_bond_0
NIC1: mlx5_bond_1
NIC2: mlx5_bond_2
NIC3: mlx5_bond_3
ulimit soft: 204800
Machine Info
Environment Variables
export NCCL_IB_DISABLE=0
export NCCL_SOCKET_IFNAME=bond0
export NCCL_NET_PLUGIN=''
export NCCL_IB_GID_INDEX=3
export NCCL_IB_TIMEOUT=22
export NCCL_IB_RETRY_CNT=7
export NCCL_IB_SL=5
export NCCL_IB_TC=136
export NCCL_IB_HCA="mlx5_bond_0,mlx5_bond_1,mlx5_bond_2,mlx5_bond_3"
export NCCL_DEBUG=WARN
export NCCL_SET_THREAD_NAME=1
export NCCL_IB_QPS_PER_CONNECTION=8
export NCCL_SET_THREAD_NAME=1
export NCCL_DEBUG_SUBSYS=INIT,TUNING,GRAPH
export NVSHMEM_DEBUG=WARN
export NVSHMEM_ENABLE_NIC_PE_MAPPING=1
export NVSHMEM_HCA_PE_MAPPING="mlx5_bond_0:1:2,mlx5_bond_1:1:2,mlx5_bond_2:1:2,mlx5_bond_3:1:2"
export NVSHMEM_IBRC_ROCE_LAG_PORT_SELECTION=3
export NVSHMEM_IB_GID_INDEX=3
export NVSHMEM_IB_TRAFFIC_CLASS=64
export GLOO_SOCKET_IFNAME=eth0
Gdrcopy Info
$ LD_LIBRARY_PATH=/opt/gdrcopy/lib /opt/gdrcopy/bin/gdrcopy_copybw
GPU id:0; name: NVIDIA H20; Bus id: 0000:08:00
GPU id:1; name: NVIDIA H20; Bus id: 0000:7e:00
GPU id:2; name: NVIDIA H20; Bus id: 0000:a2:00
GPU id:3; name: NVIDIA H20; Bus id: 0000:c6:00
GPU id:4; name: NVIDIA H20; Bus id: 0001:09:00
GPU id:5; name: NVIDIA H20; Bus id: 0001:7f:00
GPU id:6; name: NVIDIA H20; Bus id: 0001:a3:00
GPU id:7; name: NVIDIA H20; Bus id: 0001:c7:00
selecting device 0
testing size: 131072
rounded size: 131072
gpu alloc fn: cuMemAlloc
device ptr: 7f825d400000
map_d_ptr: 0x7f848c0a6000
info.va: 7f825d400000
info.mapped_size: 131072
info.page_size: 65536
info.mapped: 1
info.wc_mapping: 1
page offset: 0
user-space pointer:0x7f848c0a6000
writing test, size=131072 offset=0 num_iters=10000
write BW: 17993.9MB/s
reading test, size=131072 offset=0 num_iters=100
read BW: 716.355MB/s
unmapping buffer
unpinning buffer
closing gdrdrv
Nvshmem Info
$ LD_LIBRARY_PATH=/opt/nvshmem/lib/ /opt/nvshmem/bin/nvshmem-info -a
NVSHMEM v3.2.5
Build Information:
CUDA API 12040
CUDA Driver 12040
Build Timestamp Apr 2 2025 14:47:31
Build Variables
NVSHMEM_DEBUG=OFF NVSHMEM_DEVEL=OFF NVSHMEM_DEFAULT_PMI2=OFF
NVSHMEM_DEFAULT_PMIX=OFF NVSHMEM_DEFAULT_UCX=OFF
NVSHMEM_ENABLE_ALL_DEVICE_INLINING=OFF NVSHMEM_GPU_COLL_USE_LDST=OFF
NVSHMEM_IBGDA_SUPPORT=ON NVSHMEM_IBGDA_SUPPORT_GPUMEM_ONLY=OFF
NVSHMEM_IBDEVX_SUPPORT=OFF NVSHMEM_IBRC_SUPPORT=ON
NVSHMEM_MPI_SUPPORT=OFF NVSHMEM_NVTX=ON NVSHMEM_PMIX_SUPPORT=OFF
NVSHMEM_SHMEM_SUPPORT=OFF NVSHMEM_TEST_STATIC_LIB=OFF
NVSHMEM_TIMEOUT_DEVICE_POLLING=OFF NVSHMEM_TRACE=OFF NVSHMEM_UCX_SUPPORT=OFF
NVSHMEM_USE_DLMALLOC=OFF NVSHMEM_USE_NCCL=OFF NVSHMEM_USE_GDRCOPY=ON
NVSHMEM_VERBOSE=OFF CUDA_HOME=/usr/local/cuda GDRCOPY_HOME=/opt/gdrcopy
LIBFABRIC_HOME=/usr/local/libfabric MPI_HOME=/usr/local/ompi
NCCL_HOME=/usr/local/nccl NVSHMEM_PREFIX=/usr/local/nvshmem PMIX_HOME=/usr
SHMEM_HOME=/usr/local/ompi UCX_HOME=/usr/local/ucx
Standard options:
NVSHMEM_VERSION false (type: bool, default: false)
Print library version at startup
NVSHMEM_INFO false (type: bool, default: false)
Print environment variable options at startup
NVSHMEM_DISABLE_NVLS false (type: bool, default: false)
Disable NVLS SHARP resources for collectives, even if available for platform
NVSHMEM_SYMMETRIC_SIZE 1073741824 (type: size, default: 1073741824)
Specifies the size (in bytes) of the symmetric heap memory per PE. The
size is implementation-defined and must be at least as large as the integer
ceiling of the product of the numeric prefix and the scaling factor. The
character suffixes for the scaling factor are as follows:
* k or K multiplies by 2^10 (kibibytes)
* m or M multiplies by 2^20 (mebibytes)
* g or G multiplies by 2^30 (gibibytes)
* t or T multiplies by 2^40 (tebibytes)
For example, string '20m' is equivalent to the integer value 20971520, or 20
mebibytes. Similarly the string '3.1M' is equivalent to the integer value
3250586. Only one multiplier is recognized and any characters following the
multiplier are ignored, so '20kk' will not produce the same result as '20m'.
Usage of string '.5m' will yield the same result as the string '0.5m'.
An invalid value for NVSHMEM_SYMMETRIC_SIZE is an error, which the NVSHMEM
library shall report by either returning a nonzero value from
nvshmem_init_thread or causing program termination.
NVSHMEM_DEBUG "WARN" (type: string, default: "")
Set to enable debugging messages.
Optional values: VERSION, WARN, INFO, ABORT, TRACE
Bootstrap options:
NVSHMEM_BOOTSTRAP "PMI" (type: string, default: "PMI")
Name of the default bootstrap that should be used to initialize NVSHMEM.
Allowed values: PMI, MPI, SHMEM, plugin, UID
NVSHMEM_BOOTSTRAP_PMI "PMI" (type: string, default: "PMI")
Name of the PMI bootstrap that should be used to initialize NVSHMEM.
Allowed values: PMI, PMI-2, PMIX
NVSHMEM_BOOTSTRAP_PLUGIN "" (type: string, default: "")
Absolute path to or name of the bootstrap plugin file to load when
NVSHMEM_BOOTSTRAP=plugin is specified
NVSHMEM_BOOTSTRAP_MPI_PLUGIN "nvshmem_bootstrap_mpi.so.3" (type: string, default: "nvshmem_bootstrap_mpi.so.3")
Absolute path to or name of the MPI bootstrap plugin file.
NVSHMEM will search for the plugin based on linux linker priorities. See man
dlopen
NVSHMEM_BOOTSTRAP_SHMEM_PLUGIN "nvshmem_bootstrap_shmem.so.3" (type: string, default: "nvshmem_bootstrap_shmem.so.3")
Absolute path to or name of the SHMEM bootstrap plugin file.
NVSHMEM will search for the plugin based on linux linker priorities. See man
dlopen
NVSHMEM_BOOTSTRAP_PMI_PLUGIN "nvshmem_bootstrap_pmi.so.3" (type: string, default: "nvshmem_bootstrap_pmi.so.3")
Absolute path to or name of the PMI bootstrap plugin file.
NVSHMEM will search for the plugin based on linux linker priorities. See man
dlopen
NVSHMEM_BOOTSTRAP_PMI2_PLUGIN "nvshmem_bootstrap_pmi2.so.3" (type: string, default: "nvshmem_bootstrap_pmi2.so.3")
Absolute path to or name of the PMI-2 bootstrap plugin file.
NVSHMEM will search for the plugin based on linux linker priorities. See man
dlopen
NVSHMEM_BOOTSTRAP_PMIX_PLUGIN "nvshmem_bootstrap_pmix.so.3" (type: string, default: "nvshmem_bootstrap_pmix.so.3")
Absolute path to or name of the PMIx bootstrap plugin file.
NVSHMEM will search for the plugin based on linux linker priorities. See man
dlopen
NVSHMEM_BOOTSTRAP_UID_PLUGIN "nvshmem_bootstrap_uid.so.3" (type: string, default: "nvshmem_bootstrap_uid.so.3")
Absolute path to or name of the UID bootstrap plugin file.
NVSHMEM will search for the plugin based on linux linker priorities. See man
dlopen
Additional options:
NVSHMEM_CUDA_PATH "" (type: string, default: "")
Path to directory containing libcuda.so (for use when not in default location)
NVSHMEM_DEBUG_ATTACH_DELAY 0 (type: int, default: 0)
Delay (in seconds) during the first call to NVSHMEM_INIT to allow for attaching
a debuggger (Default 0)
NVSHMEM_DEBUG_FILE "" (type: string, default: "")
Debugging output filename, may contain %h for hostname and %p for pid
NVSHMEM_MAX_TEAMS 32 (type: long, default: 32)
Maximum number of simultaneous teams allowed
NVSHMEM_MAX_MEMORY_PER_GPU 137438953472 (type: size, default: 137438953472)
Maximum memory per GPU
NVSHMEM_DISABLE_CUDA_VMM false (type: bool, default: false)
Disable use of CUDA VMM for P2P memory mapping. By default, CUDA VMM is enabled
on x86 and disabled on P9. CUDA VMM feature in NVSHMEM requires CUDA RT version
and CUDA Driver version to be greater than or equal to 11.3.
NVSHMEM_DISABLE_P2P false (type: bool, default: false)
Disable P2P connectivity of GPUs even when available
NVSHMEM_IGNORE_CUDA_MPS_ACTIVE_THREAD_PERCENTAGE false (type: bool, default: false)
When doing Multi-Process Per GPU (MPG) run, full API support is available only
if sum of CUDA_MPS_ACTIVE_THREAD_PERCENTAGE of processes running on a GPU is <=
100%. Through this variable, user can request NVSHMEM runtime to ignore the
active thread percentage and allow full MPG support. Users enable it at their
own risk as NVSHMEM might deadlock.
NVSHMEM_CUMEM_GRANULARITY 536870912 (type: size, default: 536870912)
Granularity for cuMemAlloc/cuMemCreate
NVSHMEM_PROXY_REQUEST_BATCH_MAX 32 (type: int, default: 32)
Maxmum number of requests that the proxy thread processes in a single iteration
of the progress loop.
Collectives options:
NVSHMEM_DISABLE_NCCL false (type: bool, default: false)
Disable use of NCCL for collective operations
NVSHMEM_BARRIER_DISSEM_KVAL 2 (type: int, default: 2)
Radix of the dissemination algorithm used for barriers
NVSHMEM_BARRIER_TG_DISSEM_KVAL 2 (type: int, default: 2)
Radix of the dissemination algorithm used for thread group barriers
NVSHMEM_FCOLLECT_LL_THRESHOLD 2048 (type: size, default: 2048)
Message size threshold up to which fcollect LL algo will be used
NVSHMEM_REDUCE_SCRATCH_SIZE 524288 (type: size, default: 524288)
Amount of symmetric heap memory (minimum 16B, multiple of 8B) reserved by
runtime for every team to implement reduce and reducescatter collectives
NVSHMEM_BCAST_ALGO 0 (type: int, default: 0)
Broadcast algorithm to be used.
* 0 - use default algorithm selection strategy
NVSHMEM_REDMAXLOC_ALGO 1 (type: int, default: 1)
Reduction algorithm to be used for MAXLOC operation.
* 1 - default, flag alltoall algorithm
* 2 - flat reduce + flat bcast
* 3 - topo-aware two-level reduce + topo-aware bcast
Transport options:
NVSHMEM_REMOTE_TRANSPORT "ibrc" (type: string, default: "ibrc")
Selected transport for remote operations: ibrc, ucx, libfabric, ibdevx, none
NVSHMEM_ENABLE_NIC_PE_MAPPING true (type: bool, default: false)
When not set or set to 0, a PE is assigned the NIC on the node that is closest
to it by distance. When set to 1, NVSHMEM either assigns NICs to PEs on a
round-robin basis or uses NVSHMEM_HCA_PE_MAPPING or NVSHMEM_HCA_LIST when they
are specified.
NVSHMEM_DISABLE_LOCAL_ONLY_PROXY false (type: bool, default: false)
When running on an NVLink-only configuaration (No-IB, No-UCX), completely
disable the proxy thread. This will disable device side global exit and device
side wait timeout polling (enabled by NVSHMEM_TIMEOUT_DEVICE_POLLING build-time
variable) because these are processed by the proxy thread.
NVSHMEM_IB_ENABLE_IBGDA false (type: bool, default: false)
Set to enable GPU-initiated communication transport.
Hidden options:
NVSHMEM_INFO_HIDDEN true (type: bool, default: false)
Print hidden environment variable options at startup
NVSHMEM_DISABLE_NVLS_SHARING true (type: bool, default: true)
Disable NVLS SHARP resource sharing for user-defined teams
NVSHMEM_HEAP_KIND "DEVICE" (type: string, default: "DEVICE")
Specify the memory kind used by the NVSHMEM symmetric heap.
Allowed values: VIDMEM, SYSMEM
NVSHMEM_ENABLE_RAIL_OPT false (type: bool, default: false)
Enable Rail Optimization when heap is in SYSMEM
NVSHMEM_BOOTSTRAP_TWO_STAGE false (type: bool, default: false)
Ignore CUDA device setting during initialization,forcing two-stage
initialization
NVSHMEM_DEBUG_SUBSYS "" (type: string, default: "")
Comma separated list of debugging message sources. Prefix with '^' to exclude.
Values: INIT, COLL, P2P, PROXY, TRANSPORT, MEM, BOOTSTRAP, TOPO, UTIL, ALL
NVSHMEM_ENABLE_ERROR_CHECKS false (type: bool, default: false)
Enable error checks
NVSHMEM_DISABLE_MNNVL false (type: bool, default: false)
Disable MNNVL connectivity for GPUs even when available
NVSHMEM_CUMEM_HANDLE_TYPE "FILE_DESCRIPTOR" (type: string, default: "FILE_DESCRIPTOR")
Handle type for cuMemCreate. Supported are - FABRIC or FILE_DESCRIPTOR
NVSHMEM_BYPASS_ACCESSIBILITY_CHECK false (type: bool, default: false)
Bypass peer GPU accessbility checks
NVSHMEM_FCOLLECT_NTHREADS 512 (type: int, default: 512)
Sets number of threads per block for fcollect collective.
By default, if no env is set, default value is min(max_occupancy per CTA, msg
size per PE).
If env is specified, value overrides the default irrespective of max occupancy
per CTA
NVSHMEM_REDUCESCATTER_NTHREADS 512 (type: int, default: 512)
Sets number of threads per block for reducescatter collective.
By default, if no env is set, default value is min(max_occupancy per CTA, msg
size per PE).
If env is specified, value overrides the default irrespective of max occupancy
per CTA
NVSHMEM_MAX_CTAS 1 (type: int, default: 1)
Sets number of blocks per grid for host onstream collective.
By default, if no env is set, default value to 1 CTA
If env is specified, value overrides the default value
NVSHMEM_REDUCE_RECEXCH_KVAL 2 (type: int, default: 2)
Radix of the recursive exchange reduction algorithm
NVSHMEM_FCOLLECT_LL128_THRESHOLD 0 (type: size, default: 0)
Message size threshold up to which the fcollect LL128 algo will be used.
LL128 will be used only when FCOLLECT_LL_THRESHOLD < size
NVSHMEM_FCOLLECT_NVLS_THRESHOLD 16777216 (type: size, default: 16777216)
Message size threshold up to which fcollect NVLS algo will be used
NVSHMEM_REDUCESCATTER_NVLS_THRESHOLD 16777216 (type: size, default: 16777216)
Message size threshold up to which reducescatter NVLS algo will be used
NVSHMEM_BCAST_TREE_KVAL 2 (type: int, default: 2)
Radix of the broadcast tree algorithm
NVSHMEM_FCOLLECT_ALGO 0 (type: int, default: 0)
Fcollect algorithm to be used.
* 0 - use default algorithm selection strategy
NVSHMEM_REDUCE_ALGO 0 (type: int, default: 0)
Allreduce algorithm to be used.
* 0/1 - use default algorithm selection strategy
NVSHMEM_REDUCE_NVLS_THRESHOLD 2048 (type: int, default: 2048)
Message size threshold up to which allreduce one-shot algo will be used
NVSHMEM_REDUCESCATTER_ALGO 0 (type: int, default: 0)
Reduce Scatter algorithm to be used.
* 0 - use default algorithm selection strategy
NVSHMEM_ASSERT_ATOMICS_SYNC false (type: bool, default: false)
Bypass flush on wait_until at target
NVSHMEM_BYPASS_FLUSH false (type: bool, default: false)
Bypass flush in proxy when enforcing consistency
NVTX options:
NVSHMEM_NVTX "off" (type: string, default: "off")
Set to enable NVTX instrumentation. Accepts a comma separated list of
instrumentation groups. By default the NVTX instrumentation is disabled.
init : library setup
alloc : memory management
launch : kernel launch routines
coll : collective communications
wait : blocking point-to-point synchronization
wait_on_stream : point-to-point synchronization (on stream)
test : non-blocking point-to-point synchronization
memorder : memory ordering (quiet, fence)
quiet_on_stream : nvshmemx_quiet_on_stream
atomic_fetch : fetching atomic memory operations
atomic_set : non-fetchong atomic memory operations
rma_blocking : blocking remote memory access operations
rma_nonblocking : non-blocking remote memory access operations
proxy : activity of the proxy thread
common : init,alloc,launch,coll,memorder,wait,atomic_fetch,rma_blocking,proxy
all : all groups
off : disable all NVTX instrumentation
Metadata
Metadata
Assignees
Labels
No labels