deepseek-ai/DeepEP
1
0
Fork 0
mirror of https://github.com/deepseek-ai/DeepEP.git synced 2026-11-05 00:47:32 +00:00
Code Issues 308 Projects Packages Wiki Activity

[GH-ISSUE #192] deepep compilation error #162

New issue
Closed
opened 2026-10-26 09:55:24 +00:00 by bsakaguchi · 10 comments
bsakaguchi commented 2026-10-26 09:55:24 +00:00
Owner
Copy link

Originally created by @zhangying098 on GitHub (Jun 7, 2025).
Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/192

Could you please advise how I might solve this error that occurred while installing the latest deepep?

root@cce-9ssovpy9-hvdenhfd:/sgl-workspace/DeepEP# NVSHMEM_DIR=/sgl-workspace/nvshmem/install/ python3 setup.py build
NVSHMEM directory: /sgl-workspace/nvshmem/install/
running build
running build_py
copying deep_ep/buffer.py -> build/lib.linux-x86_64-cpython-310/deep_ep
copying deep_ep/utils.py -> build/lib.linux-x86_64-cpython-310/deep_ep
copying deep_ep/__init__.py -> build/lib.linux-x86_64-cpython-310/deep_ep
running build_ext
/usr/local/lib/python3.10/dist-packages/torch/utils/cpp_extension.py:458: UserWarning: There are no x86_64-linux-gnu-g++ version bounds defined for CUDA version 12.4
  warnings.warn(f'There are no {compiler_name} version bounds defined for CUDA version {cuda_str_version}')
building 'deep_ep_cpp' extension
Emitting ninja build file /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/build.ninja...
Compiling objects...
Allowing ninja to set a default number of workers... (overridable by setting the environment variable MAX_JOBS=N)
[1/6] /usr/local/cuda/bin/nvcc --generate-dependencies-with-compile --dependency-output /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/kernels/runtime.o.d -Icsrc/ -I/sgl-workspace/nvshmem/install//include -I/usr/local/lib/python3.10/dist-packages/torch/include -I/usr/local/lib/python3.10/dist-packages/torch/include/torch/csrc/api/include -I/usr/local/lib/python3.10/dist-packages/torch/include/TH -I/usr/local/lib/python3.10/dist-packages/torch/include/THC -I/usr/local/cuda/include -I/usr/include/python3.10 -c -c /sgl-workspace/DeepEP/csrc/kernels/runtime.cu -o /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/kernels/runtime.o -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr --compiler-options ''"'"'-fPIC'"'"'' -O3 -Xcompiler -O3 -rdc=true --ptxas-options=--register-usage-level=10 -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=deep_ep_cpp -D_GLIBCXX_USE_CXX11_ABI=0 -gencode=arch=compute_90,code=sm_90 -std=c++17
[2/6] /usr/local/cuda/bin/nvcc --generate-dependencies-with-compile --dependency-output /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/kernels/intranode.o.d -Icsrc/ -I/sgl-workspace/nvshmem/install//include -I/usr/local/lib/python3.10/dist-packages/torch/include -I/usr/local/lib/python3.10/dist-packages/torch/include/torch/csrc/api/include -I/usr/local/lib/python3.10/dist-packages/torch/include/TH -I/usr/local/lib/python3.10/dist-packages/torch/include/THC -I/usr/local/cuda/include -I/usr/include/python3.10 -c -c /sgl-workspace/DeepEP/csrc/kernels/intranode.cu -o /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/kernels/intranode.o -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr --compiler-options ''"'"'-fPIC'"'"'' -O3 -Xcompiler -O3 -rdc=true --ptxas-options=--register-usage-level=10 -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=deep_ep_cpp -D_GLIBCXX_USE_CXX11_ABI=0 -gencode=arch=compute_90,code=sm_90 -std=c++17
FAILED: /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/kernels/intranode.o 
/usr/local/cuda/bin/nvcc --generate-dependencies-with-compile --dependency-output /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/kernels/intranode.o.d -Icsrc/ -I/sgl-workspace/nvshmem/install//include -I/usr/local/lib/python3.10/dist-packages/torch/include -I/usr/local/lib/python3.10/dist-packages/torch/include/torch/csrc/api/include -I/usr/local/lib/python3.10/dist-packages/torch/include/TH -I/usr/local/lib/python3.10/dist-packages/torch/include/THC -I/usr/local/cuda/include -I/usr/include/python3.10 -c -c /sgl-workspace/DeepEP/csrc/kernels/intranode.cu -o /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/kernels/intranode.o -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr --compiler-options ''"'"'-fPIC'"'"'' -O3 -Xcompiler -O3 -rdc=true --ptxas-options=--register-usage-level=10 -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=deep_ep_cpp -D_GLIBCXX_USE_CXX11_ABI=0 -gencode=arch=compute_90,code=sm_90 -std=c++17
ptxas /tmp/tmpxft_00007f75_00000000-6_intranode.ptx, line 4771; error   : State space incorrect for instruction 'cp.async.bulk'
ptxas /tmp/tmpxft_00007f75_00000000-6_intranode.ptx, line 4814; error   : State space incorrect for instruction 'cp.async.bulk'
ptxas /tmp/tmpxft_00007f75_00000000-6_intranode.ptx, line 6577; error   : State space incorrect for instruction 'cp.async.bulk'
ptxas /tmp/tmpxft_00007f75_00000000-6_intranode.ptx, line 6620; error   : State space incorrect for instruction 'cp.async.bulk'
ptxas /tmp/tmpxft_00007f75_00000000-6_intranode.ptx, line 8383; error   : State space incorrect for instruction 'cp.async.bulk'
ptxas /tmp/tmpxft_00007f75_00000000-6_intranode.ptx, line 8426; error   : State space incorrect for instruction 'cp.async.bulk'
ptxas fatal   : Ptx assembly aborted due to errors
[3/6] /usr/local/cuda/bin/nvcc --generate-dependencies-with-compile --dependency-output /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/kernels/internode_ll.o.d -Icsrc/ -I/sgl-workspace/nvshmem/install//include -I/usr/local/lib/python3.10/dist-packages/torch/include -I/usr/local/lib/python3.10/dist-packages/torch/include/torch/csrc/api/include -I/usr/local/lib/python3.10/dist-packages/torch/include/TH -I/usr/local/lib/python3.10/dist-packages/torch/include/THC -I/usr/local/cuda/include -I/usr/include/python3.10 -c -c /sgl-workspace/DeepEP/csrc/kernels/internode_ll.cu -o /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/kernels/internode_ll.o -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr --compiler-options ''"'"'-fPIC'"'"'' -O3 -Xcompiler -O3 -rdc=true --ptxas-options=--register-usage-level=10 -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=deep_ep_cpp -D_GLIBCXX_USE_CXX11_ABI=0 -gencode=arch=compute_90,code=sm_90 -std=c++17
[4/6] c++ -MMD -MF /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/deep_ep.o.d -Wno-unused-result -Wsign-compare -DNDEBUG -g -fwrapv -O2 -Wall -g -fstack-protector-strong -Wformat -Werror=format-security -g -fwrapv -O2 -fPIC -Icsrc/ -I/sgl-workspace/nvshmem/install//include -I/usr/local/lib/python3.10/dist-packages/torch/include -I/usr/local/lib/python3.10/dist-packages/torch/include/torch/csrc/api/include -I/usr/local/lib/python3.10/dist-packages/torch/include/TH -I/usr/local/lib/python3.10/dist-packages/torch/include/THC -I/usr/local/cuda/include -I/usr/include/python3.10 -c -c /sgl-workspace/DeepEP/csrc/deep_ep.cpp -o /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/deep_ep.o -O3 -Wno-deprecated-declarations -Wno-unused-variable -Wno-sign-compare -Wno-reorder -Wno-attributes -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=deep_ep_cpp -D_GLIBCXX_USE_CXX11_ABI=0 -std=c++17
[5/6] /usr/local/cuda/bin/nvcc --generate-dependencies-with-compile --dependency-output /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/kernels/internode.o.d -Icsrc/ -I/sgl-workspace/nvshmem/install//include -I/usr/local/lib/python3.10/dist-packages/torch/include -I/usr/local/lib/python3.10/dist-packages/torch/include/torch/csrc/api/include -I/usr/local/lib/python3.10/dist-packages/torch/include/TH -I/usr/local/lib/python3.10/dist-packages/torch/include/THC -I/usr/local/cuda/include -I/usr/include/python3.10 -c -c /sgl-workspace/DeepEP/csrc/kernels/internode.cu -o /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/kernels/internode.o -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr --compiler-options ''"'"'-fPIC'"'"'' -O3 -Xcompiler -O3 -rdc=true --ptxas-options=--register-usage-level=10 -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=deep_ep_cpp -D_GLIBCXX_USE_CXX11_ABI=0 -gencode=arch=compute_90,code=sm_90 -std=c++17
ninja: build stopped: subcommand failed.
Traceback (most recent call last):
  File "/usr/local/lib/python3.10/dist-packages/torch/utils/cpp_extension.py", line 2209, in _run_ninja_build
    subprocess.run(
  File "/usr/lib/python3.10/subprocess.py", line 526, in run
    raise CalledProcessError(retcode, process.args,
subprocess.CalledProcessError: Command '['ninja', '-v']' returned non-zero exit status 1.

The above exception was the direct cause of the following exception:

Traceback (most recent call last):
  File "/sgl-workspace/DeepEP/setup.py", line 45, in <module>
    setuptools.setup(
  File "/usr/local/lib/python3.10/dist-packages/setuptools/__init__.py", line 115, in setup
    return distutils.core.setup(**attrs)
  File "/usr/local/lib/python3.10/dist-packages/setuptools/_distutils/core.py", line 186, in setup
    return run_commands(dist)
  File "/usr/local/lib/python3.10/dist-packages/setuptools/_distutils/core.py", line 202, in run_commands
    dist.run_commands()
  File "/usr/local/lib/python3.10/dist-packages/setuptools/_distutils/dist.py", line 1002, in run_commands
    self.run_command(cmd)
  File "/usr/local/lib/python3.10/dist-packages/setuptools/dist.py", line 1102, in run_command
    super().run_command(command)
  File "/usr/local/lib/python3.10/dist-packages/setuptools/_distutils/dist.py", line 1021, in run_command
    cmd_obj.run()
  File "/usr/local/lib/python3.10/dist-packages/setuptools/_distutils/command/build.py", line 135, in run
    self.run_command(cmd_name)
  File "/usr/local/lib/python3.10/dist-packages/setuptools/_distutils/cmd.py", line 357, in run_command
    self.distribution.run_command(command)
  File "/usr/local/lib/python3.10/dist-packages/setuptools/dist.py", line 1102, in run_command
    super().run_command(command)
  File "/usr/local/lib/python3.10/dist-packages/setuptools/_distutils/dist.py", line 1021, in run_command
    cmd_obj.run()
  File "/usr/local/lib/python3.10/dist-packages/setuptools/command/build_ext.py", line 96, in run
    _build_ext.run(self)
  File "/usr/local/lib/python3.10/dist-packages/setuptools/_distutils/command/build_ext.py", line 368, in run
    self.build_extensions()
  File "/usr/local/lib/python3.10/dist-packages/torch/utils/cpp_extension.py", line 900, in build_extensions
    build_ext.build_extensions(self)
  File "/usr/local/lib/python3.10/dist-packages/setuptools/_distutils/command/build_ext.py", line 484, in build_extensions
    self._build_extensions_serial()
  File "/usr/local/lib/python3.10/dist-packages/setuptools/_distutils/command/build_ext.py", line 510, in _build_extensions_serial
    self.build_extension(ext)
  File "/usr/local/lib/python3.10/dist-packages/setuptools/command/build_ext.py", line 261, in build_extension
    _build_ext.build_extension(self, ext)
  File "/usr/local/lib/python3.10/dist-packages/setuptools/_distutils/command/build_ext.py", line 565, in build_extension
    objects = self.compiler.compile(
  File "/usr/local/lib/python3.10/dist-packages/torch/utils/cpp_extension.py", line 713, in unix_wrap_ninja_compile
    _write_ninja_file_and_compile_objects(
  File "/usr/local/lib/python3.10/dist-packages/torch/utils/cpp_extension.py", line 1869, in _write_ninja_file_and_compile_objects
    _run_ninja_build(
  File "/usr/local/lib/python3.10/dist-packages/torch/utils/cpp_extension.py", line 2225, in _run_ninja_build
    raise RuntimeError(message) from e
RuntimeError: Error compiling objects for extension
Originally created by @zhangying098 on GitHub (Jun 7, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/192 **Could you please advise how I might solve this error that occurred while installing the latest deepep?** ``` root@cce-9ssovpy9-hvdenhfd:/sgl-workspace/DeepEP# NVSHMEM_DIR=/sgl-workspace/nvshmem/install/ python3 setup.py build NVSHMEM directory: /sgl-workspace/nvshmem/install/ running build running build_py copying deep_ep/buffer.py -> build/lib.linux-x86_64-cpython-310/deep_ep copying deep_ep/utils.py -> build/lib.linux-x86_64-cpython-310/deep_ep copying deep_ep/__init__.py -> build/lib.linux-x86_64-cpython-310/deep_ep running build_ext /usr/local/lib/python3.10/dist-packages/torch/utils/cpp_extension.py:458: UserWarning: There are no x86_64-linux-gnu-g++ version bounds defined for CUDA version 12.4 warnings.warn(f'There are no {compiler_name} version bounds defined for CUDA version {cuda_str_version}') building 'deep_ep_cpp' extension Emitting ninja build file /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/build.ninja... Compiling objects... Allowing ninja to set a default number of workers... (overridable by setting the environment variable MAX_JOBS=N) [1/6] /usr/local/cuda/bin/nvcc --generate-dependencies-with-compile --dependency-output /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/kernels/runtime.o.d -Icsrc/ -I/sgl-workspace/nvshmem/install//include -I/usr/local/lib/python3.10/dist-packages/torch/include -I/usr/local/lib/python3.10/dist-packages/torch/include/torch/csrc/api/include -I/usr/local/lib/python3.10/dist-packages/torch/include/TH -I/usr/local/lib/python3.10/dist-packages/torch/include/THC -I/usr/local/cuda/include -I/usr/include/python3.10 -c -c /sgl-workspace/DeepEP/csrc/kernels/runtime.cu -o /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/kernels/runtime.o -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr --compiler-options ''"'"'-fPIC'"'"'' -O3 -Xcompiler -O3 -rdc=true --ptxas-options=--register-usage-level=10 -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=deep_ep_cpp -D_GLIBCXX_USE_CXX11_ABI=0 -gencode=arch=compute_90,code=sm_90 -std=c++17 [2/6] /usr/local/cuda/bin/nvcc --generate-dependencies-with-compile --dependency-output /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/kernels/intranode.o.d -Icsrc/ -I/sgl-workspace/nvshmem/install//include -I/usr/local/lib/python3.10/dist-packages/torch/include -I/usr/local/lib/python3.10/dist-packages/torch/include/torch/csrc/api/include -I/usr/local/lib/python3.10/dist-packages/torch/include/TH -I/usr/local/lib/python3.10/dist-packages/torch/include/THC -I/usr/local/cuda/include -I/usr/include/python3.10 -c -c /sgl-workspace/DeepEP/csrc/kernels/intranode.cu -o /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/kernels/intranode.o -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr --compiler-options ''"'"'-fPIC'"'"'' -O3 -Xcompiler -O3 -rdc=true --ptxas-options=--register-usage-level=10 -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=deep_ep_cpp -D_GLIBCXX_USE_CXX11_ABI=0 -gencode=arch=compute_90,code=sm_90 -std=c++17 FAILED: /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/kernels/intranode.o /usr/local/cuda/bin/nvcc --generate-dependencies-with-compile --dependency-output /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/kernels/intranode.o.d -Icsrc/ -I/sgl-workspace/nvshmem/install//include -I/usr/local/lib/python3.10/dist-packages/torch/include -I/usr/local/lib/python3.10/dist-packages/torch/include/torch/csrc/api/include -I/usr/local/lib/python3.10/dist-packages/torch/include/TH -I/usr/local/lib/python3.10/dist-packages/torch/include/THC -I/usr/local/cuda/include -I/usr/include/python3.10 -c -c /sgl-workspace/DeepEP/csrc/kernels/intranode.cu -o /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/kernels/intranode.o -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr --compiler-options ''"'"'-fPIC'"'"'' -O3 -Xcompiler -O3 -rdc=true --ptxas-options=--register-usage-level=10 -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=deep_ep_cpp -D_GLIBCXX_USE_CXX11_ABI=0 -gencode=arch=compute_90,code=sm_90 -std=c++17 ptxas /tmp/tmpxft_00007f75_00000000-6_intranode.ptx, line 4771; error : State space incorrect for instruction 'cp.async.bulk' ptxas /tmp/tmpxft_00007f75_00000000-6_intranode.ptx, line 4814; error : State space incorrect for instruction 'cp.async.bulk' ptxas /tmp/tmpxft_00007f75_00000000-6_intranode.ptx, line 6577; error : State space incorrect for instruction 'cp.async.bulk' ptxas /tmp/tmpxft_00007f75_00000000-6_intranode.ptx, line 6620; error : State space incorrect for instruction 'cp.async.bulk' ptxas /tmp/tmpxft_00007f75_00000000-6_intranode.ptx, line 8383; error : State space incorrect for instruction 'cp.async.bulk' ptxas /tmp/tmpxft_00007f75_00000000-6_intranode.ptx, line 8426; error : State space incorrect for instruction 'cp.async.bulk' ptxas fatal : Ptx assembly aborted due to errors [3/6] /usr/local/cuda/bin/nvcc --generate-dependencies-with-compile --dependency-output /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/kernels/internode_ll.o.d -Icsrc/ -I/sgl-workspace/nvshmem/install//include -I/usr/local/lib/python3.10/dist-packages/torch/include -I/usr/local/lib/python3.10/dist-packages/torch/include/torch/csrc/api/include -I/usr/local/lib/python3.10/dist-packages/torch/include/TH -I/usr/local/lib/python3.10/dist-packages/torch/include/THC -I/usr/local/cuda/include -I/usr/include/python3.10 -c -c /sgl-workspace/DeepEP/csrc/kernels/internode_ll.cu -o /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/kernels/internode_ll.o -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr --compiler-options ''"'"'-fPIC'"'"'' -O3 -Xcompiler -O3 -rdc=true --ptxas-options=--register-usage-level=10 -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=deep_ep_cpp -D_GLIBCXX_USE_CXX11_ABI=0 -gencode=arch=compute_90,code=sm_90 -std=c++17 [4/6] c++ -MMD -MF /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/deep_ep.o.d -Wno-unused-result -Wsign-compare -DNDEBUG -g -fwrapv -O2 -Wall -g -fstack-protector-strong -Wformat -Werror=format-security -g -fwrapv -O2 -fPIC -Icsrc/ -I/sgl-workspace/nvshmem/install//include -I/usr/local/lib/python3.10/dist-packages/torch/include -I/usr/local/lib/python3.10/dist-packages/torch/include/torch/csrc/api/include -I/usr/local/lib/python3.10/dist-packages/torch/include/TH -I/usr/local/lib/python3.10/dist-packages/torch/include/THC -I/usr/local/cuda/include -I/usr/include/python3.10 -c -c /sgl-workspace/DeepEP/csrc/deep_ep.cpp -o /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/deep_ep.o -O3 -Wno-deprecated-declarations -Wno-unused-variable -Wno-sign-compare -Wno-reorder -Wno-attributes -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=deep_ep_cpp -D_GLIBCXX_USE_CXX11_ABI=0 -std=c++17 [5/6] /usr/local/cuda/bin/nvcc --generate-dependencies-with-compile --dependency-output /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/kernels/internode.o.d -Icsrc/ -I/sgl-workspace/nvshmem/install//include -I/usr/local/lib/python3.10/dist-packages/torch/include -I/usr/local/lib/python3.10/dist-packages/torch/include/torch/csrc/api/include -I/usr/local/lib/python3.10/dist-packages/torch/include/TH -I/usr/local/lib/python3.10/dist-packages/torch/include/THC -I/usr/local/cuda/include -I/usr/include/python3.10 -c -c /sgl-workspace/DeepEP/csrc/kernels/internode.cu -o /sgl-workspace/DeepEP/build/temp.linux-x86_64-cpython-310/csrc/kernels/internode.o -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr --compiler-options ''"'"'-fPIC'"'"'' -O3 -Xcompiler -O3 -rdc=true --ptxas-options=--register-usage-level=10 -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=deep_ep_cpp -D_GLIBCXX_USE_CXX11_ABI=0 -gencode=arch=compute_90,code=sm_90 -std=c++17 ninja: build stopped: subcommand failed. Traceback (most recent call last): File "/usr/local/lib/python3.10/dist-packages/torch/utils/cpp_extension.py", line 2209, in _run_ninja_build subprocess.run( File "/usr/lib/python3.10/subprocess.py", line 526, in run raise CalledProcessError(retcode, process.args, subprocess.CalledProcessError: Command '['ninja', '-v']' returned non-zero exit status 1. The above exception was the direct cause of the following exception: Traceback (most recent call last): File "/sgl-workspace/DeepEP/setup.py", line 45, in <module> setuptools.setup( File "/usr/local/lib/python3.10/dist-packages/setuptools/__init__.py", line 115, in setup return distutils.core.setup(**attrs) File "/usr/local/lib/python3.10/dist-packages/setuptools/_distutils/core.py", line 186, in setup return run_commands(dist) File "/usr/local/lib/python3.10/dist-packages/setuptools/_distutils/core.py", line 202, in run_commands dist.run_commands() File "/usr/local/lib/python3.10/dist-packages/setuptools/_distutils/dist.py", line 1002, in run_commands self.run_command(cmd) File "/usr/local/lib/python3.10/dist-packages/setuptools/dist.py", line 1102, in run_command super().run_command(command) File "/usr/local/lib/python3.10/dist-packages/setuptools/_distutils/dist.py", line 1021, in run_command cmd_obj.run() File "/usr/local/lib/python3.10/dist-packages/setuptools/_distutils/command/build.py", line 135, in run self.run_command(cmd_name) File "/usr/local/lib/python3.10/dist-packages/setuptools/_distutils/cmd.py", line 357, in run_command self.distribution.run_command(command) File "/usr/local/lib/python3.10/dist-packages/setuptools/dist.py", line 1102, in run_command super().run_command(command) File "/usr/local/lib/python3.10/dist-packages/setuptools/_distutils/dist.py", line 1021, in run_command cmd_obj.run() File "/usr/local/lib/python3.10/dist-packages/setuptools/command/build_ext.py", line 96, in run _build_ext.run(self) File "/usr/local/lib/python3.10/dist-packages/setuptools/_distutils/command/build_ext.py", line 368, in run self.build_extensions() File "/usr/local/lib/python3.10/dist-packages/torch/utils/cpp_extension.py", line 900, in build_extensions build_ext.build_extensions(self) File "/usr/local/lib/python3.10/dist-packages/setuptools/_distutils/command/build_ext.py", line 484, in build_extensions self._build_extensions_serial() File "/usr/local/lib/python3.10/dist-packages/setuptools/_distutils/command/build_ext.py", line 510, in _build_extensions_serial self.build_extension(ext) File "/usr/local/lib/python3.10/dist-packages/setuptools/command/build_ext.py", line 261, in build_extension _build_ext.build_extension(self, ext) File "/usr/local/lib/python3.10/dist-packages/setuptools/_distutils/command/build_ext.py", line 565, in build_extension objects = self.compiler.compile( File "/usr/local/lib/python3.10/dist-packages/torch/utils/cpp_extension.py", line 713, in unix_wrap_ninja_compile _write_ninja_file_and_compile_objects( File "/usr/local/lib/python3.10/dist-packages/torch/utils/cpp_extension.py", line 1869, in _write_ninja_file_and_compile_objects _run_ninja_build( File "/usr/local/lib/python3.10/dist-packages/torch/utils/cpp_extension.py", line 2225, in _run_ninja_build raise RuntimeError(message) from e RuntimeError: Error compiling objects for extension ```
bsakaguchi commented 2026-10-26 09:55:26 +00:00
Author
Owner
Copy link

@zhangying098 commented on GitHub (Jun 7, 2025):

Through binary method, I determined that this problem was caused by this pr. https://github.com/deepseek-ai/DeepEP/pull/191

<!-- gh-comment-id:2952373082 --> @zhangying098 commented on GitHub (Jun 7, 2025): Through binary method, I determined that this problem was caused by this pr. [https://github.com/deepseek-ai/DeepEP/pull/191](https://github.com/deepseek-ai/DeepEP/pull/191)
bsakaguchi commented 2026-10-26 09:55:26 +00:00
Author
Owner
Copy link

@LyricZhao commented on GitHub (Jun 9, 2025):

Which CUDA version are you using?

<!-- gh-comment-id:2954422876 --> @LyricZhao commented on GitHub (Jun 9, 2025): Which CUDA version are you using?
bsakaguchi commented 2026-10-26 09:55:27 +00:00
Author
Owner
Copy link

@haok1402 commented on GitHub (Jun 9, 2025):

I also encountered the same issue, under CUDA Version: 12.4.

<!-- gh-comment-id:2954463866 --> @haok1402 commented on GitHub (Jun 9, 2025): I also encountered the same issue, under CUDA Version: 12.4.
bsakaguchi commented 2026-10-26 09:55:27 +00:00
Author
Owner
Copy link

@MARD1NO commented on GitHub (Jun 9, 2025):

Which CUDA version are you using?

I also encounter this error in cuda12.4, when I change to use cutlass's tma copy ptx, it works...:

__device__ __forceinline__ void tma_load_1d(const void* smem_ptr, const void* gmem_ptr, uint64_t* mbar_ptr, int num_bytes,
                                            bool evict_first = true) {
#if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900))
    auto mbar_int_ptr = static_cast<uint32_t>(__cvta_generic_to_shared(mbar_ptr));
    auto smem_int_ptr  = static_cast<uint32_t>(__cvta_generic_to_shared(smem_ptr));
    // const auto cache_hint = evict_first ? kEvictFirst : kEvictNormal;
    // asm volatile("cp.async.bulk.shared::cta.global.mbarrier::complete_tx::bytes.L2::cache_hint [%0], [%1], %2, [%3], %4;\n"
    //              :: "r"(smem_int_ptr), "l"(gmem_ptr), "r"(num_bytes), "r"(mbar_int_ptr), "l"(cache_hint) : "memory");

    asm volatile("cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];\n"
                     :
                     : "r"(smem_int_ptr), "l"(gmem_ptr), "r"(num_bytes), "r"(mbar_int_ptr)
                     : "memory");
#endif 
}
<!-- gh-comment-id:2954478700 --> @MARD1NO commented on GitHub (Jun 9, 2025): > Which CUDA version are you using? I also encounter this error in cuda12.4, when I change to use cutlass's tma copy ptx, it works...: ```cpp __device__ __forceinline__ void tma_load_1d(const void* smem_ptr, const void* gmem_ptr, uint64_t* mbar_ptr, int num_bytes, bool evict_first = true) { #if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)) auto mbar_int_ptr = static_cast<uint32_t>(__cvta_generic_to_shared(mbar_ptr)); auto smem_int_ptr = static_cast<uint32_t>(__cvta_generic_to_shared(smem_ptr)); // const auto cache_hint = evict_first ? kEvictFirst : kEvictNormal; // asm volatile("cp.async.bulk.shared::cta.global.mbarrier::complete_tx::bytes.L2::cache_hint [%0], [%1], %2, [%3], %4;\n" // :: "r"(smem_int_ptr), "l"(gmem_ptr), "r"(num_bytes), "r"(mbar_int_ptr), "l"(cache_hint) : "memory"); asm volatile("cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];\n" : : "r"(smem_int_ptr), "l"(gmem_ptr), "r"(num_bytes), "r"(mbar_int_ptr) : "memory"); #endif } ```
bsakaguchi commented 2026-10-26 09:55:27 +00:00
Author
Owner
Copy link

@LyricZhao commented on GitHub (Jun 9, 2025):

Fixed in #194: Support for .shared::cta as destination state space is introduced in PTX ISA version 8.6 (> some CUDA version). See https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk.

<!-- gh-comment-id:2954495367 --> @LyricZhao commented on GitHub (Jun 9, 2025): Fixed in #194: Support for .shared::cta as destination state space is introduced in PTX ISA version 8.6 (> some CUDA version). See https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk.
bsakaguchi commented 2026-10-26 09:55:28 +00:00
Author
Owner
Copy link

@yiakwy-xpu-ml-framework-team commented on GitHub (Jun 10, 2025):

Yes @MARD1NO the previous code has problem with wrong TMA instruction.

<!-- gh-comment-id:2957716861 --> @yiakwy-xpu-ml-framework-team commented on GitHub (Jun 10, 2025): Yes @MARD1NO the previous code has problem with wrong TMA instruction.
bsakaguchi commented 2026-10-26 09:55:28 +00:00
Author
Owner
Copy link

@fzyzcjy commented on GitHub (Jun 14, 2025):

+1 I also meet this bug today

<!-- gh-comment-id:2972661646 --> @fzyzcjy commented on GitHub (Jun 14, 2025): +1 I also meet this bug today
bsakaguchi commented 2026-10-26 09:55:28 +00:00
Author
Owner
Copy link

@ayrnb commented on GitHub (Jun 16, 2025):

+1

<!-- gh-comment-id:2975765570 --> @ayrnb commented on GitHub (Jun 16, 2025): +1
bsakaguchi commented 2026-10-26 09:55:28 +00:00
Author
Owner
Copy link

@LyricZhao commented on GitHub (Jun 16, 2025):

@ayrnb @fzyzcjy Does #194 fix this issue?

<!-- gh-comment-id:2975832117 --> @LyricZhao commented on GitHub (Jun 16, 2025): @ayrnb @fzyzcjy Does #194 fix this issue?
bsakaguchi commented 2026-10-26 09:55:28 +00:00
Author
Owner
Copy link

@fzyzcjy commented on GitHub (Jun 17, 2025):

Yes I can compile now

<!-- gh-comment-id:2978737465 --> @fzyzcjy commented on GitHub (Jun 17, 2025): Yes I can compile now
Sign in to join this conversation.
No Branch/Tag specified
Branches Tags
main
hybrid-ep
chore/remove-clang-format-yapf
v2.5
nvDev
try-fix-cu128
mori-ep
antgroup-opt
compat-old-torch
add-linter
eager-rdma
v1.2.1
Labels
Clear labels   
pull-request
Mirrored from GitHub Pull Request

No labels
pull-request
Milestone
Clear milestone
No items
No milestone
Projects
Clear projects
No items
No project
Assignees
Clear assignees
No assignees
1 participant
Notifications
Due date
The due date is invalid or out of range. Please use the format "yyyy-mm-dd".

No due date set.

Dependencies

No dependencies set.

Reference
deepseek-ai/DeepEP#162
No description provided.