Ep api design - Build Infra dependencies

.dockerignore

@@ -0,0 +1,46 @@
+# Exclude bulky / host-specific paths from the docker build context.
+# Submodules are still copied (they're needed for BUILD_NVEP=1).
+
+# Python venvs / build trees
+.venv/
+venv/
+build/
+build_nvep/
+dist/
+*.egg-info/
+**/__pycache__/
+**/*.pyc
+
+# meson subproject caches inside submodules (will be regenerated in-container).
+# Note: we MUST keep the *.wrap files; Docker's .dockerignore matcher applies
+# the same pattern to files and dirs irrespective of a trailing slash, so we
+# spell out each materialized subproject dir name explicitly.
+3rdparty/nixl/subprojects/abseil-cpp-*/
+3rdparty/nixl/subprojects/asio-*/
+3rdparty/nixl/subprojects/prometheus-cpp/
+3rdparty/nixl/subprojects/taskflow-*/
+3rdparty/nixl/subprojects/tomlplusplus-*/
+3rdparty/nixl/subprojects/packagecache/
+3rdparty/nixl/subprojects/.wraplock
+
+# git internals (not needed; submodules already at correct commits via COPY)
+.git/
+3rdparty/*/.git/
+
+# CI + local-dev clutter
+.github/
+.devcontainer/
+.cursor/
+.claude/
+.pre-commit-config.yaml
+*.swp
+
+# already-built wheels, pip caches
+flashinfer-whl/
+.cache/
+.pytest_cache/
+.mypy_cache/
+
+# Docs build outputs
+docs/_build/
+docs/tutorials/generated/

.gitignore

@@ -201,3 +201,8 @@ cython_debug/
 .cursor/
 docs/tutorials/generated/
 docs/sg_execution_times.rst
+
+# moe_ep build artifacts (BUILD_NVEP=1 in-tree build of NIXL-EP + NCCL-EP)
+build_nvep/
+flashinfer/moe_ep/*/_libs/
+flashinfer/moe_ep/*/_vendored/

.gitmodules

@@ -7,3 +7,9 @@
 [submodule "3rdparty/cccl"]
 	path = 3rdparty/cccl
 	url = https://github.com/NVIDIA/cccl.git
+[submodule "3rdparty/nixl"]
+	path = 3rdparty/nixl
+	url = https://github.com/ai-dynamo/nixl.git
+[submodule "3rdparty/nccl"]
+	path = 3rdparty/nccl
+	url = https://github.com/NVIDIA/nccl.git

3rdparty_patches/nixl/0001-meson-add-blackwell-arches.patch

@@ -0,0 +1,65 @@
+From: FlashInfer build infra
+Subject: [PATCH] meson: emit nixl_ep SASS for Hopper + Blackwell
+
+The upstream meson rule pins `-arch=sm_90`, which overrides any global
+-gencode flags and ships an sm_90-only `nixl_ep_cpp.so`. FlashInfer needs
+a single .so that runs natively on H100 (sm_90), B200 (sm_100), and B300
+(sm_103), plus sm_90 PTX for forward-compat onto future arches.
+
+This patch:
+  1. Replaces the single `-arch=sm_90` flag in examples/device/ep with
+     an explicit multi-gencode list (compile-side).
+  2. Extends the project-wide nvcc_flags / nvcc_flags_link in the
+     top-level meson.build so the `-rdc=true` device link step also
+     covers sm_100 / sm_103. Without (2), nvlink discards the sm_100 /
+     sm_103 cubins generated by (1) during the final device link and
+     the resulting nixl_ep_cpp.so ends up sm_90-only.
+
+Targets the NIXL pin 05e4243f (tag v1.1.0).
+
+diff --git a/examples/device/ep/meson.build b/examples/device/ep/meson.build
+index a9ba19f..065dd13 100644
+--- a/examples/device/ep/meson.build
++++ b/examples/device/ep/meson.build
+@@ -92,7 +92,9 @@ nixl_ep_cuda_args = [
+     '-DHAVE_CUDA',
+     '-DTORCH_EXTENSION_NAME=nixl_ep_cpp',
+     '--expt-relaxed-constexpr',  # Allow calling constexpr __host__ functions from __device__ functions
+-    '-arch=sm_90',  # Only compile for sm90 (overrides global -gencode flags)
++    '-gencode=arch=compute_90,code=[sm_90,compute_90]',  # H100 SASS + PTX forward-compat
++    '-gencode=arch=compute_100,code=sm_100',  # B200 native SASS
++    '-gencode=arch=compute_103,code=sm_103',  # B300 native SASS
+     '--ptxas-options=--register-usage-level=10',  # Allow more register usage (matches setup.py)
+     '-Xcompiler', '-Wno-deprecated-declarations',
+     '-Xcompiler', '-Wno-unused-variable',
+diff --git a/meson.build b/meson.build
+index 2812ce8..b96da85 100644
+--- a/meson.build
++++ b/meson.build
+@@ -202,6 +202,15 @@ if cuda_dep.found()
+         nvcc_flags += ['-gencode', 'arch=compute_80,code=sm_80']
+     endif
+     nvcc_flags += ['-gencode', 'arch=compute_90,code=sm_90']
++    if get_option('build_nixl_ep')
++        # FlashInfer: nixl_ep targets H100 + B200 + B300. The compile-side
++        # gencode for those arches is set in examples/device/ep/meson.build
++        # via 0001-meson-add-blackwell-arches.patch; we also need them on
++        # the project-wide device-link step so nvlink doesn't drop sm_100 /
++        # sm_103 SASS during the final -rdc=true device link.
++        nvcc_flags += ['-gencode', 'arch=compute_100,code=sm_100']
++        nvcc_flags += ['-gencode', 'arch=compute_103,code=sm_103']
++    endif
+     add_project_arguments(nvcc_flags, language: 'cuda')
+
+     # Refer to https://mesonbuild.com/Cuda-module.html
+@@ -214,6 +223,10 @@ if cuda_dep.found()
+         nvcc_flags_link += ['-gencode=arch=compute_80,code=sm_80']
+     endif
+     nvcc_flags_link += ['-gencode=arch=compute_90,code=sm_90']
++    if get_option('build_nixl_ep')
++        nvcc_flags_link += ['-gencode=arch=compute_100,code=sm_100']
++        nvcc_flags_link += ['-gencode=arch=compute_103,code=sm_103']
++    endif
+     add_project_link_arguments(nvcc_flags_link, language: 'cuda')
+     message('nvcc version: ' + nvcc.version())
+     if nvcc.version().version_compare('>=12.8')

3rdparty_patches/nixl/0002-ep-only-build.patch

@@ -0,0 +1,104 @@
+From: FlashInfer build infra
+Subject: [PATCH] meson: add nixl_ep_only option to skip parent libnixl build
+
+When `nixl_ep_only=true`, the parent NIXL library (and tests, headers,
+non-EP examples) is not built. The EP example links against an
+externally-provided libnixl.so (typically the nixl-cu13 pip wheel),
+mirroring the wheel-driven build path used by FlashInfer's NCCL-EP.
+
+Two new meson options:
+  - nixl_ep_only (bool, default false): skip subdir('src')/subdir('test')
+    and install_headers; build only examples/device/ep.
+  - nixl_wheel_lib_dir (string, default ''): when nixl_ep_only=true, the
+    directory containing libnixl.so (used via cc.find_library()).
+
+Targets the NIXL pin 05e4243f (tag v1.1.0).
+
+diff --git a/examples/device/ep/meson.build b/examples/device/ep/meson.build
+index a9ba19f..a02d583 100644
+--- a/examples/device/ep/meson.build
++++ b/examples/device/ep/meson.build
+@@ -52,8 +52,21 @@ if not pybind_dep.found()
+     subdir_done()
+ endif
+
+-nixl_dep = declare_dependency(link_with: nixl_lib, include_directories: nixl_inc_dirs)
+-nixl_lib_dir = join_paths(meson.project_build_root(), 'src', 'core')
++if get_option('nixl_ep_only')
++    cc_ep = meson.get_compiler('cpp')
++    wheel_lib_dir = get_option('nixl_wheel_lib_dir')
++    if wheel_lib_dir == ''
++        error('nixl_ep_only=true requires -Dnixl_wheel_lib_dir=<path-containing-libnixl.so>')
++    endif
++    libnixl_external = cc_ep.find_library('nixl', dirs: [wheel_lib_dir], required: true)
++    nixl_dep = declare_dependency(dependencies: libnixl_external, include_directories: nixl_inc_dirs)
++    nixl_lib_dir = wheel_lib_dir
++    nixl_ep_link_with = []
++else
++    nixl_dep = declare_dependency(link_with: nixl_lib, include_directories: nixl_inc_dirs)
++    nixl_lib_dir = join_paths(meson.project_build_root(), 'src', 'core')
++    nixl_ep_link_with = [nixl_lib]
++endif
+
+ ucx_build_deps = []
+ if ucx_dep.found()
+@@ -147,7 +160,7 @@ nixl_ep_ext = py.extension_module('nixl_ep_cpp',
+     include_directories: nixl_ep_inc_dirs,
+     cpp_args: nixl_ep_cpp_args,
+     cuda_args: nixl_ep_cuda_args,
+-    link_with: [nixl_lib],
++    link_with: nixl_ep_link_with,
+     build_rpath: nixl_ep_rpath,
+     install_rpath: nixl_ep_install_rpath,
+     override_options: nixl_ep_override_options,
+diff --git a/meson.build b/meson.build
+index 2812ce8..46519e2 100644
+--- a/meson.build
++++ b/meson.build
+@@ -366,18 +366,27 @@ nixl_gpu_inc_dirs = include_directories('src/api/gpu/ucx')
+ plugins_inc_dirs = include_directories('src/plugins')
+ utils_inc_dirs = include_directories('src/utils')
+
+-subdir('src')
+-if get_option('build_tests') and get_option('buildtype') != 'release'
+-  subdir('test')
++if not get_option('nixl_ep_only')
++    subdir('src')
++    if get_option('build_tests') and get_option('buildtype') != 'release'
++        subdir('test')
++    endif
+ endif
+
+ # nixl_ep currently lives under examples/device/ep. Build that subtree when
+ # either full examples are requested or nixl_ep is explicitly requested.
+-if get_option('build_examples') or get_option('build_nixl_ep')
++# In nixl_ep_only mode, skip the rest of examples/ (cpp/ etc. depend on
++# nixl_lib from subdir('src')) and route straight to examples/device/ep.
++if get_option('nixl_ep_only')
++    if not get_option('build_nixl_ep')
++        error('nixl_ep_only=true requires -Dbuild_nixl_ep=true')
++    endif
++    subdir('examples/device/ep')
++elif get_option('build_examples') or get_option('build_nixl_ep')
+     subdir('examples')
+ endif
+
+-if get_option('install_headers')
++if get_option('install_headers') and not get_option('nixl_ep_only')
+   install_headers('src/api/cpp/nixl.h', install_dir: prefix_inc)
+   install_headers('src/api/cpp/nixl_types.h', install_dir: prefix_inc)
+   install_headers('src/api/cpp/nixl_params.h', install_dir: prefix_inc)
+diff --git a/meson_options.txt b/meson_options.txt
+index 4db4845..5b3750a 100644
+--- a/meson_options.txt
++++ b/meson_options.txt
+@@ -37,3 +37,9 @@ option('build_tests', type: 'boolean', value: true, description: 'Build all test
+ option('build_examples', type: 'boolean', value: true, description: 'Build all examples')
+ option('build_nixl_ep', type: 'boolean', value: false, description: 'Build nixl_ep example (requires sm_90)')
+ option('test_all_plugins', type: 'boolean', value: false, description: 'Testing all plugins in addition to the mocks..')
++
++# FlashInfer wheel-driven build: skip the parent libnixl build entirely and
++# link the EP example against an externally-provided libnixl.so (typically
++# from the nixl-cu13 pip wheel).
++option('nixl_ep_only', type: 'boolean', value: false, description: 'Skip parent libnixl build; build only examples/device/ep linked against an external libnixl.so')
++option('nixl_wheel_lib_dir', type: 'string', value: '', description: 'Path to directory containing external libnixl.so (used when nixl_ep_only=true)')