基于AI Agent开发范式的国产GPU大模型推理算子库优化/Fused MoE Baseline入门：快速跑通最小闭环教程.md

@@ -129,7 +129,7 @@ EOF
 
 ### 步骤 2：进入项目目录
 
-**目标：**进入本模块所需的源码目录。
+**目标：**进入本模块所需的源码目录：https://www.gitlink.org.cn/metax-maca/op_optimization/tree/master/%E5%9F%BA%E4%BA%8EAI%20Agent%E5%BC%80%E5%8F%91%E8%8C%83%E5%BC%8F%E7%9A%84%E5%9B%BD%E4%BA%A7GPU%E5%A4%A7%E6%A8%A1%E5%9E%8B%E6%8E%A8%E7%90%86%E7%AE%97%E5%AD%90%E5%BA%93%E4%BC%98%E5%8C%96%2Fbaselines%2Ffused_moe
 
 **操作：**切换到指定项目路径。
 

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/.vscode/settings.json

@@ -0,0 +1,3 @@
+{
+    "cmake.sourceDirectory": "/root/Project/fusedmoe_v2/standalone/fused_moe_i8_tn"
+}

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/query/client-vscode/query.json

@@ -0,0 +1 @@
+{"requests":[{"kind":"cache","version":2},{"kind":"codemodel","version":2},{"kind":"toolchains","version":1},{"kind":"cmakeFiles","version":1}]}

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/cache-v2-ea2ef11d05674d96d761.json

@@ -0,0 +1,471 @@
+{
+	"entries" : 
+	[
+		{
+			"name" : "CMAKE_BUILD_TYPE",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "No help, variable specified on the command line."
+				}
+			],
+			"type" : "STRING",
+			"value" : "Debug"
+		},
+		{
+			"name" : "CMAKE_CACHEFILE_DIR",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "This is the directory where this CMakeCache.txt was created"
+				}
+			],
+			"type" : "INTERNAL",
+			"value" : "/root/Project/fusedmoe/build"
+		},
+		{
+			"name" : "CMAKE_CACHE_MAJOR_VERSION",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "Major version of cmake used to create the current loaded cache"
+				}
+			],
+			"type" : "INTERNAL",
+			"value" : "3"
+		},
+		{
+			"name" : "CMAKE_CACHE_MINOR_VERSION",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "Minor version of cmake used to create the current loaded cache"
+				}
+			],
+			"type" : "INTERNAL",
+			"value" : "28"
+		},
+		{
+			"name" : "CMAKE_CACHE_PATCH_VERSION",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "Patch version of cmake used to create the current loaded cache"
+				}
+			],
+			"type" : "INTERNAL",
+			"value" : "3"
+		},
+		{
+			"name" : "CMAKE_COMMAND",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "Path to CMake executable."
+				}
+			],
+			"type" : "INTERNAL",
+			"value" : "/usr/bin/cmake"
+		},
+		{
+			"name" : "CMAKE_CPACK_COMMAND",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "Path to cpack program executable."
+				}
+			],
+			"type" : "INTERNAL",
+			"value" : "/usr/bin/cpack"
+		},
+		{
+			"name" : "CMAKE_CTEST_COMMAND",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "Path to ctest program executable."
+				}
+			],
+			"type" : "INTERNAL",
+			"value" : "/usr/bin/ctest"
+		},
+		{
+			"name" : "CMAKE_CXX_COMPILER",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "No help, variable specified on the command line."
+				}
+			],
+			"type" : "FILEPATH",
+			"value" : "/usr/bin/g++"
+		},
+		{
+			"name" : "CMAKE_C_COMPILER",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "No help, variable specified on the command line."
+				}
+			],
+			"type" : "FILEPATH",
+			"value" : "/usr/bin/gcc"
+		},
+		{
+			"name" : "CMAKE_EXPORT_COMPILE_COMMANDS",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "No help, variable specified on the command line."
+				}
+			],
+			"type" : "BOOL",
+			"value" : "TRUE"
+		},
+		{
+			"name" : "CMAKE_EXTRA_GENERATOR",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "Name of external makefile project generator."
+				}
+			],
+			"type" : "INTERNAL",
+			"value" : ""
+		},
+		{
+			"name" : "CMAKE_FIND_PACKAGE_REDIRECTS_DIR",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "Value Computed by CMake."
+				}
+			],
+			"type" : "STATIC",
+			"value" : "/root/Project/fusedmoe/build/CMakeFiles/pkgRedirects"
+		},
+		{
+			"name" : "CMAKE_GENERATOR",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "Name of generator."
+				}
+			],
+			"type" : "INTERNAL",
+			"value" : "Ninja"
+		},
+		{
+			"name" : "CMAKE_GENERATOR_INSTANCE",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "Generator instance identifier."
+				}
+			],
+			"type" : "INTERNAL",
+			"value" : ""
+		},
+		{
+			"name" : "CMAKE_GENERATOR_PLATFORM",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "Name of generator platform."
+				}
+			],
+			"type" : "INTERNAL",
+			"value" : ""
+		},
+		{
+			"name" : "CMAKE_GENERATOR_TOOLSET",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "Name of generator toolset."
+				}
+			],
+			"type" : "INTERNAL",
+			"value" : ""
+		},
+		{
+			"name" : "CMAKE_HOME_DIRECTORY",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "Source directory with the top level CMakeLists.txt file for this project"
+				}
+			],
+			"type" : "INTERNAL",
+			"value" : "/root/Project/fusedmoe/standalone/fused_moe_i8_tn"
+		},
+		{
+			"name" : "CMAKE_INSTALL_PREFIX",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "Install path prefix, prepended onto install directories."
+				}
+			],
+			"type" : "PATH",
+			"value" : "/usr/local"
+		},
+		{
+			"name" : "CMAKE_INSTALL_SO_NO_EXE",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "Install .so files without execute permission."
+				}
+			],
+			"type" : "INTERNAL",
+			"value" : "1"
+		},
+		{
+			"name" : "CMAKE_MAKE_PROGRAM",
+			"properties" : 
+			[
+				{
+					"name" : "ADVANCED",
+					"value" : "1"
+				},
+				{
+					"name" : "HELPSTRING",
+					"value" : "Program used to build from build.ninja files."
+				}
+			],
+			"type" : "FILEPATH",
+			"value" : "/opt/conda/bin/ninja"
+		},
+		{
+			"name" : "CMAKE_NUMBER_OF_MAKEFILES",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "number of local generators"
+				}
+			],
+			"type" : "INTERNAL",
+			"value" : "1"
+		},
+		{
+			"name" : "CMAKE_PLATFORM_INFO_INITIALIZED",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "Platform information initialized"
+				}
+			],
+			"type" : "INTERNAL",
+			"value" : "1"
+		},
+		{
+			"name" : "CMAKE_PROJECT_DESCRIPTION",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "Value Computed by CMake"
+				}
+			],
+			"type" : "STATIC",
+			"value" : ""
+		},
+		{
+			"name" : "CMAKE_PROJECT_HOMEPAGE_URL",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "Value Computed by CMake"
+				}
+			],
+			"type" : "STATIC",
+			"value" : ""
+		},
+		{
+			"name" : "CMAKE_PROJECT_NAME",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "Value Computed by CMake"
+				}
+			],
+			"type" : "STATIC",
+			"value" : "fused_moe_i8_tn"
+		},
+		{
+			"name" : "CMAKE_ROOT",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "Path to CMake installation."
+				}
+			],
+			"type" : "INTERNAL",
+			"value" : "/usr/share/cmake-3.28"
+		},
+		{
+			"name" : "CMAKE_SKIP_INSTALL_RPATH",
+			"properties" : 
+			[
+				{
+					"name" : "ADVANCED",
+					"value" : "1"
+				},
+				{
+					"name" : "HELPSTRING",
+					"value" : "If set, runtime paths are not added when installing shared libraries, but are added when building."
+				}
+			],
+			"type" : "BOOL",
+			"value" : "NO"
+		},
+		{
+			"name" : "CMAKE_SKIP_RPATH",
+			"properties" : 
+			[
+				{
+					"name" : "ADVANCED",
+					"value" : "1"
+				},
+				{
+					"name" : "HELPSTRING",
+					"value" : "If set, runtime paths are not added when using shared libraries."
+				}
+			],
+			"type" : "BOOL",
+			"value" : "NO"
+		},
+		{
+			"name" : "CMAKE_UNAME",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "uname command"
+				}
+			],
+			"type" : "INTERNAL",
+			"value" : "/usr/bin/uname"
+		},
+		{
+			"name" : "CMAKE_VERBOSE_MAKEFILE",
+			"properties" : 
+			[
+				{
+					"name" : "ADVANCED",
+					"value" : "1"
+				},
+				{
+					"name" : "HELPSTRING",
+					"value" : "If this value is on, makefiles will be generated without the .SILENT directive, and all commands will be echoed to the console during the make.  This is useful for debugging only. With Visual Studio IDE projects all commands are done without /nologo."
+				}
+			],
+			"type" : "BOOL",
+			"value" : "FALSE"
+		},
+		{
+			"name" : "MACA_PATH",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "Path to MACA SDK"
+				}
+			],
+			"type" : "PATH",
+			"value" : "/opt/maca"
+		},
+		{
+			"name" : "MXCC",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "Path to a program."
+				}
+			],
+			"type" : "FILEPATH",
+			"value" : "/opt/maca/mxgpu_llvm/bin/mxcc"
+		},
+		{
+			"name" : "_CMAKE_LINKER_PUSHPOP_STATE_SUPPORTED",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "linker supports push/pop state"
+				}
+			],
+			"type" : "INTERNAL",
+			"value" : "FALSE"
+		},
+		{
+			"name" : "fused_moe_i8_tn_BINARY_DIR",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "Value Computed by CMake"
+				}
+			],
+			"type" : "STATIC",
+			"value" : "/root/Project/fusedmoe/build"
+		},
+		{
+			"name" : "fused_moe_i8_tn_IS_TOP_LEVEL",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "Value Computed by CMake"
+				}
+			],
+			"type" : "STATIC",
+			"value" : "ON"
+		},
+		{
+			"name" : "fused_moe_i8_tn_SOURCE_DIR",
+			"properties" : 
+			[
+				{
+					"name" : "HELPSTRING",
+					"value" : "Value Computed by CMake"
+				}
+			],
+			"type" : "STATIC",
+			"value" : "/root/Project/fusedmoe/standalone/fused_moe_i8_tn"
+		}
+	],
+	"kind" : "cache",
+	"version" : 
+	{
+		"major" : 2,
+		"minor" : 0
+	}
+}

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/cmakeFiles-v1-7899829d23c1c1ae3e98.json

@@ -0,0 +1,73 @@
+{
+	"inputs" : 
+	[
+		{
+			"path" : "CMakeLists.txt"
+		},
+		{
+			"isCMake" : true,
+			"isExternal" : true,
+			"path" : "/usr/share/cmake-3.28/Modules/CMakeDetermineSystem.cmake"
+		},
+		{
+			"isCMake" : true,
+			"isExternal" : true,
+			"path" : "/usr/share/cmake-3.28/Modules/CMakeSystem.cmake.in"
+		},
+		{
+			"isGenerated" : true,
+			"path" : "/root/Project/fusedmoe/build/CMakeFiles/3.28.3/CMakeSystem.cmake"
+		},
+		{
+			"isCMake" : true,
+			"isExternal" : true,
+			"path" : "/usr/share/cmake-3.28/Modules/CMakeNinjaFindMake.cmake"
+		},
+		{
+			"isCMake" : true,
+			"isExternal" : true,
+			"path" : "/usr/share/cmake-3.28/Modules/CMakeSystemSpecificInitialize.cmake"
+		},
+		{
+			"isCMake" : true,
+			"isExternal" : true,
+			"path" : "/usr/share/cmake-3.28/Modules/Platform/Linux-Initialize.cmake"
+		},
+		{
+			"isCMake" : true,
+			"isExternal" : true,
+			"path" : "/usr/share/cmake-3.28/Modules/CMakeSystemSpecificInformation.cmake"
+		},
+		{
+			"isCMake" : true,
+			"isExternal" : true,
+			"path" : "/usr/share/cmake-3.28/Modules/CMakeGenericSystem.cmake"
+		},
+		{
+			"isCMake" : true,
+			"isExternal" : true,
+			"path" : "/usr/share/cmake-3.28/Modules/CMakeInitializeConfigs.cmake"
+		},
+		{
+			"isCMake" : true,
+			"isExternal" : true,
+			"path" : "/usr/share/cmake-3.28/Modules/Platform/Linux.cmake"
+		},
+		{
+			"isCMake" : true,
+			"isExternal" : true,
+			"path" : "/usr/share/cmake-3.28/Modules/Platform/UnixPaths.cmake"
+		}
+	],
+	"kind" : "cmakeFiles",
+	"paths" : 
+	{
+		"build" : "/root/Project/fusedmoe/build",
+		"source" : "/root/Project/fusedmoe/standalone/fused_moe_i8_tn"
+	},
+	"version" : 
+	{
+		"major" : 1,
+		"minor" : 0
+	}
+}

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/codemodel-v2-7cd0b1b00876e71f364f.json

@@ -0,0 +1,69 @@
+{
+	"configurations" : 
+	[
+		{
+			"directories" : 
+			[
+				{
+					"build" : ".",
+					"jsonFile" : "directory-.-Debug-f5ebdc15457944623624.json",
+					"minimumCMakeVersion" : 
+					{
+						"string" : "3.20"
+					},
+					"projectIndex" : 0,
+					"source" : ".",
+					"targetIndexes" : 
+					[
+						0,
+						1
+					]
+				}
+			],
+			"name" : "Debug",
+			"projects" : 
+			[
+				{
+					"directoryIndexes" : 
+					[
+						0
+					],
+					"name" : "fused_moe_i8_tn",
+					"targetIndexes" : 
+					[
+						0,
+						1
+					]
+				}
+			],
+			"targets" : 
+			[
+				{
+					"directoryIndex" : 0,
+					"id" : "build_fused_moe_i8_tn_example::@6890427a1f51a3e7e1df",
+					"jsonFile" : "target-build_fused_moe_i8_tn_example-Debug-a97a299baa6c6c6d83d0.json",
+					"name" : "build_fused_moe_i8_tn_example",
+					"projectIndex" : 0
+				},
+				{
+					"directoryIndex" : 0,
+					"id" : "run::@6890427a1f51a3e7e1df",
+					"jsonFile" : "target-run-Debug-0d66e135afa1376e0f20.json",
+					"name" : "run",
+					"projectIndex" : 0
+				}
+			]
+		}
+	],
+	"kind" : "codemodel",
+	"paths" : 
+	{
+		"build" : "/root/Project/fusedmoe/build",
+		"source" : "/root/Project/fusedmoe/standalone/fused_moe_i8_tn"
+	},
+	"version" : 
+	{
+		"major" : 2,
+		"minor" : 6
+	}
+}

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/directory-.-Debug-f5ebdc15457944623624.json

@@ -0,0 +1,14 @@
+{
+	"backtraceGraph" : 
+	{
+		"commands" : [],
+		"files" : [],
+		"nodes" : []
+	},
+	"installers" : [],
+	"paths" : 
+	{
+		"build" : ".",
+		"source" : "."
+	}
+}

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/index-2026-05-29T05-34-43-0302.json

@@ -0,0 +1,132 @@
+{
+	"cmake" : 
+	{
+		"generator" : 
+		{
+			"multiConfig" : false,
+			"name" : "Ninja"
+		},
+		"paths" : 
+		{
+			"cmake" : "/usr/bin/cmake",
+			"cpack" : "/usr/bin/cpack",
+			"ctest" : "/usr/bin/ctest",
+			"root" : "/usr/share/cmake-3.28"
+		},
+		"version" : 
+		{
+			"isDirty" : false,
+			"major" : 3,
+			"minor" : 28,
+			"patch" : 3,
+			"string" : "3.28.3",
+			"suffix" : ""
+		}
+	},
+	"objects" : 
+	[
+		{
+			"jsonFile" : "codemodel-v2-7cd0b1b00876e71f364f.json",
+			"kind" : "codemodel",
+			"version" : 
+			{
+				"major" : 2,
+				"minor" : 6
+			}
+		},
+		{
+			"jsonFile" : "cache-v2-ea2ef11d05674d96d761.json",
+			"kind" : "cache",
+			"version" : 
+			{
+				"major" : 2,
+				"minor" : 0
+			}
+		},
+		{
+			"jsonFile" : "cmakeFiles-v1-7899829d23c1c1ae3e98.json",
+			"kind" : "cmakeFiles",
+			"version" : 
+			{
+				"major" : 1,
+				"minor" : 0
+			}
+		},
+		{
+			"jsonFile" : "toolchains-v1-8ae3cf416ede58af34e6.json",
+			"kind" : "toolchains",
+			"version" : 
+			{
+				"major" : 1,
+				"minor" : 0
+			}
+		}
+	],
+	"reply" : 
+	{
+		"client-vscode" : 
+		{
+			"query.json" : 
+			{
+				"requests" : 
+				[
+					{
+						"kind" : "cache",
+						"version" : 2
+					},
+					{
+						"kind" : "codemodel",
+						"version" : 2
+					},
+					{
+						"kind" : "toolchains",
+						"version" : 1
+					},
+					{
+						"kind" : "cmakeFiles",
+						"version" : 1
+					}
+				],
+				"responses" : 
+				[
+					{
+						"jsonFile" : "cache-v2-ea2ef11d05674d96d761.json",
+						"kind" : "cache",
+						"version" : 
+						{
+							"major" : 2,
+							"minor" : 0
+						}
+					},
+					{
+						"jsonFile" : "codemodel-v2-7cd0b1b00876e71f364f.json",
+						"kind" : "codemodel",
+						"version" : 
+						{
+							"major" : 2,
+							"minor" : 6
+						}
+					},
+					{
+						"jsonFile" : "toolchains-v1-8ae3cf416ede58af34e6.json",
+						"kind" : "toolchains",
+						"version" : 
+						{
+							"major" : 1,
+							"minor" : 0
+						}
+					},
+					{
+						"jsonFile" : "cmakeFiles-v1-7899829d23c1c1ae3e98.json",
+						"kind" : "cmakeFiles",
+						"version" : 
+						{
+							"major" : 1,
+							"minor" : 0
+						}
+					}
+				]
+			}
+		}
+	}
+}

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/target-build_fused_moe_i8_tn_example-Debug-a97a299baa6c6c6d83d0.json

@@ -0,0 +1,73 @@
+{
+	"backtrace" : 1,
+	"backtraceGraph" : 
+	{
+		"commands" : 
+		[
+			"add_custom_target"
+		],
+		"files" : 
+		[
+			"CMakeLists.txt"
+		],
+		"nodes" : 
+		[
+			{
+				"file" : 0
+			},
+			{
+				"command" : 0,
+				"file" : 0,
+				"line" : 38,
+				"parent" : 0
+			}
+		]
+	},
+	"id" : "build_fused_moe_i8_tn_example::@6890427a1f51a3e7e1df",
+	"name" : "build_fused_moe_i8_tn_example",
+	"paths" : 
+	{
+		"build" : ".",
+		"source" : "."
+	},
+	"sourceGroups" : 
+	[
+		{
+			"name" : "",
+			"sourceIndexes" : 
+			[
+				0
+			]
+		},
+		{
+			"name" : "CMake Rules",
+			"sourceIndexes" : 
+			[
+				1,
+				2
+			]
+		}
+	],
+	"sources" : 
+	[
+		{
+			"backtrace" : 1,
+			"isGenerated" : true,
+			"path" : "/root/Project/fusedmoe/build/CMakeFiles/build_fused_moe_i8_tn_example",
+			"sourceGroupIndex" : 0
+		},
+		{
+			"backtrace" : 0,
+			"isGenerated" : true,
+			"path" : "/root/Project/fusedmoe/build/CMakeFiles/build_fused_moe_i8_tn_example.rule",
+			"sourceGroupIndex" : 1
+		},
+		{
+			"backtrace" : 0,
+			"isGenerated" : true,
+			"path" : "/root/Project/fusedmoe/build/fused_moe_i8_tn_example.rule",
+			"sourceGroupIndex" : 1
+		}
+	],
+	"type" : "UTILITY"
+}

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/target-run-Debug-0d66e135afa1376e0f20.json

@@ -0,0 +1,73 @@
+{
+	"backtrace" : 1,
+	"backtraceGraph" : 
+	{
+		"commands" : 
+		[
+			"add_custom_target"
+		],
+		"files" : 
+		[
+			"CMakeLists.txt"
+		],
+		"nodes" : 
+		[
+			{
+				"file" : 0
+			},
+			{
+				"command" : 0,
+				"file" : 0,
+				"line" : 40,
+				"parent" : 0
+			}
+		]
+	},
+	"id" : "run::@6890427a1f51a3e7e1df",
+	"name" : "run",
+	"paths" : 
+	{
+		"build" : ".",
+		"source" : "."
+	},
+	"sourceGroups" : 
+	[
+		{
+			"name" : "",
+			"sourceIndexes" : 
+			[
+				0
+			]
+		},
+		{
+			"name" : "CMake Rules",
+			"sourceIndexes" : 
+			[
+				1,
+				2
+			]
+		}
+	],
+	"sources" : 
+	[
+		{
+			"backtrace" : 1,
+			"isGenerated" : true,
+			"path" : "/root/Project/fusedmoe/build/CMakeFiles/run",
+			"sourceGroupIndex" : 0
+		},
+		{
+			"backtrace" : 0,
+			"isGenerated" : true,
+			"path" : "/root/Project/fusedmoe/build/CMakeFiles/run.rule",
+			"sourceGroupIndex" : 1
+		},
+		{
+			"backtrace" : 0,
+			"isGenerated" : true,
+			"path" : "/root/Project/fusedmoe/build/fused_moe_i8_tn_example.rule",
+			"sourceGroupIndex" : 1
+		}
+	],
+	"type" : "UTILITY"
+}

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/toolchains-v1-8ae3cf416ede58af34e6.json

@@ -0,0 +1,18 @@
+{
+	"kind" : "toolchains",
+	"toolchains" : 
+	[
+		{
+			"compiler" : 
+			{
+				"implicit" : {}
+			},
+			"language" : "NONE"
+		}
+	],
+	"version" : 
+	{
+		"major" : 1,
+		"minor" : 0
+	}
+}

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeCache.txt

@@ -0,0 +1,127 @@
+# This is the CMakeCache file.
+# For build in directory: /root/Project/fusedmoe/build
+# It was generated by CMake: /usr/bin/cmake
+# You can edit this file to change values found and used by cmake.
+# If you do not want to change any of the values, simply exit the editor.
+# If you do want to change a value, simply edit, save, and exit the editor.
+# The syntax for the file is as follows:
+# KEY:TYPE=VALUE
+# KEY is the name of a variable in the cache.
+# TYPE is a hint to GUIs for the type of VALUE, DO NOT EDIT TYPE!.
+# VALUE is the current value for the KEY.
+
+########################
+# EXTERNAL cache entries
+########################
+
+//No help, variable specified on the command line.
+CMAKE_BUILD_TYPE:STRING=Debug
+
+//No help, variable specified on the command line.
+CMAKE_CXX_COMPILER:FILEPATH=/usr/bin/g++
+
+//No help, variable specified on the command line.
+CMAKE_C_COMPILER:FILEPATH=/usr/bin/gcc
+
+//No help, variable specified on the command line.
+CMAKE_EXPORT_COMPILE_COMMANDS:BOOL=TRUE
+
+//Value Computed by CMake.
+CMAKE_FIND_PACKAGE_REDIRECTS_DIR:STATIC=/root/Project/fusedmoe/build/CMakeFiles/pkgRedirects
+
+//Install path prefix, prepended onto install directories.
+CMAKE_INSTALL_PREFIX:PATH=/usr/local
+
+//Program used to build from build.ninja files.
+CMAKE_MAKE_PROGRAM:FILEPATH=/opt/conda/bin/ninja
+
+//Value Computed by CMake
+CMAKE_PROJECT_DESCRIPTION:STATIC=
+
+//Value Computed by CMake
+CMAKE_PROJECT_HOMEPAGE_URL:STATIC=
+
+//Value Computed by CMake
+CMAKE_PROJECT_NAME:STATIC=fused_moe_i8_tn
+
+//If set, runtime paths are not added when installing shared libraries,
+// but are added when building.
+CMAKE_SKIP_INSTALL_RPATH:BOOL=NO
+
+//If set, runtime paths are not added when using shared libraries.
+CMAKE_SKIP_RPATH:BOOL=NO
+
+//If this value is on, makefiles will be generated without the
+// .SILENT directive, and all commands will be echoed to the console
+// during the make.  This is useful for debugging only. With Visual
+// Studio IDE projects all commands are done without /nologo.
+CMAKE_VERBOSE_MAKEFILE:BOOL=FALSE
+
+//Path to MACA SDK
+MACA_PATH:PATH=/opt/maca
+
+//Path to a program.
+MXCC:FILEPATH=/opt/maca/mxgpu_llvm/bin/mxcc
+
+//Value Computed by CMake
+fused_moe_i8_tn_BINARY_DIR:STATIC=/root/Project/fusedmoe/build
+
+//Value Computed by CMake
+fused_moe_i8_tn_IS_TOP_LEVEL:STATIC=ON
+
+//Value Computed by CMake
+fused_moe_i8_tn_SOURCE_DIR:STATIC=/root/Project/fusedmoe/standalone/fused_moe_i8_tn
+
+
+########################
+# INTERNAL cache entries
+########################
+
+//This is the directory where this CMakeCache.txt was created
+CMAKE_CACHEFILE_DIR:INTERNAL=/root/Project/fusedmoe/build
+//Major version of cmake used to create the current loaded cache
+CMAKE_CACHE_MAJOR_VERSION:INTERNAL=3
+//Minor version of cmake used to create the current loaded cache
+CMAKE_CACHE_MINOR_VERSION:INTERNAL=28
+//Patch version of cmake used to create the current loaded cache
+CMAKE_CACHE_PATCH_VERSION:INTERNAL=3
+//Path to CMake executable.
+CMAKE_COMMAND:INTERNAL=/usr/bin/cmake
+//Path to cpack program executable.
+CMAKE_CPACK_COMMAND:INTERNAL=/usr/bin/cpack
+//Path to ctest program executable.
+CMAKE_CTEST_COMMAND:INTERNAL=/usr/bin/ctest
+//Name of external makefile project generator.
+CMAKE_EXTRA_GENERATOR:INTERNAL=
+//Name of generator.
+CMAKE_GENERATOR:INTERNAL=Ninja
+//Generator instance identifier.
+CMAKE_GENERATOR_INSTANCE:INTERNAL=
+//Name of generator platform.
+CMAKE_GENERATOR_PLATFORM:INTERNAL=
+//Name of generator toolset.
+CMAKE_GENERATOR_TOOLSET:INTERNAL=
+//Source directory with the top level CMakeLists.txt file for this
+// project
+CMAKE_HOME_DIRECTORY:INTERNAL=/root/Project/fusedmoe/standalone/fused_moe_i8_tn
+//Install .so files without execute permission.
+CMAKE_INSTALL_SO_NO_EXE:INTERNAL=1
+//ADVANCED property for variable: CMAKE_MAKE_PROGRAM
+CMAKE_MAKE_PROGRAM-ADVANCED:INTERNAL=1
+//number of local generators
+CMAKE_NUMBER_OF_MAKEFILES:INTERNAL=1
+//Platform information initialized
+CMAKE_PLATFORM_INFO_INITIALIZED:INTERNAL=1
+//Path to CMake installation.
+CMAKE_ROOT:INTERNAL=/usr/share/cmake-3.28
+//ADVANCED property for variable: CMAKE_SKIP_INSTALL_RPATH
+CMAKE_SKIP_INSTALL_RPATH-ADVANCED:INTERNAL=1
+//ADVANCED property for variable: CMAKE_SKIP_RPATH
+CMAKE_SKIP_RPATH-ADVANCED:INTERNAL=1
+//uname command
+CMAKE_UNAME:INTERNAL=/usr/bin/uname
+//ADVANCED property for variable: CMAKE_VERBOSE_MAKEFILE
+CMAKE_VERBOSE_MAKEFILE-ADVANCED:INTERNAL=1
+//linker supports push/pop state
+_CMAKE_LINKER_PUSHPOP_STATE_SUPPORTED:INTERNAL=FALSE
+

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeFiles/3.28.3/CMakeSystem.cmake

@@ -0,0 +1,15 @@
+set(CMAKE_HOST_SYSTEM "Linux-5.15.0-58-generic")
+set(CMAKE_HOST_SYSTEM_NAME "Linux")
+set(CMAKE_HOST_SYSTEM_VERSION "5.15.0-58-generic")
+set(CMAKE_HOST_SYSTEM_PROCESSOR "x86_64")
+
+
+
+set(CMAKE_SYSTEM "Linux-5.15.0-58-generic")
+set(CMAKE_SYSTEM_NAME "Linux")
+set(CMAKE_SYSTEM_VERSION "5.15.0-58-generic")
+set(CMAKE_SYSTEM_PROCESSOR "x86_64")
+
+set(CMAKE_CROSSCOMPILING "FALSE")
+
+set(CMAKE_SYSTEM_LOADED 1)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeFiles/CMakeConfigureLog.yaml

@@ -0,0 +1,11 @@
+
+---
+events:
+  -
+    kind: "message-v1"
+    backtrace:
+      - "/usr/share/cmake-3.28/Modules/CMakeDetermineSystem.cmake:233 (message)"
+      - "CMakeLists.txt:3 (project)"
+    message: |
+      The system is: Linux - 5.15.0-58-generic - x86_64
+...

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeFiles/TargetDirectories.txt

@@ -0,0 +1,4 @@
+/root/Project/fusedmoe/build/CMakeFiles/build_fused_moe_i8_tn_example.dir
+/root/Project/fusedmoe/build/CMakeFiles/run.dir
+/root/Project/fusedmoe/build/CMakeFiles/edit_cache.dir
+/root/Project/fusedmoe/build/CMakeFiles/rebuild_cache.dir

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeFiles/cmake.check_cache

@@ -0,0 +1 @@
+# This file is generated by cmake for dependency checking of the CMakeCache.txt file

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeFiles/rules.ninja

@@ -0,0 +1,45 @@
+# CMAKE generated file: DO NOT EDIT!
+# Generated by "Ninja" Generator, CMake Version 3.28
+
+# This file contains all the rules used to get the outputs files
+# built from the input files.
+# It is included in the main 'build.ninja'.
+
+# =============================================================================
+# Project: fused_moe_i8_tn
+# Configurations: Debug
+# =============================================================================
+# =============================================================================
+
+#############################################
+# Rule for running custom commands.
+
+rule CUSTOM_COMMAND
+  command = $COMMAND
+  description = $DESC
+
+
+#############################################
+# Rule for re-running cmake.
+
+rule RERUN_CMAKE
+  command = /usr/bin/cmake --regenerate-during-build -S/root/Project/fusedmoe/standalone/fused_moe_i8_tn -B/root/Project/fusedmoe/build
+  description = Re-running CMake...
+  generator = 1
+
+
+#############################################
+# Rule for cleaning all built files.
+
+rule CLEAN
+  command = /opt/conda/bin/ninja $FILE_ARG -t clean $TARGETS
+  description = Cleaning all built files...
+
+
+#############################################
+# Rule for printing all primary targets available.
+
+rule HELP
+  command = /opt/conda/bin/ninja -t targets
+  description = All primary targets available:
+

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/build.ninja

@@ -0,0 +1,146 @@
+# CMAKE generated file: DO NOT EDIT!
+# Generated by "Ninja" Generator, CMake Version 3.28
+
+# This file contains all the build statements describing the
+# compilation DAG.
+
+# =============================================================================
+# Write statements declared in CMakeLists.txt:
+# 
+# Which is the root file.
+# =============================================================================
+
+# =============================================================================
+# Project: fused_moe_i8_tn
+# Configurations: Debug
+# =============================================================================
+
+#############################################
+# Minimal version of Ninja required by this file
+
+ninja_required_version = 1.5
+
+
+#############################################
+# Set configuration variable for custom commands.
+
+CONFIGURATION = Debug
+# =============================================================================
+# Include auxiliary files.
+
+
+#############################################
+# Include rules file.
+
+include CMakeFiles/rules.ninja
+
+# =============================================================================
+
+#############################################
+# Logical path to working directory; prefix for absolute paths.
+
+cmake_ninja_workdir = /root/Project/fusedmoe/build/
+
+#############################################
+# Utility command for build_fused_moe_i8_tn_example
+
+build build_fused_moe_i8_tn_example: phony CMakeFiles/build_fused_moe_i8_tn_example fused_moe_i8_tn_example
+
+
+#############################################
+# Utility command for run
+
+build run: phony CMakeFiles/run fused_moe_i8_tn_example
+
+
+#############################################
+# Utility command for edit_cache
+
+build CMakeFiles/edit_cache.util: CUSTOM_COMMAND
+  COMMAND = cd /root/Project/fusedmoe/build && /usr/bin/cmake -E echo No\ interactive\ CMake\ dialog\ available.
+  DESC = No interactive CMake dialog available...
+  restat = 1
+
+build edit_cache: phony CMakeFiles/edit_cache.util
+
+
+#############################################
+# Utility command for rebuild_cache
+
+build CMakeFiles/rebuild_cache.util: CUSTOM_COMMAND
+  COMMAND = cd /root/Project/fusedmoe/build && /usr/bin/cmake --regenerate-during-build -S/root/Project/fusedmoe/standalone/fused_moe_i8_tn -B/root/Project/fusedmoe/build
+  DESC = Running CMake to regenerate build system...
+  pool = console
+  restat = 1
+
+build rebuild_cache: phony CMakeFiles/rebuild_cache.util
+
+
+#############################################
+# Phony custom command for CMakeFiles/build_fused_moe_i8_tn_example
+
+build CMakeFiles/build_fused_moe_i8_tn_example | ${cmake_ninja_workdir}CMakeFiles/build_fused_moe_i8_tn_example: phony fused_moe_i8_tn_example
+
+
+#############################################
+# Custom command for fused_moe_i8_tn_example
+
+build fused_moe_i8_tn_example | ${cmake_ninja_workdir}fused_moe_i8_tn_example: CUSTOM_COMMAND /root/Project/fusedmoe/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_example.cpp
+  COMMAND = cd /root/Project/fusedmoe/standalone/fused_moe_i8_tn && /opt/maca/mxgpu_llvm/bin/mxcc -std=c++17 -xmaca -I\"/root/Project/fusedmoe/standalone/fused_moe_i8_tn/src\" -I\"/opt/maca/include\" /root/Project/fusedmoe/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_example.cpp -L\"/opt/maca/lib\" -lmcruntime -o /root/Project/fusedmoe/build/fused_moe_i8_tn_example
+  DESC = Generating fused_moe_i8_tn_example
+  restat = 1
+
+
+#############################################
+# Custom command for CMakeFiles/run
+
+build CMakeFiles/run | ${cmake_ninja_workdir}CMakeFiles/run: CUSTOM_COMMAND fused_moe_i8_tn_example
+  COMMAND = cd /root/Project/fusedmoe/build && /root/Project/fusedmoe/build/fused_moe_i8_tn_example
+  pool = console
+
+# =============================================================================
+# Target aliases.
+
+# =============================================================================
+# Folder targets.
+
+# =============================================================================
+
+#############################################
+# Folder: /root/Project/fusedmoe/build
+
+build all: phony build_fused_moe_i8_tn_example
+
+# =============================================================================
+# Built-in targets
+
+
+#############################################
+# Re-run CMake if any of its inputs changed.
+
+build build.ninja: RERUN_CMAKE | /root/Project/fusedmoe/standalone/fused_moe_i8_tn/CMakeLists.txt /usr/share/cmake-3.28/Modules/CMakeDetermineSystem.cmake /usr/share/cmake-3.28/Modules/CMakeGenericSystem.cmake /usr/share/cmake-3.28/Modules/CMakeInitializeConfigs.cmake /usr/share/cmake-3.28/Modules/CMakeNinjaFindMake.cmake /usr/share/cmake-3.28/Modules/CMakeSystem.cmake.in /usr/share/cmake-3.28/Modules/CMakeSystemSpecificInformation.cmake /usr/share/cmake-3.28/Modules/CMakeSystemSpecificInitialize.cmake /usr/share/cmake-3.28/Modules/Platform/Linux-Initialize.cmake /usr/share/cmake-3.28/Modules/Platform/Linux.cmake /usr/share/cmake-3.28/Modules/Platform/UnixPaths.cmake CMakeCache.txt CMakeFiles/3.28.3/CMakeSystem.cmake
+  pool = console
+
+
+#############################################
+# A missing CMake input file is not an error.
+
+build /root/Project/fusedmoe/standalone/fused_moe_i8_tn/CMakeLists.txt /usr/share/cmake-3.28/Modules/CMakeDetermineSystem.cmake /usr/share/cmake-3.28/Modules/CMakeGenericSystem.cmake /usr/share/cmake-3.28/Modules/CMakeInitializeConfigs.cmake /usr/share/cmake-3.28/Modules/CMakeNinjaFindMake.cmake /usr/share/cmake-3.28/Modules/CMakeSystem.cmake.in /usr/share/cmake-3.28/Modules/CMakeSystemSpecificInformation.cmake /usr/share/cmake-3.28/Modules/CMakeSystemSpecificInitialize.cmake /usr/share/cmake-3.28/Modules/Platform/Linux-Initialize.cmake /usr/share/cmake-3.28/Modules/Platform/Linux.cmake /usr/share/cmake-3.28/Modules/Platform/UnixPaths.cmake CMakeCache.txt CMakeFiles/3.28.3/CMakeSystem.cmake: phony
+
+
+#############################################
+# Clean all the built files.
+
+build clean: CLEAN
+
+
+#############################################
+# Print all primary targets available.
+
+build help: HELP
+
+
+#############################################
+# Make the all target the default.
+
+default all

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/cmake_install.cmake

@@ -0,0 +1,49 @@
+# Install script for directory: /root/Project/fusedmoe/standalone/fused_moe_i8_tn
+
+# Set the install prefix
+if(NOT DEFINED CMAKE_INSTALL_PREFIX)
+  set(CMAKE_INSTALL_PREFIX "/usr/local")
+endif()
+string(REGEX REPLACE "/$" "" CMAKE_INSTALL_PREFIX "${CMAKE_INSTALL_PREFIX}")
+
+# Set the install configuration name.
+if(NOT DEFINED CMAKE_INSTALL_CONFIG_NAME)
+  if(BUILD_TYPE)
+    string(REGEX REPLACE "^[^A-Za-z0-9_]+" ""
+           CMAKE_INSTALL_CONFIG_NAME "${BUILD_TYPE}")
+  else()
+    set(CMAKE_INSTALL_CONFIG_NAME "Debug")
+  endif()
+  message(STATUS "Install configuration: \"${CMAKE_INSTALL_CONFIG_NAME}\"")
+endif()
+
+# Set the component getting installed.
+if(NOT CMAKE_INSTALL_COMPONENT)
+  if(COMPONENT)
+    message(STATUS "Install component: \"${COMPONENT}\"")
+    set(CMAKE_INSTALL_COMPONENT "${COMPONENT}")
+  else()
+    set(CMAKE_INSTALL_COMPONENT)
+  endif()
+endif()
+
+# Install shared libraries without execute permission?
+if(NOT DEFINED CMAKE_INSTALL_SO_NO_EXE)
+  set(CMAKE_INSTALL_SO_NO_EXE "1")
+endif()
+
+# Is this installation the result of a crosscompile?
+if(NOT DEFINED CMAKE_CROSSCOMPILING)
+  set(CMAKE_CROSSCOMPILING "FALSE")
+endif()
+
+if(CMAKE_INSTALL_COMPONENT)
+  set(CMAKE_INSTALL_MANIFEST "install_manifest_${CMAKE_INSTALL_COMPONENT}.txt")
+else()
+  set(CMAKE_INSTALL_MANIFEST "install_manifest.txt")
+endif()
+
+string(REPLACE ";" "\n" CMAKE_INSTALL_MANIFEST_CONTENT
+       "${CMAKE_INSTALL_MANIFEST_FILES}")
+file(WRITE "/root/Project/fusedmoe/build/${CMAKE_INSTALL_MANIFEST}"
+     "${CMAKE_INSTALL_MANIFEST_CONTENT}")

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/scripts/build_fused_moe_i8_tn_pybind.sh

@@ -0,0 +1,52 @@
+#!/usr/bin/env bash
+
+set -euo pipefail
+
+MACA_PATH="${MACA_PATH:-/opt/maca}"
+PYTHON_BIN="${PYTHON_BIN:-/opt/conda/bin/python}"
+ROOT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)"
+
+export MACA_PATH
+export LD_LIBRARY_PATH="$MACA_PATH/mxgpu_llvm/lib:$MACA_PATH/lib:${LD_LIBRARY_PATH:-}"
+
+BUILD_DIR="$ROOT_DIR/standalone/fused_moe_i8_tn/build"
+SO="$BUILD_DIR/fused_moe_i8_tn_pybind.so"
+SRC="$ROOT_DIR/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_pybind.cpp"
+
+mkdir -p "$BUILD_DIR"
+
+PYTHON_INCLUDE="$("${PYTHON_BIN}" -c "import sysconfig; print(sysconfig.get_path('include'))")"
+PYTHON_LIB="$("${PYTHON_BIN}" -c "import sysconfig; print(sysconfig.get_config_var('LIBDIR'))")"
+PYTHON_LDLIB="$("${PYTHON_BIN}" -c "import sysconfig; print(sysconfig.get_config_var('LDLIBRARY'))")"
+
+PYBIND11_INCLUDE="$ROOT_DIR/third_party/pybind11/include"
+
+"$MACA_PATH/mxgpu_llvm/bin/mxcc" \
+    -std=c++17 \
+    -O2 \
+    -c \
+    -xmaca \
+    -fPIC \
+    -I"$ROOT_DIR/include" \
+    -I"$MACA_PATH/include" \
+    -I"$PYTHON_INCLUDE" \
+    -I"$PYBIND11_INCLUDE" \
+    "$SRC" \
+    -o "$BUILD_DIR/fused_moe_i8_tn_pybind.o"
+
+
+g++ \
+    -std=c++17 \
+    -O2 \
+    -shared \
+    -fPIC \
+    "$BUILD_DIR/fused_moe_i8_tn_pybind.o" \
+    -L"$MACA_PATH/lib" \
+    -L"/opt/conda/lib" \
+    -Wl,-rpath,"/opt/conda/lib" \
+    -l:"libpython3.12.so" \
+    -lmcruntime \
+    -lmccompiler \
+    -o "$SO"
+
+echo "[SUCCESS] $SO"

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/scripts/run_fused_moe_i8_tn_benchmark.sh

@@ -0,0 +1,12 @@
+#!/usr/bin/env bash
+
+set -euo pipefail
+
+ROOT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)"
+PYTHON_BIN="${PYTHON_BIN:-/opt/conda/bin/python}"
+MACA_PATH="${MACA_PATH:-/opt/maca-20260318}"
+
+export MACA_PATH
+export LD_LIBRARY_PATH="$MACA_PATH/mxgpu_llvm/lib:$MACA_PATH/lib:${LD_LIBRARY_PATH:-}"
+
+"$PYTHON_BIN" "$ROOT_DIR/standalone/fused_moe_i8_tn/python/benchmark_fused_moe_i8_tn.py" "$@"

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/scripts/run_fused_moe_i8_tn_pybind_test.sh

@@ -0,0 +1,12 @@
+#!/usr/bin/env bash
+
+set -euo pipefail
+
+ROOT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)"
+PYTHON_BIN="${PYTHON_BIN:-/opt/conda/bin/python}"
+MACA_PATH="${MACA_PATH:-/opt/maca-20260318}"
+
+export MACA_PATH
+export LD_LIBRARY_PATH="$MACA_PATH/mxgpu_llvm/lib:$MACA_PATH/lib:${LD_LIBRARY_PATH:-}"
+
+"$PYTHON_BIN" "$ROOT_DIR/standalone/fused_moe_i8_tn/python/test_fused_moe_i8_tn_pybind.py" "$@"

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_bf16_tn/CMakeLists.txt

@@ -0,0 +1,45 @@
+cmake_minimum_required(VERSION 3.20)
+
+project(fused_moe_bf16_tn LANGUAGES NONE)
+
+set(MACA_PATH "$ENV{MACA_PATH}" CACHE PATH "Path to MACA SDK")
+if(NOT MACA_PATH)
+  set(MACA_PATH "/opt/maca")
+endif()
+
+find_program(MXCC
+  NAMES mxcc
+  PATHS "${MACA_PATH}/mxgpu_llvm/bin"
+  NO_DEFAULT_PATH)
+
+if(NOT MXCC)
+  message(FATAL_ERROR "mxcc not found under ${MACA_PATH}/mxgpu_llvm/bin")
+endif()
+
+set(EXAMPLE_ROOT "${CMAKE_CURRENT_SOURCE_DIR}")
+get_filename_component(ROOT_DIR "${EXAMPLE_ROOT}/../.." ABSOLUTE)
+set(SRC "${EXAMPLE_ROOT}/src/fused_moe_bf16_tn_example.cpp")
+set(BIN "${CMAKE_CURRENT_BINARY_DIR}/fused_moe_bf16_tn_example")
+
+add_custom_command(
+  OUTPUT "${BIN}"
+  COMMAND "${MXCC}"
+          -std=c++17
+          -xmaca
+          -I"${ROOT_DIR}/include"
+          -I"${MACA_PATH}/include"
+          "${SRC}"
+          -L"${MACA_PATH}/lib"
+          -lmcruntime
+          -o "${BIN}"
+  DEPENDS "${SRC}"
+  WORKING_DIRECTORY "${ROOT_DIR}"
+  VERBATIM)
+
+add_custom_target(build_fused_moe_bf16_tn_example ALL DEPENDS "${BIN}")
+
+add_custom_target(
+  run
+  COMMAND "${BIN}"
+  DEPENDS "${BIN}"
+  USES_TERMINAL)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_bf16_tn/Makefile

@@ -0,0 +1,28 @@
+MACA_PATH ?= /opt/maca
+MXCC := $(MACA_PATH)/mxgpu_llvm/bin/mxcc
+ROOT_DIR := $(abspath $(CURDIR)/../..)
+BUILD_DIR := $(CURDIR)/build
+SRC := $(CURDIR)/src/fused_moe_bf16_tn_example.cpp
+BIN := $(BUILD_DIR)/fused_moe_bf16_tn_example
+
+.PHONY: all build run clean
+
+all: build
+
+build: $(BIN)
+
+$(BIN): $(SRC)
+	mkdir -p $(BUILD_DIR)
+	$(MXCC) -std=c++17 -xmaca \
+		-I$(ROOT_DIR)/include \
+		-I$(MACA_PATH)/include \
+		$(SRC) \
+		-L$(MACA_PATH)/lib \
+		-lmcruntime \
+		-o $(BIN)
+
+run: $(BIN)
+	$(BIN)
+
+clean:
+	rm -rf $(BUILD_DIR)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_bf16_tn/src/fused_moe_bf16_tn_example.cpp

@@ -0,0 +1,246 @@
+#include <mc_runtime.h>
+
+#include <algorithm>
+#include <cmath>
+#include <cstdint>
+#include <cstdlib>
+#include <iostream>
+#include <vector>
+
+#include "mctlass/bfloat16.h"
+#include "mctlass/frontend_op/gemm_config.h"
+#include "mctlass/frontend_op/mctlass_moe_gemm.h"
+
+namespace {
+
+using Bf16 = maca_bfloat16;
+using LayoutA = mctlass::layout::RowMajor;
+using LayoutB = mctlass::layout::ColumnMajor;
+using LayoutC = mctlass::layout::RowMajor;
+using GemmOp = mctlassMoeGemm<Bf16, LayoutA, Bf16, LayoutB, Bf16, LayoutC, float>;
+
+constexpr int kNumExperts = 2;
+constexpr int kNumTokens = 256;
+constexpr int kTopK = 1;
+constexpr int kEM = 256;
+constexpr int kN = 128;
+constexpr int kK = 64;
+constexpr int kTileM = 128;
+
+void check_mc(mcError_t status, const char *expr) {
+    if (status != mcSuccess) {
+        std::cerr << expr << " failed: " << mcGetErrorString(status) << '\n';
+        std::exit(EXIT_FAILURE);
+    }
+}
+
+void check_mctlass(mctlass::Status status, const char *expr) {
+    if (status != mctlass::Status::kSuccess) {
+        std::cerr << expr << " failed: " << mctlass::mctlassGetStatusString(status) << '\n';
+        std::exit(EXIT_FAILURE);
+    }
+}
+
+Bf16 float_to_bf16(float value) {
+    return mctlass::bfloat16_t(value).to_bfloat();
+}
+
+float bf16_to_float(Bf16 value) {
+    return static_cast<float>(mctlass::bfloat16_t(value));
+}
+
+void fill_inputs(std::vector<Bf16> &a,
+                 std::vector<Bf16> &b_col_major,
+                 std::vector<float> &moe_weights,
+                 std::vector<int> &token_ids,
+                 std::vector<int> &expert_ids) {
+    a.resize(static_cast<size_t>(kNumTokens) * kK);
+    b_col_major.resize(static_cast<size_t>(kNumExperts) * kN * kK);
+    moe_weights.resize(kEM);
+    token_ids.resize(kEM);
+    expert_ids = {0, 1};
+
+    for (int row = 0; row < kNumTokens; ++row) {
+        for (int kk = 0; kk < kK; ++kk) {
+            const float value = ((row * 11 + kk * 5 + 7) % 29 - 14) * 0.125f;
+            a[static_cast<size_t>(row) * kK + kk] = float_to_bf16(value);
+        }
+        token_ids[row] = row;
+        moe_weights[row] = 0.5f + 0.03125f * static_cast<float>(row % 7);
+    }
+
+    for (int expert = 0; expert < kNumExperts; ++expert) {
+        for (int col = 0; col < kN; ++col) {
+            for (int kk = 0; kk < kK; ++kk) {
+                const float value = ((expert * 13 + col * 3 + kk * 7 + 1) % 31 - 15) * 0.0625f;
+                b_col_major[(static_cast<size_t>(expert) * kN + col) * kK + kk] = float_to_bf16(value);
+            }
+        }
+    }
+}
+
+std::vector<Bf16> reference_fused_moe(const std::vector<Bf16> &a,
+                                      const std::vector<Bf16> &b_col_major,
+                                      const std::vector<float> &moe_weights,
+                                      const std::vector<int> &token_ids,
+                                      const std::vector<int> &expert_ids) {
+    std::vector<Bf16> out(static_cast<size_t>(kEM) * kN, float_to_bf16(0.0f));
+
+    for (int routed_row = 0; routed_row < kEM; ++routed_row) {
+        const int token = token_ids[routed_row];
+        const int tile_idx = routed_row / kTileM;
+        const int expert = expert_ids[tile_idx];
+        const float moe_weight = moe_weights[routed_row];
+
+        for (int col = 0; col < kN; ++col) {
+            float acc = 0.0f;
+            for (int kk = 0; kk < kK; ++kk) {
+                acc += bf16_to_float(a[static_cast<size_t>(token) * kK + kk]) *
+                       bf16_to_float(b_col_major[(static_cast<size_t>(expert) * kN + col) * kK + kk]);
+            }
+            out[static_cast<size_t>(routed_row) * kN + col] = float_to_bf16(acc * moe_weight);
+        }
+    }
+
+    return out;
+}
+
+bool validate_result(const std::vector<Bf16> &got, const std::vector<Bf16> &expected) {
+    constexpr float kTolerance = 2e-2f;
+    size_t mismatch_count = 0;
+    size_t first_bad = 0;
+    float max_abs = 0.0f;
+
+    for (size_t i = 0; i < got.size(); ++i) {
+        const float got_f = bf16_to_float(got[i]);
+        const float exp_f = bf16_to_float(expected[i]);
+        const float abs_err = std::fabs(got_f - exp_f);
+        max_abs = std::max(max_abs, abs_err);
+        if (abs_err > kTolerance) {
+            if (mismatch_count == 0) {
+                first_bad = i;
+            }
+            ++mismatch_count;
+        }
+    }
+
+    if (mismatch_count != 0) {
+        const int row = static_cast<int>(first_bad / kN);
+        const int col = static_cast<int>(first_bad % kN);
+        std::cerr << "fused_moe_bf16_tn failed"
+                  << ": mismatches=" << mismatch_count
+                  << ", first mismatch at (" << row << ", " << col << ")"
+                  << ", got=" << bf16_to_float(got[first_bad])
+                  << ", expected=" << bf16_to_float(expected[first_bad])
+                  << ", max_abs=" << max_abs << '\n';
+        return false;
+    }
+
+    std::cout << "fused_moe_bf16_tn passed"
+              << ": rows=" << kEM
+              << ", topk=" << kTopK
+              << ", N=" << kN
+              << ", K=" << kK
+              << ", sample C[0]=" << bf16_to_float(got.front())
+              << ", C[last]=" << bf16_to_float(got.back())
+              << ", max_abs=" << max_abs << '\n';
+    return true;
+}
+
+}  // namespace
+
+int main() {
+    int device_count = 0;
+    check_mc(mcGetDeviceCount(&device_count), "mcGetDeviceCount");
+    if (device_count <= 0) {
+        std::cerr << "No MACA device is visible.\n";
+        return EXIT_FAILURE;
+    }
+    check_mc(mcSetDevice(0), "mcSetDevice");
+
+    std::vector<Bf16> host_a;
+    std::vector<Bf16> host_b;
+    std::vector<float> host_moe_weights;
+    std::vector<int> host_token_ids;
+    std::vector<int> host_expert_ids;
+    std::vector<int32_t> host_num_tokens_post_padded(1, kEM);
+    std::vector<Bf16> host_output(static_cast<size_t>(kEM) * kN, float_to_bf16(0.0f));
+
+    fill_inputs(host_a, host_b, host_moe_weights, host_token_ids, host_expert_ids);
+    const std::vector<Bf16> expected =
+        reference_fused_moe(host_a, host_b, host_moe_weights, host_token_ids, host_expert_ids);
+
+    Bf16 *dev_a = nullptr;
+    Bf16 *dev_b = nullptr;
+    float *dev_moe_weights = nullptr;
+    int *dev_token_ids = nullptr;
+    int *dev_expert_ids = nullptr;
+    int32_t *dev_num_tokens_post_padded = nullptr;
+    Bf16 *dev_c = nullptr;
+
+    check_mc(mcMalloc(reinterpret_cast<void **>(&dev_a), host_a.size() * sizeof(Bf16)), "mcMalloc(dev_a)");
+    check_mc(mcMalloc(reinterpret_cast<void **>(&dev_b), host_b.size() * sizeof(Bf16)), "mcMalloc(dev_b)");
+    check_mc(mcMalloc(reinterpret_cast<void **>(&dev_moe_weights), host_moe_weights.size() * sizeof(float)),
+             "mcMalloc(dev_moe_weights)");
+    check_mc(mcMalloc(reinterpret_cast<void **>(&dev_token_ids), host_token_ids.size() * sizeof(int)),
+             "mcMalloc(dev_token_ids)");
+    check_mc(mcMalloc(reinterpret_cast<void **>(&dev_expert_ids), host_expert_ids.size() * sizeof(int)),
+             "mcMalloc(dev_expert_ids)");
+    check_mc(mcMalloc(reinterpret_cast<void **>(&dev_num_tokens_post_padded),
+                      host_num_tokens_post_padded.size() * sizeof(int32_t)),
+             "mcMalloc(dev_num_tokens_post_padded)");
+    check_mc(mcMalloc(reinterpret_cast<void **>(&dev_c), host_output.size() * sizeof(Bf16)), "mcMalloc(dev_c)");
+
+    check_mc(mcMemcpy(dev_a, host_a.data(), host_a.size() * sizeof(Bf16), mcMemcpyHostToDevice), "mcMemcpy(dev_a)");
+    check_mc(mcMemcpy(dev_b, host_b.data(), host_b.size() * sizeof(Bf16), mcMemcpyHostToDevice), "mcMemcpy(dev_b)");
+    check_mc(mcMemcpy(dev_moe_weights,
+                      host_moe_weights.data(),
+                      host_moe_weights.size() * sizeof(float),
+                      mcMemcpyHostToDevice),
+             "mcMemcpy(dev_moe_weights)");
+    check_mc(mcMemcpy(dev_token_ids, host_token_ids.data(), host_token_ids.size() * sizeof(int), mcMemcpyHostToDevice),
+             "mcMemcpy(dev_token_ids)");
+    check_mc(mcMemcpy(dev_expert_ids,
+                      host_expert_ids.data(),
+                      host_expert_ids.size() * sizeof(int),
+                      mcMemcpyHostToDevice),
+             "mcMemcpy(dev_expert_ids)");
+    check_mc(mcMemcpy(dev_num_tokens_post_padded,
+                      host_num_tokens_post_padded.data(),
+                      host_num_tokens_post_padded.size() * sizeof(int32_t),
+                      mcMemcpyHostToDevice),
+             "mcMemcpy(dev_num_tokens_post_padded)");
+    check_mc(mcMemset(dev_c, 0, host_output.size() * sizeof(Bf16)), "mcMemset(dev_c)");
+
+    GemmOp gemm_op;
+    typename GemmOp::Arguments args(
+        mctlass::gemm::GemmUniversalMode::kGemm,
+        mctlass::gemm::BatchedGemmCoord(kEM, kN, kK, kNumExperts),
+        typename GemmOp::epilogueParams(dev_moe_weights),
+        dev_a,
+        dev_b,
+        dev_c,
+        typename GemmOp::moeParams(dev_token_ids,
+                                   dev_expert_ids,
+                                   dev_num_tokens_post_padded,
+                                   kEM,
+                                   kTopK,
+                                   true));
+
+    check_mctlass(gemm_op(args, nullptr, nullptr), "gemm_op");
+    check_mc(mcDeviceSynchronize(), "mcDeviceSynchronize");
+    check_mc(mcGetLastError(), "mcGetLastError");
+
+    check_mc(mcMemcpy(host_output.data(), dev_c, host_output.size() * sizeof(Bf16), mcMemcpyDeviceToHost),
+             "mcMemcpy(host_output)");
+
+    check_mc(mcFree(dev_a), "mcFree(dev_a)");
+    check_mc(mcFree(dev_b), "mcFree(dev_b)");
+    check_mc(mcFree(dev_moe_weights), "mcFree(dev_moe_weights)");
+    check_mc(mcFree(dev_token_ids), "mcFree(dev_token_ids)");
+    check_mc(mcFree(dev_expert_ids), "mcFree(dev_expert_ids)");
+    check_mc(mcFree(dev_num_tokens_post_padded), "mcFree(dev_num_tokens_post_padded)");
+    check_mc(mcFree(dev_c), "mcFree(dev_c)");
+
+    return validate_result(host_output, expected) ? EXIT_SUCCESS : EXIT_FAILURE;
+}

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/CMakeLists.txt

@@ -0,0 +1,44 @@
+cmake_minimum_required(VERSION 3.20)
+
+project(fused_moe_i8_tn LANGUAGES NONE)
+
+set(MACA_PATH "$ENV{MACA_PATH}" CACHE PATH "Path to MACA SDK")
+if(NOT MACA_PATH)
+  set(MACA_PATH "/opt/maca")
+endif()
+
+find_program(MXCC
+  NAMES mxcc
+  PATHS "${MACA_PATH}/mxgpu_llvm/bin"
+  NO_DEFAULT_PATH)
+
+if(NOT MXCC)
+  message(FATAL_ERROR "mxcc not found under ${MACA_PATH}/mxgpu_llvm/bin")
+endif()
+
+set(EXAMPLE_ROOT "${CMAKE_CURRENT_SOURCE_DIR}")
+set(SRC "${EXAMPLE_ROOT}/src/fused_moe_i8_tn_example.cpp")
+set(BIN "${CMAKE_CURRENT_BINARY_DIR}/fused_moe_i8_tn_example")
+
+add_custom_command(
+  OUTPUT "${BIN}"
+  COMMAND "${MXCC}"
+          -std=c++17
+          -xmaca
+          -I"${EXAMPLE_ROOT}/src"
+          -I"${MACA_PATH}/include"
+          "${SRC}"
+          -L"${MACA_PATH}/lib"
+          -lmcruntime
+          -o "${BIN}"
+  DEPENDS "${SRC}"
+  WORKING_DIRECTORY "${EXAMPLE_ROOT}"
+  VERBATIM)
+
+add_custom_target(build_fused_moe_i8_tn_example ALL DEPENDS "${BIN}")
+
+add_custom_target(
+  run
+  COMMAND "${BIN}"
+  DEPENDS "${BIN}"
+  USES_TERMINAL)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/Makefile

@@ -0,0 +1,27 @@
+MACA_PATH ?= /opt/maca
+MXCC := $(MACA_PATH)/mxgpu_llvm/bin/mxcc
+BUILD_DIR := $(CURDIR)/build
+SRC := $(CURDIR)/src/fused_moe_i8_tn_example.cpp
+BIN := $(BUILD_DIR)/fused_moe_i8_tn_example
+
+.PHONY: all build run clean
+
+all: build
+
+build: $(BIN)
+
+$(BIN): $(SRC)
+	mkdir -p $(BUILD_DIR)
+	$(MXCC) -std=c++17 -xmaca \
+		-I$(CURDIR)/src \
+		-I$(MACA_PATH)/include \
+		$(SRC) \
+		-L$(MACA_PATH)/lib \
+		-lmcruntime \
+		-o $(BIN)
+
+run: $(BIN)
+	$(BIN)
+
+clean:
+	rm -rf $(BUILD_DIR)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/benchmark_results.md

@@ -0,0 +1,30 @@
+# Fused MoE i8 TN Benchmark Results
+
+Remote run environment:
+
+- host: `10.2.118.21`
+- repo: `/home/acl_dnn/mcTlass`
+- binary: `standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_example`
+
+Benchmark command used:
+
+```bash
+MCTLASS_MOE_WARMUP=2 MCTLASS_MOE_ITERS=5 ./standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_example
+```
+
+Results:
+
+```text
+Benchmark config: warmup=2, iters=5
+fused_moe_i8_tn_topk1 passed: rows=256, topk=1, N=128, K=128, sample C[0]=0.695312, C[last]=-0.445312, max_abs=0
+fused_moe_i8_tn_topk1 benchmark: avg_ms=0.0137216, TOPS=0.611343, warmup=2, iters=5
+fused_moe_i8_tn_topk2 passed: rows=512, topk=2, N=128, K=128, sample C[0]=-0.578125, C[last]=-0.498047, max_abs=0
+fused_moe_i8_tn_topk2 benchmark: avg_ms=0.0121856, TOPS=1.37681, warmup=2, iters=5
+fused_moe_i8_tn_topk3 passed: rows=384, topk=3, N=128, K=128, sample C[0]=-1.08594, C[last]=-0.335938, max_abs=0
+fused_moe_i8_tn_topk3 benchmark: avg_ms=0.0124416, TOPS=1.01136, warmup=2, iters=5
+```
+
+Notes:
+
+- current numbers are from a short sanity benchmark, not a long stabilized run
+- default code path still uses `warmup=20` and `iters=100` if env vars are not set

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/fused-moe-i8-tn-pybind-triton-guide.md

@@ -0,0 +1,214 @@
+# fused_moe_i8_tn Python / Triton / Pybind 使用说明
+
+## 1. 目标
+
+当前仓库已经为 `standalone/fused_moe_i8_tn` 提供了三套可用于 Python 层验证与对比的实现：
+
+1. `pybind`
+   调用 MACA / MCTLASS C++ kernel，通过 Python 扩展模块暴露给 Python。
+2. `triton`
+   使用 Triton 实现的 `fused moe` 路径，结构参考 vLLM 的 `fused_moe_kernel`。
+3. `reference`
+   纯 Python / NumPy 参考实现，用于结果校验，不用于性能。
+
+这三套后端都已经接入统一测试和统一 benchmark 脚本，可以直接按后端切换。
+
+## 2. 相关文件
+
+### 2.1 C++ / Pybind
+
+- [standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_runner.h](/home/zguo/mcTlass/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_runner.h)
+- [standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_pybind.cpp](/home/zguo/mcTlass/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_pybind.cpp)
+- [scripts/build_fused_moe_i8_tn_pybind.sh](/home/zguo/mcTlass/scripts/build_fused_moe_i8_tn_pybind.sh)
+
+### 2.2 Triton
+
+- [standalone/fused_moe_i8_tn/python/fused_moe_i8_tn_triton.py](/home/zguo/mcTlass/standalone/fused_moe_i8_tn/python/fused_moe_i8_tn_triton.py)
+
+### 2.3 测试与性能
+
+- [standalone/fused_moe_i8_tn/python/test_fused_moe_i8_tn_pybind.py](/home/zguo/mcTlass/standalone/fused_moe_i8_tn/python/test_fused_moe_i8_tn_pybind.py)
+- [standalone/fused_moe_i8_tn/python/benchmark_fused_moe_i8_tn.py](/home/zguo/mcTlass/standalone/fused_moe_i8_tn/python/benchmark_fused_moe_i8_tn.py)
+- [scripts/run_fused_moe_i8_tn_pybind_test.sh](/home/zguo/mcTlass/scripts/run_fused_moe_i8_tn_pybind_test.sh)
+- [scripts/run_fused_moe_i8_tn_benchmark.sh](/home/zguo/mcTlass/scripts/run_fused_moe_i8_tn_benchmark.sh)
+
+## 3. 默认远端环境
+
+当前脚本已经默认适配远端环境：
+
+```bash
+PYTHON_BIN=/home/wtliu/miniforge3/envs/py310/bin/python
+MACA_PATH=/opt/maca-20260318
+LD_LIBRARY_PATH=$MACA_PATH/mxgpu_llvm/lib:$MACA_PATH/lib:$LD_LIBRARY_PATH
+```
+
+其中：
+
+1. `py310` 环境中已确认存在：
+   - `torch`
+   - `triton`
+   - `numpy`
+2. `torch` 实际版本为：
+   - `2.8.0+metax3.6.0.5`
+3. `triton` 实际版本为：
+   - `3.0.0`
+
+如果你需要切换 Python 环境，也可以在命令前覆盖：
+
+```bash
+PYTHON_BIN=/path/to/python bash scripts/run_fused_moe_i8_tn_pybind_test.sh --backend triton
+```
+
+## 4. Pybind 编译
+
+远端进入仓库后，编译 `pybind` 模块：
+
+```bash
+cd /home/acl_dnn/mcTlass
+bash scripts/build_fused_moe_i8_tn_pybind.sh
+```
+
+编译成功后，会生成：
+
+```bash
+standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_pybind.cpython-310-x86_64-linux-gnu.so
+```
+
+## 5. 正确性测试
+
+### 5.1 只测 pybind
+
+```bash
+cd /home/acl_dnn/mcTlass
+bash scripts/run_fused_moe_i8_tn_pybind_test.sh --backend pybind
+```
+
+### 5.2 只测 triton
+
+```bash
+cd /home/acl_dnn/mcTlass
+bash scripts/run_fused_moe_i8_tn_pybind_test.sh --backend triton
+```
+
+### 5.3 跑全部后端
+
+```bash
+cd /home/acl_dnn/mcTlass
+bash scripts/run_fused_moe_i8_tn_pybind_test.sh --backend all
+```
+
+支持的后端选项：
+
+- `pybind`
+- `triton`
+- `reference`
+- `all`
+
+## 6. 性能测试
+
+### 6.1 只测 triton
+
+```bash
+cd /home/acl_dnn/mcTlass
+bash scripts/run_fused_moe_i8_tn_benchmark.sh --backend triton --warmup 5 --iters 20
+```
+
+### 6.2 跑全部后端
+
+```bash
+cd /home/acl_dnn/mcTlass
+bash scripts/run_fused_moe_i8_tn_benchmark.sh --backend all --warmup 5 --iters 20
+```
+
+参数说明：
+
+- `--warmup`
+  预热次数
+- `--iters`
+  正式计时次数
+
+## 7. 当前 Triton 实现说明
+
+当前 Triton 后端不是最初那版“逐行加载、逐行计算”的简化实现，而是已经改成参考 vLLM `fused_moe_kernel` 的结构。
+
+核心特征包括：
+
+1. 使用 `sorted routed rows`
+2. 使用 `block expert ids`
+3. 使用 `num_tokens_post_padded`
+4. 使用 grouped `pid_m / pid_n` 调度方式
+5. 使用 `MUL_ROUTED_WEIGHT`
+6. 使用 `int8_w8a8 + per_channel_quant` 这条固定路径
+
+当前为了适配本仓库现有 `fused_moe_i8_tn` 输入定义，做了以下简化：
+
+1. `HAS_BIAS=False`
+2. `use_int8_w8a8=True`
+3. `use_fp8_w8a8=False`
+4. `use_int8_w8a16=False`
+5. `group_k=0, group_n=0`
+6. `naive_block_assignment=False`
+
+因此它的目的目前是：
+
+1. 让 Triton 路径和当前 `fused_moe_i8_tn` 测试数据对齐
+2. 保持与 vLLM fused moe kernel 结构尽量接近
+3. 在 MetaX + Triton 环境中可实际运行
+
+## 8. 当前对比结论
+
+### 8.1 正确性
+
+远端实际验证结果表明：
+
+1. `triton` 和 `reference` 的输出一致
+2. `pybind` 也通过校验，但由于其输出走 BF16 路径，和 `reference/triton` 相比会存在 BF16 舍入差
+
+这属于当前实现预期行为，不是错误。
+
+### 8.2 性能
+
+远端使用：
+
+```bash
+bash scripts/run_fused_moe_i8_tn_benchmark.sh --backend all --warmup 5 --iters 20
+```
+
+得到的结果为：
+
+```text
+pybind:fused_moe_i8_tn_topk1 benchmark: avg_ms=0.109936, TOPS=0.076305
+pybind:fused_moe_i8_tn_topk2 benchmark: avg_ms=0.126724, TOPS=0.132392
+pybind:fused_moe_i8_tn_topk3 benchmark: avg_ms=0.113708, TOPS=0.110660
+
+triton:fused_moe_i8_tn_topk1 benchmark: avg_ms=9.841149, TOPS=0.000852
+triton:fused_moe_i8_tn_topk2 benchmark: avg_ms=9.924612, TOPS=0.001690
+triton:fused_moe_i8_tn_topk3 benchmark: avg_ms=9.938139, TOPS=0.001266
+
+reference:fused_moe_i8_tn_topk1 benchmark: avg_ms=717.996218, TOPS=0.000012
+reference:fused_moe_i8_tn_topk2 benchmark: avg_ms=1437.042784, TOPS=0.000012
+reference:fused_moe_i8_tn_topk3 benchmark: avg_ms=1077.384005, TOPS=0.000012
+```
+
+当前结论：
+
+1. `pybind` 最快
+2. `triton` 明显快于 `reference`
+3. `triton` 仍显著慢于 `pybind`
+
+也就是说，当前 Triton 路径已经具备：
+
+1. 结构正确
+2. 数值正确
+3. 能在远端真实运行
+
+但它还不是性能优化完成版。
+
+## 9. 推荐后续工作
+
+如果后续继续优化，优先建议做：
+
+1. 分析 MetaX Triton backend 的 `tl.dot(int8, int8)` lowering 是否真正命中高效硬件路径
+2. 调整 `BLOCK_SIZE_K / BLOCK_SIZE_N / GROUP_SIZE_M`
+3. 进一步减少 Python 侧 routing / packing 的额外开销
+4. 对齐 C++ kernel 的 tile 组织方式，逐步缩小 `pybind` 与 `triton` 的性能差距

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/benchmark_fused_moe_i8_tn.py

@@ -0,0 +1,63 @@
+import argparse
+import time
+
+from test_fused_moe_i8_tn_pybind import fill_inputs, resolve_backends
+
+
+K_N = 128
+K_K = 128
+
+
+def compute_tops(rows: int, cols: int, k_dim: int, avg_ms: float) -> float:
+    if avg_ms <= 0.0:
+        return 0.0
+    operations = 2.0 * float(rows) * float(cols) * float(k_dim)
+    return operations / (avg_ms * 1.0e9)
+
+
+def benchmark_backend_case(backend: str, backend_fn, tag: str, num_tokens: int, topk: int, em: int, tile_experts: list[int], warmup: int, iters: int):
+    a, b, scale_a, scale_b, moe_weights, token_ids, expert_ids = fill_inputs(num_tokens, topk, tile_experts)
+
+    for _ in range(warmup):
+        backend_fn(a, b, scale_a, scale_b, moe_weights, token_ids, expert_ids, topk)
+
+    start = time.perf_counter()
+    for _ in range(iters):
+        backend_fn(a, b, scale_a, scale_b, moe_weights, token_ids, expert_ids, topk)
+    elapsed_s = time.perf_counter() - start
+
+    avg_ms = elapsed_s * 1000.0 / iters
+    tops = compute_tops(em, K_N, K_K, avg_ms)
+    print(
+        f"{backend}:{tag} benchmark: avg_ms={avg_ms:.6f}, "
+        f"TOPS={tops:.6f}, warmup={warmup}, iters={iters}"
+    )
+
+
+def parse_args():
+    parser = argparse.ArgumentParser()
+    parser.add_argument(
+        "--backend",
+        choices=("pybind", "triton", "reference", "all"),
+        default="pybind",
+    )
+    parser.add_argument("--warmup", type=int, default=5)
+    parser.add_argument("--iters", type=int, default=20)
+    return parser.parse_args()
+
+
+def main():
+    args = parse_args()
+    cases = [
+        ("fused_moe_i8_tn_topk1", 256, 1, 256, [0, 1]),
+        ("fused_moe_i8_tn_topk2", 256, 2, 512, [0, 1, 1, 0]),
+        ("fused_moe_i8_tn_topk3", 128, 3, 384, [0, 1, 0]),
+    ]
+    backends = resolve_backends(args.backend)
+    for backend, backend_fn in backends:
+        for case in cases:
+            benchmark_backend_case(backend, backend_fn, *case, warmup=args.warmup, iters=args.iters)
+
+
+if __name__ == "__main__":
+    main()

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/fused_moe_i8_tn_triton.py

@@ -0,0 +1,368 @@
+from __future__ import annotations
+
+from typing import Any
+
+import numpy as np
+
+
+K_TILE_M = 128
+BLOCK_SIZE_M = 128
+BLOCK_SIZE_N = 128
+BLOCK_SIZE_K = 32
+GROUP_SIZE_M = 8
+
+_TRITON = None
+_TL = None
+triton = None
+tl = None
+
+
+def _require_triton_runtime() -> tuple[Any, Any, Any]:
+    global _TRITON, _TL, triton, tl
+    try:
+        import torch
+        import triton as triton_mod
+        import triton.language as tl_mod
+    except ImportError as exc:
+        raise RuntimeError("torch and triton are required for the Triton backend") from exc
+
+    if not torch.cuda.is_available():
+        raise RuntimeError("the Triton backend currently requires a CUDA-capable torch runtime")
+
+    _TRITON = triton_mod
+    _TL = tl_mod
+    triton = triton_mod
+    tl = tl_mod
+    return torch, triton_mod, tl_mod
+
+
+def triton_backend_available() -> tuple[bool, str]:
+    try:
+        _require_triton_runtime()
+    except RuntimeError as exc:
+        return False, str(exc)
+    return True, ""
+
+
+def _build_blocked_routing(token_ids: np.ndarray, expert_ids: np.ndarray, num_experts: int, block_size_m: int):
+    total_rows = int(token_ids.shape[0])
+    routed_rows_per_expert = [[] for _ in range(num_experts)]
+
+    for routed_row in range(total_rows):
+        tile_idx = routed_row // K_TILE_M
+        expert = int(expert_ids[tile_idx])
+        routed_rows_per_expert[expert].append(routed_row)
+
+    sorted_routed_rows: list[int] = []
+    block_expert_ids: list[int] = []
+    invalid_row = total_rows
+
+    for expert, rows in enumerate(routed_rows_per_expert):
+        if not rows:
+            continue
+        sorted_routed_rows.extend(rows)
+        padded = (-len(rows)) % block_size_m
+        if padded:
+            sorted_routed_rows.extend([invalid_row] * padded)
+        block_count = (len(rows) + padded) // block_size_m
+        block_expert_ids.extend([expert] * block_count)
+
+    num_tokens_post_padded = len(sorted_routed_rows)
+    return (
+        np.asarray(sorted_routed_rows, dtype=np.int32),
+        np.asarray(block_expert_ids, dtype=np.int32),
+        np.asarray([num_tokens_post_padded], dtype=np.int32),
+    )
+
+
+def _ensure_triton_symbols():
+    if _TRITON is None or _TL is None:
+        _require_triton_runtime()
+    return _TRITON, _TL
+
+
+def _get_fused_moe_kernel():
+    triton, tl = _ensure_triton_symbols()
+
+    @triton.jit
+    def _fused_moe_kernel(
+        a_ptr,
+        b_ptr,
+        c_ptr,
+        b_bias_ptr,
+        scale_a_ptr,
+        scale_b_ptr,
+        moe_weights_ptr,
+        sorted_routed_rows_ptr,
+        block_expert_ids_ptr,
+        num_tokens_post_padded_ptr,
+        n_dim,
+        k_dim,
+        em,
+        num_valid_tokens,
+        stride_am,
+        stride_ak,
+        stride_be,
+        stride_bk,
+        stride_bn,
+        stride_cm,
+        stride_cn,
+        stride_asm,
+        stride_ask,
+        stride_bse,
+        stride_bsk,
+        stride_bsn,
+        stride_bbe,
+        stride_bbn,
+        group_n: tl.constexpr,
+        group_k: tl.constexpr,
+        naive_block_assignment: tl.constexpr,
+        BLOCK_SIZE_M: tl.constexpr,
+        BLOCK_SIZE_N: tl.constexpr,
+        BLOCK_SIZE_K: tl.constexpr,
+        GROUP_SIZE_M: tl.constexpr,
+        SPLIT_K: tl.constexpr,
+        MUL_ROUTED_WEIGHT: tl.constexpr,
+        top_k: tl.constexpr,
+        compute_type: tl.constexpr,
+        use_fp8_w8a8: tl.constexpr,
+        use_int8_w8a8: tl.constexpr,
+        use_int8_w8a16: tl.constexpr,
+        per_channel_quant: tl.constexpr,
+        HAS_BIAS: tl.constexpr,
+    ):
+        pid = tl.program_id(axis=0)
+        num_pid_m = tl.cdiv(em, BLOCK_SIZE_M)
+        num_pid_n = tl.cdiv(n_dim, BLOCK_SIZE_N)
+        num_pid_in_group = GROUP_SIZE_M * num_pid_n
+        group_id = pid // num_pid_in_group
+        first_pid_m = group_id * GROUP_SIZE_M
+        group_size_m = tl.minimum(num_pid_m - first_pid_m, GROUP_SIZE_M)
+        pid_m = first_pid_m + ((pid % num_pid_in_group) % group_size_m)
+        pid_n = (pid % num_pid_in_group) // group_size_m
+
+        offs = tl.arange(0, BLOCK_SIZE_M).to(tl.int64)
+        num_tokens_post_padded = tl.load(num_tokens_post_padded_ptr)
+        if pid_m * BLOCK_SIZE_M >= num_tokens_post_padded:
+            return
+
+        if not naive_block_assignment:
+            offs_token_id = pid_m * BLOCK_SIZE_M + offs
+            offs_token = tl.load(sorted_routed_rows_ptr + offs_token_id)
+        else:
+            offs_token = tl.where(
+                offs == 0,
+                pid_m,
+                num_valid_tokens,
+            )
+
+        offs_token = offs_token.to(tl.int64)
+        token_mask = offs_token < num_valid_tokens
+
+        off_experts = tl.load(block_expert_ids_ptr + pid_m).to(tl.int64)
+        if off_experts == -1:
+            zero_acc = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=compute_type)
+            zero_offs_cn = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
+            zero_c_ptrs = c_ptr + stride_cm * offs_token[:, None] + stride_cn * zero_offs_cn[None, :]
+            zero_c_mask = token_mask[:, None] & (zero_offs_cn[None, :] < n_dim)
+            tl.store(zero_c_ptrs, zero_acc, mask=zero_c_mask)
+            return
+
+        offs_bn = (pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N).to(tl.int64)) % n_dim
+        offs_k = tl.arange(0, BLOCK_SIZE_K)
+        a_ptrs = a_ptr + (offs_token[:, None] // top_k * stride_am + offs_k[None, :] * stride_ak)
+        b_ptrs = b_ptr + off_experts * stride_be + (offs_k[:, None] * stride_bk + offs_bn[None, :] * stride_bn)
+
+        if use_int8_w8a16:
+            b_scale_ptrs = scale_b_ptr + off_experts * stride_bse + offs_bn[None, :] * stride_bsn
+            b_scale = tl.load(b_scale_ptrs)
+
+        if use_fp8_w8a8 or use_int8_w8a8:
+            if group_k > 0 and group_n > 0:
+                a_scale_ptrs = scale_a_ptr + (offs_token // top_k) * stride_asm
+                offs_bsn = offs_bn // group_n
+                b_scale_ptrs = scale_b_ptr + off_experts * stride_bse + offs_bsn * stride_bsn
+            elif per_channel_quant:
+                b_scale_ptrs = scale_b_ptr + off_experts * stride_bse + offs_bn[None, :] * stride_bsn
+                b_scale = tl.load(b_scale_ptrs)
+                a_scale_ptrs = scale_a_ptr + (offs_token // top_k) * stride_asm
+                a_scale = tl.load(a_scale_ptrs, mask=token_mask, other=0.0)[:, None]
+            else:
+                a_scale = tl.load(scale_a_ptr)
+                b_scale = tl.load(scale_b_ptr + off_experts)
+
+        if HAS_BIAS:
+            bias_ptrs = b_bias_ptr + off_experts * stride_bbe + offs_bn * stride_bbn
+            bias = tl.load(bias_ptrs, mask=(offs_bn < n_dim), other=0.0)
+
+        accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
+        for k in range(0, tl.cdiv(k_dim, BLOCK_SIZE_K)):
+            a = tl.load(
+                a_ptrs,
+                mask=token_mask[:, None] & (offs_k[None, :] < k_dim - k * BLOCK_SIZE_K),
+                other=0.0,
+            )
+            b = tl.load(b_ptrs, mask=offs_k[:, None] < k_dim - k * BLOCK_SIZE_K, other=0.0)
+            if use_int8_w8a16:
+                accumulator = tl.dot(a, b.to(compute_type), acc=accumulator)
+            elif use_fp8_w8a8 or use_int8_w8a8:
+                if group_k > 0 and group_n > 0:
+                    k_start = k * BLOCK_SIZE_K
+                    offs_ks = k_start // group_k
+                    a_scale = tl.load(a_scale_ptrs + offs_ks * stride_ask, mask=token_mask, other=0.0)
+                    b_scale = tl.load(b_scale_ptrs + offs_ks * stride_bsk)
+                    accumulator += tl.dot(a, b) * a_scale[:, None] * b_scale[None, :]
+                else:
+                    if use_fp8_w8a8:
+                        accumulator = tl.dot(a, b, acc=accumulator)
+                    else:
+                        accumulator += tl.dot(a, b)
+            else:
+                accumulator += tl.dot(a, b)
+
+            a_ptrs += BLOCK_SIZE_K * stride_ak
+            b_ptrs += BLOCK_SIZE_K * stride_bk
+
+        if use_int8_w8a16:
+            accumulator = accumulator * b_scale
+        elif (use_fp8_w8a8 or use_int8_w8a8) and not (group_k > 0 and group_n > 0):
+            accumulator = accumulator * a_scale * b_scale
+
+        if HAS_BIAS:
+            accumulator += bias[None, :]
+
+        if MUL_ROUTED_WEIGHT:
+            moe_weight = tl.load(
+                moe_weights_ptr + offs_token,
+                mask=token_mask,
+                other=0,
+            )
+            accumulator *= moe_weight[:, None]
+
+        accumulator = accumulator.to(compute_type)
+
+        offs_cn = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
+        c_ptrs = c_ptr + stride_cm * offs_token[:, None] + stride_cn * offs_cn[None, :]
+        c_mask = token_mask[:, None] & (offs_cn[None, :] < n_dim)
+        tl.store(c_ptrs, accumulator, mask=c_mask)
+
+    return _fused_moe_kernel
+
+
+def run_fused_moe_i8_tn_triton(
+    a: np.ndarray,
+    b_col_major: np.ndarray,
+    scale_a: np.ndarray,
+    scale_b: np.ndarray,
+    moe_weights: np.ndarray,
+    token_ids: np.ndarray,
+    expert_ids: np.ndarray,
+    topk: int,
+    device: str = "cuda",
+) -> np.ndarray:
+    # Adapted from vLLM's fused MoE Triton path:
+    # https://github.com/vllm-project/vllm/blob/main/vllm/model_executor/layers/fused_moe/fused_moe.py
+    torch, triton, tl = _require_triton_runtime()
+
+    if a.ndim != 2:
+        raise ValueError("a must be a 2D array")
+    if b_col_major.ndim != 3:
+        raise ValueError("b_col_major must be a 3D array")
+    if scale_a.ndim != 1:
+        raise ValueError("scale_a must be a 1D array")
+    if scale_b.ndim != 2:
+        raise ValueError("scale_b must be a 2D array")
+    if moe_weights.ndim != 1:
+        raise ValueError("moe_weights must be a 1D array")
+    if token_ids.ndim != 1:
+        raise ValueError("token_ids must be a 1D array")
+    if expert_ids.ndim != 1:
+        raise ValueError("expert_ids must be a 1D array")
+    if topk <= 0:
+        raise ValueError("topk must be > 0")
+
+    num_tokens, k_dim = a.shape
+    num_experts, n_dim, b_k = b_col_major.shape
+    total_rows = moe_weights.shape[0]
+
+    if b_k != k_dim:
+        raise ValueError("B K dimension must match A K dimension")
+    if scale_a.shape[0] != num_tokens:
+        raise ValueError("scale_a size mismatch")
+    if scale_b.shape != (num_experts, n_dim):
+        raise ValueError("scale_b shape mismatch")
+    if token_ids.shape[0] != total_rows:
+        raise ValueError("token_ids size mismatch")
+    if total_rows != num_tokens * topk:
+        raise ValueError("moe_weights size must equal num_tokens * topk")
+    if total_rows % K_TILE_M != 0:
+        raise ValueError("num_tokens * topk must be a multiple of 128")
+    if expert_ids.shape[0] != total_rows // K_TILE_M:
+        raise ValueError("expert_ids size mismatch")
+
+    sorted_routed_rows, block_expert_ids, num_tokens_post_padded = _build_blocked_routing(
+        token_ids, expert_ids, num_experts, BLOCK_SIZE_M
+    )
+
+    a_t = torch.as_tensor(np.ascontiguousarray(a), device=device, dtype=torch.int8)
+    b_t = torch.as_tensor(np.ascontiguousarray(b_col_major), device=device, dtype=torch.int8)
+    scale_a_t = torch.as_tensor(np.ascontiguousarray(scale_a), device=device, dtype=torch.float32)
+    scale_b_t = torch.as_tensor(np.ascontiguousarray(scale_b), device=device, dtype=torch.float32)
+    moe_weights_t = torch.as_tensor(np.ascontiguousarray(moe_weights), device=device, dtype=torch.float32)
+    sorted_routed_rows_t = torch.as_tensor(sorted_routed_rows, device=device, dtype=torch.int32)
+    block_expert_ids_t = torch.as_tensor(block_expert_ids, device=device, dtype=torch.int32)
+    num_tokens_post_padded_t = torch.as_tensor(num_tokens_post_padded, device=device, dtype=torch.int32)
+    dummy_bias_t = torch.zeros((num_experts, n_dim), device=device, dtype=torch.float32)
+    out_t = torch.empty((total_rows, n_dim), device=device, dtype=torch.float32)
+
+    fused_moe_kernel = _get_fused_moe_kernel()
+
+    grid = (triton.cdiv(int(num_tokens_post_padded[0]), BLOCK_SIZE_M) * triton.cdiv(n_dim, BLOCK_SIZE_N),)
+    fused_moe_kernel[grid](
+        a_t,
+        b_t,
+        out_t,
+        dummy_bias_t,
+        scale_a_t,
+        scale_b_t,
+        moe_weights_t,
+        sorted_routed_rows_t,
+        block_expert_ids_t,
+        num_tokens_post_padded_t,
+        n_dim,
+        k_dim,
+        int(num_tokens_post_padded[0]),
+        total_rows,
+        a_t.stride(0),
+        a_t.stride(1),
+        b_t.stride(0),
+        b_t.stride(2),
+        b_t.stride(1),
+        out_t.stride(0),
+        out_t.stride(1),
+        scale_a_t.stride(0),
+        0,
+        scale_b_t.stride(0),
+        0,
+        scale_b_t.stride(1),
+        dummy_bias_t.stride(0),
+        dummy_bias_t.stride(1),
+        group_n=0,
+        group_k=0,
+        naive_block_assignment=False,
+        BLOCK_SIZE_M=BLOCK_SIZE_M,
+        BLOCK_SIZE_N=BLOCK_SIZE_N,
+        BLOCK_SIZE_K=BLOCK_SIZE_K,
+        GROUP_SIZE_M=GROUP_SIZE_M,
+        SPLIT_K=1,
+        MUL_ROUTED_WEIGHT=True,
+        top_k=topk,
+        compute_type=tl.float32,
+        use_fp8_w8a8=False,
+        use_int8_w8a8=True,
+        use_int8_w8a16=False,
+        per_channel_quant=True,
+        HAS_BIAS=False,
+    )
+
+    return out_t.cpu().numpy()

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/test_fused_moe_i8_tn_pybind.py

@@ -0,0 +1,154 @@
+import argparse
+import importlib.util
+import os
+from pathlib import Path
+
+import numpy as np
+
+from fused_moe_i8_tn_triton import run_fused_moe_i8_tn_triton, triton_backend_available
+
+
+K_NUM_EXPERTS = 2
+K_TILE_M = 128
+K_N = 128
+K_K = 128
+
+
+def load_extension():
+    build_dir = Path(__file__).resolve().parents[1] / "build"
+    candidates = sorted(build_dir.glob("fused_moe_i8_tn_pybind*.so"))
+    if not candidates:
+        raise FileNotFoundError(f"no built extension found under {build_dir}")
+    module_path = candidates[0]
+    spec = importlib.util.spec_from_file_location("fused_moe_i8_tn_pybind", module_path)
+    module = importlib.util.module_from_spec(spec)
+    assert spec.loader is not None
+    spec.loader.exec_module(module)
+    return module
+
+
+def fill_inputs(num_tokens: int, topk: int, tile_experts: list[int]):
+    total_rows = num_tokens * topk
+
+    a = np.empty((num_tokens, K_K), dtype=np.int8)
+    scale_a = np.empty((num_tokens,), dtype=np.float32)
+    moe_weights = np.empty((total_rows,), dtype=np.float32)
+    token_ids = np.empty((total_rows,), dtype=np.int32)
+    expert_ids = np.asarray(tile_experts, dtype=np.int32)
+
+    for row in range(num_tokens):
+        for kk in range(K_K):
+            a[row, kk] = ((row * 13 + kk * 7 + topk * 5 + 3) % 11) - 5
+        scale_a[row] = 0.125 + 0.015625 * ((row + topk) % 7)
+
+    for routed_row in range(total_rows):
+        token_ids[routed_row] = routed_row
+        moe_weights[routed_row] = 0.5 + 0.03125 * ((routed_row + topk) % 5)
+
+    b = np.empty((K_NUM_EXPERTS, K_N, K_K), dtype=np.int8)
+    scale_b = np.empty((K_NUM_EXPERTS, K_N), dtype=np.float32)
+    for expert in range(K_NUM_EXPERTS):
+        for col in range(K_N):
+            scale_b[expert, col] = 0.25 + 0.03125 * ((expert * 3 + col + topk) % 9)
+            for kk in range(K_K):
+                b[expert, col, kk] = ((expert * 17 + col * 5 + kk * 3 + topk) % 9) - 4
+
+    return a, b, scale_a, scale_b, moe_weights, token_ids, expert_ids
+
+
+def reference_fused_moe(a, b, scale_a, scale_b, moe_weights, token_ids, expert_ids, topk):
+    total_rows = token_ids.shape[0]
+    out = np.zeros((total_rows, K_N), dtype=np.float32)
+    for routed_row in range(total_rows):
+        token = token_ids[routed_row] // topk
+        tile_idx = routed_row // K_TILE_M
+        expert = expert_ids[tile_idx]
+        row_scale = scale_a[token] * moe_weights[routed_row]
+        for col in range(K_N):
+            acc = 0
+            for kk in range(K_K):
+                acc += int(a[token, kk]) * int(b[expert, col, kk])
+            out[routed_row, col] = np.float32(acc * row_scale * scale_b[expert, col])
+    return out
+
+
+def run_pybind_backend(module, a, b, scale_a, scale_b, moe_weights, token_ids, expert_ids, topk):
+    return module.run_fused_moe_i8_tn(
+        a,
+        b,
+        scale_a,
+        scale_b.reshape(-1),
+        moe_weights,
+        token_ids,
+        expert_ids,
+        topk,
+        K_NUM_EXPERTS,
+        int(os.environ.get("MCTLASS_PY_DEVICE_ID", "0")),
+    )
+
+
+def run_reference_backend(a, b, scale_a, scale_b, moe_weights, token_ids, expert_ids, topk):
+    return reference_fused_moe(a, b, scale_a, scale_b, moe_weights, token_ids, expert_ids, topk)
+
+
+def run_case(backend: str, backend_fn, tag: str, num_tokens: int, topk: int, em: int, tile_experts: list[int]):
+    assert em == num_tokens * topk
+    a, b, scale_a, scale_b, moe_weights, token_ids, expert_ids = fill_inputs(num_tokens, topk, tile_experts)
+    expected = reference_fused_moe(a, b, scale_a, scale_b, moe_weights, token_ids, expert_ids, topk)
+    got = backend_fn(a, b, scale_a, scale_b, moe_weights, token_ids, expert_ids, topk)
+    np.testing.assert_allclose(got, expected, rtol=0.0, atol=1e-2)
+    print(
+        f"{backend}:{tag} passed: rows={got.shape[0]}, cols={got.shape[1]}, "
+        f"sample C[0]={got.reshape(-1)[0]}, C[last]={got.reshape(-1)[-1]}"
+    )
+
+
+def resolve_backends(requested_backend: str):
+    backends = []
+
+    if requested_backend in {"pybind", "all"}:
+        module = load_extension()
+        backends.append(("pybind", lambda a, b, scale_a, scale_b, moe_weights, token_ids, expert_ids, topk: run_pybind_backend(
+            module, a, b, scale_a, scale_b, moe_weights, token_ids, expert_ids, topk
+        )))
+
+    if requested_backend in {"reference", "all"}:
+        backends.append(("reference", run_reference_backend))
+
+    if requested_backend in {"triton", "all"}:
+        available, reason = triton_backend_available()
+        if available:
+            backends.append(("triton", run_fused_moe_i8_tn_triton))
+        elif requested_backend == "triton":
+            raise RuntimeError(f"Triton backend is unavailable: {reason}")
+        else:
+            print(f"skip triton backend: {reason}")
+
+    return backends
+
+
+def parse_args():
+    parser = argparse.ArgumentParser()
+    parser.add_argument(
+        "--backend",
+        choices=("pybind", "triton", "reference", "all"),
+        default=os.environ.get("MCTLASS_FUSED_MOE_BACKEND", "pybind"),
+    )
+    return parser.parse_args()
+
+
+def main():
+    args = parse_args()
+    cases = [
+        ("fused_moe_i8_tn_topk1", 256, 1, 256, [0, 1]),
+        ("fused_moe_i8_tn_topk2", 256, 2, 512, [0, 1, 1, 0]),
+        ("fused_moe_i8_tn_topk3", 128, 3, 384, [0, 1, 0]),
+    ]
+    backends = resolve_backends(args.backend)
+    for backend, backend_fn in backends:
+        for case in cases:
+            run_case(backend, backend_fn, *case)
+
+
+if __name__ == "__main__":
+    main()

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_example.cpp

@@ -0,0 +1,324 @@
+#include <mc_runtime.h>
+
+#include <cmath>
+#include <cstdint>
+#include <cstdlib>
+#include <iostream>
+#include <string>
+#include <vector>
+
+#include "fused_moe_i8_tn_runner.h"
+
+namespace {
+
+constexpr int kNumExperts = fused_moe_i8_tn::kDefaultNumExperts;
+constexpr int kTileM = fused_moe_i8_tn::kDefaultTileM;
+constexpr int kN = fused_moe_i8_tn::kDefaultN;
+constexpr int kK = fused_moe_i8_tn::kDefaultK;
+constexpr int kDefaultWarmupIterations = 20;
+constexpr int kDefaultMeasuredIterations = 100;
+
+struct CaseConfig {
+    const char *tag;
+    int num_tokens;
+    int topk;
+    int em;
+    std::vector<int> tile_experts;
+};
+
+struct BenchmarkConfig {
+    int warmup_iterations = 0;
+    int measured_iterations = 0;
+};
+
+struct BenchmarkResult {
+    float avg_ms = 0.0f;
+    double tops = 0.0;
+    int warmup_iterations = 0;
+    int measured_iterations = 0;
+};
+
+void check_mc(mcError_t status, const char *expr) {
+    if (status != mcSuccess) {
+        std::cerr << expr << " failed: " << mcGetErrorString(status) << '\n';
+        std::exit(EXIT_FAILURE);
+    }
+}
+
+int read_env_int(const char *name, int default_value) {
+    const char *value = std::getenv(name);
+    if (!value || !value[0]) {
+        return default_value;
+    }
+
+    char *end = nullptr;
+    long parsed = std::strtol(value, &end, 10);
+    if (end == value || (end && *end != '\0') || parsed < 0) {
+        std::cerr << "Ignoring invalid " << name << '=' << value
+                  << ", using default " << default_value << '\n';
+        return default_value;
+    }
+    return static_cast<int>(parsed);
+}
+
+BenchmarkConfig make_benchmark_config() {
+    BenchmarkConfig config;
+    config.warmup_iterations = read_env_int("MCTLASS_MOE_WARMUP", kDefaultWarmupIterations);
+    config.measured_iterations = read_env_int("MCTLASS_MOE_ITERS", kDefaultMeasuredIterations);
+    if (config.measured_iterations <= 0) {
+        std::cerr << "MCTLASS_MOE_ITERS must be > 0\n";
+        std::exit(EXIT_FAILURE);
+    }
+    return config;
+}
+
+double compute_tops(int m, int n, int k, float avg_ms) {
+    if (avg_ms <= 0.0f) {
+        return 0.0;
+    }
+    const double operations = 2.0 * static_cast<double>(m) * static_cast<double>(n) * static_cast<double>(k);
+    return operations / (static_cast<double>(avg_ms) * 1.0e9);
+}
+
+void fill_inputs(const CaseConfig &cfg,
+                 std::vector<int8_t> &a,
+                 std::vector<int8_t> &b_col_major,
+                 std::vector<float> &scale_a,
+                 std::vector<float> &scale_b,
+                 std::vector<float> &moe_weights,
+                 std::vector<int> &token_ids,
+                 std::vector<int> &expert_ids) {
+    const int total_rows = cfg.num_tokens * cfg.topk;
+
+    a.resize(static_cast<size_t>(cfg.num_tokens) * kK);
+    scale_a.resize(cfg.num_tokens);
+    moe_weights.resize(total_rows);
+    token_ids.resize(total_rows);
+    expert_ids = cfg.tile_experts;
+
+    for (int row = 0; row < cfg.num_tokens; ++row) {
+        for (int kk = 0; kk < kK; ++kk) {
+            a[static_cast<size_t>(row) * kK + kk] =
+                static_cast<int8_t>(((row * 13 + kk * 7 + cfg.topk * 5 + 3) % 11) - 5);
+        }
+        scale_a[row] = 0.125f + 0.015625f * static_cast<float>((row + cfg.topk) % 7);
+    }
+
+    for (int routed_row = 0; routed_row < total_rows; ++routed_row) {
+        token_ids[routed_row] = routed_row;
+        moe_weights[routed_row] = 0.5f + 0.03125f * static_cast<float>((routed_row + cfg.topk) % 5);
+    }
+
+    b_col_major.resize(static_cast<size_t>(kNumExperts) * kN * kK);
+    scale_b.resize(static_cast<size_t>(kNumExperts) * kN);
+    for (int expert = 0; expert < kNumExperts; ++expert) {
+        for (int col = 0; col < kN; ++col) {
+            scale_b[static_cast<size_t>(expert) * kN + col] =
+                0.25f + 0.03125f * static_cast<float>((expert * 3 + col + cfg.topk) % 9);
+            for (int kk = 0; kk < kK; ++kk) {
+                const int value = ((expert * 17 + col * 5 + kk * 3 + cfg.topk) % 9) - 4;
+                b_col_major[(static_cast<size_t>(expert) * kN + col) * kK + kk] =
+                    static_cast<int8_t>(value);
+            }
+        }
+    }
+}
+
+std::vector<fused_moe_i8_tn::BFloat16> reference_fused_moe(const CaseConfig &cfg,
+                                                           const std::vector<int8_t> &a,
+                                                           const std::vector<int8_t> &b_col_major,
+                                                           const std::vector<float> &scale_a,
+                                                           const std::vector<float> &scale_b,
+                                                           const std::vector<float> &moe_weights,
+                                                           const std::vector<int> &token_ids,
+                                                           const std::vector<int> &expert_ids) {
+    const int total_rows = cfg.num_tokens * cfg.topk;
+    std::vector<fused_moe_i8_tn::BFloat16> out(static_cast<size_t>(total_rows) * kN,
+                                               fused_moe_i8_tn::float_to_bf16(0.0f));
+
+    for (int routed_row = 0; routed_row < total_rows; ++routed_row) {
+        const int token = token_ids[routed_row] / cfg.topk;
+        const int tile_idx = routed_row / kTileM;
+        const int expert = expert_ids[tile_idx];
+        const float row_scale = scale_a[token] * moe_weights[routed_row];
+
+        for (int col = 0; col < kN; ++col) {
+            int32_t acc = 0;
+            for (int kk = 0; kk < kK; ++kk) {
+                const int32_t lhs = static_cast<int32_t>(a[static_cast<size_t>(token) * kK + kk]);
+                const int32_t rhs =
+                    static_cast<int32_t>(b_col_major[(static_cast<size_t>(expert) * kN + col) * kK + kk]);
+                acc += lhs * rhs;
+            }
+            const float scaled = static_cast<float>(acc) * row_scale *
+                                 scale_b[static_cast<size_t>(expert) * kN + col];
+            out[static_cast<size_t>(routed_row) * kN + col] = fused_moe_i8_tn::float_to_bf16(scaled);
+        }
+    }
+
+    return out;
+}
+
+bool validate_result(const std::vector<fused_moe_i8_tn::BFloat16> &got,
+                     const std::vector<fused_moe_i8_tn::BFloat16> &expected,
+                     const CaseConfig &cfg) {
+    size_t mismatch_count = 0;
+    size_t first_bad = 0;
+    float max_abs = 0.0f;
+    const int total_rows = cfg.num_tokens * cfg.topk;
+
+    for (size_t i = 0; i < got.size(); ++i) {
+        const float got_f = fused_moe_i8_tn::bf16_to_float(got[i]);
+        const float exp_f = fused_moe_i8_tn::bf16_to_float(expected[i]);
+        const float abs_err = std::fabs(got_f - exp_f);
+        max_abs = std::max(max_abs, abs_err);
+        if (abs_err > 1e-2f) {
+            if (mismatch_count == 0) {
+                first_bad = i;
+            }
+            ++mismatch_count;
+        }
+    }
+
+    if (mismatch_count != 0) {
+        const int row = static_cast<int>(first_bad / kN);
+        const int col = static_cast<int>(first_bad % kN);
+        std::cerr << cfg.tag << " failed"
+                  << ": mismatches=" << mismatch_count
+                  << ", first mismatch at (" << row << ", " << col << ")"
+                  << ", got=" << fused_moe_i8_tn::bf16_to_float(got[first_bad])
+                  << ", expected=" << fused_moe_i8_tn::bf16_to_float(expected[first_bad])
+                  << ", max_abs=" << max_abs << '\n';
+        return false;
+    }
+
+    std::cout << cfg.tag
+              << " passed"
+              << ": rows=" << total_rows
+              << ", topk=" << cfg.topk
+              << ", N=" << kN
+              << ", K=" << kK
+              << ", sample C[0]=" << fused_moe_i8_tn::bf16_to_float(got.front())
+              << ", C[last]=" << fused_moe_i8_tn::bf16_to_float(got.back())
+              << ", max_abs=" << max_abs << '\n';
+    return true;
+}
+
+template <typename LaunchFn>
+BenchmarkResult run_benchmark(LaunchFn &&launch, int m, int n, int k, const BenchmarkConfig &config) {
+    for (int iter = 0; iter < config.warmup_iterations; ++iter) {
+        launch();
+    }
+    check_mc(mcDeviceSynchronize(), "mcDeviceSynchronize(warmup)");
+
+    mcEvent_t start;
+    mcEvent_t stop;
+    check_mc(mcEventCreate(&start), "mcEventCreate(start)");
+    check_mc(mcEventCreate(&stop), "mcEventCreate(stop)");
+
+    check_mc(mcEventRecord(start, nullptr), "mcEventRecord(start)");
+    for (int iter = 0; iter < config.measured_iterations; ++iter) {
+        launch();
+    }
+    check_mc(mcEventRecord(stop, nullptr), "mcEventRecord(stop)");
+    check_mc(mcEventSynchronize(stop), "mcEventSynchronize(stop)");
+
+    float elapsed_ms = 0.0f;
+    check_mc(mcEventElapsedTime(&elapsed_ms, start, stop), "mcEventElapsedTime");
+    check_mc(mcEventDestroy(start), "mcEventDestroy(start)");
+    check_mc(mcEventDestroy(stop), "mcEventDestroy(stop)");
+
+    BenchmarkResult result;
+    result.warmup_iterations = config.warmup_iterations;
+    result.measured_iterations = config.measured_iterations;
+    result.avg_ms = elapsed_ms / static_cast<float>(config.measured_iterations);
+    result.tops = compute_tops(m, n, k, result.avg_ms);
+    return result;
+}
+
+bool run_case(const CaseConfig &cfg, const BenchmarkConfig &benchmark_config) {
+    std::vector<int8_t> host_a;
+    std::vector<int8_t> host_b;
+    std::vector<float> host_scale_a;
+    std::vector<float> host_scale_b;
+    std::vector<float> host_moe_weights;
+    std::vector<int> host_token_ids;
+    std::vector<int> host_expert_ids;
+
+    fill_inputs(cfg,
+                host_a,
+                host_b,
+                host_scale_a,
+                host_scale_b,
+                host_moe_weights,
+                host_token_ids,
+                host_expert_ids);
+
+    const std::vector<fused_moe_i8_tn::BFloat16> expected = reference_fused_moe(
+        cfg, host_a, host_b, host_scale_a, host_scale_b, host_moe_weights, host_token_ids, host_expert_ids);
+
+    fused_moe_i8_tn::HostInputs inputs;
+    inputs.num_tokens = cfg.num_tokens;
+    inputs.topk = cfg.topk;
+    inputs.em = cfg.em;
+    inputs.num_experts = kNumExperts;
+    inputs.n = kN;
+    inputs.k = kK;
+    inputs.a = host_a;
+    inputs.b_col_major = host_b;
+    inputs.scale_a = host_scale_a;
+    inputs.scale_b = host_scale_b;
+    inputs.moe_weights = host_moe_weights;
+    inputs.token_ids = host_token_ids;
+    inputs.expert_ids = host_expert_ids;
+
+    const fused_moe_i8_tn::RunResult run_result = fused_moe_i8_tn::run_fused_moe_i8_tn(inputs);
+    const bool valid = validate_result(run_result.output, expected, cfg);
+    if (!valid) {
+        return false;
+    }
+
+    const BenchmarkResult benchmark = run_benchmark([&]() { fused_moe_i8_tn::run_fused_moe_i8_tn(inputs); },
+                                                    cfg.em,
+                                                    kN,
+                                                    kK,
+                                                    benchmark_config);
+    std::cout << cfg.tag
+              << " benchmark"
+              << ": avg_ms=" << benchmark.avg_ms
+              << ", TOPS=" << benchmark.tops
+              << ", warmup=" << benchmark.warmup_iterations
+              << ", iters=" << benchmark.measured_iterations
+              << '\n';
+    return true;
+}
+
+}  // namespace
+
+int main() {
+    int device_count = 0;
+    check_mc(mcGetDeviceCount(&device_count), "mcGetDeviceCount");
+    if (device_count <= 0) {
+        std::cerr << "No MACA device is visible.\n";
+        return EXIT_FAILURE;
+    }
+    check_mc(mcSetDevice(0), "mcSetDevice");
+
+    const BenchmarkConfig benchmark_config = make_benchmark_config();
+    std::cout << "Benchmark config"
+              << ": warmup=" << benchmark_config.warmup_iterations
+              << ", iters=" << benchmark_config.measured_iterations
+              << '\n';
+
+    const std::vector<CaseConfig> cases = {
+        {"fused_moe_i8_tn_topk1", 256, 1, 256, {0, 1}},
+        {"fused_moe_i8_tn_topk2", 256, 2, 512, {0, 1, 1, 0}},
+        {"fused_moe_i8_tn_topk3", 128, 3, 384, {0, 1, 0}},
+    };
+
+    bool ok = true;
+    for (const CaseConfig &cfg : cases) {
+        ok &= run_case(cfg, benchmark_config);
+    }
+    return ok ? EXIT_SUCCESS : EXIT_FAILURE;
+}

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_kernel.h

@@ -0,0 +1,629 @@
+#pragma once
+
+#include <algorithm>
+#include <cstdint>
+
+#include <cute/tensor.hpp>
+
+#include "fused_moe_i8_tn_macros.h"
+#include "fused_moe_i8_tn_types.h"
+
+namespace fused_moe_i8_tn {
+
+using ElementA = int8_t;
+using ElementB = int8_t;
+using ElementC = BFloat16;
+using ElementAccumulator = int32_t;
+using ElementCompute = float;
+
+using INT1 = __NATIVE_VECTOR__(1, int32_t);
+using INT4 = __NATIVE_VECTOR__(4, int32_t);
+using FLOAT2 = __NATIVE_VECTOR__(2, float);
+using FLOAT4 = __NATIVE_VECTOR__(4, float);
+using LdgType = __NATIVE_VECTOR__(4, int32_t);
+using StsType = LdgType;
+using LdsType = LdgType;
+using StgType = __NATIVE_VECTOR__(2, uint);
+using Tc = maca_bfloat16;
+
+constexpr int kTileM = 128;
+constexpr int kTileN = 128;
+constexpr int kTileK = 128;
+constexpr int kThreadCount = 256;
+constexpr int kWaveSize = 64;
+constexpr int kWaveNum = kThreadCount / kWaveSize;
+constexpr int kWaveM = 4;
+constexpr int kWaveN = kWaveNum / kWaveM;
+constexpr int kLdgSize = sizeof(LdgType) * kThreadCount;
+constexpr int kMNPerLdg = kLdgSize / kTileK;
+constexpr int kLdgSizePerWave = kLdgSize / kWaveNum;
+constexpr int kSizeA = kTileM * kTileK * sizeof(ElementA);
+constexpr int kSizeB = kTileN * kTileK * sizeof(ElementB);
+constexpr int kLdgNumA = kSizeA / kLdgSize;
+constexpr int kLdgNumB = kSizeB / kLdgSize;
+constexpr int kLdsNumA = kSizeA / (kLdgSizePerWave * kWaveM);
+constexpr int kLdsNumB = kSizeB / (kLdgSizePerWave * kWaveN);
+constexpr int kStsNumA = kLdgNumA;
+constexpr int kStsNumB = kLdgNumB;
+constexpr int kMmaM = kTileM / 16 / kWaveM;
+constexpr int kMmaN = kTileN / 16 / kWaveN;
+constexpr int kMmaK = kTileK / 16;
+constexpr int kRowCSize = 8;
+constexpr int kOutputCount = 16;
+constexpr int kSmemSize = kSizeA + kSizeB;
+
+template <bool IsTopkLog2>
+struct DirectMoeKernel {
+    static constexpr bool kIsTopkLog2 = IsTopkLog2;
+    using EpilogueOutputOp = fused_moe_i8_tn::EpilogueOutputOp;
+
+    struct Arguments {
+        BatchedGemmCoord problem_size;
+        typename EpilogueOutputOp::Params output_op;
+        void const *ptr_A;
+        void const *ptr_B;
+        void *ptr_C;
+        MoeParams moe_params;
+
+        FUSED_MOE_HOST_DEVICE
+        Arguments() : ptr_A(nullptr), ptr_B(nullptr), ptr_C(nullptr) {}
+
+        FUSED_MOE_HOST_DEVICE
+        Arguments(BatchedGemmCoord problem_size_,
+                  typename EpilogueOutputOp::Params output_op_,
+                  void const *ptr_A_,
+                  void const *ptr_B_,
+                  void *ptr_C_,
+                  MoeParams moe_params_)
+            : problem_size(problem_size_),
+              output_op(output_op_),
+              ptr_A(ptr_A_),
+              ptr_B(ptr_B_),
+              ptr_C(ptr_C_),
+              moe_params(moe_params_) {}
+    };
+};
+
+template <bool IsTopkLog2>
+__global__ void direct_moe_kernel(typename DirectMoeKernel<IsTopkLog2>::Arguments args) {
+    using namespace cute;
+
+#define MMA_STAGE_MNKX2(m, n, k)                                                                    \
+    accum[m][n] = FUSED_MOE_BUILTIN_MMA_16X16X16_I8(a[m][k], b[n][k], accum[m][n]);                \
+    accum[m][n] = FUSED_MOE_BUILTIN_MMA_16X16X16_I8(a[m][k + 1], b[n][k + 1], accum[m][n])
+
+#define LDG_A_STAGE_I(ldgi)                                                                         \
+    A[ldgi] = __builtin_mxc_ldg_b128_predicator(Aaddr + ldg_a_offs_m[ldgi] + ldg_k,                 \
+                                                0,                                                  \
+                                                true,                                               \
+                                                true,                                               \
+                                                false,                                              \
+                                                false,                                              \
+                                                rowA_mask[ldgi],                                    \
+                                                1,                                                  \
+                                                MACA_ICMP_EQ)
+
+#define LDG_B_STAGE_I(ldgi)                                                                         \
+    B[ldgi] = __builtin_mxc_ldg_b128(&(gB(ldg_n[ldgi], ldg_k, tile_k)),                             \
+                                     0,                                                             \
+                                     -1,                                                            \
+                                     true,                                                          \
+                                     true,                                                          \
+                                     false,                                                         \
+                                     false)
+
+#define LDS_A_B128(rowi, coli) FUSED_MOE_LDS(a[rowi][coli * 4], sA(lds_row_A[rowi], lds_col[coli]), LdsType)
+#define LDS_B_B128(rowi, coli) FUSED_MOE_LDS(b[rowi][coli * 4], sB(lds_row_B[rowi], lds_col[coli]), LdsType)
+
+#define CVT_F32_TO_BF16(dst, src0, src1)                                                            \
+    src0 = ((src0 >> 16) & 1) + src0 + 0x7fff;                                                      \
+    src1 = ((src1 >> 16) & 1) + src1 + 0x7fff;                                                      \
+    dst = __builtin_mxc_byte_perm(src0, src1, 0x03020706)
+
+    int *token_ids_ptr = args.moe_params.token_ids;
+    int *expert_ids_ptr = args.moe_params.expert_ids;
+    int *num_tokens_post_padded_ptr = args.moe_params.num_tokens_post_padded_ptr;
+
+    int num_tokens_post_padded = num_tokens_post_padded_ptr[0];
+    int tid = threadIdx.x;
+    int bidx = blockIdx.x + blockIdx.z * gridDim.x;
+    int bidy = blockIdx.y;
+    int wave = tid / kWaveSize;
+    int lane = tid % kWaveSize;
+
+    if (bidx * kTileM >= num_tokens_post_padded) {
+        return;
+    }
+
+    EpilogueOutputOp output_op(args.output_op);
+
+    __shared__ int8_t smem_data[kSmemSize];
+    int8_t *smem_A = smem_data;
+    int8_t *smem_B = smem_A + kSizeA;
+
+    int group_idx = expert_ids_ptr[bidx];
+    int prev_m = bidx * kTileM;
+    ElementB *Baddr = (ElementB *)args.ptr_B + uint64_t(group_idx) * args.problem_size.n() * args.problem_size.k();
+
+    Tensor mB = make_tensor(make_gmem_ptr((ElementB *)Baddr),
+                            make_shape(args.problem_size.n(), args.problem_size.k()),
+                            make_stride(args.problem_size.k(), Int<1>{}));
+    Tensor gB = local_tile(mB, make_tile(Int<kTileN>{}, Int<kTileK>{}), make_coord(bidy, _));
+
+    LdgType A[kLdgNumA], B[kLdgNumB];
+    int k_head = (args.problem_size.k() - 1) % kTileK + 1;
+    int col_limit = min(kTileN, args.problem_size.n() - bidy * kTileN);
+    int ldg_n[kLdgNumB], ldg_a_offs_m[kLdgNumA];
+    bool rowA_mask[kLdgNumA];
+    int ldg_m_base = tid / 8;
+    int ldg_n_base = tid / 8 * kLdgNumB;
+    int ldg_k = (lane % 8) * 16;
+    int num_tile_k = size<2>(gB);
+
+    ElementA *Aaddr = (ElementA *)args.ptr_A + (num_tile_k - 1) * kTileK;
+
+#pragma unroll
+    for (uint32_t ldgi = 0; ldgi < kLdgNumA; ++ldgi) {
+        int idx_row_a = ldg_m_base + kMNPerLdg * ldgi;
+        reinterpret_cast<INT1 *>(&ldg_a_offs_m)[ldgi] =
+            __builtin_mxc_ldg_b32(token_ids_ptr + idx_row_a + prev_m, 0, -1, true, true, false, false);
+    }
+#pragma unroll
+    for (uint32_t ldgi = 0; ldgi < kLdgNumB; ++ldgi) {
+        ldg_n[ldgi] = min(ldg_n_base + ldgi, col_limit - 1);
+        B[ldgi] = __builtin_mxc_ldg_b128_predicator(&(gB(ldg_n[ldgi], ldg_k, num_tile_k - 1)),
+                                                    0,
+                                                    true,
+                                                    true,
+                                                    false,
+                                                    false,
+                                                    ldg_k,
+                                                    k_head,
+                                                    MACA_ICMP_SLT);
+    }
+#pragma unroll
+    for (uint32_t ldgi = 0; ldgi < kLdgNumA; ++ldgi) {
+        rowA_mask[ldgi] = ldg_a_offs_m[ldgi] < args.problem_size.m();
+        if constexpr (IsTopkLog2) {
+            ldg_a_offs_m[ldgi] = (ldg_a_offs_m[ldgi] >> args.moe_params.topk_bits) * args.problem_size.k();
+        } else {
+            ldg_a_offs_m[ldgi] = (ldg_a_offs_m[ldgi] / args.moe_params.topk) * args.problem_size.k();
+        }
+        A[ldgi] = __builtin_mxc_ldg_b128_predicator(Aaddr + ldg_a_offs_m[ldgi] + ldg_k,
+                                                    0,
+                                                    true,
+                                                    true,
+                                                    false,
+                                                    false,
+                                                    (ldg_k < k_head) && rowA_mask[ldgi],
+                                                    1,
+                                                    MACA_ICMP_EQ);
+    }
+
+    Tensor sA = make_tensor(make_smem_ptr((ElementA *)smem_A),
+                            make_shape(Int<kTileM>{}, Int<kTileK>{}),
+                            make_stride(Int<kTileK>{}, Int<1>{}));
+    Tensor sB = make_tensor(make_smem_ptr((ElementB *)smem_B),
+                            make_shape(Int<kTileN>{}, Int<kTileK>{}),
+                            make_stride(Int<kTileK>{}, Int<1>{}));
+
+    int sts_rowA[kStsNumA], sts_rowB[kStsNumB];
+    int sts_col = (((tid / 8) + (tid % 8)) % 8) * 16;
+#pragma unroll
+    for (uint32_t i = 0; i < kStsNumB; ++i) {
+        sts_rowB[i] = tid / 8 + kMNPerLdg * i;
+        FUSED_MOE_STS(sB(sts_rowB[i], sts_col), B[i], StsType);
+    }
+#pragma unroll
+    for (uint32_t i = 0; i < kStsNumA; ++i) {
+        sts_rowA[i] = wave * 32 + lane / 8 + i * 8;
+    }
+    FUSED_MOE_STS(sA(sts_rowA[0], sts_col), A[0], StsType);
+    FUSED_MOE_STS(sA(sts_rowA[1], sts_col), A[1], StsType);
+
+    INT4 accum[kMmaM][kMmaN] = {0};
+    int32_t a[kMmaM][kMmaK], b[kMmaN][kMmaK];
+    int lds_row_A[2], lds_row_B[8], lds_col[2];
+
+#pragma unroll
+    for (int i = 0; i < 2; ++i) {
+        lds_col[i] = (((tid % 16) + (lane / 16) + 4 * i) % 8) * 16;
+        lds_row_A[i] = (tid % 16) + wave * 32 + 16 * i;
+    }
+#pragma unroll
+    for (int i = 0; i < 8; ++i) {
+        lds_row_B[i] = (tid % 16) + 16 * i;
+    }
+
+    __syncthreadshared();
+
+    LDS_A_B128(0, 0);
+    LDS_B_B128(0, 0);
+    LDS_B_B128(1, 0);
+    LDS_B_B128(2, 0);
+    LDS_B_B128(3, 0);
+
+    int loop_tile_k = size<2>(gB) - 1;
+    Aaddr = (ElementA *)args.ptr_A;
+    for (uint32_t tile_k = 0; tile_k < loop_tile_k; ++tile_k) {
+        LDG_B_STAGE_I(0);
+        LDG_B_STAGE_I(1);
+        MMA_STAGE_MNKX2(0, 0, 0);
+        LDS_B_B128(4, 0);
+        MMA_STAGE_MNKX2(0, 0, 2);
+        LDS_B_B128(5, 0);
+        MMA_STAGE_MNKX2(0, 1, 0);
+        LDS_B_B128(6, 0);
+        LDG_B_STAGE_I(2);
+        MMA_STAGE_MNKX2(0, 1, 2);
+        LDS_B_B128(7, 0);
+        MMA_STAGE_MNKX2(0, 2, 0);
+        LDG_B_STAGE_I(3);
+        MMA_STAGE_MNKX2(0, 2, 2);
+        MMA_STAGE_MNKX2(0, 3, 0);
+        LDG_A_STAGE_I(0);
+        MMA_STAGE_MNKX2(0, 3, 2);
+        LDG_A_STAGE_I(1);
+
+        MMA_STAGE_MNKX2(0, 4, 0);
+        LDS_A_B128(0, 1);
+        MMA_STAGE_MNKX2(0, 4, 2);
+        LDS_B_B128(0, 1);
+        MMA_STAGE_MNKX2(0, 5, 0);
+        LDS_B_B128(1, 1);
+        MMA_STAGE_MNKX2(0, 5, 2);
+        LDS_B_B128(2, 1);
+        MMA_STAGE_MNKX2(0, 6, 0);
+        LDS_B_B128(3, 1);
+        MMA_STAGE_MNKX2(0, 6, 2);
+        MMA_STAGE_MNKX2(0, 7, 0);
+        MMA_STAGE_MNKX2(0, 7, 2);
+
+        LDS_B_B128(4, 1);
+        MMA_STAGE_MNKX2(0, 0, 4);
+        LDS_B_B128(5, 1);
+        MMA_STAGE_MNKX2(0, 0, 6);
+        LDS_B_B128(6, 1);
+        MMA_STAGE_MNKX2(0, 1, 4);
+        LDS_B_B128(7, 1);
+        MMA_STAGE_MNKX2(0, 1, 6);
+        MMA_STAGE_MNKX2(0, 2, 4);
+        MMA_STAGE_MNKX2(0, 2, 6);
+        FUSED_MOE_STS(sA(sts_rowA[2], sts_col), A[2], StsType);
+        MMA_STAGE_MNKX2(0, 3, 4);
+        MMA_STAGE_MNKX2(0, 3, 6);
+        FUSED_MOE_STS(sA(sts_rowA[3], sts_col), A[3], StsType);
+
+        MMA_STAGE_MNKX2(0, 4, 4);
+        LDG_A_STAGE_I(2);
+        MMA_STAGE_MNKX2(0, 4, 6);
+        LDG_A_STAGE_I(3);
+        MMA_STAGE_MNKX2(0, 5, 4);
+        MMA_STAGE_MNKX2(0, 5, 6);
+        MMA_STAGE_MNKX2(0, 6, 4);
+        LDS_A_B128(1, 0);
+        MMA_STAGE_MNKX2(0, 6, 6);
+        MMA_STAGE_MNKX2(0, 7, 4);
+        Aaddr += kTileK;
+        MMA_STAGE_MNKX2(0, 7, 6);
+
+        __syncthreadshared();
+        MMA_STAGE_MNKX2(1, 0, 0);
+        LDS_A_B128(1, 1);
+        MMA_STAGE_MNKX2(1, 0, 2);
+        MMA_STAGE_MNKX2(1, 1, 0);
+        MMA_STAGE_MNKX2(1, 1, 2);
+        MMA_STAGE_MNKX2(1, 2, 0);
+        MMA_STAGE_MNKX2(1, 2, 2);
+        MMA_STAGE_MNKX2(1, 3, 0);
+        MMA_STAGE_MNKX2(1, 3, 2);
+
+        MMA_STAGE_MNKX2(1, 4, 0);
+        FUSED_MOE_STS(sB(sts_rowB[0], sts_col), B[0], StsType);
+        MMA_STAGE_MNKX2(1, 4, 2);
+        MMA_STAGE_MNKX2(1, 5, 0);
+        MMA_STAGE_MNKX2(1, 5, 2);
+        FUSED_MOE_STS(sB(sts_rowB[1], sts_col), B[1], StsType);
+        MMA_STAGE_MNKX2(1, 6, 0);
+        MMA_STAGE_MNKX2(1, 6, 2);
+        MMA_STAGE_MNKX2(1, 7, 0);
+        FUSED_MOE_STS(sB(sts_rowB[2], sts_col), B[2], StsType);
+        MMA_STAGE_MNKX2(1, 7, 2);
+
+        MMA_STAGE_MNKX2(1, 0, 4);
+        MMA_STAGE_MNKX2(1, 0, 6);
+        FUSED_MOE_STS(sB(sts_rowB[3], sts_col), B[3], StsType);
+        MMA_STAGE_MNKX2(1, 1, 4);
+        MMA_STAGE_MNKX2(1, 1, 6);
+        MMA_STAGE_MNKX2(1, 2, 4);
+        FUSED_MOE_STS(sA(sts_rowA[0], sts_col), A[0], StsType);
+        MMA_STAGE_MNKX2(1, 2, 6);
+        MMA_STAGE_MNKX2(1, 3, 4);
+        MMA_STAGE_MNKX2(1, 3, 6);
+        FUSED_MOE_STS(sA(sts_rowA[1], sts_col), A[1], StsType);
+
+        MMA_STAGE_MNKX2(1, 4, 4);
+        MMA_STAGE_MNKX2(1, 4, 6);
+        MMA_STAGE_MNKX2(1, 5, 4);
+        __syncthreadshared();
+        MMA_STAGE_MNKX2(1, 5, 6);
+        LDS_A_B128(0, 0);
+        LDS_B_B128(0, 0);
+        MMA_STAGE_MNKX2(1, 6, 4);
+        LDS_B_B128(1, 0);
+        MMA_STAGE_MNKX2(1, 6, 6);
+        LDS_B_B128(2, 0);
+        MMA_STAGE_MNKX2(1, 7, 4);
+        LDS_B_B128(3, 0);
+        MMA_STAGE_MNKX2(1, 7, 6);
+    }
+
+    int rowC[kRowCSize];
+    MMA_STAGE_MNKX2(0, 0, 0);
+    LDS_B_B128(4, 0);
+    MMA_STAGE_MNKX2(0, 0, 2);
+    LDS_B_B128(5, 0);
+    MMA_STAGE_MNKX2(0, 1, 0);
+    LDS_B_B128(6, 0);
+    MMA_STAGE_MNKX2(0, 1, 2);
+    LDS_B_B128(7, 0);
+    MMA_STAGE_MNKX2(0, 2, 0);
+    int token_row_m = prev_m + ((lane / 16) % 2) * 4 + wave * 8 + (lane / 32) * 32;
+    MMA_STAGE_MNKX2(0, 2, 2);
+    MMA_STAGE_MNKX2(0, 3, 0);
+    MMA_STAGE_MNKX2(0, 3, 2);
+
+#pragma unroll
+    for (int j = 0; j < 4; ++j) {
+        *(reinterpret_cast<INT1 *>(&rowC) + j) =
+            __builtin_mxc_ldg_b32(token_ids_ptr + token_row_m + j, 0, -1, true, true, false, false);
+    }
+
+    MMA_STAGE_MNKX2(0, 4, 0);
+    LDS_A_B128(0, 1);
+    MMA_STAGE_MNKX2(0, 4, 2);
+    LDS_B_B128(0, 1);
+    MMA_STAGE_MNKX2(0, 5, 0);
+    LDS_B_B128(1, 1);
+    MMA_STAGE_MNKX2(0, 5, 2);
+    LDS_B_B128(2, 1);
+    MMA_STAGE_MNKX2(0, 6, 0);
+    LDS_B_B128(3, 1);
+    MMA_STAGE_MNKX2(0, 6, 2);
+    MMA_STAGE_MNKX2(0, 7, 0);
+    MMA_STAGE_MNKX2(0, 7, 2);
+
+    LDS_B_B128(4, 1);
+    MMA_STAGE_MNKX2(0, 0, 4);
+    LDS_B_B128(5, 1);
+    MMA_STAGE_MNKX2(0, 0, 6);
+    LDS_B_B128(6, 1);
+    MMA_STAGE_MNKX2(0, 1, 4);
+    LDS_B_B128(7, 1);
+    MMA_STAGE_MNKX2(0, 1, 6);
+    MMA_STAGE_MNKX2(0, 2, 4);
+    FUSED_MOE_STS(sA(sts_rowA[2], sts_col), A[2], StsType);
+    MMA_STAGE_MNKX2(0, 2, 6);
+    MMA_STAGE_MNKX2(0, 3, 4);
+    MMA_STAGE_MNKX2(0, 3, 6);
+    FUSED_MOE_STS(sA(sts_rowA[3], sts_col), A[3], StsType);
+
+    MMA_STAGE_MNKX2(0, 4, 4);
+    MMA_STAGE_MNKX2(0, 4, 6);
+    MMA_STAGE_MNKX2(0, 5, 4);
+    MMA_STAGE_MNKX2(0, 5, 6);
+    MMA_STAGE_MNKX2(0, 6, 4);
+    LDS_A_B128(1, 0);
+    MMA_STAGE_MNKX2(0, 6, 6);
+    MMA_STAGE_MNKX2(0, 7, 4);
+    MMA_STAGE_MNKX2(0, 7, 6);
+
+#pragma unroll
+    for (int j = 0; j < 4; ++j) {
+        *(reinterpret_cast<INT1 *>(&rowC) + 4 + j) =
+            __builtin_mxc_ldg_b32(token_ids_ptr + token_row_m + 64 + j, 0, -1, true, true, false, false);
+    }
+
+    MMA_STAGE_MNKX2(1, 0, 0);
+    MMA_STAGE_MNKX2(1, 0, 2);
+    MMA_STAGE_MNKX2(1, 1, 0);
+    MMA_STAGE_MNKX2(1, 1, 2);
+    MMA_STAGE_MNKX2(1, 2, 0);
+    MMA_STAGE_MNKX2(1, 2, 2);
+    MMA_STAGE_MNKX2(1, 3, 0);
+    MMA_STAGE_MNKX2(1, 3, 2);
+
+    MMA_STAGE_MNKX2(1, 4, 0);
+    MMA_STAGE_MNKX2(1, 4, 2);
+    LDS_A_B128(1, 1);
+    MMA_STAGE_MNKX2(1, 5, 0);
+    MMA_STAGE_MNKX2(1, 5, 2);
+    MMA_STAGE_MNKX2(1, 6, 0);
+    MMA_STAGE_MNKX2(1, 6, 2);
+    MMA_STAGE_MNKX2(1, 7, 0);
+    MMA_STAGE_MNKX2(1, 7, 2);
+
+    MMA_STAGE_MNKX2(1, 0, 4);
+    MMA_STAGE_MNKX2(1, 0, 6);
+    MMA_STAGE_MNKX2(1, 1, 4);
+    MMA_STAGE_MNKX2(1, 1, 6);
+    MMA_STAGE_MNKX2(1, 2, 4);
+    MMA_STAGE_MNKX2(1, 2, 6);
+    MMA_STAGE_MNKX2(1, 3, 4);
+    MMA_STAGE_MNKX2(1, 3, 6);
+
+    MMA_STAGE_MNKX2(1, 4, 4);
+    MMA_STAGE_MNKX2(1, 4, 6);
+    MMA_STAGE_MNKX2(1, 5, 4);
+    MMA_STAGE_MNKX2(1, 5, 6);
+    MMA_STAGE_MNKX2(1, 6, 4);
+    MMA_STAGE_MNKX2(1, 6, 6);
+    MMA_STAGE_MNKX2(1, 7, 4);
+    MMA_STAGE_MNKX2(1, 7, 6);
+
+    INT4 output[kOutputCount];
+#pragma unroll
+    for (uint32_t i = 0; i < 2; ++i) {
+#pragma unroll
+        for (uint32_t j = 0; j < 4; ++j) {
+            output[i * 8 + 2 * j][0] = accum[i][0][j];
+            output[i * 8 + 2 * j][1] = accum[i][2][j];
+            output[i * 8 + 2 * j][2] = accum[i][4][j];
+            output[i * 8 + 2 * j][3] = accum[i][6][j];
+            output[i * 8 + 2 * j + 1][0] = accum[i][1][j];
+            output[i * 8 + 2 * j + 1][1] = accum[i][3][j];
+            output[i * 8 + 2 * j + 1][2] = accum[i][5][j];
+            output[i * 8 + 2 * j + 1][3] = accum[i][7][j];
+        }
+    }
+
+    int colC[2];
+    bool colC_mask[2];
+    colC[0] = (tid % 16) * 4;
+    colC[1] = colC[0] + 64;
+    colC_mask[0] = colC[0] < col_limit;
+    colC_mask[1] = colC[1] < col_limit;
+
+    float weights[2][4], a_scale[2][4];
+    FLOAT4 b_scale[2];
+
+#pragma unroll
+    for (uint32_t i = 0; i < 2; ++i) {
+#pragma unroll
+        for (uint32_t j = 0; j < 4; ++j) {
+            if (output_op.MUL_WEIGHTS) {
+                const void *moe_weights_ptr = output_op.moe_weights_ + rowC[i * 4 + j];
+                *(reinterpret_cast<INT1 *>(&weights[i]) + j) =
+                    __builtin_mxc_ldg_b32_predicator(const_cast<void *>(moe_weights_ptr),
+                                                     0,
+                                                     true,
+                                                     true,
+                                                     false,
+                                                     false,
+                                                     rowC[i * 4 + j],
+                                                     args.problem_size.m(),
+                                                     MACA_ICMP_SLT);
+            }
+
+            int row_a_scale;
+            if constexpr (IsTopkLog2) {
+                row_a_scale = (rowC[i * 4 + j] >> args.moe_params.topk_bits);
+            } else {
+                row_a_scale = (rowC[i * 4 + j] / args.moe_params.topk);
+            }
+
+            const void *scale_a_ptr = output_op.scale_a_ + row_a_scale;
+            *(reinterpret_cast<INT1 *>(&a_scale[i]) + j) =
+                __builtin_mxc_ldg_b32_predicator(const_cast<void *>(scale_a_ptr),
+                                                 0,
+                                                 true,
+                                                 true,
+                                                 false,
+                                                 false,
+                                                 rowC[i * 4 + j],
+                                                 args.problem_size.m(),
+                                                 MACA_ICMP_SLT);
+        }
+    }
+
+#pragma unroll
+    for (uint32_t i = 0; i < 2; ++i) {
+        const void *scale_b_ptr =
+            (const float *)output_op.scale_b_ + group_idx * args.problem_size.n() + bidy * kTileN + colC[i];
+        b_scale[i] = __builtin_mxc_ldg_b128_predicator(const_cast<void *>(scale_b_ptr),
+                                                       0,
+                                                       true,
+                                                       true,
+                                                       false,
+                                                       false,
+                                                       colC_mask[i],
+                                                       1,
+                                                       MACA_ICMP_EQ);
+    }
+
+    Tc *Caddr = (Tc *)args.ptr_C + bidy * kTileN;
+    FLOAT2 zero2 = {0.f, 0.f};
+    StgType tempC;
+
+#pragma unroll
+    for (uint32_t i = 0; i < 2; ++i) {
+#pragma unroll
+        for (uint32_t j = 0; j < 4; ++j) {
+            float out[8];
+            out[0] = output[i * 8 + 2 * j][0];
+            out[1] = output[i * 8 + 2 * j][1];
+            out[2] = output[i * 8 + 2 * j][2];
+            out[3] = output[i * 8 + 2 * j][3];
+            out[4] = output[i * 8 + 2 * j + 1][0];
+            out[5] = output[i * 8 + 2 * j + 1][1];
+            out[6] = output[i * 8 + 2 * j + 1][2];
+            out[7] = output[i * 8 + 2 * j + 1][3];
+
+            if (output_op.MUL_WEIGHTS) {
+                a_scale[i][j] *= weights[i][j];
+            }
+
+            FLOAT2 a_scale_f2 = {a_scale[i][j], a_scale[i][j]};
+            FLOAT2 scale[4];
+            scale[0] = __builtin_mxc_pk_fma_f32(reinterpret_cast<FLOAT2 *>(&b_scale[0])[0], a_scale_f2, zero2);
+            scale[1] = __builtin_mxc_pk_fma_f32(reinterpret_cast<FLOAT2 *>(&b_scale[0])[1], a_scale_f2, zero2);
+            scale[2] = __builtin_mxc_pk_fma_f32(reinterpret_cast<FLOAT2 *>(&b_scale[1])[0], a_scale_f2, zero2);
+            scale[3] = __builtin_mxc_pk_fma_f32(reinterpret_cast<FLOAT2 *>(&b_scale[1])[1], a_scale_f2, zero2);
+            *reinterpret_cast<FLOAT2 *>(&out[0]) =
+                __builtin_mxc_pk_fma_f32(*reinterpret_cast<FLOAT2 *>(&out[0]), scale[0], zero2);
+            *reinterpret_cast<FLOAT2 *>(&out[2]) =
+                __builtin_mxc_pk_fma_f32(*reinterpret_cast<FLOAT2 *>(&out[2]), scale[1], zero2);
+            *reinterpret_cast<FLOAT2 *>(&out[4]) =
+                __builtin_mxc_pk_fma_f32(*reinterpret_cast<FLOAT2 *>(&out[4]), scale[2], zero2);
+            *reinterpret_cast<FLOAT2 *>(&out[6]) =
+                __builtin_mxc_pk_fma_f32(*reinterpret_cast<FLOAT2 *>(&out[6]), scale[3], zero2);
+
+            CVT_F32_TO_BF16(tempC[0], reinterpret_cast<uint *>(&out)[0], reinterpret_cast<uint *>(&out)[1]);
+            CVT_F32_TO_BF16(tempC[1], reinterpret_cast<uint *>(&out)[2], reinterpret_cast<uint *>(&out)[3]);
+            __builtin_mxc_stg_b64_predicator(Caddr + rowC[i * 4 + j] * args.problem_size.n() + colC[0],
+                                             0,
+                                             *(reinterpret_cast<uint64_t *>(&tempC)),
+                                             true,
+                                             false,
+                                             false,
+                                             (rowC[i * 4 + j] < args.problem_size.m()) && colC_mask[0],
+                                             1,
+                                             MACA_ICMP_EQ);
+
+            CVT_F32_TO_BF16(tempC[0], reinterpret_cast<uint *>(&out)[4], reinterpret_cast<uint *>(&out)[5]);
+            CVT_F32_TO_BF16(tempC[1], reinterpret_cast<uint *>(&out)[6], reinterpret_cast<uint *>(&out)[7]);
+            __builtin_mxc_stg_b64_predicator(Caddr + rowC[i * 4 + j] * args.problem_size.n() + colC[1],
+                                             0,
+                                             *(reinterpret_cast<uint64_t *>(&tempC)),
+                                             true,
+                                             false,
+                                             false,
+                                             (rowC[i * 4 + j] < args.problem_size.m()) && colC_mask[1],
+                                             1,
+                                             MACA_ICMP_EQ);
+        }
+    }
+}
+
+template <bool IsTopkLog2>
+using DirectMoeGemmKernel = DirectMoeKernel<IsTopkLog2>;
+
+template <typename Kernel>
+inline dim3 get_grid_shape(typename Kernel::Arguments const &args) {
+    const int grid_m = (args.moe_params.EM + kTileM - 1) / kTileM;
+    const int group_bidx = std::max(1, std::min(8, (args.problem_size.m() / args.problem_size.batch() + kTileM - 1) / kTileM));
+    const int grid_x = std::min(grid_m, group_bidx);
+    const int grid_z = (grid_m + grid_x - 1) / grid_x;
+    const int grid_y = (args.problem_size.n() + kTileN - 1) / kTileN;
+    return dim3(grid_x, grid_y, grid_z);
+}
+
+template <typename Kernel>
+inline Status launch(typename Kernel::Arguments const &args, mcStream_t stream = nullptr) {
+    dim3 const block(kThreadCount, 1, 1);
+    dim3 const grid = get_grid_shape<Kernel>(args);
+    direct_moe_kernel<Kernel::kIsTopkLog2><<<grid, block, 0, stream>>>(args);
+    return Status::kSuccess;
+}
+
+}  // namespace fused_moe_i8_tn

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_macros.h

@@ -0,0 +1,21 @@
+#pragma once
+
+#include "fused_moe_i8_tn_types.h"
+
+#define FUSED_MOE_CP_ASYNC_FENC() asm(";--------------")
+
+#define FUSED_MOE_LDS(dst, src, type_)                                                              \
+    FUSED_MOE_CP_ASYNC_FENC();                                                                      \
+    *reinterpret_cast<type_ *>(&(dst)) = *reinterpret_cast<type_ *>(&(src));                        \
+    FUSED_MOE_CP_ASYNC_FENC()
+
+#define FUSED_MOE_STS(dst, src, type_)                                                              \
+    FUSED_MOE_CP_ASYNC_FENC();                                                                      \
+    *reinterpret_cast<type_ *>(&(dst)) = *reinterpret_cast<type_ *>(&(src));                        \
+    FUSED_MOE_CP_ASYNC_FENC()
+
+#if defined(__MACA_ARCH__) && (__MACA_ARCH__ == 1000 || __MACA_ARCH__ == 1089)
+#define FUSED_MOE_BUILTIN_MMA_16X16X16_I8(a, b, c) __builtin_mxc_mma_16x16x16i8(a, b, c)
+#else
+#define FUSED_MOE_BUILTIN_MMA_16X16X16_I8(a, b, c) 0
+#endif

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_pybind.cpp

@@ -0,0 +1,101 @@
+#include <Python.h>
+
+#include <pybind11/numpy.h>
+#include <pybind11/pybind11.h>
+
+#include <cstdint>
+#include <cstring>
+#include <stdexcept>
+#include <vector>
+
+#include "fused_moe_i8_tn_runner.h"
+
+namespace py = pybind11;
+
+namespace {
+
+template <typename T>
+std::vector<T> copy_1d_array(const py::array &array, const char *name) {
+    auto view = py::array_t<T, py::array::c_style | py::array::forcecast>::ensure(array);
+    if (!view) {
+        throw std::invalid_argument(std::string("failed to cast ") + name);
+    }
+    if (view.ndim() != 1) {
+        throw std::invalid_argument(std::string(name) + " must be a 1D array");
+    }
+    const T *ptr = static_cast<const T *>(view.data());
+    return std::vector<T>(ptr, ptr + view.size());
+}
+
+std::vector<int8_t> copy_2d_int8_array(const py::array &array, const char *name) {
+    auto view = py::array_t<int8_t, py::array::c_style | py::array::forcecast>::ensure(array);
+    if (!view) {
+        throw std::invalid_argument(std::string("failed to cast ") + name);
+    }
+    if (view.ndim() != 2) {
+        throw std::invalid_argument(std::string(name) + " must be a 2D array");
+    }
+    const int8_t *ptr = static_cast<const int8_t *>(view.data());
+    return std::vector<int8_t>(ptr, ptr + view.size());
+}
+
+py::array_t<float> run_fused_moe_pybind(const py::array &a,
+                                        const py::array &b_col_major,
+                                        const py::array &scale_a,
+                                        const py::array &scale_b,
+                                        const py::array &moe_weights,
+                                        const py::array &token_ids,
+                                        const py::array &expert_ids,
+                                        int topk,
+                                        int num_experts = fused_moe_i8_tn::kDefaultNumExperts,
+                                        int device_id = 0) {
+    auto a_view = py::array_t<int8_t, py::array::c_style | py::array::forcecast>::ensure(a);
+    auto b_view = py::array_t<int8_t, py::array::c_style | py::array::forcecast>::ensure(b_col_major);
+    if (!a_view || a_view.ndim() != 2) {
+        throw std::invalid_argument("a must be a 2D int8 array");
+    }
+    if (!b_view || b_view.ndim() != 3) {
+        throw std::invalid_argument("b_col_major must be a 3D int8 array");
+    }
+
+    fused_moe_i8_tn::HostInputs inputs;
+    inputs.num_tokens = static_cast<int>(a_view.shape(0));
+    inputs.k = static_cast<int>(a_view.shape(1));
+    inputs.num_experts = num_experts;
+    inputs.n = static_cast<int>(b_view.shape(1));
+    inputs.topk = topk;
+    inputs.em = inputs.num_tokens * inputs.topk;
+    inputs.a = copy_2d_int8_array(a, "a");
+    inputs.b_col_major.assign(static_cast<const int8_t *>(b_view.data()),
+                              static_cast<const int8_t *>(b_view.data()) + b_view.size());
+    inputs.scale_a = copy_1d_array<float>(scale_a, "scale_a");
+    inputs.scale_b = copy_1d_array<float>(scale_b, "scale_b");
+    inputs.moe_weights = copy_1d_array<float>(moe_weights, "moe_weights");
+    inputs.token_ids = copy_1d_array<int>(token_ids, "token_ids");
+    inputs.expert_ids = copy_1d_array<int>(expert_ids, "expert_ids");
+
+    fused_moe_i8_tn::RunResult result = fused_moe_i8_tn::run_fused_moe_i8_tn(inputs, device_id);
+    std::vector<float> output = fused_moe_i8_tn::bf16_vector_to_float(result.output);
+
+    py::array_t<float> out({result.rows, result.cols});
+    std::memcpy(out.mutable_data(), output.data(), output.size() * sizeof(float));
+    return out;
+}
+
+}  // namespace
+
+PYBIND11_MODULE(fused_moe_i8_tn_pybind, m) {
+    m.doc() = "Pybind wrapper for standalone fused_moe_i8_tn";
+    m.def("run_fused_moe_i8_tn",
+          &run_fused_moe_pybind,
+          py::arg("a"),
+          py::arg("b_col_major"),
+          py::arg("scale_a"),
+          py::arg("scale_b"),
+          py::arg("moe_weights"),
+          py::arg("token_ids"),
+          py::arg("expert_ids"),
+          py::arg("topk"),
+          py::arg("num_experts") = fused_moe_i8_tn::kDefaultNumExperts,
+          py::arg("device_id") = 0);
+}

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_runner.h

@@ -0,0 +1,331 @@
+#pragma once
+
+#include <mc_runtime.h>
+
+#include <cmath>
+#include <cstdint>
+#include <cstring>
+#include <stdexcept>
+#include <string>
+#include <vector>
+
+#include "fused_moe_i8_tn_kernel.h"
+#include "fused_moe_i8_tn_types.h"
+
+namespace fused_moe_i8_tn {
+
+constexpr int kDefaultNumExperts = 2;
+constexpr int kDefaultTileM = 128;
+constexpr int kDefaultN = 128;
+constexpr int kDefaultK = 128;
+
+struct HostInputs {
+    int num_tokens = 0;
+    int topk = 0;
+    int em = 0;
+    int num_experts = kDefaultNumExperts;
+    int n = kDefaultN;
+    int k = kDefaultK;
+    std::vector<int8_t> a;
+    std::vector<int8_t> b_col_major;
+    std::vector<float> scale_a;
+    std::vector<float> scale_b;
+    std::vector<float> moe_weights;
+    std::vector<int> token_ids;
+    std::vector<int> expert_ids;
+};
+
+struct RunResult {
+    int rows = 0;
+    int cols = 0;
+    std::vector<BFloat16> output;
+};
+
+inline float bf16_to_float(BFloat16 value) {
+    return static_cast<float>(value);
+}
+
+inline BFloat16 float_to_bf16(float value) {
+    return BFloat16(value);
+}
+
+inline std::vector<float> bf16_vector_to_float(const std::vector<BFloat16> &input) {
+    std::vector<float> out(input.size());
+    for (size_t i = 0; i < input.size(); ++i) {
+        out[i] = bf16_to_float(input[i]);
+    }
+    return out;
+}
+
+inline void check_mc_or_throw(mcError_t status, const char *expr) {
+    if (status != mcSuccess) {
+        throw std::runtime_error(std::string(expr) + " failed: " + mcGetErrorString(status));
+    }
+}
+
+inline void check_status_or_throw(Status status, const char *expr) {
+    if (status != Status::kSuccess) {
+        throw std::runtime_error(std::string(expr) + " failed: " + get_status_string(status));
+    }
+}
+
+inline bool is_log2_value(int value) {
+    return value > 0 && ((value & (value - 1)) == 0);
+}
+
+inline void validate_host_inputs(const HostInputs &inputs) {
+    if (inputs.num_tokens <= 0) {
+        throw std::invalid_argument("num_tokens must be > 0");
+    }
+    if (inputs.topk <= 0) {
+        throw std::invalid_argument("topk must be > 0");
+    }
+    if (inputs.n != kDefaultN) {
+        throw std::invalid_argument("only N=128 is supported by fused_moe_i8_tn");
+    }
+    if (inputs.k != kDefaultK) {
+        throw std::invalid_argument("only K=128 is supported by fused_moe_i8_tn");
+    }
+    if (inputs.em <= 0) {
+        throw std::invalid_argument("em must be > 0");
+    }
+    if (inputs.em != inputs.num_tokens * inputs.topk) {
+        throw std::invalid_argument("em must equal num_tokens * topk");
+    }
+    if (inputs.em % kDefaultTileM != 0) {
+        throw std::invalid_argument("em must be a multiple of 128");
+    }
+    if (inputs.num_experts <= 0) {
+        throw std::invalid_argument("num_experts must be > 0");
+    }
+
+    const size_t total_rows = static_cast<size_t>(inputs.em);
+    const size_t expected_a = static_cast<size_t>(inputs.num_tokens) * inputs.k;
+    const size_t expected_b = static_cast<size_t>(inputs.num_experts) * inputs.n * inputs.k;
+    const size_t expected_scale_a = static_cast<size_t>(inputs.num_tokens);
+    const size_t expected_scale_b = static_cast<size_t>(inputs.num_experts) * inputs.n;
+    const size_t expected_moe_weights = total_rows;
+    const size_t expected_token_ids = total_rows;
+    const size_t expected_expert_ids = static_cast<size_t>(inputs.em / kDefaultTileM);
+
+    if (inputs.a.size() != expected_a) {
+        throw std::invalid_argument("A size mismatch");
+    }
+    if (inputs.b_col_major.size() != expected_b) {
+        throw std::invalid_argument("B size mismatch");
+    }
+    if (inputs.scale_a.size() != expected_scale_a) {
+        throw std::invalid_argument("scale_a size mismatch");
+    }
+    if (inputs.scale_b.size() != expected_scale_b) {
+        throw std::invalid_argument("scale_b size mismatch");
+    }
+    if (inputs.moe_weights.size() != expected_moe_weights) {
+        throw std::invalid_argument("moe_weights size mismatch");
+    }
+    if (inputs.token_ids.size() != expected_token_ids) {
+        throw std::invalid_argument("token_ids size mismatch");
+    }
+    if (inputs.expert_ids.size() != expected_expert_ids) {
+        throw std::invalid_argument("expert_ids size mismatch");
+    }
+
+    for (int expert : inputs.expert_ids) {
+        if (expert < 0 || expert >= inputs.num_experts) {
+            throw std::invalid_argument("expert_ids contains out-of-range expert index");
+        }
+    }
+}
+
+template <typename GemmKernel>
+inline typename GemmKernel::Arguments make_direct_arguments(const HostInputs &inputs,
+                                                            int total_rows,
+                                                            int8_t *dev_a,
+                                                            int8_t *dev_b,
+                                                            float *dev_scale_a,
+                                                            float *dev_scale_b,
+                                                            float *dev_moe_weights,
+                                                            int *dev_token_ids,
+                                                            int *dev_expert_ids,
+                                                            int32_t *dev_num_tokens_post_padded,
+                                                            BFloat16 *dev_c) {
+    return typename GemmKernel::Arguments(
+        BatchedGemmCoord(total_rows, inputs.n, inputs.k, inputs.num_experts),
+        typename GemmKernel::EpilogueOutputOp::Params(dev_scale_a, dev_scale_b, dev_moe_weights),
+        dev_a,
+        dev_b,
+        dev_c,
+        MoeParams(dev_token_ids, dev_expert_ids, dev_num_tokens_post_padded, inputs.em, inputs.topk, true));
+}
+
+inline RunResult run_fused_moe_i8_tn(const HostInputs &inputs, int device_id = 0) {
+    validate_host_inputs(inputs);
+
+    int device_count = 0;
+    check_mc_or_throw(mcGetDeviceCount(&device_count), "mcGetDeviceCount");
+    if (device_count <= 0) {
+        throw std::runtime_error("No MACA device is visible.");
+    }
+    if (device_id < 0 || device_id >= device_count) {
+        throw std::invalid_argument("device_id is out of range");
+    }
+    check_mc_or_throw(mcSetDevice(device_id), "mcSetDevice");
+
+    const int total_rows = inputs.em;
+    std::vector<int32_t> host_num_tokens_post_padded(1, inputs.em);
+    std::vector<BFloat16> host_output(static_cast<size_t>(total_rows) * inputs.n, float_to_bf16(0.0f));
+
+    int8_t *dev_a = nullptr;
+    int8_t *dev_b = nullptr;
+    float *dev_scale_a = nullptr;
+    float *dev_scale_b = nullptr;
+    float *dev_moe_weights = nullptr;
+    int *dev_token_ids = nullptr;
+    int *dev_expert_ids = nullptr;
+    int32_t *dev_num_tokens_post_padded = nullptr;
+    BFloat16 *dev_c = nullptr;
+
+    try {
+        check_mc_or_throw(mcMalloc(reinterpret_cast<void **>(&dev_a), inputs.a.size() * sizeof(int8_t)), "mcMalloc(dev_a)");
+        check_mc_or_throw(mcMalloc(reinterpret_cast<void **>(&dev_b), inputs.b_col_major.size() * sizeof(int8_t)), "mcMalloc(dev_b)");
+        check_mc_or_throw(mcMalloc(reinterpret_cast<void **>(&dev_scale_a), inputs.scale_a.size() * sizeof(float)),
+                          "mcMalloc(dev_scale_a)");
+        check_mc_or_throw(mcMalloc(reinterpret_cast<void **>(&dev_scale_b), inputs.scale_b.size() * sizeof(float)),
+                          "mcMalloc(dev_scale_b)");
+        check_mc_or_throw(mcMalloc(reinterpret_cast<void **>(&dev_moe_weights), inputs.moe_weights.size() * sizeof(float)),
+                          "mcMalloc(dev_moe_weights)");
+        check_mc_or_throw(mcMalloc(reinterpret_cast<void **>(&dev_token_ids), inputs.token_ids.size() * sizeof(int)),
+                          "mcMalloc(dev_token_ids)");
+        check_mc_or_throw(mcMalloc(reinterpret_cast<void **>(&dev_expert_ids), inputs.expert_ids.size() * sizeof(int)),
+                          "mcMalloc(dev_expert_ids)");
+        check_mc_or_throw(mcMalloc(reinterpret_cast<void **>(&dev_num_tokens_post_padded),
+                                   host_num_tokens_post_padded.size() * sizeof(int32_t)),
+                          "mcMalloc(dev_num_tokens_post_padded)");
+        check_mc_or_throw(mcMalloc(reinterpret_cast<void **>(&dev_c), host_output.size() * sizeof(BFloat16)), "mcMalloc(dev_c)");
+
+        check_mc_or_throw(mcMemcpy(dev_a, inputs.a.data(), inputs.a.size() * sizeof(int8_t), mcMemcpyHostToDevice),
+                          "mcMemcpy(dev_a)");
+        check_mc_or_throw(mcMemcpy(dev_b,
+                                   inputs.b_col_major.data(),
+                                   inputs.b_col_major.size() * sizeof(int8_t),
+                                   mcMemcpyHostToDevice),
+                          "mcMemcpy(dev_b)");
+        check_mc_or_throw(mcMemcpy(dev_scale_a,
+                                   inputs.scale_a.data(),
+                                   inputs.scale_a.size() * sizeof(float),
+                                   mcMemcpyHostToDevice),
+                          "mcMemcpy(dev_scale_a)");
+        check_mc_or_throw(mcMemcpy(dev_scale_b,
+                                   inputs.scale_b.data(),
+                                   inputs.scale_b.size() * sizeof(float),
+                                   mcMemcpyHostToDevice),
+                          "mcMemcpy(dev_scale_b)");
+        check_mc_or_throw(mcMemcpy(dev_moe_weights,
+                                   inputs.moe_weights.data(),
+                                   inputs.moe_weights.size() * sizeof(float),
+                                   mcMemcpyHostToDevice),
+                          "mcMemcpy(dev_moe_weights)");
+        check_mc_or_throw(mcMemcpy(dev_token_ids,
+                                   inputs.token_ids.data(),
+                                   inputs.token_ids.size() * sizeof(int),
+                                   mcMemcpyHostToDevice),
+                          "mcMemcpy(dev_token_ids)");
+        check_mc_or_throw(mcMemcpy(dev_expert_ids,
+                                   inputs.expert_ids.data(),
+                                   inputs.expert_ids.size() * sizeof(int),
+                                   mcMemcpyHostToDevice),
+                          "mcMemcpy(dev_expert_ids)");
+        check_mc_or_throw(mcMemcpy(dev_num_tokens_post_padded,
+                                   host_num_tokens_post_padded.data(),
+                                   host_num_tokens_post_padded.size() * sizeof(int32_t),
+                                   mcMemcpyHostToDevice),
+                          "mcMemcpy(dev_num_tokens_post_padded)");
+        check_mc_or_throw(mcMemset(dev_c, 0, host_output.size() * sizeof(BFloat16)), "mcMemset(dev_c)");
+
+        const bool topk_is_log2 = is_log2_value(inputs.topk);
+        if (topk_is_log2) {
+            using GemmKernel = DirectMoeKernel<true>;
+            auto const args = make_direct_arguments<GemmKernel>(inputs,
+                                                                total_rows,
+                                                                dev_a,
+                                                                dev_b,
+                                                                dev_scale_a,
+                                                                dev_scale_b,
+                                                                dev_moe_weights,
+                                                                dev_token_ids,
+                                                                dev_expert_ids,
+                                                                dev_num_tokens_post_padded,
+                                                                dev_c);
+            check_status_or_throw(launch<GemmKernel>(args), "direct_moe_kernel");
+        } else {
+            using GemmKernel = DirectMoeKernel<false>;
+            auto const args = make_direct_arguments<GemmKernel>(inputs,
+                                                                total_rows,
+                                                                dev_a,
+                                                                dev_b,
+                                                                dev_scale_a,
+                                                                dev_scale_b,
+                                                                dev_moe_weights,
+                                                                dev_token_ids,
+                                                                dev_expert_ids,
+                                                                dev_num_tokens_post_padded,
+                                                                dev_c);
+            check_status_or_throw(launch<GemmKernel>(args), "direct_moe_kernel");
+        }
+
+        check_mc_or_throw(mcDeviceSynchronize(), "mcDeviceSynchronize");
+        check_mc_or_throw(mcGetLastError(), "mcGetLastError");
+        check_mc_or_throw(mcMemcpy(host_output.data(),
+                                   dev_c,
+                                   host_output.size() * sizeof(BFloat16),
+                                   mcMemcpyDeviceToHost),
+                          "mcMemcpy(host_output)");
+    } catch (...) {
+        if (dev_a) {
+            mcFree(dev_a);
+        }
+        if (dev_b) {
+            mcFree(dev_b);
+        }
+        if (dev_scale_a) {
+            mcFree(dev_scale_a);
+        }
+        if (dev_scale_b) {
+            mcFree(dev_scale_b);
+        }
+        if (dev_moe_weights) {
+            mcFree(dev_moe_weights);
+        }
+        if (dev_token_ids) {
+            mcFree(dev_token_ids);
+        }
+        if (dev_expert_ids) {
+            mcFree(dev_expert_ids);
+        }
+        if (dev_num_tokens_post_padded) {
+            mcFree(dev_num_tokens_post_padded);
+        }
+        if (dev_c) {
+            mcFree(dev_c);
+        }
+        throw;
+    }
+
+    check_mc_or_throw(mcFree(dev_a), "mcFree(dev_a)");
+    check_mc_or_throw(mcFree(dev_b), "mcFree(dev_b)");
+    check_mc_or_throw(mcFree(dev_scale_a), "mcFree(dev_scale_a)");
+    check_mc_or_throw(mcFree(dev_scale_b), "mcFree(dev_scale_b)");
+    check_mc_or_throw(mcFree(dev_moe_weights), "mcFree(dev_moe_weights)");
+    check_mc_or_throw(mcFree(dev_token_ids), "mcFree(dev_token_ids)");
+    check_mc_or_throw(mcFree(dev_expert_ids), "mcFree(dev_expert_ids)");
+    check_mc_or_throw(mcFree(dev_num_tokens_post_padded), "mcFree(dev_num_tokens_post_padded)");
+    check_mc_or_throw(mcFree(dev_c), "mcFree(dev_c)");
+
+    RunResult result;
+    result.rows = total_rows;
+    result.cols = inputs.n;
+    result.output = std::move(host_output);
+    return result;
+}
+
+}  // namespace fused_moe_i8_tn

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_types.h

@@ -0,0 +1,167 @@
+#pragma once
+
+#include <cstdint>
+#include <cstring>
+
+#include <maca_bfloat16.h>
+#include <mc_runtime_api.h>
+#include <mc_runtime_types.h>
+
+namespace fused_moe_i8_tn {
+
+#if defined(__MXCC__) || (defined(__clang__) && defined(__MACA__))
+#define FUSED_MOE_HOST_DEVICE __forceinline__ __device__ __host__
+#define FUSED_MOE_DEVICE __forceinline__ __device__
+#else
+#define FUSED_MOE_HOST_DEVICE inline
+#define FUSED_MOE_DEVICE inline
+#endif
+
+enum class Status {
+    kSuccess,
+    kErrorInternal,
+};
+
+inline const char *get_status_string(Status status) {
+    switch (status) {
+        case Status::kSuccess:
+            return "Success";
+        case Status::kErrorInternal:
+            return "Error Internal";
+    }
+    return "Invalid status";
+}
+
+struct alignas(2) BFloat16 {
+    uint16_t storage;
+
+    FUSED_MOE_HOST_DEVICE
+    BFloat16() : storage(0) {}
+
+    FUSED_MOE_HOST_DEVICE
+    explicit BFloat16(float x) {
+#if defined(__MACA_ARCH__)
+        auto tmp = __float2bfloat16(x);
+        storage = reinterpret_cast<uint16_t const &>(tmp);
+#else
+        uint32_t bits;
+        std::memcpy(&bits, &x, sizeof(bits));
+        bits += ((bits >> 16) & 1) + 0x7fff;
+        storage = static_cast<uint16_t>(bits >> 16);
+#endif
+    }
+
+    FUSED_MOE_HOST_DEVICE
+    operator float() const {
+#if defined(__MACA_ARCH__)
+        __maca_bfloat16_raw raw;
+        raw.x = storage;
+        return __bfloat162float(__maca_bfloat16(raw));
+#else
+        uint32_t bits = static_cast<uint32_t>(storage) << 16;
+        float out;
+        std::memcpy(&out, &bits, sizeof(out));
+        return out;
+#endif
+    }
+};
+
+struct BatchedGemmCoord {
+    int m_;
+    int n_;
+    int k_;
+    int batch_;
+
+    FUSED_MOE_HOST_DEVICE
+    BatchedGemmCoord() : m_(0), n_(0), k_(0), batch_(0) {}
+
+    FUSED_MOE_HOST_DEVICE
+    BatchedGemmCoord(int m, int n, int k, int batch) : m_(m), n_(n), k_(k), batch_(batch) {}
+
+    FUSED_MOE_HOST_DEVICE
+    int m() const { return m_; }
+
+    FUSED_MOE_HOST_DEVICE
+    int n() const { return n_; }
+
+    FUSED_MOE_HOST_DEVICE
+    int k() const { return k_; }
+
+    FUSED_MOE_HOST_DEVICE
+    int batch() const { return batch_; }
+};
+
+struct MoeParams {
+    int *token_ids;
+    int *expert_ids;
+    int *num_tokens_post_padded_ptr;
+    int32_t EM;
+    int32_t topk;
+    bool mul_weight;
+    int topk_bits;
+
+    FUSED_MOE_HOST_DEVICE
+    MoeParams()
+        : token_ids(nullptr),
+          expert_ids(nullptr),
+          num_tokens_post_padded_ptr(nullptr),
+          EM(0),
+          topk(0),
+          mul_weight(false),
+          topk_bits(0) {}
+
+    FUSED_MOE_HOST_DEVICE
+    MoeParams(int *token_ids_,
+              int *expert_ids_,
+              int *num_tokens_post_padded_ptr_,
+              int EM_,
+              int topk_,
+              bool mul_weight_)
+        : token_ids(token_ids_),
+          expert_ids(expert_ids_),
+          num_tokens_post_padded_ptr(num_tokens_post_padded_ptr_),
+          EM(EM_),
+          topk(topk_),
+          mul_weight(mul_weight_),
+          topk_bits(0) {
+        int num = topk_;
+        while (num >>= 1) {
+            ++topk_bits;
+        }
+    }
+};
+
+struct EpilogueOutputOp {
+    using ElementOutput = BFloat16;
+    using ElementCompute = float;
+    static constexpr int kCount = 2;
+    static constexpr bool MUL_WEIGHTS = true;
+
+    struct Params {
+        ElementCompute const *scale_a;
+        ElementCompute const *scale_b;
+        ElementCompute const *moe_weights;
+
+        FUSED_MOE_HOST_DEVICE
+        Params() : scale_a(nullptr), scale_b(nullptr), moe_weights(nullptr) {}
+
+        FUSED_MOE_HOST_DEVICE
+        Params(ElementCompute const *scale_a_,
+               ElementCompute const *scale_b_,
+               ElementCompute const *moe_weights_)
+            : scale_a(scale_a_), scale_b(scale_b_), moe_weights(moe_weights_) {}
+    };
+
+    ElementCompute const *scale_a_;
+    ElementCompute const *scale_b_;
+    ElementCompute const *moe_weights_;
+
+    FUSED_MOE_HOST_DEVICE
+    EpilogueOutputOp() : scale_a_(nullptr), scale_b_(nullptr), moe_weights_(nullptr) {}
+
+    FUSED_MOE_HOST_DEVICE
+    explicit EpilogueOutputOp(Params const &params)
+        : scale_a_(params.scale_a), scale_b_(params.scale_b), moe_weights_(params.moe_weights) {}
+};
+
+}  // namespace fused_moe_i8_tn

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/i8_tn_256x256x128_raw_arrays/CMakeLists.txt

@@ -0,0 +1,46 @@
+cmake_minimum_required(VERSION 3.20)
+
+project(i8_tn_256x256x128_raw_arrays LANGUAGES NONE)
+
+set(MACA_PATH "$ENV{MACA_PATH}" CACHE PATH "Path to MACA SDK")
+if(NOT MACA_PATH)
+  set(MACA_PATH "/opt/maca")
+endif()
+
+find_program(MXCC
+  NAMES mxcc
+  PATHS "${MACA_PATH}/mxgpu_llvm/bin"
+  NO_DEFAULT_PATH)
+
+if(NOT MXCC)
+  message(FATAL_ERROR "mxcc not found under ${MACA_PATH}/mxgpu_llvm/bin")
+endif()
+
+set(STANDALONE_ROOT "${CMAKE_CURRENT_SOURCE_DIR}")
+set(SRC "${STANDALONE_ROOT}/src/test.cpp")
+set(BIN "${CMAKE_CURRENT_BINARY_DIR}/i8_tn_256x256x128_raw_arrays_test")
+
+add_custom_command(
+  OUTPUT "${BIN}"
+  COMMAND "${MXCC}"
+          -std=c++17
+          -xmaca
+          -I"${STANDALONE_ROOT}/include"
+          -I"${MACA_PATH}/include"
+          "${SRC}"
+          -L"${MACA_PATH}/lib"
+          -lmcruntime
+          -o "${BIN}"
+  DEPENDS
+          "${SRC}"
+          "${STANDALONE_ROOT}/include/standalone_maca_kernel_utils.hpp"
+          "${STANDALONE_ROOT}/include/gemm_i8_tn_256x256x128_raw_arrays.hpp"
+  VERBATIM)
+
+add_custom_target(build_i8_tn_256x256x128_raw_arrays_test ALL DEPENDS "${BIN}")
+
+add_custom_target(
+  run
+  COMMAND "${BIN}"
+  DEPENDS "${BIN}"
+  USES_TERMINAL)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/i8_tn_256x256x128_raw_arrays/include/gemm_i8_tn_256x256x128_raw_arrays.hpp

@@ -0,0 +1,368 @@
+#pragma once
+
+#include "standalone_maca_kernel_utils.hpp"
+
+namespace standalone_i8_tn_256x256x128_raw_arrays {
+
+struct KernelConfig {
+    static constexpr int kTileM = 256;
+    static constexpr int kTileN = 256;
+    static constexpr int kTileK = 128;
+    static constexpr int kThreadCount = 512;
+    static constexpr int kWaveSize = 64;
+    static constexpr int kWaveCount = kThreadCount / kWaveSize;
+    static constexpr int kRowsPerWaveGroup = 64;
+    static constexpr int kColsPerWaveGroup = 128;
+    static constexpr int kRowsPerMicroTile = 16;
+    static constexpr int kColsPerMicroTile = 32;
+    static constexpr int kColumnGroupsPerWave = 2;
+    static constexpr int kColBlocksPerWaveGroup = 2;
+    static constexpr int kRowBlocksPerWave = 4;
+    static constexpr int kOutputVectorsPerMicroTile = 4;
+    static constexpr int kOutputVectorsPerColBlock = kRowBlocksPerWave * kOutputVectorsPerMicroTile;
+    static constexpr int kElementsPer128b = 16;
+    static constexpr int kSharedBytesA = kTileM * kTileK;
+    static constexpr int kSharedBytesB = kTileN * kTileK;
+    static constexpr int kSmemSize = kSharedBytesA + kSharedBytesB;
+};
+
+using StoreVector = __NATIVE_VECTOR__(2, int32_t);
+using LdsTypeI8Mma = __NATIVE_VECTOR__(4, int32_t);
+using AbTypeI8Mma = int32_t;
+using AccumTypeI8Mma = __NATIVE_VECTOR__(4, int32_t);
+
+struct ThreadContext {
+    int a_store_offset[4];
+    int a_row_local[4];
+    int a_load_k[4];
+    int b_store_offset[4];
+    int b_col_local[4];
+    int b_load_k[4];
+    int a_lds_offset[KernelConfig::kRowBlocksPerWave][2];
+    int b_lds_offset[KernelConfig::kColumnGroupsPerWave][KernelConfig::kColBlocksPerWaveGroup][2][2];
+};
+
+__forceinline__ __device__ int swizzled_slot8(int row_or_col, int q) {
+    return q ^ (row_or_col & 7);
+}
+
+__forceinline__ __device__ void clear_accumulators(
+    AccumTypeI8Mma (&accum)[KernelConfig::kColumnGroupsPerWave]
+                            [KernelConfig::kColBlocksPerWaveGroup]
+                            [KernelConfig::kRowBlocksPerWave][2]) {
+#pragma unroll
+    for (int column_group = 0; column_group < KernelConfig::kColumnGroupsPerWave; ++column_group) {
+#pragma unroll
+        for (int col_block = 0; col_block < KernelConfig::kColBlocksPerWaveGroup; ++col_block) {
+#pragma unroll
+            for (int row_block = 0; row_block < KernelConfig::kRowBlocksPerWave; ++row_block) {
+#pragma unroll
+                for (int half = 0; half < 2; ++half) {
+#pragma unroll
+                    for (int i = 0; i < KernelConfig::kOutputVectorsPerMicroTile; ++i) {
+                        accum[column_group][col_block][row_block][half][i] = 0;
+                    }
+                }
+            }
+        }
+    }
+}
+
+__forceinline__ __device__ void build_thread_context(ThreadContext &ctx,
+                                                     int tid,
+                                                     int wave_row_group,
+                                                     int wave_col_group,
+                                                     int lane16,
+                                                     int lane_q_block,
+                                                     int K) {
+    const int row_or_col_local = tid >> 3;
+    const int slot8_phys = tid & 7;
+
+#pragma unroll
+    for (int pass = 0; pass < 4; ++pass) {
+        const int row_or_col = pass * 64 + row_or_col_local;
+        const int q = slot8_phys ^ (row_or_col & 7);
+        ctx.a_store_offset[pass] = row_or_col * 128 + slot8_phys * KernelConfig::kElementsPer128b;
+        ctx.a_row_local[pass] = row_or_col;
+        ctx.a_load_k[pass] = q * KernelConfig::kElementsPer128b;
+        ctx.b_store_offset[pass] = row_or_col * 128 + slot8_phys * KernelConfig::kElementsPer128b;
+        ctx.b_col_local[pass] = row_or_col;
+        ctx.b_load_k[pass] = q * KernelConfig::kElementsPer128b;
+    }
+
+    const int q_values[2] = {lane_q_block, lane_q_block + 4};
+
+#pragma unroll
+    for (int row_block = 0; row_block < KernelConfig::kRowBlocksPerWave; ++row_block) {
+        const int a_row =
+            wave_row_group * KernelConfig::kRowsPerWaveGroup + row_block * KernelConfig::kRowsPerMicroTile + lane16;
+#pragma unroll
+        for (int half = 0; half < 2; ++half) {
+            const int slot8_phys_local = swizzled_slot8(a_row, q_values[half]);
+            ctx.a_lds_offset[row_block][half] = a_row * 128 + slot8_phys_local * KernelConfig::kElementsPer128b;
+        }
+    }
+
+    const int lds_k_b = (lane_q_block ^ (lane16 & 3)) * KernelConfig::kElementsPer128b;
+#pragma unroll
+    for (int column_group = 0; column_group < KernelConfig::kColumnGroupsPerWave; ++column_group) {
+#pragma unroll
+        for (int col_block = 0; col_block < KernelConfig::kColBlocksPerWaveGroup; ++col_block) {
+            const int b_chunk = wave_col_group * 4 + column_group * 2 + col_block;
+            const int b_col0 = b_chunk * KernelConfig::kColsPerMicroTile + lane16;
+            const int b_col1 = b_col0 + 16;
+            const int cols[2] = {b_col0, b_col1};
+#pragma unroll
+            for (int half = 0; half < 2; ++half) {
+#pragma unroll
+                for (int which = 0; which < 2; ++which) {
+                    const int row_or_col = cols[which];
+                    const int slot8_phys_local = swizzled_slot8(row_or_col, q_values[half]);
+                    ctx.b_lds_offset[column_group][col_block][half][which] =
+                        row_or_col * 128 + slot8_phys_local * KernelConfig::kElementsPer128b;
+                }
+            }
+        }
+    }
+}
+
+template <int Pass>
+__forceinline__ __device__ void load_a_pass(int8_t *smem_a,
+                                            const int8_t *a_ptr,
+                                            int K,
+                                            int global_row_base,
+                                            int k_tile_base,
+                                            ThreadContext const &ctx) {
+    const int row = global_row_base + ctx.a_row_local[Pass];
+    __builtin_mxc_ldg_b128_bsm(
+        smem_a + ctx.a_store_offset[Pass],
+        const_cast<void *>(reinterpret_cast<void const *>(
+            a_ptr + static_cast<size_t>(row) * K + k_tile_base + ctx.a_load_k[Pass])),
+        0,
+        -1,
+        true,
+        true,
+        false,
+        true);
+}
+
+template <int Pass>
+__forceinline__ __device__ void load_b_pass(int8_t *smem_b,
+                                            const int8_t *b_ptr,
+                                            int K,
+                                            int global_col_base,
+                                            int k_tile_base,
+                                            ThreadContext const &ctx) {
+    const int col = global_col_base + ctx.b_col_local[Pass];
+    __builtin_mxc_ldg_b128_bsm(
+        smem_b + ctx.b_store_offset[Pass],
+        const_cast<void *>(reinterpret_cast<void const *>(
+            b_ptr + static_cast<size_t>(col) * K + k_tile_base + ctx.b_load_k[Pass])),
+        0,
+        -1,
+        true,
+        true,
+        false,
+        true);
+}
+
+__forceinline__ __device__ void wait_for_tile_load() {
+    standalone_arrive_gvmcnt(0);
+    __builtin_mxc_barrier_inst();
+}
+
+__forceinline__ __device__ void mma_on_pack16(AccumTypeI8Mma &accum_left,
+                                              AccumTypeI8Mma &accum_right,
+                                              LdsTypeI8Mma const &a_pack,
+                                              LdsTypeI8Mma const &b0_pack,
+                                              LdsTypeI8Mma const &b1_pack) {
+    AbTypeI8Mma const *a_frag = reinterpret_cast<AbTypeI8Mma const *>(&a_pack);
+    AbTypeI8Mma const *b0_frag = reinterpret_cast<AbTypeI8Mma const *>(&b0_pack);
+    AbTypeI8Mma const *b1_frag = reinterpret_cast<AbTypeI8Mma const *>(&b1_pack);
+#pragma unroll
+    for (int step = 0; step < 4; ++step) {
+        accum_left = STANDALONE_BUILTIN_MMA_16X16X16_I8(a_frag[step], b0_frag[step], accum_left);
+        accum_right = STANDALONE_BUILTIN_MMA_16X16X16_I8(a_frag[step], b1_frag[step], accum_right);
+    }
+}
+
+__forceinline__ __device__ void consume_full_k128_from_shared(
+    AccumTypeI8Mma (&accum)[KernelConfig::kColumnGroupsPerWave]
+                            [KernelConfig::kColBlocksPerWaveGroup]
+                            [KernelConfig::kRowBlocksPerWave][2],
+    int8_t const *smem_a,
+    int8_t const *smem_b,
+    ThreadContext const &ctx) {
+    LdsTypeI8Mma b0_pack_low[KernelConfig::kColumnGroupsPerWave][KernelConfig::kColBlocksPerWaveGroup];
+    LdsTypeI8Mma b1_pack_low[KernelConfig::kColumnGroupsPerWave][KernelConfig::kColBlocksPerWaveGroup];
+    LdsTypeI8Mma b0_pack_high[KernelConfig::kColumnGroupsPerWave][KernelConfig::kColBlocksPerWaveGroup];
+    LdsTypeI8Mma b1_pack_high[KernelConfig::kColumnGroupsPerWave][KernelConfig::kColBlocksPerWaveGroup];
+
+#pragma unroll
+    for (int column_group = 0; column_group < KernelConfig::kColumnGroupsPerWave; ++column_group) {
+#pragma unroll
+        for (int col_block = 0; col_block < KernelConfig::kColBlocksPerWaveGroup; ++col_block) {
+            STANDALONE_LDS(
+                b0_pack_low[column_group][col_block],
+                *const_cast<int8_t *>(smem_b + ctx.b_lds_offset[column_group][col_block][0][0]),
+                LdsTypeI8Mma);
+            STANDALONE_LDS(
+                b1_pack_low[column_group][col_block],
+                *const_cast<int8_t *>(smem_b + ctx.b_lds_offset[column_group][col_block][0][1]),
+                LdsTypeI8Mma);
+            STANDALONE_LDS(
+                b0_pack_high[column_group][col_block],
+                *const_cast<int8_t *>(smem_b + ctx.b_lds_offset[column_group][col_block][1][0]),
+                LdsTypeI8Mma);
+            STANDALONE_LDS(
+                b1_pack_high[column_group][col_block],
+                *const_cast<int8_t *>(smem_b + ctx.b_lds_offset[column_group][col_block][1][1]),
+                LdsTypeI8Mma);
+        }
+    }
+
+#pragma unroll
+    for (int row_block = 0; row_block < KernelConfig::kRowBlocksPerWave; ++row_block) {
+        LdsTypeI8Mma a_pack_low;
+        LdsTypeI8Mma a_pack_high;
+        STANDALONE_LDS(a_pack_low, *const_cast<int8_t *>(smem_a + ctx.a_lds_offset[row_block][0]), LdsTypeI8Mma);
+        STANDALONE_LDS(a_pack_high, *const_cast<int8_t *>(smem_a + ctx.a_lds_offset[row_block][1]), LdsTypeI8Mma);
+
+#pragma unroll
+        for (int column_group = 0; column_group < KernelConfig::kColumnGroupsPerWave; ++column_group) {
+#pragma unroll
+            for (int col_block = 0; col_block < KernelConfig::kColBlocksPerWaveGroup; ++col_block) {
+                mma_on_pack16(
+                    accum[column_group][col_block][row_block][0],
+                    accum[column_group][col_block][row_block][1],
+                    a_pack_low,
+                    b0_pack_low[column_group][col_block],
+                    b1_pack_low[column_group][col_block]);
+                mma_on_pack16(
+                    accum[column_group][col_block][row_block][0],
+                    accum[column_group][col_block][row_block][1],
+                    a_pack_high,
+                    b0_pack_high[column_group][col_block],
+                    b1_pack_high[column_group][col_block]);
+            }
+        }
+    }
+}
+
+__forceinline__ __device__ void store_accumulators_to_global(
+    int32_t *D,
+    int M,
+    int N,
+    int bidx,
+    int bidy,
+    int bidz,
+    AccumTypeI8Mma const (&accum)[KernelConfig::kColumnGroupsPerWave]
+                                  [KernelConfig::kColBlocksPerWaveGroup]
+                                  [KernelConfig::kRowBlocksPerWave][2]) {
+    const int row_limit = ((M - bidy * KernelConfig::kTileM) < KernelConfig::kTileM)
+                              ? (M - bidy * KernelConfig::kTileM)
+                              : KernelConfig::kTileM;
+    const int col_limit = ((N - bidx * KernelConfig::kTileN) < KernelConfig::kTileN)
+                              ? (N - bidx * KernelConfig::kTileN)
+                              : KernelConfig::kTileN;
+    const int tid = threadIdx.x;
+    const int wave_idx = tid / KernelConfig::kWaveSize;
+    const int lane_idx = tid % KernelConfig::kWaveSize;
+    const int lane_col = lane_idx % 16;
+    const int lane_row_group = lane_idx / 16;
+    const int wave_row_base = (wave_idx / 2) * KernelConfig::kRowsPerWaveGroup;
+    const int wave_col_base = (wave_idx % 2) * KernelConfig::kColsPerWaveGroup;
+
+    int32_t *d_ptr = D + static_cast<size_t>(bidz) * M * N;
+
+#pragma unroll
+    for (int column_group = 0; column_group < KernelConfig::kColumnGroupsPerWave; ++column_group) {
+#pragma unroll
+        for (int col_block = 0; col_block < KernelConfig::kColBlocksPerWaveGroup; ++col_block) {
+            const int microtile_col_base =
+                wave_col_base + column_group * 64 + col_block * KernelConfig::kColsPerMicroTile;
+            const int col0 = microtile_col_base + lane_col;
+            const int col1 = col0 + 16;
+
+#pragma unroll
+            for (int row_block = 0; row_block < KernelConfig::kRowBlocksPerWave; ++row_block) {
+                const int output_base =
+                    column_group * (KernelConfig::kColBlocksPerWaveGroup * KernelConfig::kOutputVectorsPerColBlock) +
+                    col_block * KernelConfig::kOutputVectorsPerColBlock +
+                    row_block * KernelConfig::kOutputVectorsPerMicroTile;
+#pragma unroll
+                for (int i = 0; i < KernelConfig::kOutputVectorsPerMicroTile; ++i) {
+                    const int row = wave_row_base + row_block * KernelConfig::kRowsPerMicroTile +
+                                    lane_row_group * 4 + i;
+                    if (row >= row_limit) {
+                        continue;
+                    }
+                    const size_t row_offset = static_cast<size_t>(bidy * KernelConfig::kTileM + row) * N +
+                                              bidx * KernelConfig::kTileN;
+                    if (col0 < col_limit) {
+                        d_ptr[row_offset + col0] = accum[column_group][col_block][row_block][0][i];
+                    }
+                    if (col1 < col_limit) {
+                        d_ptr[row_offset + col1] = accum[column_group][col_block][row_block][1][i];
+                    }
+                }
+            }
+        }
+    }
+}
+
+__global__ void gemm_i8_tn_256x256x128_raw_arrays_kernel(const int8_t *A,
+                                                         const int8_t *B,
+                                                         int32_t *D,
+                                                         int M,
+                                                         int N,
+                                                         int K) {
+    __shared__ int8_t smem_data[KernelConfig::kSmemSize];
+
+    int8_t *smem_a = smem_data;
+    int8_t *smem_b = smem_data + KernelConfig::kSharedBytesA;
+
+    const int tid = threadIdx.x;
+    const int bidx = blockIdx.y;
+    const int bidy = blockIdx.x;
+    const int bidz = blockIdx.z;
+
+    const int wave_idx = tid / KernelConfig::kWaveSize;
+    const int lane_idx = tid % KernelConfig::kWaveSize;
+    const int wave_row_group = wave_idx / 2;
+    const int wave_col_group = wave_idx % 2;
+    const int lane16 = lane_idx % 16;
+    const int lane_q_block = lane_idx / 16;
+    const int global_row_base = bidy * KernelConfig::kTileM;
+    const int global_col_base = bidx * KernelConfig::kTileN;
+
+    const int8_t *a_ptr = A + static_cast<size_t>(bidz) * M * K;
+    const int8_t *b_ptr = B + static_cast<size_t>(bidz) * N * K;
+
+    AccumTypeI8Mma accum[KernelConfig::kColumnGroupsPerWave]
+                        [KernelConfig::kColBlocksPerWaveGroup]
+                        [KernelConfig::kRowBlocksPerWave][2];
+    ThreadContext ctx;
+    clear_accumulators(accum);
+    build_thread_context(ctx, tid, wave_row_group, wave_col_group, lane16, lane_q_block, K);
+
+    for (int k_tile = 0; k_tile < K; k_tile += KernelConfig::kTileK) {
+        load_a_pass<0>(smem_a, a_ptr, K, global_row_base, k_tile, ctx);
+        load_a_pass<1>(smem_a, a_ptr, K, global_row_base, k_tile, ctx);
+        load_a_pass<2>(smem_a, a_ptr, K, global_row_base, k_tile, ctx);
+        load_a_pass<3>(smem_a, a_ptr, K, global_row_base, k_tile, ctx);
+
+        load_b_pass<0>(smem_b, b_ptr, K, global_col_base, k_tile, ctx);
+        load_b_pass<1>(smem_b, b_ptr, K, global_col_base, k_tile, ctx);
+        load_b_pass<2>(smem_b, b_ptr, K, global_col_base, k_tile, ctx);
+        load_b_pass<3>(smem_b, b_ptr, K, global_col_base, k_tile, ctx);
+
+        wait_for_tile_load();
+        consume_full_k128_from_shared(accum, smem_a, smem_b, ctx);
+        __syncthreadshared();
+    }
+
+    store_accumulators_to_global(D, M, N, bidx, bidy, bidz, accum);
+}
+
+}  // namespace standalone_i8_tn_256x256x128_raw_arrays

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/i8_tn_256x256x128_raw_arrays/include/standalone_maca_kernel_utils.hpp

@@ -0,0 +1,18 @@
+#pragma once
+
+#include <mc_runtime.h>
+
+#define standalone_arrive_gvmcnt(count) __builtin_mxc_arrive(64 + count)
+
+#if defined(__MACA_ARCH__) && (__MACA_ARCH__ == 1000 || __MACA_ARCH__ == 1089)
+#define STANDALONE_BUILTIN_MMA_16X16X16_I8(a, b, c) __builtin_mxc_mma_16x16x16i8(a, b, c)
+#else
+#define STANDALONE_BUILTIN_MMA_16X16X16_I8(a, b, c) 0
+#endif
+
+#define standalone_cp_async_fenc() asm(";--------------")
+
+#define STANDALONE_LDS(dst, src, ldstype)                                                          \
+    standalone_cp_async_fenc();                                                                    \
+    *reinterpret_cast<ldstype *>(&(dst)) = *reinterpret_cast<ldstype *>(&(src));                  \
+    standalone_cp_async_fenc()

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/i8_tn_256x256x128_raw_arrays/src/test.cpp

@@ -0,0 +1,229 @@
+#include <mc_runtime.h>
+
+#include <cstdint>
+#include <cstdlib>
+#include <iomanip>
+#include <iostream>
+#include <vector>
+
+#include "../include/gemm_i8_tn_256x256x128_raw_arrays.hpp"
+
+namespace {
+
+using namespace standalone_i8_tn_256x256x128_raw_arrays;
+
+struct BenchmarkResult {
+    float avg_ms = 0.0f;
+    double tflops = 0.0;
+    int warmup_iterations = 0;
+    int measured_iterations = 0;
+};
+
+void check_mc(mcError_t status, const char *expr) {
+    if (status != mcSuccess) {
+        std::cerr << expr << " failed: " << mcGetErrorString(status) << '\n';
+        std::exit(EXIT_FAILURE);
+    }
+}
+
+void fill_row_major_a(std::vector<int8_t> &a, int m, int k) {
+    a.resize(static_cast<size_t>(m) * k);
+    for (int row = 0; row < m; ++row) {
+        for (int kk = 0; kk < k; ++kk) {
+            a[static_cast<size_t>(row) * k + kk] =
+                static_cast<int8_t>(((row * 13 + kk * 7 + 5) % 9) - 4);
+        }
+    }
+}
+
+void fill_col_major_b(std::vector<int8_t> &b, int n, int k) {
+    b.resize(static_cast<size_t>(n) * k);
+    for (int col = 0; col < n; ++col) {
+        for (int kk = 0; kk < k; ++kk) {
+            b[static_cast<size_t>(col) * k + kk] =
+                static_cast<int8_t>(((kk * 11 + col * 5 + 3) % 7) - 3);
+        }
+    }
+}
+
+std::vector<int32_t> reference_gemm_tn(const std::vector<int8_t> &a,
+                                       const std::vector<int8_t> &b_col_major,
+                                       int m,
+                                       int n,
+                                       int k) {
+    std::vector<int32_t> out(static_cast<size_t>(m) * n, 0);
+    for (int row = 0; row < m; ++row) {
+        for (int col = 0; col < n; ++col) {
+            int32_t acc = 0;
+            for (int kk = 0; kk < k; ++kk) {
+                acc += static_cast<int32_t>(a[static_cast<size_t>(row) * k + kk]) *
+                       static_cast<int32_t>(b_col_major[static_cast<size_t>(col) * k + kk]);
+            }
+            out[static_cast<size_t>(row) * n + col] = acc;
+        }
+    }
+    return out;
+}
+
+bool run_case(int m, int n, int k, const char *tag) {
+    std::vector<int8_t> host_a;
+    std::vector<int8_t> host_b;
+    std::vector<int32_t> host_d(static_cast<size_t>(m) * n, -1);
+    fill_row_major_a(host_a, m, k);
+    fill_col_major_b(host_b, n, k);
+    const std::vector<int32_t> reference = reference_gemm_tn(host_a, host_b, m, n, k);
+
+    int8_t *dev_a = nullptr;
+    int8_t *dev_b = nullptr;
+    int32_t *dev_d = nullptr;
+    check_mc(mcMalloc(reinterpret_cast<void **>(&dev_a), host_a.size() * sizeof(int8_t)), "mcMalloc(dev_a)");
+    check_mc(mcMalloc(reinterpret_cast<void **>(&dev_b), host_b.size() * sizeof(int8_t)), "mcMalloc(dev_b)");
+    check_mc(mcMalloc(reinterpret_cast<void **>(&dev_d), host_d.size() * sizeof(int32_t)), "mcMalloc(dev_d)");
+
+    check_mc(mcMemcpy(dev_a, host_a.data(), host_a.size() * sizeof(int8_t), mcMemcpyHostToDevice), "mcMemcpy(dev_a)");
+    check_mc(mcMemcpy(dev_b, host_b.data(), host_b.size() * sizeof(int8_t), mcMemcpyHostToDevice), "mcMemcpy(dev_b)");
+    check_mc(mcMemset(dev_d, 0, host_d.size() * sizeof(int32_t)), "mcMemset(dev_d)");
+
+    dim3 grid((m + KernelConfig::kTileM - 1) / KernelConfig::kTileM,
+              (n + KernelConfig::kTileN - 1) / KernelConfig::kTileN,
+              1);
+    gemm_i8_tn_256x256x128_raw_arrays_kernel<<<grid, KernelConfig::kThreadCount>>>(dev_a, dev_b, dev_d, m, n, k);
+    check_mc(mcDeviceSynchronize(), "mcDeviceSynchronize");
+    check_mc(mcGetLastError(), "mcGetLastError");
+
+    check_mc(mcMemcpy(host_d.data(), dev_d, host_d.size() * sizeof(int32_t), mcMemcpyDeviceToHost), "mcMemcpy(host_d)");
+
+    check_mc(mcFree(dev_a), "mcFree(dev_a)");
+    check_mc(mcFree(dev_b), "mcFree(dev_b)");
+    check_mc(mcFree(dev_d), "mcFree(dev_d)");
+
+    size_t mismatch_count = 0;
+    size_t first_bad = 0;
+    for (size_t i = 0; i < host_d.size(); ++i) {
+        if (host_d[i] != reference[i]) {
+            if (mismatch_count == 0) {
+                first_bad = i;
+            }
+            ++mismatch_count;
+        }
+    }
+
+    if (mismatch_count != 0) {
+        const int row = static_cast<int>(first_bad / n);
+        const int col = static_cast<int>(first_bad % n);
+        std::cerr << tag << " validation failed. mismatches=" << mismatch_count
+                  << ", first mismatch at (" << row << ", " << col << ")"
+                  << ", got=" << host_d[first_bad]
+                  << ", expected=" << reference[first_bad] << '\n';
+        return false;
+    }
+
+    std::cout << tag << " passed: M=" << m
+              << ", N=" << n
+              << ", K=" << k
+              << ", sample D[0]=" << host_d[0]
+              << ", D[last]=" << host_d.back() << '\n';
+    return true;
+}
+
+double compute_tflops(int m, int n, int k, float avg_ms) {
+    if (avg_ms <= 0.0f) {
+        return 0.0;
+    }
+    const double operations = 2.0 * static_cast<double>(m) * n * k;
+    return operations / (static_cast<double>(avg_ms) * 1.0e9);
+}
+
+BenchmarkResult run_benchmark(int m, int n, int k, int warmup_iterations, int measured_iterations) {
+    std::vector<int8_t> host_a;
+    std::vector<int8_t> host_b;
+    std::vector<int32_t> host_d(static_cast<size_t>(m) * n, 0);
+    fill_row_major_a(host_a, m, k);
+    fill_col_major_b(host_b, n, k);
+
+    int8_t *dev_a = nullptr;
+    int8_t *dev_b = nullptr;
+    int32_t *dev_d = nullptr;
+    check_mc(mcMalloc(reinterpret_cast<void **>(&dev_a), host_a.size() * sizeof(int8_t)), "mcMalloc(dev_a)");
+    check_mc(mcMalloc(reinterpret_cast<void **>(&dev_b), host_b.size() * sizeof(int8_t)), "mcMalloc(dev_b)");
+    check_mc(mcMalloc(reinterpret_cast<void **>(&dev_d), host_d.size() * sizeof(int32_t)), "mcMalloc(dev_d)");
+
+    check_mc(mcMemcpy(dev_a, host_a.data(), host_a.size() * sizeof(int8_t), mcMemcpyHostToDevice), "mcMemcpy(dev_a)");
+    check_mc(mcMemcpy(dev_b, host_b.data(), host_b.size() * sizeof(int8_t), mcMemcpyHostToDevice), "mcMemcpy(dev_b)");
+    check_mc(mcMemset(dev_d, 0, host_d.size() * sizeof(int32_t)), "mcMemset(dev_d)");
+
+    dim3 grid((m + KernelConfig::kTileM - 1) / KernelConfig::kTileM,
+              (n + KernelConfig::kTileN - 1) / KernelConfig::kTileN,
+              1);
+
+    for (int iter = 0; iter < warmup_iterations; ++iter) {
+        gemm_i8_tn_256x256x128_raw_arrays_kernel<<<grid, KernelConfig::kThreadCount>>>(dev_a, dev_b, dev_d, m, n, k);
+    }
+    check_mc(mcDeviceSynchronize(), "mcDeviceSynchronize(warmup)");
+
+    mcEvent_t start;
+    mcEvent_t stop;
+    check_mc(mcEventCreate(&start), "mcEventCreate(start)");
+    check_mc(mcEventCreate(&stop), "mcEventCreate(stop)");
+    check_mc(mcEventRecord(start), "mcEventRecord(start)");
+    for (int iter = 0; iter < measured_iterations; ++iter) {
+        gemm_i8_tn_256x256x128_raw_arrays_kernel<<<grid, KernelConfig::kThreadCount>>>(dev_a, dev_b, dev_d, m, n, k);
+    }
+    check_mc(mcEventRecord(stop), "mcEventRecord(stop)");
+    check_mc(mcEventSynchronize(stop), "mcEventSynchronize(stop)");
+    check_mc(mcGetLastError(), "mcGetLastError");
+
+    float elapsed_ms = 0.0f;
+    check_mc(mcEventElapsedTime(&elapsed_ms, start, stop), "mcEventElapsedTime");
+    check_mc(mcEventDestroy(start), "mcEventDestroy(start)");
+    check_mc(mcEventDestroy(stop), "mcEventDestroy(stop)");
+
+    check_mc(mcFree(dev_a), "mcFree(dev_a)");
+    check_mc(mcFree(dev_b), "mcFree(dev_b)");
+    check_mc(mcFree(dev_d), "mcFree(dev_d)");
+
+    BenchmarkResult result;
+    result.warmup_iterations = warmup_iterations;
+    result.measured_iterations = measured_iterations;
+    result.avg_ms = elapsed_ms / static_cast<float>(measured_iterations);
+    result.tflops = compute_tflops(m, n, k, result.avg_ms);
+    return result;
+}
+
+}  // namespace
+
+int main() {
+    constexpr int kExactM = 2048;
+    constexpr int kExactN = 2048;
+    constexpr int kExactK = 2048;
+    constexpr int kBenchM = 2048;
+    constexpr int kBenchN = 2048;
+    constexpr int kBenchK = 2048;
+    constexpr int kWarmupIterations = 3;
+    constexpr int kMeasuredIterations = 10;
+
+    int device_count = 0;
+    check_mc(mcGetDeviceCount(&device_count), "mcGetDeviceCount");
+    if (device_count <= 0) {
+        std::cerr << "No MACA device is visible.\n";
+        return EXIT_FAILURE;
+    }
+    check_mc(mcSetDevice(0), "mcSetDevice");
+
+    if (!run_case(kExactM, kExactN, kExactK, "standalone_i8_tn_256x256x128_raw_arrays_exact")) {
+        return EXIT_FAILURE;
+    }
+
+    const BenchmarkResult benchmark =
+        run_benchmark(kBenchM, kBenchN, kBenchK, kWarmupIterations, kMeasuredIterations);
+    std::cout << std::fixed << std::setprecision(3)
+              << "standalone_i8_tn_256x256x128_raw_arrays benchmark: M=" << kBenchM
+              << ", N=" << kBenchN
+              << ", K=" << kBenchK
+              << ", avg_ms=" << benchmark.avg_ms
+              << ", TFLOPS=" << benchmark.tflops
+              << ", warmup=" << benchmark.warmup_iterations
+              << ", iters=" << benchmark.measured_iterations << '\n';
+
+    return EXIT_SUCCESS;
+}

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/attr.h

@@ -0,0 +1,730 @@
+/*
+    pybind11/attr.h: Infrastructure for processing custom
+    type and function attributes
+
+    Copyright (c) 2016 Wenzel Jakob <wenzel.jakob@epfl.ch>
+
+    All rights reserved. Use of this source code is governed by a
+    BSD-style license that can be found in the LICENSE file.
+*/
+
+#pragma once
+
+#include "detail/common.h"
+#include "cast.h"
+#include "trampoline_self_life_support.h"
+
+#include <functional>
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+
+/// \addtogroup annotations
+/// @{
+
+/// Annotation for methods
+struct is_method {
+    handle class_;
+    explicit is_method(const handle &c) : class_(c) {}
+};
+
+/// Annotation for setters
+struct is_setter {};
+
+/// Annotation for operators
+struct is_operator {};
+
+/// Annotation for classes that cannot be subclassed
+struct is_final {};
+
+/// Annotation for parent scope
+struct scope {
+    handle value;
+    explicit scope(const handle &s) : value(s) {}
+};
+
+/// Annotation for documentation
+struct doc {
+    const char *value;
+    explicit doc(const char *value) : value(value) {}
+};
+
+/// Annotation for function names
+struct name {
+    const char *value;
+    explicit name(const char *value) : value(value) {}
+};
+
+/// Annotation indicating that a function is an overload associated with a given "sibling"
+struct sibling {
+    handle value;
+    explicit sibling(const handle &value) : value(value.ptr()) {}
+};
+
+/// Annotation indicating that a class derives from another given type
+template <typename T>
+struct base {
+
+    PYBIND11_DEPRECATED(
+        "base<T>() was deprecated in favor of specifying 'T' as a template argument to class_")
+    base() = default;
+};
+
+/// Keep patient alive while nurse lives
+template <size_t Nurse, size_t Patient>
+struct keep_alive {};
+
+/// Annotation indicating that a class is involved in a multiple inheritance relationship
+struct multiple_inheritance {};
+
+/// Annotation which enables dynamic attributes, i.e. adds `__dict__` to a class
+struct dynamic_attr {};
+
+/// Annotation which enables the buffer protocol for a type
+struct buffer_protocol {};
+
+/// Annotation which enables releasing the GIL before calling the C++ destructor of wrapped
+/// instances (pybind/pybind11#1446).
+struct release_gil_before_calling_cpp_dtor {};
+
+/// Annotation which requests that a special metaclass is created for a type
+struct metaclass {
+    handle value;
+
+    PYBIND11_DEPRECATED("py::metaclass() is no longer required. It's turned on by default now.")
+    metaclass() = default;
+
+    /// Override pybind11's default metaclass
+    explicit metaclass(handle value) : value(value) {}
+};
+
+/// Specifies a custom callback with signature `void (PyHeapTypeObject*)` that
+/// may be used to customize the Python type.
+///
+/// The callback is invoked immediately before `PyType_Ready`.
+///
+/// Note: This is an advanced interface, and uses of it may require changes to
+/// work with later versions of pybind11.  You may wish to consult the
+/// implementation of `make_new_python_type` in `detail/classes.h` to understand
+/// the context in which the callback will be run.
+struct custom_type_setup {
+    using callback = std::function<void(PyHeapTypeObject *heap_type)>;
+
+    explicit custom_type_setup(callback value) : value(std::move(value)) {}
+
+    callback value;
+};
+
+/// Annotation that marks a class as local to the module:
+struct module_local {
+    const bool value;
+    constexpr explicit module_local(bool v = true) : value(v) {}
+};
+
+/// Annotation to mark enums as an arithmetic type
+struct arithmetic {};
+
+/// Mark a function for addition at the beginning of the existing overload chain instead of the end
+struct prepend {};
+
+/** \rst
+    A call policy which places one or more guard variables (``Ts...``) around the function call.
+
+    For example, this definition:
+
+    .. code-block:: cpp
+
+        m.def("foo", foo, py::call_guard<T>());
+
+    is equivalent to the following pseudocode:
+
+    .. code-block:: cpp
+
+        m.def("foo", [](args...) {
+            T scope_guard;
+            return foo(args...); // forwarded arguments
+        });
+ \endrst */
+template <typename... Ts>
+struct call_guard;
+
+template <>
+struct call_guard<> {
+    using type = detail::void_type;
+};
+
+template <typename T>
+struct call_guard<T> {
+    static_assert(std::is_default_constructible<T>::value,
+                  "The guard type must be default constructible");
+
+    using type = T;
+};
+
+template <typename T, typename... Ts>
+struct call_guard<T, Ts...> {
+    struct type {
+        T guard{}; // Compose multiple guard types with left-to-right default-constructor order
+        typename call_guard<Ts...>::type next{};
+    };
+};
+
+/// @} annotations
+
+PYBIND11_NAMESPACE_BEGIN(detail)
+/* Forward declarations */
+enum op_id : int;
+enum op_type : int;
+struct undefined_t;
+template <op_id id, op_type ot, typename L = undefined_t, typename R = undefined_t>
+struct op_;
+void keep_alive_impl(size_t Nurse, size_t Patient, function_call &call, handle ret);
+
+/// Internal data structure which holds metadata about a keyword argument
+struct argument_record {
+    const char *name;  ///< Argument name
+    const char *descr; ///< Human-readable version of the argument value
+    handle value;      ///< Associated Python object
+    bool convert : 1;  ///< True if the argument is allowed to convert when loading
+    bool none : 1;     ///< True if None is allowed when loading
+
+    argument_record(const char *name, const char *descr, handle value, bool convert, bool none)
+        : name(name), descr(descr), value(value), convert(convert), none(none) {}
+};
+
+/// Internal data structure which holds metadata about a bound function (signature, overloads,
+/// etc.)
+#define PYBIND11_DETAIL_FUNCTION_RECORD_ABI_ID "v1" // PLEASE UPDATE if the struct is changed.
+struct function_record {
+    function_record()
+        : is_constructor(false), is_new_style_constructor(false), is_stateless(false),
+          is_operator(false), is_method(false), is_setter(false), has_args(false),
+          has_kwargs(false), prepend(false) {}
+
+    /// Function name
+    char *name = nullptr; /* why no C++ strings? They generate heavier code.. */
+
+    // User-specified documentation string
+    char *doc = nullptr;
+
+    /// Human-readable version of the function signature
+    char *signature = nullptr;
+
+    /// List of registered keyword arguments
+    std::vector<argument_record> args;
+
+    /// Pointer to lambda function which converts arguments and performs the actual call
+    handle (*impl)(function_call &) = nullptr;
+
+    /// Storage for the wrapped function pointer and captured data, if any
+    void *data[3] = {};
+
+    /// Pointer to custom destructor for 'data' (if needed)
+    void (*free_data)(function_record *ptr) = nullptr;
+
+    /// Return value policy associated with this function
+    return_value_policy policy = return_value_policy::automatic;
+
+    /// True if name == '__init__'
+    bool is_constructor : 1;
+
+    /// True if this is a new-style `__init__` defined in `detail/init.h`
+    bool is_new_style_constructor : 1;
+
+    /// True if this is a stateless function pointer
+    bool is_stateless : 1;
+
+    /// True if this is an operator (__add__), etc.
+    bool is_operator : 1;
+
+    /// True if this is a method
+    bool is_method : 1;
+
+    /// True if this is a setter
+    bool is_setter : 1;
+
+    /// True if the function has a '*args' argument
+    bool has_args : 1;
+
+    /// True if the function has a '**kwargs' argument
+    bool has_kwargs : 1;
+
+    /// True if this function is to be inserted at the beginning of the overload resolution chain
+    bool prepend : 1;
+
+    /// Number of arguments (including py::args and/or py::kwargs, if present)
+    std::uint16_t nargs;
+
+    /// Number of leading positional arguments, which are terminated by a py::args or py::kwargs
+    /// argument or by a py::kw_only annotation.
+    std::uint16_t nargs_pos = 0;
+
+    /// Number of leading arguments (counted in `nargs`) that are positional-only
+    std::uint16_t nargs_pos_only = 0;
+
+    /// Python method object
+    PyMethodDef *def = nullptr;
+
+    /// Python handle to the parent scope (a class or a module)
+    handle scope;
+
+    /// Python handle to the sibling function representing an overload chain
+    handle sibling;
+
+    /// Pointer to next overload
+    function_record *next = nullptr;
+};
+// The main purpose of this macro is to make it easy to pin-point the critically related code
+// sections.
+#define PYBIND11_ENSURE_PRECONDITION_FOR_FUNCTIONAL_H_PERFORMANCE_OPTIMIZATIONS(...)              \
+    static_assert(                                                                                \
+        __VA_ARGS__,                                                                              \
+        "Violation of precondition for pybind11/functional.h performance optimizations!")
+
+/// Special data structure which (temporarily) holds metadata about a bound class
+struct type_record {
+    PYBIND11_NOINLINE type_record()
+        : multiple_inheritance(false), dynamic_attr(false), buffer_protocol(false),
+          module_local(false), is_final(false), release_gil_before_calling_cpp_dtor(false) {}
+
+    /// Handle to the parent scope
+    handle scope;
+
+    /// Name of the class
+    const char *name = nullptr;
+
+    // Pointer to RTTI type_info data structure
+    const std::type_info *type = nullptr;
+
+    /// How large is the underlying C++ type?
+    size_t type_size = 0;
+
+    /// What is the alignment of the underlying C++ type?
+    size_t type_align = 0;
+
+    /// How large is the type's holder?
+    size_t holder_size = 0;
+
+    /// The global operator new can be overridden with a class-specific variant
+    void *(*operator_new)(size_t) = nullptr;
+
+    /// Function pointer to class_<..>::init_instance
+    void (*init_instance)(instance *, const void *) = nullptr;
+
+    /// Function pointer to class_<..>::dealloc
+    void (*dealloc)(detail::value_and_holder &) = nullptr;
+
+    /// Function pointer for casting alias class (aka trampoline) pointer to
+    /// trampoline_self_life_support pointer. Sidesteps cross-DSO RTTI issues
+    /// on platforms like macOS (see PR #5728 for details).
+    get_trampoline_self_life_support_fn get_trampoline_self_life_support
+        = [](void *) -> trampoline_self_life_support * { return nullptr; };
+
+    /// List of base classes of the newly created type
+    list bases;
+
+    /// Optional docstring
+    const char *doc = nullptr;
+
+    /// Custom metaclass (optional)
+    handle metaclass;
+
+    /// Custom type setup.
+    custom_type_setup::callback custom_type_setup_callback;
+
+    /// Multiple inheritance marker
+    bool multiple_inheritance : 1;
+
+    /// Does the class manage a __dict__?
+    bool dynamic_attr : 1;
+
+    /// Does the class implement the buffer protocol?
+    bool buffer_protocol : 1;
+
+    /// Is the class definition local to the module shared object?
+    bool module_local : 1;
+
+    /// Is the class inheritable from python classes?
+    bool is_final : 1;
+
+    /// Solves pybind/pybind11#1446
+    bool release_gil_before_calling_cpp_dtor : 1;
+
+    holder_enum_t holder_enum_v = holder_enum_t::undefined;
+
+    PYBIND11_NOINLINE void add_base(const std::type_info &base, void *(*caster)(void *) ) {
+        auto *base_info = detail::get_type_info(base, false);
+        if (!base_info) {
+            std::string tname(base.name());
+            detail::clean_type_id(tname);
+            pybind11_fail("generic_type: type \"" + std::string(name)
+                          + "\" referenced unknown base type \"" + tname + "\"");
+        }
+
+        // SMART_HOLDER_BAKEIN_FOLLOW_ON: Refine holder compatibility checks.
+        bool this_has_unique_ptr_holder = (holder_enum_v == holder_enum_t::std_unique_ptr);
+        bool base_has_unique_ptr_holder
+            = (base_info->holder_enum_v == holder_enum_t::std_unique_ptr);
+        if (this_has_unique_ptr_holder != base_has_unique_ptr_holder) {
+            std::string tname(base.name());
+            detail::clean_type_id(tname);
+            pybind11_fail("generic_type: type \"" + std::string(name) + "\" "
+                          + (this_has_unique_ptr_holder ? "does not have" : "has")
+                          + " a non-default holder type while its base \"" + tname + "\" "
+                          + (base_has_unique_ptr_holder ? "does not" : "does"));
+        }
+
+        bases.append(reinterpret_cast<PyObject *>(base_info->type));
+
+#ifdef PYBIND11_BACKWARD_COMPATIBILITY_TP_DICTOFFSET
+        dynamic_attr |= base_info->type->tp_dictoffset != 0;
+#else
+        dynamic_attr |= (PyType_GetFlags(base_info->type) & Py_TPFLAGS_MANAGED_DICT) != 0;
+#endif
+
+        if (caster) {
+            base_info->implicit_casts.emplace_back(type, caster);
+        }
+    }
+};
+
+inline function_call::function_call(const function_record &f, handle p) : func(f), parent(p) {
+    args.reserve(f.nargs);
+    args_convert.reserve(f.nargs);
+}
+
+/// Tag for a new-style `__init__` defined in `detail/init.h`
+struct is_new_style_constructor {};
+
+/**
+ * Partial template specializations to process custom attributes provided to
+ * cpp_function_ and class_. These are either used to initialize the respective
+ * fields in the type_record and function_record data structures or executed at
+ * runtime to deal with custom call policies (e.g. keep_alive).
+ */
+template <typename T, typename SFINAE = void>
+struct process_attribute;
+
+template <typename T>
+struct process_attribute_default {
+    /// Default implementation: do nothing
+    static void init(const T &, function_record *) {}
+    static void init(const T &, type_record *) {}
+    static void precall(function_call &) {}
+    static void postcall(function_call &, handle) {}
+};
+
+/// Process an attribute specifying the function's name
+template <>
+struct process_attribute<name> : process_attribute_default<name> {
+    static void init(const name &n, function_record *r) { r->name = const_cast<char *>(n.value); }
+};
+
+/// Process an attribute specifying the function's docstring
+template <>
+struct process_attribute<doc> : process_attribute_default<doc> {
+    static void init(const doc &n, function_record *r) { r->doc = const_cast<char *>(n.value); }
+};
+
+/// Process an attribute specifying the function's docstring (provided as a C-style string)
+template <>
+struct process_attribute<const char *> : process_attribute_default<const char *> {
+    static void init(const char *d, function_record *r) { r->doc = const_cast<char *>(d); }
+    static void init(const char *d, type_record *r) { r->doc = d; }
+};
+template <>
+struct process_attribute<char *> : process_attribute<const char *> {};
+
+/// Process an attribute indicating the function's return value policy
+template <>
+struct process_attribute<return_value_policy> : process_attribute_default<return_value_policy> {
+    static void init(const return_value_policy &p, function_record *r) { r->policy = p; }
+};
+
+/// Process an attribute which indicates that this is an overloaded function associated with a
+/// given sibling
+template <>
+struct process_attribute<sibling> : process_attribute_default<sibling> {
+    static void init(const sibling &s, function_record *r) { r->sibling = s.value; }
+};
+
+/// Process an attribute which indicates that this function is a method
+template <>
+struct process_attribute<is_method> : process_attribute_default<is_method> {
+    static void init(const is_method &s, function_record *r) {
+        r->is_method = true;
+        r->scope = s.class_;
+    }
+};
+
+/// Process an attribute which indicates that this function is a setter
+template <>
+struct process_attribute<is_setter> : process_attribute_default<is_setter> {
+    static void init(const is_setter &, function_record *r) { r->is_setter = true; }
+};
+
+/// Process an attribute which indicates the parent scope of a method
+template <>
+struct process_attribute<scope> : process_attribute_default<scope> {
+    static void init(const scope &s, function_record *r) { r->scope = s.value; }
+};
+
+/// Process an attribute which indicates that this function is an operator
+template <>
+struct process_attribute<is_operator> : process_attribute_default<is_operator> {
+    static void init(const is_operator &, function_record *r) { r->is_operator = true; }
+};
+
+template <>
+struct process_attribute<is_new_style_constructor>
+    : process_attribute_default<is_new_style_constructor> {
+    static void init(const is_new_style_constructor &, function_record *r) {
+        r->is_new_style_constructor = true;
+    }
+};
+
+inline void check_kw_only_arg(const arg &a, function_record *r) {
+    if (r->args.size() > r->nargs_pos && (!a.name || a.name[0] == '\0')) {
+        pybind11_fail("arg(): cannot specify an unnamed argument after a kw_only() annotation or "
+                      "args() argument");
+    }
+}
+
+inline void append_self_arg_if_needed(function_record *r) {
+    if (r->is_method && r->args.empty()) {
+        r->args.emplace_back("self", nullptr, handle(), /*convert=*/true, /*none=*/false);
+    }
+}
+
+/// Process a keyword argument attribute (*without* a default value)
+template <>
+struct process_attribute<arg> : process_attribute_default<arg> {
+    static void init(const arg &a, function_record *r) {
+        append_self_arg_if_needed(r);
+        r->args.emplace_back(a.name, nullptr, handle(), !a.flag_noconvert, a.flag_none);
+
+        check_kw_only_arg(a, r);
+    }
+};
+
+/// Process a keyword argument attribute (*with* a default value)
+template <>
+struct process_attribute<arg_v> : process_attribute_default<arg_v> {
+    static void init(const arg_v &a, function_record *r) {
+        if (r->is_method && r->args.empty()) {
+            r->args.emplace_back(
+                "self", /*descr=*/nullptr, /*parent=*/handle(), /*convert=*/true, /*none=*/false);
+        }
+
+        if (!a.value) {
+#if defined(PYBIND11_DETAILED_ERROR_MESSAGES)
+            std::string descr("'");
+            if (a.name) {
+                descr += std::string(a.name) + ": ";
+            }
+            descr += a.type + "'";
+            if (r->is_method) {
+                if (r->name) {
+                    descr += " in method '" + (std::string) str(r->scope) + "."
+                             + (std::string) r->name + "'";
+                } else {
+                    descr += " in method of '" + (std::string) str(r->scope) + "'";
+                }
+            } else if (r->name) {
+                descr += " in function '" + (std::string) r->name + "'";
+            }
+            pybind11_fail("arg(): could not convert default argument " + descr
+                          + " into a Python object (type not registered yet?)");
+#else
+            pybind11_fail("arg(): could not convert default argument "
+                          "into a Python object (type not registered yet?). "
+                          "#define PYBIND11_DETAILED_ERROR_MESSAGES or compile in debug mode for "
+                          "more information.");
+#endif
+        }
+        r->args.emplace_back(a.name, a.descr, a.value.inc_ref(), !a.flag_noconvert, a.flag_none);
+
+        check_kw_only_arg(a, r);
+    }
+};
+
+/// Process a keyword-only-arguments-follow pseudo argument
+template <>
+struct process_attribute<kw_only> : process_attribute_default<kw_only> {
+    static void init(const kw_only &, function_record *r) {
+        append_self_arg_if_needed(r);
+        if (r->has_args && r->nargs_pos != static_cast<std::uint16_t>(r->args.size())) {
+            pybind11_fail("Mismatched args() and kw_only(): they must occur at the same relative "
+                          "argument location (or omit kw_only() entirely)");
+        }
+        r->nargs_pos = static_cast<std::uint16_t>(r->args.size());
+    }
+};
+
+/// Process a positional-only-argument maker
+template <>
+struct process_attribute<pos_only> : process_attribute_default<pos_only> {
+    static void init(const pos_only &, function_record *r) {
+        append_self_arg_if_needed(r);
+        r->nargs_pos_only = static_cast<std::uint16_t>(r->args.size());
+        if (r->nargs_pos_only > r->nargs_pos) {
+            pybind11_fail("pos_only(): cannot follow a py::args() argument");
+        }
+        // It also can't follow a kw_only, but a static_assert in pybind11.h checks that
+    }
+};
+
+/// Process a parent class attribute.  Single inheritance only (class_ itself already guarantees
+/// that)
+template <typename T>
+struct process_attribute<T, enable_if_t<is_pyobject<T>::value>>
+    : process_attribute_default<handle> {
+    static void init(const handle &h, type_record *r) { r->bases.append(h); }
+};
+
+/// Process a parent class attribute (deprecated, does not support multiple inheritance)
+template <typename T>
+struct process_attribute<base<T>> : process_attribute_default<base<T>> {
+    static void init(const base<T> &, type_record *r) { r->add_base(typeid(T), nullptr); }
+};
+
+/// Process a multiple inheritance attribute
+template <>
+struct process_attribute<multiple_inheritance> : process_attribute_default<multiple_inheritance> {
+    static void init(const multiple_inheritance &, type_record *r) {
+        r->multiple_inheritance = true;
+    }
+};
+
+template <>
+struct process_attribute<dynamic_attr> : process_attribute_default<dynamic_attr> {
+    static void init(const dynamic_attr &, type_record *r) { r->dynamic_attr = true; }
+};
+
+template <>
+struct process_attribute<custom_type_setup> {
+    static void init(const custom_type_setup &value, type_record *r) {
+        r->custom_type_setup_callback = value.value;
+    }
+};
+
+template <>
+struct process_attribute<is_final> : process_attribute_default<is_final> {
+    static void init(const is_final &, type_record *r) { r->is_final = true; }
+};
+
+template <>
+struct process_attribute<buffer_protocol> : process_attribute_default<buffer_protocol> {
+    static void init(const buffer_protocol &, type_record *r) { r->buffer_protocol = true; }
+};
+
+template <>
+struct process_attribute<metaclass> : process_attribute_default<metaclass> {
+    static void init(const metaclass &m, type_record *r) { r->metaclass = m.value; }
+};
+
+template <>
+struct process_attribute<module_local> : process_attribute_default<module_local> {
+    static void init(const module_local &l, type_record *r) { r->module_local = l.value; }
+};
+
+template <>
+struct process_attribute<release_gil_before_calling_cpp_dtor>
+    : process_attribute_default<release_gil_before_calling_cpp_dtor> {
+    static void init(const release_gil_before_calling_cpp_dtor &, type_record *r) {
+        r->release_gil_before_calling_cpp_dtor = true;
+    }
+};
+
+/// Process a 'prepend' attribute, putting this at the beginning of the overload chain
+template <>
+struct process_attribute<prepend> : process_attribute_default<prepend> {
+    static void init(const prepend &, function_record *r) { r->prepend = true; }
+};
+
+/// Process an 'arithmetic' attribute for enums (does nothing here)
+template <>
+struct process_attribute<arithmetic> : process_attribute_default<arithmetic> {};
+
+template <typename... Ts>
+struct process_attribute<call_guard<Ts...>> : process_attribute_default<call_guard<Ts...>> {};
+
+/**
+ * Process a keep_alive call policy -- invokes keep_alive_impl during the
+ * pre-call handler if both Nurse, Patient != 0 and use the post-call handler
+ * otherwise
+ */
+template <size_t Nurse, size_t Patient>
+struct process_attribute<keep_alive<Nurse, Patient>>
+    : public process_attribute_default<keep_alive<Nurse, Patient>> {
+    template <size_t N = Nurse, size_t P = Patient, enable_if_t<N != 0 && P != 0, int> = 0>
+    static void precall(function_call &call) {
+        keep_alive_impl(Nurse, Patient, call, handle());
+    }
+    template <size_t N = Nurse, size_t P = Patient, enable_if_t<N != 0 && P != 0, int> = 0>
+    static void postcall(function_call &, handle) {}
+    template <size_t N = Nurse, size_t P = Patient, enable_if_t<N == 0 || P == 0, int> = 0>
+    static void precall(function_call &) {}
+    template <size_t N = Nurse, size_t P = Patient, enable_if_t<N == 0 || P == 0, int> = 0>
+    static void postcall(function_call &call, handle ret) {
+        keep_alive_impl(Nurse, Patient, call, ret);
+    }
+};
+
+/// Recursively iterate over variadic template arguments
+template <typename... Args>
+struct process_attributes {
+    static void init(const Args &...args, function_record *r) {
+        PYBIND11_WORKAROUND_INCORRECT_MSVC_C4100(r);
+        PYBIND11_WORKAROUND_INCORRECT_GCC_UNUSED_BUT_SET_PARAMETER(r);
+        using expander = int[];
+        (void) expander{
+            0, ((void) process_attribute<typename std::decay<Args>::type>::init(args, r), 0)...};
+    }
+    static void init(const Args &...args, type_record *r) {
+        PYBIND11_WORKAROUND_INCORRECT_MSVC_C4100(r);
+        PYBIND11_WORKAROUND_INCORRECT_GCC_UNUSED_BUT_SET_PARAMETER(r);
+        using expander = int[];
+        (void) expander{0,
+                        (process_attribute<typename std::decay<Args>::type>::init(args, r), 0)...};
+    }
+    static void precall(function_call &call) {
+        PYBIND11_WORKAROUND_INCORRECT_MSVC_C4100(call);
+        using expander = int[];
+        (void) expander{0,
+                        (process_attribute<typename std::decay<Args>::type>::precall(call), 0)...};
+    }
+    static void postcall(function_call &call, handle fn_ret) {
+        PYBIND11_WORKAROUND_INCORRECT_MSVC_C4100(call, fn_ret);
+        PYBIND11_WORKAROUND_INCORRECT_GCC_UNUSED_BUT_SET_PARAMETER(fn_ret);
+        using expander = int[];
+        (void) expander{
+            0, (process_attribute<typename std::decay<Args>::type>::postcall(call, fn_ret), 0)...};
+    }
+};
+
+template <typename T>
+struct is_keep_alive : std::false_type {};
+
+template <size_t Nurse, size_t Patient>
+struct is_keep_alive<keep_alive<Nurse, Patient>> : std::true_type {};
+
+template <typename T>
+using is_call_guard = is_instantiation<call_guard, T>;
+
+/// Extract the ``type`` from the first `call_guard` in `Extras...` (or `void_type` if none found)
+template <typename... Extra>
+using extract_guard_t = typename exactly_one_t<is_call_guard, call_guard<>, Extra...>::type;
+
+/// Check the number of named arguments at compile time
+template <typename... Extra,
+          size_t named = constexpr_sum(std::is_base_of<arg, Extra>::value...),
+          size_t self = constexpr_sum(std::is_same<is_method, Extra>::value...)>
+constexpr bool expected_num_args(size_t nargs, bool has_args, bool has_kwargs) {
+    PYBIND11_WORKAROUND_INCORRECT_MSVC_C4100(nargs, has_args, has_kwargs);
+    return named == 0
+           || (self + named + static_cast<size_t>(has_args) + static_cast<size_t>(has_kwargs))
+                  == nargs;
+}
+
+PYBIND11_NAMESPACE_END(detail)
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/buffer_info.h

@@ -0,0 +1,209 @@
+/*
+    pybind11/buffer_info.h: Python buffer object interface
+
+    Copyright (c) 2016 Wenzel Jakob <wenzel.jakob@epfl.ch>
+
+    All rights reserved. Use of this source code is governed by a
+    BSD-style license that can be found in the LICENSE file.
+*/
+
+#pragma once
+
+#include "detail/common.h"
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+// Default, C-style strides
+inline std::vector<ssize_t> c_strides(const std::vector<ssize_t> &shape, ssize_t itemsize) {
+    auto ndim = shape.size();
+    std::vector<ssize_t> strides(ndim, itemsize);
+    if (ndim > 0) {
+        for (size_t i = ndim - 1; i > 0; --i) {
+            strides[i - 1] = strides[i] * shape[i];
+        }
+    }
+    return strides;
+}
+
+// F-style strides; default when constructing an array_t with `ExtraFlags & f_style`
+inline std::vector<ssize_t> f_strides(const std::vector<ssize_t> &shape, ssize_t itemsize) {
+    auto ndim = shape.size();
+    std::vector<ssize_t> strides(ndim, itemsize);
+    for (size_t i = 1; i < ndim; ++i) {
+        strides[i] = strides[i - 1] * shape[i - 1];
+    }
+    return strides;
+}
+
+template <typename T, typename SFINAE = void>
+struct compare_buffer_info;
+
+PYBIND11_NAMESPACE_END(detail)
+
+/// Information record describing a Python buffer object
+struct buffer_info {
+    void *ptr = nullptr;          // Pointer to the underlying storage
+    ssize_t itemsize = 0;         // Size of individual items in bytes
+    ssize_t size = 0;             // Total number of entries
+    std::string format;           // For homogeneous buffers, this should be set to
+                                  // format_descriptor<T>::format()
+    ssize_t ndim = 0;             // Number of dimensions
+    std::vector<ssize_t> shape;   // Shape of the tensor (1 entry per dimension)
+    std::vector<ssize_t> strides; // Number of bytes between adjacent entries
+                                  // (for each per dimension)
+    bool readonly = false;        // flag to indicate if the underlying storage may be written to
+
+    buffer_info() = default;
+
+    buffer_info(void *ptr,
+                ssize_t itemsize,
+                const std::string &format,
+                ssize_t ndim,
+                detail::any_container<ssize_t> shape_in,
+                detail::any_container<ssize_t> strides_in,
+                bool readonly = false)
+        : ptr(ptr), itemsize(itemsize), size(1), format(format), ndim(ndim),
+          shape(std::move(shape_in)), strides(std::move(strides_in)), readonly(readonly) {
+        if (ndim != static_cast<ssize_t>(shape.size())
+            || ndim != static_cast<ssize_t>(strides.size())) {
+            pybind11_fail("buffer_info: ndim doesn't match shape and/or strides length");
+        }
+        for (size_t i = 0; i < static_cast<size_t>(ndim); ++i) {
+            size *= shape[i];
+        }
+    }
+
+    template <typename T>
+    buffer_info(T *ptr,
+                detail::any_container<ssize_t> shape_in,
+                detail::any_container<ssize_t> strides_in,
+                bool readonly = false)
+        : buffer_info(private_ctr_tag(),
+                      ptr,
+                      sizeof(T),
+                      format_descriptor<T>::format(),
+                      static_cast<ssize_t>(shape_in->size()),
+                      std::move(shape_in),
+                      std::move(strides_in),
+                      readonly) {}
+
+    buffer_info(void *ptr,
+                ssize_t itemsize,
+                const std::string &format,
+                ssize_t size,
+                bool readonly = false)
+        : buffer_info(ptr, itemsize, format, 1, {size}, {itemsize}, readonly) {}
+
+    template <typename T>
+    buffer_info(T *ptr, ssize_t size, bool readonly = false)
+        : buffer_info(ptr, sizeof(T), format_descriptor<T>::format(), size, readonly) {}
+
+    template <typename T>
+    buffer_info(const T *ptr, ssize_t size, bool readonly = true)
+        : buffer_info(
+              const_cast<T *>(ptr), sizeof(T), format_descriptor<T>::format(), size, readonly) {}
+
+    explicit buffer_info(Py_buffer *view, bool ownview = true)
+        : buffer_info(
+              view->buf,
+              view->itemsize,
+              view->format,
+              view->ndim,
+              {view->shape, view->shape + view->ndim},
+              /* Though buffer::request() requests PyBUF_STRIDES, ctypes objects
+               * ignore this flag and return a view with NULL strides.
+               * When strides are NULL, build them manually.  */
+              view->strides
+                  ? std::vector<ssize_t>(view->strides, view->strides + view->ndim)
+                  : detail::c_strides({view->shape, view->shape + view->ndim}, view->itemsize),
+              (view->readonly != 0)) {
+        // NOLINTNEXTLINE(cppcoreguidelines-prefer-member-initializer)
+        this->m_view = view;
+        // NOLINTNEXTLINE(cppcoreguidelines-prefer-member-initializer)
+        this->ownview = ownview;
+    }
+
+    buffer_info(const buffer_info &) = delete;
+    buffer_info &operator=(const buffer_info &) = delete;
+
+    buffer_info(buffer_info &&other) noexcept { (*this) = std::move(other); }
+
+    buffer_info &operator=(buffer_info &&rhs) noexcept {
+        ptr = rhs.ptr;
+        itemsize = rhs.itemsize;
+        size = rhs.size;
+        format = std::move(rhs.format);
+        ndim = rhs.ndim;
+        shape = std::move(rhs.shape);
+        strides = std::move(rhs.strides);
+        std::swap(m_view, rhs.m_view);
+        std::swap(ownview, rhs.ownview);
+        readonly = rhs.readonly;
+        return *this;
+    }
+
+    ~buffer_info() {
+        if (m_view && ownview) {
+            PyBuffer_Release(m_view);
+            delete m_view;
+        }
+    }
+
+    Py_buffer *view() const { return m_view; }
+    Py_buffer *&view() { return m_view; }
+
+    /* True if the buffer item type is equivalent to `T`. */
+    // To define "equivalent" by example:
+    // `buffer_info::item_type_is_equivalent_to<int>(b)` and
+    // `buffer_info::item_type_is_equivalent_to<long>(b)` may both be true
+    // on some platforms, but `int` and `unsigned` will never be equivalent.
+    // For the ground truth, please inspect `detail::compare_buffer_info<>`.
+    template <typename T>
+    bool item_type_is_equivalent_to() const {
+        return detail::compare_buffer_info<T>::compare(*this);
+    }
+
+private:
+    struct private_ctr_tag {};
+
+    buffer_info(private_ctr_tag,
+                void *ptr,
+                ssize_t itemsize,
+                const std::string &format,
+                ssize_t ndim,
+                detail::any_container<ssize_t> &&shape_in,
+                detail::any_container<ssize_t> &&strides_in,
+                bool readonly)
+        : buffer_info(
+              ptr, itemsize, format, ndim, std::move(shape_in), std::move(strides_in), readonly) {}
+
+    Py_buffer *m_view = nullptr;
+    bool ownview = false;
+};
+
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+template <typename T, typename SFINAE>
+struct compare_buffer_info {
+    static bool compare(const buffer_info &b) {
+        // NOLINTNEXTLINE(bugprone-sizeof-expression) Needed for `PyObject *`
+        return b.format == format_descriptor<T>::format() && b.itemsize == (ssize_t) sizeof(T);
+    }
+};
+
+template <typename T>
+struct compare_buffer_info<T, detail::enable_if_t<std::is_integral<T>::value>> {
+    static bool compare(const buffer_info &b) {
+        return static_cast<size_t>(b.itemsize) == sizeof(T)
+               && (b.format == format_descriptor<T>::value
+                   || ((sizeof(T) == sizeof(long))
+                       && b.format == (std::is_unsigned<T>::value ? "L" : "l"))
+                   || ((sizeof(T) == sizeof(size_t))
+                       && b.format == (std::is_unsigned<T>::value ? "N" : "n")));
+    }
+};
+
+PYBIND11_NAMESPACE_END(detail)
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/chrono.h

@@ -0,0 +1,227 @@
+/*
+    pybind11/chrono.h: Transparent conversion between std::chrono and python's datetime
+
+    Copyright (c) 2016 Trent Houliston <trent@houliston.me> and
+                       Wenzel Jakob <wenzel.jakob@epfl.ch>
+
+    All rights reserved. Use of this source code is governed by a
+    BSD-style license that can be found in the LICENSE file.
+*/
+
+#pragma once
+
+#include "pybind11.h"
+
+#include <chrono>
+#include <cmath>
+#include <ctime>
+#include <datetime.h>
+#include <mutex>
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+template <typename type>
+class duration_caster {
+public:
+    using rep = typename type::rep;
+    using period = typename type::period;
+
+    // signed 25 bits required by the standard.
+    using days = std::chrono::duration<int_least32_t, std::ratio<86400>>;
+
+    bool load(handle src, bool) {
+        using namespace std::chrono;
+
+        // Lazy initialise the PyDateTime import
+        if (!PyDateTimeAPI) {
+            PyDateTime_IMPORT;
+        }
+
+        if (!src) {
+            return false;
+        }
+        // If invoked with datetime.delta object
+        if (PyDelta_Check(src.ptr())) {
+            value = type(duration_cast<duration<rep, period>>(
+                days(PyDateTime_DELTA_GET_DAYS(src.ptr()))
+                + seconds(PyDateTime_DELTA_GET_SECONDS(src.ptr()))
+                + microseconds(PyDateTime_DELTA_GET_MICROSECONDS(src.ptr()))));
+            return true;
+        }
+        // If invoked with a float we assume it is seconds and convert
+        if (PyFloat_Check(src.ptr())) {
+            value = type(duration_cast<duration<rep, period>>(
+                duration<double>(PyFloat_AsDouble(src.ptr()))));
+            return true;
+        }
+        return false;
+    }
+
+    // If this is a duration just return it back
+    static const std::chrono::duration<rep, period> &
+    get_duration(const std::chrono::duration<rep, period> &src) {
+        return src;
+    }
+    static const std::chrono::duration<rep, period> &
+    get_duration(const std::chrono::duration<rep, period> &&) = delete;
+
+    // If this is a time_point get the time_since_epoch
+    template <typename Clock>
+    static std::chrono::duration<rep, period>
+    get_duration(const std::chrono::time_point<Clock, std::chrono::duration<rep, period>> &src) {
+        return src.time_since_epoch();
+    }
+
+    static handle cast(const type &src, return_value_policy /* policy */, handle /* parent */) {
+        using namespace std::chrono;
+
+        // Use overloaded function to get our duration from our source
+        // Works out if it is a duration or time_point and get the duration
+        auto d = get_duration(src);
+
+        // Lazy initialise the PyDateTime import
+        if (!PyDateTimeAPI) {
+            PyDateTime_IMPORT;
+        }
+
+        // Declare these special duration types so the conversions happen with the correct
+        // primitive types (int)
+        using dd_t = duration<int, std::ratio<86400>>;
+        using ss_t = duration<int, std::ratio<1>>;
+        using us_t = duration<int, std::micro>;
+
+        auto dd = duration_cast<dd_t>(d);
+        auto subd = d - dd;
+        auto ss = duration_cast<ss_t>(subd);
+        auto us = duration_cast<us_t>(subd - ss);
+        return PyDelta_FromDSU(dd.count(), ss.count(), us.count());
+    }
+
+    PYBIND11_TYPE_CASTER(type, const_name("datetime.timedelta"));
+};
+
+inline std::tm *localtime_thread_safe(const std::time_t *time, std::tm *buf) {
+#if (defined(__STDC_LIB_EXT1__) && defined(__STDC_WANT_LIB_EXT1__)) || defined(_MSC_VER)
+    if (localtime_s(buf, time))
+        return nullptr;
+    return buf;
+#else
+    static std::mutex mtx;
+    std::lock_guard<std::mutex> lock(mtx);
+    std::tm *tm_ptr = std::localtime(time);
+    if (tm_ptr != nullptr) {
+        *buf = *tm_ptr;
+    }
+    return tm_ptr;
+#endif
+}
+
+// This is for casting times on the system clock into datetime.datetime instances
+template <typename Duration>
+class type_caster<std::chrono::time_point<std::chrono::system_clock, Duration>> {
+public:
+    using type = std::chrono::time_point<std::chrono::system_clock, Duration>;
+    bool load(handle src, bool) {
+        using namespace std::chrono;
+
+        // Lazy initialise the PyDateTime import
+        if (!PyDateTimeAPI) {
+            PyDateTime_IMPORT;
+        }
+
+        if (!src) {
+            return false;
+        }
+
+        std::tm cal;
+        microseconds msecs;
+
+        if (PyDateTime_Check(src.ptr())) {
+            cal.tm_sec = PyDateTime_DATE_GET_SECOND(src.ptr());
+            cal.tm_min = PyDateTime_DATE_GET_MINUTE(src.ptr());
+            cal.tm_hour = PyDateTime_DATE_GET_HOUR(src.ptr());
+            cal.tm_mday = PyDateTime_GET_DAY(src.ptr());
+            cal.tm_mon = PyDateTime_GET_MONTH(src.ptr()) - 1;
+            cal.tm_year = PyDateTime_GET_YEAR(src.ptr()) - 1900;
+            cal.tm_isdst = -1;
+            msecs = microseconds(PyDateTime_DATE_GET_MICROSECOND(src.ptr()));
+        } else if (PyDate_Check(src.ptr())) {
+            cal.tm_sec = 0;
+            cal.tm_min = 0;
+            cal.tm_hour = 0;
+            cal.tm_mday = PyDateTime_GET_DAY(src.ptr());
+            cal.tm_mon = PyDateTime_GET_MONTH(src.ptr()) - 1;
+            cal.tm_year = PyDateTime_GET_YEAR(src.ptr()) - 1900;
+            cal.tm_isdst = -1;
+            msecs = microseconds(0);
+        } else if (PyTime_Check(src.ptr())) {
+            cal.tm_sec = PyDateTime_TIME_GET_SECOND(src.ptr());
+            cal.tm_min = PyDateTime_TIME_GET_MINUTE(src.ptr());
+            cal.tm_hour = PyDateTime_TIME_GET_HOUR(src.ptr());
+            cal.tm_mday = 1;  // This date (day, month, year) = (1, 0, 70)
+            cal.tm_mon = 0;   // represents 1-Jan-1970, which is the first
+            cal.tm_year = 70; // earliest available date for Python's datetime
+            cal.tm_isdst = -1;
+            msecs = microseconds(PyDateTime_TIME_GET_MICROSECOND(src.ptr()));
+        } else {
+            return false;
+        }
+
+        value = time_point_cast<Duration>(system_clock::from_time_t(std::mktime(&cal)) + msecs);
+        return true;
+    }
+
+    static handle cast(const std::chrono::time_point<std::chrono::system_clock, Duration> &src,
+                       return_value_policy /* policy */,
+                       handle /* parent */) {
+        using namespace std::chrono;
+
+        // Lazy initialise the PyDateTime import
+        if (!PyDateTimeAPI) {
+            PyDateTime_IMPORT;
+        }
+
+        // Get out microseconds, and make sure they are positive, to avoid bug in eastern
+        // hemisphere time zones (cfr. https://github.com/pybind/pybind11/issues/2417)
+        using us_t = duration<int, std::micro>;
+        auto us = duration_cast<us_t>(src.time_since_epoch() % seconds(1));
+        if (us.count() < 0) {
+            us += duration_cast<us_t>(seconds(1));
+        }
+
+        // Subtract microseconds BEFORE `system_clock::to_time_t`, because:
+        // > If std::time_t has lower precision, it is implementation-defined whether the value is
+        // rounded or truncated. (https://en.cppreference.com/w/cpp/chrono/system_clock/to_time_t)
+        std::time_t tt
+            = system_clock::to_time_t(time_point_cast<system_clock::duration>(src - us));
+
+        std::tm localtime;
+        std::tm *localtime_ptr = localtime_thread_safe(&tt, &localtime);
+        if (!localtime_ptr) {
+            throw cast_error("Unable to represent system_clock in local time");
+        }
+        return PyDateTime_FromDateAndTime(localtime.tm_year + 1900,
+                                          localtime.tm_mon + 1,
+                                          localtime.tm_mday,
+                                          localtime.tm_hour,
+                                          localtime.tm_min,
+                                          localtime.tm_sec,
+                                          us.count());
+    }
+    PYBIND11_TYPE_CASTER(type, const_name("datetime.datetime"));
+};
+
+// Other clocks that are not the system clock are not measured as datetime.datetime objects
+// since they are not measured on calendar time. So instead we just make them timedeltas
+// Or if they have passed us a time as a float we convert that
+template <typename Clock, typename Duration>
+class type_caster<std::chrono::time_point<Clock, Duration>>
+    : public duration_caster<std::chrono::time_point<Clock, Duration>> {};
+
+template <typename Rep, typename Period>
+class type_caster<std::chrono::duration<Rep, Period>>
+    : public duration_caster<std::chrono::duration<Rep, Period>> {};
+
+PYBIND11_NAMESPACE_END(detail)
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/common.h

@@ -0,0 +1,2 @@
+#include "detail/common.h"
+#warning "Including 'common.h' is deprecated. It will be removed in v3.0. Use 'pybind11.h'."

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/complex.h

@@ -0,0 +1,95 @@
+/*
+    pybind11/complex.h: Complex number support
+
+    Copyright (c) 2016 Wenzel Jakob <wenzel.jakob@epfl.ch>
+
+    All rights reserved. Use of this source code is governed by a
+    BSD-style license that can be found in the LICENSE file.
+*/
+
+#pragma once
+
+#include "pybind11.h"
+
+#include <complex>
+
+/// glibc defines I as a macro which breaks things, e.g., boost template names
+#ifdef I
+#    undef I
+#endif
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+
+template <typename T>
+struct format_descriptor<std::complex<T>, detail::enable_if_t<std::is_floating_point<T>::value>> {
+    static constexpr const char c = format_descriptor<T>::c;
+    static constexpr const char value[3] = {'Z', c, '\0'};
+    static std::string format() { return std::string(value); }
+};
+
+#ifndef PYBIND11_CPP17
+
+template <typename T>
+constexpr const char
+    format_descriptor<std::complex<T>,
+                      detail::enable_if_t<std::is_floating_point<T>::value>>::value[3];
+
+#endif
+
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+template <typename T>
+struct is_fmt_numeric<std::complex<T>, detail::enable_if_t<std::is_floating_point<T>::value>> {
+    static constexpr bool value = true;
+    static constexpr int index = is_fmt_numeric<T>::index + 3;
+};
+
+template <typename T>
+class type_caster<std::complex<T>> {
+public:
+    bool load(handle src, bool convert) {
+        if (!src) {
+            return false;
+        }
+        if (!convert
+            && !(PyComplex_Check(src.ptr()) || PyFloat_Check(src.ptr())
+                 || PYBIND11_LONG_CHECK(src.ptr()))) {
+            return false;
+        }
+        handle src_or_index = src;
+        // PyPy: 7.3.7's 3.8 does not implement PyLong_*'s __index__ calls.
+        // The same logic is used in numeric_caster for ints and floats
+#if defined(PYPY_VERSION)
+        object index;
+        if (PYBIND11_INDEX_CHECK(src.ptr())) {
+            index = reinterpret_steal<object>(PyNumber_Index(src.ptr()));
+            if (!index) {
+                PyErr_Clear();
+                if (!convert)
+                    return false;
+            } else {
+                src_or_index = index;
+            }
+        }
+#endif
+        Py_complex result = PyComplex_AsCComplex(src_or_index.ptr());
+        if (result.real == -1.0 && PyErr_Occurred()) {
+            PyErr_Clear();
+            return false;
+        }
+        value = std::complex<T>((T) result.real, (T) result.imag);
+        return true;
+    }
+
+    static handle
+    cast(const std::complex<T> &src, return_value_policy /* policy */, handle /* parent */) {
+        return PyComplex_FromDoubles((double) src.real(), (double) src.imag());
+    }
+
+    PYBIND11_TYPE_CASTER(
+        std::complex<T>,
+        io_name("typing.SupportsComplex | typing.SupportsFloat | typing.SupportsIndex",
+                "complex"));
+};
+PYBIND11_NAMESPACE_END(detail)
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/conduit/README.txt

@@ -0,0 +1,15 @@
+NOTE
+----
+
+The C++ code here
+
+** only depends on <Python.h> **
+
+and nothing else.
+
+DO NOT ADD CODE WITH OTHER EXTERNAL DEPENDENCIES TO THIS DIRECTORY.
+
+Read on:
+
+pybind11_conduit_v1.h — Type-safe interoperability between different
+                        independent Python/C++ bindings systems.

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/conduit/pybind11_conduit_v1.h

@@ -0,0 +1,116 @@
+// Copyright (c) 2024 The pybind Community.
+
+/* The pybind11_conduit_v1 feature enables type-safe interoperability between
+
+* different independent Python/C++ bindings systems,
+
+* including pybind11 versions with different PYBIND11_INTERNALS_VERSION's.
+
+    * NOTE: The conduit feature
+            only covers    from-Python-to-C++ conversions, it
+            does not cover from-C++-to-Python conversions.
+            (For the latter, a different feature would have to be added.)
+
+The naming of the feature is a bit misleading:
+
+* The feature is in no way tied to pybind11 internals.
+
+* It just happens to originate from pybind11 and currently still lives there.
+
+* The only external dependency is <Python.h>.
+
+The implementation is a VERY light-weight dependency. It is designed to be
+compatible with any ISO C++11 (or higher) compiler, and does NOT require
+C++ Exception Handling to be enabled.
+
+Please see https://github.com/pybind/pybind11/pull/5296 for more background.
+
+The implementation involves a
+
+def _pybind11_conduit_v1_(
+    self,
+    pybind11_platform_abi_id: bytes,
+    cpp_type_info_capsule: capsule,
+    pointer_kind: bytes) -> capsule
+
+method that is meant to be added to Python objects wrapping C++ objects
+(e.g. pybind11::class_-wrapped types).
+
+The design of the _pybind11_conduit_v1_ feature provides two layers of
+protection against C++ ABI mismatches:
+
+* The first and most important layer is that the pybind11_platform_abi_id's
+  must match between extensions. — This will never be perfect, but is the same
+  pragmatic approach used in pybind11 since 2017
+  (https://github.com/pybind/pybind11/commit/96997a4b9d4ec3d389a570604394af5d5eee2557,
+  PYBIND11_INTERNALS_ID).
+
+* The second layer is that the typeid(std::type_info).name()'s must match
+  between extensions.
+
+The implementation below (which is shorter than this comment!), serves as a
+battle-tested specification. The main API is this one function:
+
+auto *cpp_pointer = pybind11_conduit_v1::get_type_pointer_ephemeral<YourType>(py_obj);
+
+It is meant to be a minimalistic reference implementation, intentionally
+without comprehensive error reporting. It is expected that major bindings
+systems will roll their own, compatible implementations, potentially with
+system-specific error reporting. The essential specifications all bindings
+systems need to agree on are merely:
+
+* PYBIND11_PLATFORM_ABI_ID (const char* literal).
+
+* The cpp_type_info capsule (see below: a void *ptr and a const char *name).
+
+* The cpp_conduit capsule (see below: a void *ptr and a const char *name).
+
+* "raw_pointer_ephemeral" means: the lifetime of the pointer is the lifetime
+  of the py_obj.
+
+*/
+
+// THIS MUST STAY AT THE TOP!
+#include "pybind11_platform_abi_id.h"
+
+#include <Python.h>
+#include <typeinfo>
+
+namespace pybind11_conduit_v1 {
+
+inline void *get_raw_pointer_ephemeral(PyObject *py_obj, const std::type_info *cpp_type_info) {
+    PyObject *cpp_type_info_capsule
+        = PyCapsule_New(const_cast<void *>(static_cast<const void *>(cpp_type_info)),
+                        typeid(std::type_info).name(),
+                        nullptr);
+    if (cpp_type_info_capsule == nullptr) {
+        return nullptr;
+    }
+    PyObject *cpp_conduit = PyObject_CallMethod(py_obj,
+                                                "_pybind11_conduit_v1_",
+                                                "yOy",
+                                                PYBIND11_PLATFORM_ABI_ID,
+                                                cpp_type_info_capsule,
+                                                "raw_pointer_ephemeral");
+    Py_DECREF(cpp_type_info_capsule);
+    if (cpp_conduit == nullptr) {
+        return nullptr;
+    }
+    void *raw_ptr = PyCapsule_GetPointer(cpp_conduit, cpp_type_info->name());
+    Py_DECREF(cpp_conduit);
+    if (PyErr_Occurred()) {
+        return nullptr;
+    }
+    return raw_ptr;
+}
+
+template <typename T>
+T *get_type_pointer_ephemeral(PyObject *py_obj) {
+    void *raw_ptr = get_raw_pointer_ephemeral(py_obj, &typeid(T));
+    if (raw_ptr == nullptr) {
+        return nullptr;
+    }
+    return static_cast<T *>(raw_ptr);
+}
+
+} // namespace pybind11_conduit_v1

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/conduit/pybind11_platform_abi_id.h

@@ -0,0 +1,87 @@
+#pragma once
+
+// Copyright (c) 2024 The pybind Community.
+
+// To maximize reusability:
+// DO NOT ADD CODE THAT REQUIRES C++ EXCEPTION HANDLING.
+
+#include "wrap_include_python_h.h"
+
+// Implementation details. DO NOT USE ELSEWHERE. (Unfortunately we cannot #undef them.)
+// This is duplicated here to maximize portability.
+#define PYBIND11_PLATFORM_ABI_ID_STRINGIFY(x) #x
+#define PYBIND11_PLATFORM_ABI_ID_TOSTRING(x) PYBIND11_PLATFORM_ABI_ID_STRINGIFY(x)
+
+#ifdef PYBIND11_COMPILER_TYPE
+//   // To maintain backward compatibility (see PR #5439).
+#    define PYBIND11_COMPILER_TYPE_LEADING_UNDERSCORE ""
+#else
+#    define PYBIND11_COMPILER_TYPE_LEADING_UNDERSCORE "_"
+#    if defined(__MINGW32__)
+#        define PYBIND11_COMPILER_TYPE "mingw"
+#    elif defined(__CYGWIN__)
+#        define PYBIND11_COMPILER_TYPE "gcc_cygwin"
+#    elif defined(_MSC_VER)
+#        define PYBIND11_COMPILER_TYPE "msvc"
+#    elif defined(__clang__) || defined(__GNUC__)
+#        define PYBIND11_COMPILER_TYPE "system" // Assumed compatible with system compiler.
+#    else
+#        error "Unknown PYBIND11_COMPILER_TYPE: PLEASE REVISE THIS CODE."
+#    endif
+#endif
+
+// PR #5439 made this macro obsolete. However, there are many manipulations of this macro in the
+// wild. Therefore, to maintain backward compatibility, it is kept around.
+#ifndef PYBIND11_STDLIB
+#    define PYBIND11_STDLIB ""
+#endif
+
+#ifndef PYBIND11_BUILD_ABI
+#    if defined(_MSC_VER)                 // See PR #4953.
+#        if defined(_MT) && defined(_DLL) // Corresponding to CL command line options /MD or /MDd.
+#            if (_MSC_VER) / 100 == 19
+#                define PYBIND11_BUILD_ABI "_md_mscver19"
+#            else
+#                error "Unknown major version for MSC_VER: PLEASE REVISE THIS CODE."
+#            endif
+#        elif defined(_MT) // Corresponding to CL command line options /MT or /MTd.
+#            define PYBIND11_BUILD_ABI "_mt_mscver" PYBIND11_PLATFORM_ABI_ID_TOSTRING(_MSC_VER)
+#        else
+#            if (_MSC_VER) / 100 == 19
+#                define PYBIND11_BUILD_ABI "_none_mscver19"
+#            else
+#                error "Unknown major version for MSC_VER: PLEASE REVISE THIS CODE."
+#            endif
+#        endif
+#    elif defined(_LIBCPP_ABI_VERSION) // https://libcxx.llvm.org/DesignDocs/ABIVersioning.html
+#        define PYBIND11_BUILD_ABI                                                                \
+            "_libcpp_abi" PYBIND11_PLATFORM_ABI_ID_TOSTRING(_LIBCPP_ABI_VERSION)
+#    elif defined(_GLIBCXX_USE_CXX11_ABI) // See PR #5439.
+#        if defined(__NVCOMPILER)
+//           // Assume that NVHPC is in the 1xxx ABI family.
+//           // THIS ASSUMPTION IS NOT FUTURE PROOF but apparently the best we can do.
+//           // Please let us know if there is a way to validate the assumption here.
+#        elif !defined(__GXX_ABI_VERSION)
+#            error                                                                                \
+                "Unknown platform or compiler (_GLIBCXX_USE_CXX11_ABI): PLEASE REVISE THIS CODE."
+#        endif
+#        if defined(__GXX_ABI_VERSION) && __GXX_ABI_VERSION < 1002 || __GXX_ABI_VERSION >= 2000
+#            error "Unknown platform or compiler (__GXX_ABI_VERSION): PLEASE REVISE THIS CODE."
+#        endif
+#        define PYBIND11_BUILD_ABI                                                                \
+            "_libstdcpp_gxx_abi_1xxx_use_cxx11_abi_" PYBIND11_PLATFORM_ABI_ID_TOSTRING(           \
+                _GLIBCXX_USE_CXX11_ABI)
+#    else
+#        error "Unknown platform or compiler: PLEASE REVISE THIS CODE."
+#    endif
+#endif
+
+// On MSVC, debug and release builds are not ABI-compatible!
+#if defined(_MSC_VER) && defined(_DEBUG)
+#    define PYBIND11_BUILD_TYPE "_debug"
+#else
+#    define PYBIND11_BUILD_TYPE ""
+#endif
+
+#define PYBIND11_PLATFORM_ABI_ID                                                                  \
+    PYBIND11_COMPILER_TYPE PYBIND11_STDLIB PYBIND11_BUILD_ABI PYBIND11_BUILD_TYPE

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/conduit/wrap_include_python_h.h

@@ -0,0 +1,72 @@
+#pragma once
+
+// Copyright (c) 2024 The pybind Community.
+
+// STRONG REQUIREMENT:
+//   This header is a wrapper around `#include <Python.h>`, therefore it
+//   MUST BE INCLUDED BEFORE ANY STANDARD HEADERS are included.
+// See also:
+//   https://docs.python.org/3/c-api/intro.html#include-files
+// Quoting from there:
+//   Note: Since Python may define some pre-processor definitions which affect
+//   the standard headers on some systems, you must include Python.h before
+//   any standard headers are included.
+
+// To maximize reusability:
+// DO NOT ADD CODE THAT REQUIRES C++ EXCEPTION HANDLING.
+
+// Disable linking to pythonX_d.lib on Windows in debug mode.
+#if defined(_MSC_VER) && defined(_DEBUG) && !defined(Py_DEBUG)
+// Workaround for a VS 2022 issue.
+// See https://github.com/pybind/pybind11/pull/3497 for full context.
+// NOTE: This workaround knowingly violates the Python.h include order
+//       requirement (see above).
+#    include <yvals.h>
+#    if _MSVC_STL_VERSION >= 143
+#        include <crtdefs.h>
+#    endif
+#    define PYBIND11_DEBUG_MARKER
+#    undef _DEBUG
+#endif
+
+// Don't let Python.h #define (v)snprintf as macro because they are implemented
+// properly in Visual Studio since 2015.
+#if defined(_MSC_VER)
+#    define HAVE_SNPRINTF 1
+#endif
+
+#if defined(_MSC_VER)
+#    pragma warning(push)
+#    pragma warning(disable : 4505)
+// C4505: 'PySlice_GetIndicesEx': unreferenced local function has been removed
+#endif
+
+#include <Python.h>
+#include <frameobject.h>
+#include <pythread.h>
+
+#if defined(_MSC_VER)
+#    pragma warning(pop)
+#endif
+
+#if defined(PYBIND11_DEBUG_MARKER)
+#    define _DEBUG 1
+#    undef PYBIND11_DEBUG_MARKER
+#endif
+
+// Python #defines overrides on all sorts of core functions, which
+// tends to wreak havok in C++ codebases that expect these to work
+// like regular functions (potentially with several overloads).
+#if defined(isalnum)
+#    undef isalnum
+#    undef isalpha
+#    undef islower
+#    undef isspace
+#    undef isupper
+#    undef tolower
+#    undef toupper
+#endif
+
+#if defined(copysign)
+#    undef copysign
+#endif

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/critical_section.h

@@ -0,0 +1,56 @@
+// Copyright (c) 2016-2025 The Pybind Development Team.
+// All rights reserved. Use of this source code is governed by a
+// BSD-style license that can be found in the LICENSE file.
+
+#pragma once
+
+#include "pytypes.h"
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+
+/// This does not do anything if there's a GIL. On free-threaded Python,
+/// it locks an object. This uses the CPython API, which has limits
+class scoped_critical_section {
+public:
+#ifdef Py_GIL_DISABLED
+    explicit scoped_critical_section(handle obj1, handle obj2 = handle{}) {
+        if (obj1) {
+            if (obj2) {
+                PyCriticalSection2_Begin(&section2, obj1.ptr(), obj2.ptr());
+                rank = 2;
+            } else {
+                PyCriticalSection_Begin(&section, obj1.ptr());
+                rank = 1;
+            }
+        } else if (obj2) {
+            PyCriticalSection_Begin(&section, obj2.ptr());
+            rank = 1;
+        }
+    }
+
+    ~scoped_critical_section() {
+        if (rank == 1) {
+            PyCriticalSection_End(&section);
+        } else if (rank == 2) {
+            PyCriticalSection2_End(&section2);
+        }
+    }
+#else
+    explicit scoped_critical_section(handle, handle = handle{}) {};
+    ~scoped_critical_section() = default;
+#endif
+
+    scoped_critical_section(const scoped_critical_section &) = delete;
+    scoped_critical_section &operator=(const scoped_critical_section &) = delete;
+
+private:
+#ifdef Py_GIL_DISABLED
+    int rank{0};
+    union {
+        PyCriticalSection section;
+        PyCriticalSection2 section2;
+    };
+#endif
+};
+
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/detail/argument_vector.h

@@ -0,0 +1,417 @@
+/*
+    pybind11/detail/argument_vector.h: small_vector-like containers to
+    avoid heap allocation of arguments during function call dispatch.
+
+    Copyright (c) Meta Platforms, Inc. and affiliates.
+
+    All rights reserved. Use of this source code is governed by a
+    BSD-style license that can be found in the LICENSE file.
+*/
+
+#pragma once
+
+#include <pybind11/pytypes.h>
+
+#include "common.h"
+
+#include <algorithm>
+#include <array>
+#include <cstddef>
+#include <cstdint>
+#include <cstring>
+#include <iterator>
+#include <type_traits>
+#include <utility>
+#include <vector>
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+
+PYBIND11_WARNING_DISABLE_MSVC(4127)
+
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+// Shared implementation utility for our small_vector-like containers.
+// We support C++11 and C++14, so we cannot use
+// std::variant. Union with the tag packed next to the inline
+// array's size is smaller anyway, allowing 1 extra handle of
+// inline storage for free. Compare the layouts (1 line per
+// size_t/void*, assuming a 64-bit machine):
+// With variant, total is N + 2 for N >= 2:
+// - variant tag (cannot be packed with the array size)
+// - array size (or first pointer of 3 in std::vector)
+// - N pointers of inline storage (or 2 remaining pointers of std::vector)
+// Custom union, total is N + 1 for N >= 3:
+// - variant tag & array size if applicable
+// - N pointers of inline storage (or 3 pointers of std::vector)
+//
+// NOTE: this is a low-level representational convenience; the two
+// use cases of this union are materially different and in particular
+// have different semantics for inline_array::size. All that is being
+// shared is the memory management behavior.
+template <typename ArrayT, std::size_t InlineSize, typename VectorT = ArrayT>
+union inline_array_or_vector {
+    struct inline_array {
+        bool is_inline = true;
+        std::uint32_t size = 0;
+        std::array<ArrayT, InlineSize> arr;
+    };
+    struct heap_vector {
+        bool is_inline = false;
+        std::vector<VectorT> vec;
+
+        heap_vector() = default;
+        heap_vector(std::size_t count, VectorT value) : vec(count, value) {}
+    };
+
+    inline_array iarray;
+    heap_vector hvector;
+
+    inline_array_or_vector() : iarray() {}
+
+    ~inline_array_or_vector() {
+        if (is_inline()) {
+            iarray.~inline_array();
+        } else {
+            hvector.~heap_vector();
+        }
+    }
+
+    // Disable copy ctor and assignment.
+    inline_array_or_vector(const inline_array_or_vector &) = delete;
+    inline_array_or_vector &operator=(const inline_array_or_vector &) = delete;
+
+    inline_array_or_vector(inline_array_or_vector &&rhs) noexcept {
+        if (rhs.is_inline()) {
+            new (&iarray) inline_array(std::move(rhs.iarray));
+        } else {
+            new (&hvector) heap_vector(std::move(rhs.hvector));
+        }
+        assert(is_inline() == rhs.is_inline());
+    }
+
+    inline_array_or_vector &operator=(inline_array_or_vector &&rhs) noexcept {
+        if (this == &rhs) {
+            return *this;
+        }
+
+        if (is_inline()) {
+            iarray.~inline_array();
+        } else {
+            hvector.~heap_vector();
+        }
+
+        if (rhs.is_inline()) {
+            new (&iarray) inline_array(std::move(rhs.iarray));
+        } else {
+            new (&hvector) heap_vector(std::move(rhs.hvector));
+        }
+        return *this;
+    }
+
+    bool is_inline() const {
+        // It is undefined behavior to access the inactive member of a
+        // union directly. However, it is well-defined to reinterpret_cast any
+        // pointer into a pointer to char and examine it as an array
+        // of bytes. See
+        // https://dev-discuss.pytorch.org/t/unionizing-for-profit-how-to-exploit-the-power-of-unions-in-c/444#the-memcpy-loophole-4
+        bool result = false;
+        static_assert(offsetof(inline_array, is_inline) == 0,
+                      "untagged union implementation relies on this");
+        static_assert(offsetof(heap_vector, is_inline) == 0,
+                      "untagged union implementation relies on this");
+        std::memcpy(&result, reinterpret_cast<const char *>(this), sizeof(bool));
+        return result;
+    }
+};
+
+template <typename T, std::size_t InlineSize>
+struct small_vector {
+public:
+    small_vector() = default;
+
+    // Disable copy ctor and assignment.
+    small_vector(const small_vector &) = delete;
+    small_vector &operator=(const small_vector &) = delete;
+    small_vector(small_vector &&) noexcept = default;
+    small_vector &operator=(small_vector &&) noexcept = default;
+
+    std::size_t size() const {
+        if (is_inline()) {
+            return m_repr.iarray.size;
+        }
+        return m_repr.hvector.vec.size();
+    }
+
+    T const *data() const {
+        if (is_inline()) {
+            return m_repr.iarray.arr.data();
+        }
+        return m_repr.hvector.vec.data();
+    }
+
+    T &operator[](std::size_t idx) {
+        assert(idx < size());
+        if (is_inline()) {
+            return m_repr.iarray.arr[idx];
+        }
+        return m_repr.hvector.vec[idx];
+    }
+
+    T const &operator[](std::size_t idx) const {
+        assert(idx < size());
+        if (is_inline()) {
+            return m_repr.iarray.arr[idx];
+        }
+        return m_repr.hvector.vec[idx];
+    }
+
+    void push_back(const T &x) { emplace_back(x); }
+
+    void push_back(T &&x) { emplace_back(std::move(x)); }
+
+    template <typename... Args>
+    void emplace_back(Args &&...x) {
+        if (is_inline()) {
+            auto &ha = m_repr.iarray;
+            if (ha.size == InlineSize) {
+                move_to_heap_vector_with_reserved_size(InlineSize + 1);
+                m_repr.hvector.vec.emplace_back(std::forward<Args>(x)...);
+            } else {
+                ha.arr[ha.size++] = T(std::forward<Args>(x)...);
+            }
+        } else {
+            m_repr.hvector.vec.emplace_back(std::forward<Args>(x)...);
+        }
+    }
+
+    void reserve(std::size_t sz) {
+        if (is_inline()) {
+            if (sz > InlineSize) {
+                move_to_heap_vector_with_reserved_size(sz);
+            }
+        } else {
+            reserve_slow_path(sz);
+        }
+    }
+
+private:
+    using repr_type = inline_array_or_vector<T, InlineSize>;
+    repr_type m_repr;
+
+    PYBIND11_NOINLINE void move_to_heap_vector_with_reserved_size(std::size_t reserved_size) {
+        assert(is_inline());
+        auto &ha = m_repr.iarray;
+        using heap_vector = typename repr_type::heap_vector;
+        heap_vector hv;
+        hv.vec.reserve(reserved_size);
+        static_assert(std::is_nothrow_move_constructible<T>::value,
+                      "this conversion is not exception safe");
+        static_assert(std::is_nothrow_move_constructible<heap_vector>::value,
+                      "this conversion is not exception safe");
+        std::move(ha.arr.begin(), ha.arr.begin() + ha.size, std::back_inserter(hv.vec));
+        new (&m_repr.hvector) heap_vector(std::move(hv));
+    }
+
+    PYBIND11_NOINLINE void reserve_slow_path(std::size_t sz) { m_repr.hvector.vec.reserve(sz); }
+
+    bool is_inline() const { return m_repr.is_inline(); }
+};
+
+// Container to avoid heap allocation for kRequestedInlineSize or fewer booleans.
+template <std::size_t kRequestedInlineSize>
+struct small_vector<bool, kRequestedInlineSize> {
+private:
+public:
+    small_vector() = default;
+
+    // Disable copy ctor and assignment.
+    small_vector(const small_vector &) = delete;
+    small_vector &operator=(const small_vector &) = delete;
+    small_vector(small_vector &&) noexcept = default;
+    small_vector &operator=(small_vector &&) noexcept = default;
+
+    small_vector(std::size_t count, bool value) {
+        if (count > kInlineSize) {
+            new (&m_repr.hvector) typename repr_type::heap_vector(count, value);
+        } else {
+            auto &inline_arr = m_repr.iarray;
+            inline_arr.arr.fill(value ? static_cast<std::size_t>(-1) : 0);
+            inline_arr.size = static_cast<decltype(inline_arr.size)>(count);
+        }
+    }
+
+    std::size_t size() const {
+        if (is_inline()) {
+            return m_repr.iarray.size;
+        }
+        return m_repr.hvector.vec.size();
+    }
+
+    void reserve(std::size_t sz) {
+        if (is_inline()) {
+            if (sz > kInlineSize) {
+                move_to_heap_vector_with_reserved_size(sz);
+            }
+        } else {
+            m_repr.hvector.vec.reserve(sz);
+        }
+    }
+
+    bool operator[](std::size_t idx) const {
+        if (is_inline()) {
+            return inline_index(idx);
+        }
+        assert(idx < m_repr.hvector.vec.size());
+        return m_repr.hvector.vec[idx];
+    }
+
+    void push_back(bool b) {
+        if (is_inline()) {
+            auto &ha = m_repr.iarray;
+            if (ha.size == kInlineSize) {
+                move_to_heap_vector_with_reserved_size(kInlineSize + 1);
+                push_back_slow_path(b);
+            } else {
+                assert(ha.size < kInlineSize);
+                const auto wbi = word_and_bit_index(ha.size++);
+                assert(wbi.word < kWords);
+                assert(wbi.bit < kBitsPerWord);
+                if (b) {
+                    ha.arr[wbi.word] |= (static_cast<std::size_t>(1) << wbi.bit);
+                } else {
+                    ha.arr[wbi.word] &= ~(static_cast<std::size_t>(1) << wbi.bit);
+                }
+                assert(operator[](ha.size - 1) == b);
+            }
+        } else {
+            push_back_slow_path(b);
+        }
+    }
+
+    void set(std::size_t idx, bool value = true) {
+        if (is_inline()) {
+            auto &ha = m_repr.iarray;
+            assert(ha.size < kInlineSize);
+            const auto wbi = word_and_bit_index(idx);
+            assert(wbi.word < kWords);
+            assert(wbi.bit < kBitsPerWord);
+            if (value) {
+                ha.arr[wbi.word] |= (static_cast<std::size_t>(1) << wbi.bit);
+            } else {
+                ha.arr[wbi.word] &= ~(static_cast<std::size_t>(1) << wbi.bit);
+            }
+        } else {
+            m_repr.hvector.vec[idx] = value;
+        }
+    }
+
+    void swap(small_vector &rhs) noexcept { std::swap(m_repr, rhs.m_repr); }
+
+private:
+    struct WordAndBitIndex {
+        std::size_t word;
+        std::size_t bit;
+    };
+
+    static WordAndBitIndex word_and_bit_index(std::size_t idx) {
+        return WordAndBitIndex{idx / kBitsPerWord, idx % kBitsPerWord};
+    }
+
+    bool inline_index(std::size_t idx) const {
+        const auto wbi = word_and_bit_index(idx);
+        assert(wbi.word < kWords);
+        assert(wbi.bit < kBitsPerWord);
+        return m_repr.iarray.arr[wbi.word] & (static_cast<std::size_t>(1) << wbi.bit);
+    }
+
+    PYBIND11_NOINLINE void move_to_heap_vector_with_reserved_size(std::size_t reserved_size) {
+        auto &inline_arr = m_repr.iarray;
+        using heap_vector = typename repr_type::heap_vector;
+        heap_vector hv;
+        hv.vec.reserve(reserved_size);
+        for (std::size_t ii = 0; ii < inline_arr.size; ++ii) {
+            hv.vec.push_back(inline_index(ii));
+        }
+        new (&m_repr.hvector) heap_vector(std::move(hv));
+    }
+
+    PYBIND11_NOINLINE void push_back_slow_path(bool b) { m_repr.hvector.vec.push_back(b); }
+
+    static constexpr auto kBitsPerWord = 8 * sizeof(std::size_t);
+    static constexpr auto kWords = (kRequestedInlineSize + kBitsPerWord - 1) / kBitsPerWord;
+    static constexpr auto kInlineSize = kWords * kBitsPerWord;
+
+    using repr_type = inline_array_or_vector<std::size_t, kWords, bool>;
+    repr_type m_repr;
+
+    bool is_inline() const { return m_repr.is_inline(); }
+};
+
+// Container to avoid heap allocation for N or fewer arguments.
+template <size_t N>
+using argument_vector = small_vector<handle, N>;
+
+// Container to avoid heap allocation for N or fewer booleans.
+template <size_t N>
+using args_convert_vector = small_vector<bool, N>;
+
+/// A small_vector of PyObject* that holds references and releases them on destruction.
+/// This provides explicit ownership semantics without relying on py::object's
+/// destructor, and avoids the need for reinterpret_cast when passing to vectorcall.
+template <std::size_t InlineSize>
+class ref_small_vector {
+public:
+    ref_small_vector() = default;
+
+    ~ref_small_vector() {
+        for (std::size_t i = 0; i < m_ptrs.size(); ++i) {
+            Py_XDECREF(m_ptrs[i]);
+        }
+    }
+
+    // Disable copy (prevent accidental double-decref)
+    ref_small_vector(const ref_small_vector &) = delete;
+    ref_small_vector &operator=(const ref_small_vector &) = delete;
+
+    // Move is allowed
+    ref_small_vector(ref_small_vector &&other) noexcept : m_ptrs(std::move(other.m_ptrs)) {
+        // other.m_ptrs is now empty, so its destructor won't decref anything
+    }
+
+    ref_small_vector &operator=(ref_small_vector &&other) noexcept {
+        if (this != &other) {
+            // Decref our current contents
+            for (std::size_t i = 0; i < m_ptrs.size(); ++i) {
+                Py_XDECREF(m_ptrs[i]);
+            }
+            m_ptrs = std::move(other.m_ptrs);
+        }
+        return *this;
+    }
+
+    /// Add a pointer, taking ownership (no incref, will decref on destruction)
+    void push_back_steal(PyObject *p) { m_ptrs.push_back(p); }
+
+    /// Add a pointer, borrowing (increfs now, will decref on destruction)
+    void push_back_borrow(PyObject *p) {
+        Py_XINCREF(p);
+        m_ptrs.push_back(p);
+    }
+
+    /// Add a null pointer (for PY_VECTORCALL_ARGUMENTS_OFFSET slot)
+    void push_back_null() { m_ptrs.push_back(nullptr); }
+
+    void reserve(std::size_t sz) { m_ptrs.reserve(sz); }
+
+    std::size_t size() const { return m_ptrs.size(); }
+
+    PyObject *operator[](std::size_t idx) const { return m_ptrs[idx]; }
+
+    PyObject *const *data() const { return m_ptrs.data(); }
+
+private:
+    small_vector<PyObject *, InlineSize> m_ptrs;
+};
+
+PYBIND11_NAMESPACE_END(detail)
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/detail/class.h

@@ -0,0 +1,850 @@
+/*
+    pybind11/detail/class.h: Python C API implementation details for py::class_
+
+    Copyright (c) 2017 Wenzel Jakob <wenzel.jakob@epfl.ch>
+
+    All rights reserved. Use of this source code is governed by a
+    BSD-style license that can be found in the LICENSE file.
+*/
+
+#pragma once
+
+#include <pybind11/attr.h>
+#include <pybind11/options.h>
+
+#include "exception_translation.h"
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+#if !defined(PYPY_VERSION)
+#    define PYBIND11_BUILTIN_QUALNAME
+#    define PYBIND11_SET_OLDPY_QUALNAME(obj, nameobj)
+#else
+// In PyPy, we still set __qualname__ so that we can produce reliable function type
+// signatures; in CPython this macro expands to nothing:
+#    define PYBIND11_SET_OLDPY_QUALNAME(obj, nameobj)                                             \
+        setattr((PyObject *) obj, "__qualname__", nameobj)
+#endif
+
+inline std::string get_fully_qualified_tp_name(PyTypeObject *type) {
+#if !defined(PYPY_VERSION)
+    return type->tp_name;
+#else
+    auto module_name = handle((PyObject *) type).attr("__module__").cast<std::string>();
+    if (module_name == PYBIND11_BUILTINS_MODULE)
+        return type->tp_name;
+    else
+        return std::move(module_name) + "." + type->tp_name;
+#endif
+}
+
+inline PyTypeObject *type_incref(PyTypeObject *type) {
+    Py_INCREF(type);
+    return type;
+}
+
+#if !defined(PYPY_VERSION)
+
+/// `pybind11_static_property.__get__()`: Always pass the class instead of the instance.
+extern "C" inline PyObject *pybind11_static_get(PyObject *self, PyObject * /*ob*/, PyObject *cls) {
+    return PyProperty_Type.tp_descr_get(self, cls, cls);
+}
+
+/// `pybind11_static_property.__set__()`: Just like the above `__get__()`.
+extern "C" inline int pybind11_static_set(PyObject *self, PyObject *obj, PyObject *value) {
+    PyObject *cls = PyType_Check(obj) ? obj : (PyObject *) Py_TYPE(obj);
+    return PyProperty_Type.tp_descr_set(self, cls, value);
+}
+
+// Forward declaration to use in `make_static_property_type()`
+inline void enable_dynamic_attributes(PyHeapTypeObject *heap_type);
+
+/** A `static_property` is the same as a `property` but the `__get__()` and `__set__()`
+    methods are modified to always use the object type instead of a concrete instance.
+    Return value: New reference. */
+inline PyTypeObject *make_static_property_type() {
+    constexpr auto *name = "pybind11_static_property";
+    auto name_obj = reinterpret_steal<object>(PYBIND11_FROM_STRING(name));
+
+    /* Danger zone: from now (and until PyType_Ready), make sure to
+       issue no Python C API calls which could potentially invoke the
+       garbage collector (the GC will call type_traverse(), which will in
+       turn find the newly constructed type in an invalid state) */
+    auto *heap_type = reinterpret_cast<PyHeapTypeObject *>(PyType_Type.tp_alloc(&PyType_Type, 0));
+    if (!heap_type) {
+        pybind11_fail("make_static_property_type(): error allocating type!");
+    }
+
+    heap_type->ht_name = name_obj.inc_ref().ptr();
+#    ifdef PYBIND11_BUILTIN_QUALNAME
+    heap_type->ht_qualname = name_obj.inc_ref().ptr();
+#    endif
+
+    auto *type = &heap_type->ht_type;
+    type->tp_name = name;
+    type->tp_base = type_incref(&PyProperty_Type);
+    type->tp_flags = Py_TPFLAGS_DEFAULT | Py_TPFLAGS_BASETYPE | Py_TPFLAGS_HEAPTYPE;
+    type->tp_descr_get = pybind11_static_get;
+    type->tp_descr_set = pybind11_static_set;
+
+#    if PY_VERSION_HEX >= 0x030C0000
+    // Since Python-3.12 property-derived types are required to
+    // have dynamic attributes (to set `__doc__`)
+    enable_dynamic_attributes(heap_type);
+#    endif
+
+    if (PyType_Ready(type) < 0) {
+        pybind11_fail("make_static_property_type(): failure in PyType_Ready()!");
+    }
+
+    setattr(reinterpret_cast<PyObject *>(type), "__module__", str(PYBIND11_DUMMY_MODULE_NAME));
+    PYBIND11_SET_OLDPY_QUALNAME(type, name_obj);
+
+    return type;
+}
+
+#else // PYPY
+
+/** PyPy has some issues with the above C API, so we evaluate Python code instead.
+    This function will only be called once so performance isn't really a concern.
+    Return value: New reference. */
+inline PyTypeObject *make_static_property_type() {
+    auto d = dict();
+    PyObject *result = PyRun_String(R"(\
+class pybind11_static_property(property):
+    def __get__(self, obj, cls):
+        return property.__get__(self, cls, cls)
+
+    def __set__(self, obj, value):
+        cls = obj if isinstance(obj, type) else type(obj)
+        property.__set__(self, cls, value)
+)",
+                                    Py_file_input,
+                                    d.ptr(),
+                                    d.ptr());
+    if (result == nullptr)
+        throw error_already_set();
+    Py_DECREF(result);
+    return (PyTypeObject *) d["pybind11_static_property"].cast<object>().release().ptr();
+}
+
+#endif // PYPY
+
+/** Types with static properties need to handle `Type.static_prop = x` in a specific way.
+    By default, Python replaces the `static_property` itself, but for wrapped C++ types
+    we need to call `static_property.__set__()` in order to propagate the new value to
+    the underlying C++ data structure. */
+extern "C" inline int pybind11_meta_setattro(PyObject *obj, PyObject *name, PyObject *value) {
+    // Use `_PyType_Lookup()` instead of `PyObject_GetAttr()` in order to get the raw
+    // descriptor (`property`) instead of calling `tp_descr_get` (`property.__get__()`).
+    PyObject *descr = _PyType_Lookup((PyTypeObject *) obj, name);
+
+    // The following assignment combinations are possible:
+    //   1. `Type.static_prop = value`             --> descr_set: `Type.static_prop.__set__(value)`
+    //   2. `Type.static_prop = other_static_prop` --> setattro:  replace existing `static_prop`
+    //   3. `Type.regular_attribute = value`       --> setattro:  regular attribute assignment
+    auto *const static_prop = (PyObject *) get_internals().static_property_type;
+    const auto call_descr_set = (descr != nullptr) && (value != nullptr)
+                                && (PyObject_IsInstance(descr, static_prop) != 0)
+                                && (PyObject_IsInstance(value, static_prop) == 0);
+    if (call_descr_set) {
+        // Call `static_property.__set__()` instead of replacing the `static_property`.
+#if !defined(PYPY_VERSION)
+        return Py_TYPE(descr)->tp_descr_set(descr, obj, value);
+#else
+        if (PyObject *result = PyObject_CallMethod(descr, "__set__", "OO", obj, value)) {
+            Py_DECREF(result);
+            return 0;
+        } else {
+            return -1;
+        }
+#endif
+    } else {
+        // Replace existing attribute.
+        return PyType_Type.tp_setattro(obj, name, value);
+    }
+}
+
+/**
+ * Python 3's PyInstanceMethod_Type hides itself via its tp_descr_get, which prevents aliasing
+ * methods via cls.attr("m2") = cls.attr("m1"): instead the tp_descr_get returns a plain function,
+ * when called on a class, or a PyMethod, when called on an instance.  Override that behaviour here
+ * to do a special case bypass for PyInstanceMethod_Types.
+ */
+extern "C" inline PyObject *pybind11_meta_getattro(PyObject *obj, PyObject *name) {
+    PyObject *descr = _PyType_Lookup((PyTypeObject *) obj, name);
+    if (descr && PyInstanceMethod_Check(descr)) {
+        Py_INCREF(descr);
+        return descr;
+    }
+    return PyType_Type.tp_getattro(obj, name);
+}
+
+/// metaclass `__call__` function that is used to create all pybind11 objects.
+extern "C" inline PyObject *pybind11_meta_call(PyObject *type, PyObject *args, PyObject *kwargs) {
+
+    // use the default metaclass call to create/initialize the object
+    PyObject *self = PyType_Type.tp_call(type, args, kwargs);
+    if (self == nullptr) {
+        return nullptr;
+    }
+
+    // Ensure that the base __init__ function(s) were called
+    values_and_holders vhs(self);
+    for (const auto &vh : vhs) {
+        if (!vh.holder_constructed() && !vhs.is_redundant_value_and_holder(vh)) {
+            PyErr_Format(PyExc_TypeError,
+                         "%.200s.__init__() must be called when overriding __init__",
+                         get_fully_qualified_tp_name(vh.type->type).c_str());
+            Py_DECREF(self);
+            return nullptr;
+        }
+    }
+
+    return self;
+}
+
+/// Cleanup the type-info for a pybind11-registered type.
+extern "C" inline void pybind11_meta_dealloc(PyObject *obj) {
+    with_internals_if_internals([obj](internals &internals) {
+        auto *type = (PyTypeObject *) obj;
+
+        // A pybind11-registered type will:
+        // 1) be found in internals.registered_types_py
+        // 2) have exactly one associated `detail::type_info`
+        auto found_type = internals.registered_types_py.find(type);
+        if (found_type != internals.registered_types_py.end() && found_type->second.size() == 1
+            && found_type->second[0]->type == type) {
+
+            auto *tinfo = found_type->second[0];
+            auto tindex = std::type_index(*tinfo->cpptype);
+            internals.direct_conversions.erase(tindex);
+
+            auto &local_internals = get_local_internals();
+            if (tinfo->module_local) {
+                local_internals.registered_types_cpp.erase(tinfo->cpptype);
+            } else {
+                internals.registered_types_cpp.erase(tindex);
+#if PYBIND11_INTERNALS_VERSION >= 12
+                internals.registered_types_cpp_fast.erase(tinfo->cpptype);
+                for (const std::type_info *alias : tinfo->alias_chain) {
+                    auto num_erased = internals.registered_types_cpp_fast.erase(alias);
+                    (void) num_erased;
+                    assert(num_erased > 0);
+                }
+#endif
+            }
+            internals.registered_types_py.erase(tinfo->type);
+
+            // Actually just `std::erase_if`, but that's only available in C++20
+            auto &cache = internals.inactive_override_cache;
+            for (auto it = cache.begin(), last = cache.end(); it != last;) {
+                if (it->first == (PyObject *) tinfo->type) {
+                    it = cache.erase(it);
+                } else {
+                    ++it;
+                }
+            }
+
+            delete tinfo;
+        }
+    });
+
+    PyType_Type.tp_dealloc(obj);
+}
+
+/** This metaclass is assigned by default to all pybind11 types and is required in order
+    for static properties to function correctly. Users may override this using `py::metaclass`.
+    Return value: New reference. */
+inline PyTypeObject *make_default_metaclass() {
+    constexpr auto *name = "pybind11_type";
+    auto name_obj = reinterpret_steal<object>(PYBIND11_FROM_STRING(name));
+
+    /* Danger zone: from now (and until PyType_Ready), make sure to
+       issue no Python C API calls which could potentially invoke the
+       garbage collector (the GC will call type_traverse(), which will in
+       turn find the newly constructed type in an invalid state) */
+    auto *heap_type = reinterpret_cast<PyHeapTypeObject *>(PyType_Type.tp_alloc(&PyType_Type, 0));
+    if (!heap_type) {
+        pybind11_fail("make_default_metaclass(): error allocating metaclass!");
+    }
+
+    heap_type->ht_name = name_obj.inc_ref().ptr();
+#ifdef PYBIND11_BUILTIN_QUALNAME
+    heap_type->ht_qualname = name_obj.inc_ref().ptr();
+#endif
+
+    auto *type = &heap_type->ht_type;
+    type->tp_name = name;
+    type->tp_base = type_incref(&PyType_Type);
+    type->tp_flags = Py_TPFLAGS_DEFAULT | Py_TPFLAGS_BASETYPE | Py_TPFLAGS_HEAPTYPE;
+
+    type->tp_call = pybind11_meta_call;
+
+    type->tp_setattro = pybind11_meta_setattro;
+    type->tp_getattro = pybind11_meta_getattro;
+
+    type->tp_dealloc = pybind11_meta_dealloc;
+
+    if (PyType_Ready(type) < 0) {
+        pybind11_fail("make_default_metaclass(): failure in PyType_Ready()!");
+    }
+
+    setattr(reinterpret_cast<PyObject *>(type), "__module__", str(PYBIND11_DUMMY_MODULE_NAME));
+    PYBIND11_SET_OLDPY_QUALNAME(type, name_obj);
+
+    return type;
+}
+
+/// For multiple inheritance types we need to recursively register/deregister base pointers for any
+/// base classes with pointers that are difference from the instance value pointer so that we can
+/// correctly recognize an offset base class pointer. This calls a function with any offset base
+/// ptrs.
+inline void traverse_offset_bases(void *valueptr,
+                                  const detail::type_info *tinfo,
+                                  instance *self,
+                                  bool (*f)(void * /*parentptr*/, instance * /*self*/)) {
+    for (handle h : reinterpret_borrow<tuple>(tinfo->type->tp_bases)) {
+        if (auto *parent_tinfo = get_type_info(reinterpret_cast<PyTypeObject *>(h.ptr()))) {
+            for (auto &c : parent_tinfo->implicit_casts) {
+                if (c.first == tinfo->cpptype) {
+                    auto *parentptr = c.second(valueptr);
+                    if (parentptr != valueptr) {
+                        f(parentptr, self);
+                    }
+                    traverse_offset_bases(parentptr, parent_tinfo, self, f);
+                    break;
+                }
+            }
+        }
+    }
+}
+
+#ifdef Py_GIL_DISABLED
+inline void enable_try_inc_ref(PyObject *obj) {
+#    if PY_VERSION_HEX >= 0x030E00A4
+    PyUnstable_EnableTryIncRef(obj);
+#    else
+    if (_Py_IsImmortal(obj)) {
+        return;
+    }
+    for (;;) {
+        Py_ssize_t shared = _Py_atomic_load_ssize_relaxed(&obj->ob_ref_shared);
+        if ((shared & _Py_REF_SHARED_FLAG_MASK) != 0) {
+            // Nothing to do if it's in WEAKREFS, QUEUED, or MERGED states.
+            return;
+        }
+        if (_Py_atomic_compare_exchange_ssize(
+                &obj->ob_ref_shared, &shared, shared | _Py_REF_MAYBE_WEAKREF)) {
+            return;
+        }
+    }
+#    endif
+}
+#endif
+
+inline bool register_instance_impl(void *ptr, instance *self) {
+    assert(ptr);
+#ifdef Py_GIL_DISABLED
+    enable_try_inc_ref(reinterpret_cast<PyObject *>(self));
+#endif
+    with_instance_map(ptr, [&](instance_map &instances) { instances.emplace(ptr, self); });
+    return true; // unused, but gives the same signature as the deregister func
+}
+inline bool deregister_instance_impl(void *ptr, instance *self) {
+    assert(ptr);
+    return with_instance_map(ptr, [&](instance_map &instances) {
+        auto range = instances.equal_range(ptr);
+        for (auto it = range.first; it != range.second; ++it) {
+            if (self == it->second) {
+                instances.erase(it);
+                return true;
+            }
+        }
+        return false;
+    });
+}
+
+inline void register_instance(instance *self, void *valptr, const type_info *tinfo) {
+    register_instance_impl(valptr, self);
+    if (!tinfo->simple_ancestors) {
+        traverse_offset_bases(valptr, tinfo, self, register_instance_impl);
+    }
+}
+
+inline bool deregister_instance(instance *self, void *valptr, const type_info *tinfo) {
+    bool ret = deregister_instance_impl(valptr, self);
+    if (!tinfo->simple_ancestors) {
+        traverse_offset_bases(valptr, tinfo, self, deregister_instance_impl);
+    }
+    return ret;
+}
+
+/// Instance creation function for all pybind11 types. It allocates the internal instance layout
+/// for holding C++ objects and holders.  Allocation is done lazily (the first time the instance is
+/// cast to a reference or pointer), and initialization is done by an `__init__` function.
+inline PyObject *make_new_instance(PyTypeObject *type) {
+#if defined(PYPY_VERSION)
+    // PyPy gets tp_basicsize wrong (issue 2482) under multiple inheritance when the first
+    // inherited object is a plain Python type (i.e. not derived from an extension type).  Fix it.
+    ssize_t instance_size = static_cast<ssize_t>(sizeof(instance));
+    if (type->tp_basicsize < instance_size) {
+        type->tp_basicsize = instance_size;
+    }
+#endif
+    PyObject *self = type->tp_alloc(type, 0);
+    auto *inst = reinterpret_cast<instance *>(self);
+    // Allocate the value/holder internals:
+    inst->allocate_layout();
+
+    return self;
+}
+
+/// Instance creation function for all pybind11 types. It only allocates space for the
+/// C++ object, but doesn't call the constructor -- an `__init__` function must do that.
+extern "C" inline PyObject *pybind11_object_new(PyTypeObject *type, PyObject *, PyObject *) {
+    return make_new_instance(type);
+}
+
+/// An `__init__` function constructs the C++ object. Users should provide at least one
+/// of these using `py::init` or directly with `.def(__init__, ...)`. Otherwise, the
+/// following default function will be used which simply throws an exception.
+extern "C" inline int pybind11_object_init(PyObject *self, PyObject *, PyObject *) {
+    PyTypeObject *type = Py_TYPE(self);
+    std::string msg = get_fully_qualified_tp_name(type) + ": No constructor defined!";
+    set_error(PyExc_TypeError, msg.c_str());
+    return -1;
+}
+
+inline void add_patient(PyObject *nurse, PyObject *patient) {
+    auto *instance = reinterpret_cast<detail::instance *>(nurse);
+    instance->has_patients = true;
+    Py_INCREF(patient);
+
+    with_internals([&](internals &internals) { internals.patients[nurse].push_back(patient); });
+}
+
+inline void clear_patients(PyObject *self) {
+    auto *instance = reinterpret_cast<detail::instance *>(self);
+    std::vector<PyObject *> patients;
+
+    with_internals([&](internals &internals) {
+        auto pos = internals.patients.find(self);
+
+        if (pos == internals.patients.end()) {
+            pybind11_fail(
+                "FATAL: Internal consistency check failed: Invalid clear_patients() call.");
+        }
+
+        // Clearing the patients can cause more Python code to run, which
+        // can invalidate the iterator. Extract the vector of patients
+        // from the unordered_map first.
+        patients = std::move(pos->second);
+        internals.patients.erase(pos);
+    });
+
+    instance->has_patients = false;
+    for (PyObject *&patient : patients) {
+        Py_CLEAR(patient);
+    }
+}
+
+/// Clears all internal data from the instance and removes it from registered instances in
+/// preparation for deallocation.
+inline void clear_instance(PyObject *self) {
+    auto *instance = reinterpret_cast<detail::instance *>(self);
+
+    // Deallocate any values/holders, if present:
+    for (auto &v_h : values_and_holders(instance)) {
+        if (v_h) {
+
+            // We have to deregister before we call dealloc because, for virtual MI types, we still
+            // need to be able to get the parent pointers.
+            if (v_h.instance_registered()
+                && !deregister_instance(instance, v_h.value_ptr(), v_h.type)) {
+                pybind11_fail(
+                    "pybind11_object_dealloc(): Tried to deallocate unregistered instance!");
+            }
+
+            if (instance->owned || v_h.holder_constructed()) {
+                v_h.type->dealloc(v_h);
+            }
+        } else if (v_h.holder_constructed()) {
+            v_h.type->dealloc(v_h); // Disowned instance.
+        }
+    }
+    // Deallocate the value/holder layout internals:
+    instance->deallocate_layout();
+
+    if (instance->weakrefs) {
+        PyObject_ClearWeakRefs(self);
+    }
+
+    PyObject **dict_ptr = _PyObject_GetDictPtr(self);
+    if (dict_ptr) {
+        Py_CLEAR(*dict_ptr);
+    }
+
+    if (instance->has_patients) {
+        clear_patients(self);
+    }
+}
+
+/// Instance destructor function for all pybind11 types. It calls `type_info.dealloc`
+/// to destroy the C++ object itself, while the rest is Python bookkeeping.
+extern "C" inline void pybind11_object_dealloc(PyObject *self) {
+    auto *type = Py_TYPE(self);
+
+    // If this is a GC tracked object, untrack it first
+    // Note that the track call is implicitly done by the
+    // default tp_alloc, which we never override.
+    if (PyType_HasFeature(type, Py_TPFLAGS_HAVE_GC) != 0) {
+        PyObject_GC_UnTrack(self);
+    }
+
+#if PY_VERSION_HEX >= 0x030D0000
+    // PyObject_ClearManagedDict() is available from Python 3.13+. It must be
+    // called before tp_free() because on Python 3.14+ tp_free no longer
+    // implicitly clears the managed dict, which would abandon the refcounts of
+    // objects stored in __dict__ of py::dynamic_attr() types, causing permanent
+    // memory leaks.
+    if (PyType_HasFeature(type, Py_TPFLAGS_MANAGED_DICT)) {
+        PyObject_ClearManagedDict(self);
+    }
+#endif
+
+    clear_instance(self);
+
+    type->tp_free(self);
+
+    // This was not needed before Python 3.8 (Python issue 35810)
+    // https://github.com/pybind/pybind11/issues/1946
+    Py_DECREF(type);
+}
+
+PYBIND11_WARNING_PUSH
+PYBIND11_WARNING_DISABLE_GCC("-Wredundant-decls")
+
+std::string error_string();
+
+PYBIND11_WARNING_POP
+
+/** Create the type which can be used as a common base for all classes.  This is
+    needed in order to satisfy Python's requirements for multiple inheritance.
+    Return value: New reference. */
+inline PyObject *make_object_base_type(PyTypeObject *metaclass) {
+    constexpr auto *name = "pybind11_object";
+    auto name_obj = reinterpret_steal<object>(PYBIND11_FROM_STRING(name));
+
+    /* Danger zone: from now (and until PyType_Ready), make sure to
+       issue no Python C API calls which could potentially invoke the
+       garbage collector (the GC will call type_traverse(), which will in
+       turn find the newly constructed type in an invalid state) */
+    auto *heap_type = reinterpret_cast<PyHeapTypeObject *>(metaclass->tp_alloc(metaclass, 0));
+    if (!heap_type) {
+        pybind11_fail("make_object_base_type(): error allocating type!");
+    }
+
+    heap_type->ht_name = name_obj.inc_ref().ptr();
+#ifdef PYBIND11_BUILTIN_QUALNAME
+    heap_type->ht_qualname = name_obj.inc_ref().ptr();
+#endif
+
+    auto *type = &heap_type->ht_type;
+    type->tp_name = name;
+    type->tp_base = type_incref(&PyBaseObject_Type);
+    type->tp_basicsize = static_cast<ssize_t>(sizeof(instance));
+    type->tp_flags = Py_TPFLAGS_DEFAULT | Py_TPFLAGS_BASETYPE | Py_TPFLAGS_HEAPTYPE;
+
+    type->tp_new = pybind11_object_new;
+    type->tp_init = pybind11_object_init;
+    type->tp_dealloc = pybind11_object_dealloc;
+
+    /* Support weak references (needed for the keep_alive feature) */
+    type->tp_weaklistoffset = offsetof(instance, weakrefs);
+
+    if (PyType_Ready(type) < 0) {
+        pybind11_fail("PyType_Ready failed in make_object_base_type(): " + error_string());
+    }
+
+    setattr(reinterpret_cast<PyObject *>(type), "__module__", str(PYBIND11_DUMMY_MODULE_NAME));
+    PYBIND11_SET_OLDPY_QUALNAME(type, name_obj);
+
+    assert(!PyType_HasFeature(type, Py_TPFLAGS_HAVE_GC));
+    return reinterpret_cast<PyObject *>(heap_type);
+}
+
+/// dynamic_attr: Allow the garbage collector to traverse the internal instance `__dict__`.
+extern "C" inline int pybind11_traverse(PyObject *self, visitproc visit, void *arg) {
+#if PY_VERSION_HEX >= 0x030D0000
+    int ret = PyObject_VisitManagedDict(self, visit, arg);
+    if (ret) {
+        return ret;
+    }
+#else
+    PyObject *&dict = *_PyObject_GetDictPtr(self);
+    Py_VISIT(dict);
+#endif
+// https://docs.python.org/3/c-api/typeobj.html#c.PyTypeObject.tp_traverse
+#if PY_VERSION_HEX >= 0x03090000
+    Py_VISIT(Py_TYPE(self));
+#endif
+    return 0;
+}
+
+/// dynamic_attr: Allow the GC to clear the dictionary.
+extern "C" inline int pybind11_clear(PyObject *self) {
+#if PY_VERSION_HEX >= 0x030D0000
+    PyObject_ClearManagedDict(self);
+#else
+    PyObject *&dict = *_PyObject_GetDictPtr(self);
+    Py_CLEAR(dict);
+#endif
+    return 0;
+}
+
+/// Give instances of this type a `__dict__` and opt into garbage collection.
+inline void enable_dynamic_attributes(PyHeapTypeObject *heap_type) {
+    auto *type = &heap_type->ht_type;
+    type->tp_flags |= Py_TPFLAGS_HAVE_GC;
+#ifdef PYBIND11_BACKWARD_COMPATIBILITY_TP_DICTOFFSET
+    type->tp_dictoffset = type->tp_basicsize;           // place dict at the end
+    type->tp_basicsize += (ssize_t) sizeof(PyObject *); // and allocate enough space for it
+#else
+    type->tp_flags |= Py_TPFLAGS_MANAGED_DICT;
+#endif
+    type->tp_traverse = pybind11_traverse;
+    type->tp_clear = pybind11_clear;
+
+    static PyGetSetDef getset[]
+        = {{"__dict__", PyObject_GenericGetDict, PyObject_GenericSetDict, nullptr, nullptr},
+           {nullptr, nullptr, nullptr, nullptr, nullptr}};
+    type->tp_getset = getset;
+}
+
+/// buffer_protocol: Fill in the view as specified by flags.
+extern "C" inline int pybind11_getbuffer(PyObject *obj, Py_buffer *view, int flags) {
+    // Look for a `get_buffer` implementation in this type's info or any bases (following MRO).
+    type_info *tinfo = nullptr;
+    for (auto type : reinterpret_borrow<tuple>(Py_TYPE(obj)->tp_mro)) {
+        tinfo = get_type_info((PyTypeObject *) type.ptr());
+        if (tinfo && tinfo->get_buffer) {
+            break;
+        }
+    }
+    if (view == nullptr || !tinfo || !tinfo->get_buffer) {
+        if (view) {
+            view->obj = nullptr;
+        }
+        set_error(PyExc_BufferError, "pybind11_getbuffer(): Internal error");
+        return -1;
+    }
+    std::memset(view, 0, sizeof(Py_buffer));
+    std::unique_ptr<buffer_info> info = nullptr;
+    try {
+        info.reset(tinfo->get_buffer(obj, tinfo->get_buffer_data));
+    } catch (...) {
+        try_translate_exceptions();
+        raise_from(PyExc_BufferError, "Error getting buffer");
+        return -1;
+    }
+    if (info == nullptr) {
+        pybind11_fail("FATAL UNEXPECTED SITUATION: tinfo->get_buffer() returned nullptr.");
+    }
+
+    if ((flags & PyBUF_WRITABLE) == PyBUF_WRITABLE && info->readonly) {
+        // view->obj = nullptr;  // Was just memset to 0, so not necessary
+        set_error(PyExc_BufferError, "Writable buffer requested for readonly storage");
+        return -1;
+    }
+
+    // Fill in all the information, and then downgrade as requested by the caller, or raise an
+    // error if that's not possible.
+    view->itemsize = info->itemsize;
+    view->len = view->itemsize;
+    for (auto s : info->shape) {
+        view->len *= s;
+    }
+    view->ndim = static_cast<int>(info->ndim);
+    view->shape = info->shape.data();
+    view->strides = info->strides.data();
+    view->readonly = static_cast<int>(info->readonly);
+    if ((flags & PyBUF_FORMAT) == PyBUF_FORMAT) {
+        view->format = const_cast<char *>(info->format.c_str());
+    }
+
+    // Note, all contiguity flags imply PyBUF_STRIDES and lower.
+    if ((flags & PyBUF_C_CONTIGUOUS) == PyBUF_C_CONTIGUOUS) {
+        if (PyBuffer_IsContiguous(view, 'C') == 0) {
+            std::memset(view, 0, sizeof(Py_buffer));
+            set_error(PyExc_BufferError,
+                      "C-contiguous buffer requested for discontiguous storage");
+            return -1;
+        }
+    } else if ((flags & PyBUF_F_CONTIGUOUS) == PyBUF_F_CONTIGUOUS) {
+        if (PyBuffer_IsContiguous(view, 'F') == 0) {
+            std::memset(view, 0, sizeof(Py_buffer));
+            set_error(PyExc_BufferError,
+                      "Fortran-contiguous buffer requested for discontiguous storage");
+            return -1;
+        }
+    } else if ((flags & PyBUF_ANY_CONTIGUOUS) == PyBUF_ANY_CONTIGUOUS) {
+        if (PyBuffer_IsContiguous(view, 'A') == 0) {
+            std::memset(view, 0, sizeof(Py_buffer));
+            set_error(PyExc_BufferError, "Contiguous buffer requested for discontiguous storage");
+            return -1;
+        }
+
+    } else if ((flags & PyBUF_STRIDES) != PyBUF_STRIDES) {
+        // If no strides are requested, the buffer must be C-contiguous.
+        // https://docs.python.org/3/c-api/buffer.html#contiguity-requests
+        if (PyBuffer_IsContiguous(view, 'C') == 0) {
+            std::memset(view, 0, sizeof(Py_buffer));
+            set_error(PyExc_BufferError,
+                      "C-contiguous buffer requested for discontiguous storage");
+            return -1;
+        }
+
+        view->strides = nullptr;
+
+        // Since this is a contiguous buffer, it can also pretend to be 1D.
+        if ((flags & PyBUF_ND) != PyBUF_ND) {
+            view->shape = nullptr;
+            view->ndim = 0;
+        }
+    }
+
+    // Set these after all checks so they don't leak out into the caller, and can be automatically
+    // cleaned up on error.
+    view->buf = info->ptr;
+    view->internal = info.release();
+    view->obj = obj;
+    Py_INCREF(view->obj);
+    return 0;
+}
+
+/// buffer_protocol: Release the resources of the buffer.
+extern "C" inline void pybind11_releasebuffer(PyObject *, Py_buffer *view) {
+    delete (buffer_info *) view->internal;
+}
+
+/// Give this type a buffer interface.
+inline void enable_buffer_protocol(PyHeapTypeObject *heap_type) {
+    heap_type->ht_type.tp_as_buffer = &heap_type->as_buffer;
+
+    heap_type->as_buffer.bf_getbuffer = pybind11_getbuffer;
+    heap_type->as_buffer.bf_releasebuffer = pybind11_releasebuffer;
+}
+
+/** Create a brand new Python type according to the `type_record` specification.
+    Return value: New reference. */
+inline PyObject *make_new_python_type(const type_record &rec) {
+    auto name = reinterpret_steal<object>(PYBIND11_FROM_STRING(rec.name));
+
+    auto qualname = name;
+    if (rec.scope && !PyModule_Check(rec.scope.ptr()) && hasattr(rec.scope, "__qualname__")) {
+        qualname = reinterpret_steal<object>(
+            PyUnicode_FromFormat("%U.%U", rec.scope.attr("__qualname__").ptr(), name.ptr()));
+    }
+
+    object module_ = get_module_name_if_available(rec.scope);
+    const auto *full_name = c_str(
+#if !defined(PYPY_VERSION)
+        module_ ? str(module_).cast<std::string>() + "." + rec.name :
+#endif
+                rec.name);
+
+    char *tp_doc = nullptr;
+    if (rec.doc && options::show_user_defined_docstrings()) {
+        /* Allocate memory for docstring (Python will free this later on) */
+        size_t size = std::strlen(rec.doc) + 1;
+#if PY_VERSION_HEX >= 0x030D0000
+        tp_doc = static_cast<char *>(PyMem_MALLOC(size));
+#else
+        tp_doc = (char *) PyObject_MALLOC(size);
+#endif
+        std::memcpy((void *) tp_doc, rec.doc, size);
+    }
+
+    auto &internals = get_internals();
+    auto bases = tuple(rec.bases);
+    auto *base = (bases.empty()) ? internals.instance_base : bases[0].ptr();
+
+    /* Danger zone: from now (and until PyType_Ready), make sure to
+       issue no Python C API calls which could potentially invoke the
+       garbage collector (the GC will call type_traverse(), which will in
+       turn find the newly constructed type in an invalid state) */
+    auto *metaclass = rec.metaclass.ptr() ? reinterpret_cast<PyTypeObject *>(rec.metaclass.ptr())
+                                          : internals.default_metaclass;
+
+    auto *heap_type = reinterpret_cast<PyHeapTypeObject *>(metaclass->tp_alloc(metaclass, 0));
+    if (!heap_type) {
+        pybind11_fail(std::string(rec.name) + ": Unable to create type object!");
+    }
+
+    heap_type->ht_name = name.release().ptr();
+#ifdef PYBIND11_BUILTIN_QUALNAME
+    heap_type->ht_qualname = qualname.inc_ref().ptr();
+#endif
+
+    auto *type = &heap_type->ht_type;
+    type->tp_name = full_name;
+    type->tp_doc = tp_doc;
+    type->tp_base = type_incref(reinterpret_cast<PyTypeObject *>(base));
+    type->tp_basicsize = static_cast<ssize_t>(sizeof(instance));
+    if (!bases.empty()) {
+        type->tp_bases = bases.release().ptr();
+    }
+
+    /* Don't inherit base __init__ */
+    type->tp_init = pybind11_object_init;
+
+    /* Supported protocols */
+    type->tp_as_number = &heap_type->as_number;
+    type->tp_as_sequence = &heap_type->as_sequence;
+    type->tp_as_mapping = &heap_type->as_mapping;
+    type->tp_as_async = &heap_type->as_async;
+
+    /* Flags */
+    type->tp_flags |= Py_TPFLAGS_DEFAULT | Py_TPFLAGS_HEAPTYPE;
+    if (!rec.is_final) {
+        type->tp_flags |= Py_TPFLAGS_BASETYPE;
+    }
+
+    if (rec.dynamic_attr) {
+        enable_dynamic_attributes(heap_type);
+    }
+
+    if (rec.buffer_protocol) {
+        enable_buffer_protocol(heap_type);
+    }
+
+    if (rec.custom_type_setup_callback) {
+        rec.custom_type_setup_callback(heap_type);
+    }
+
+    if (PyType_Ready(type) < 0) {
+        pybind11_fail(std::string(rec.name) + ": PyType_Ready failed: " + error_string());
+    }
+
+    assert(!rec.dynamic_attr || PyType_HasFeature(type, Py_TPFLAGS_HAVE_GC));
+
+    /* Register type with the parent scope */
+    if (rec.scope) {
+        setattr(rec.scope, rec.name, reinterpret_cast<PyObject *>(type));
+    } else {
+        Py_INCREF(type); // Keep it alive forever (reference leak)
+    }
+
+    if (module_) { // Needed by pydoc
+        setattr(reinterpret_cast<PyObject *>(type), "__module__", module_);
+    }
+
+    PYBIND11_SET_OLDPY_QUALNAME(type, qualname);
+
+    return reinterpret_cast<PyObject *>(type);
+}
+
+PYBIND11_NAMESPACE_END(detail)
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/detail/cpp_conduit.h

@@ -0,0 +1,75 @@
+// Copyright (c) 2024 The pybind Community.
+
+#pragma once
+
+#include <pybind11/pytypes.h>
+
+#include "common.h"
+#include "internals.h"
+
+#include <typeinfo>
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+// Forward declaration needed here: Refactoring opportunity.
+extern "C" inline PyObject *pybind11_object_new(PyTypeObject *type, PyObject *, PyObject *);
+
+inline bool type_is_managed_by_our_internals(PyTypeObject *type_obj) {
+#if defined(PYPY_VERSION)
+    auto &internals = get_internals();
+    return bool(internals.registered_types_py.find(type_obj)
+                != internals.registered_types_py.end());
+#else
+    return (type_obj->tp_new == pybind11_object_new);
+#endif
+}
+
+inline bool is_instance_method_of_type(PyTypeObject *type_obj, PyObject *attr_name) {
+    PyObject *descr = _PyType_Lookup(type_obj, attr_name);
+    return ((descr != nullptr) && PyInstanceMethod_Check(descr));
+}
+
+inline object try_get_cpp_conduit_method(PyObject *obj) {
+    if (PyType_Check(obj)) {
+        return object();
+    }
+    PyTypeObject *type_obj = Py_TYPE(obj);
+    str attr_name("_pybind11_conduit_v1_");
+    bool assumed_to_be_callable = false;
+    if (type_is_managed_by_our_internals(type_obj)) {
+        if (!is_instance_method_of_type(type_obj, attr_name.ptr())) {
+            return object();
+        }
+        assumed_to_be_callable = true;
+    }
+    PyObject *method = PyObject_GetAttr(obj, attr_name.ptr());
+    if (method == nullptr) {
+        PyErr_Clear();
+        return object();
+    }
+    if (!assumed_to_be_callable && PyCallable_Check(method) == 0) {
+        Py_DECREF(method);
+        return object();
+    }
+    return reinterpret_steal<object>(method);
+}
+
+inline void *try_raw_pointer_ephemeral_from_cpp_conduit(handle src,
+                                                        const std::type_info *cpp_type_info) {
+    object method = try_get_cpp_conduit_method(src.ptr());
+    if (method) {
+        capsule cpp_type_info_capsule(const_cast<void *>(static_cast<const void *>(cpp_type_info)),
+                                      typeid(std::type_info).name());
+        object cpp_conduit = method(bytes(PYBIND11_PLATFORM_ABI_ID),
+                                    cpp_type_info_capsule,
+                                    bytes("raw_pointer_ephemeral"));
+        if (isinstance<capsule>(cpp_conduit)) {
+            return reinterpret_borrow<capsule>(cpp_conduit).get_pointer();
+        }
+    }
+    return nullptr;
+}
+
+PYBIND11_NAMESPACE_END(detail)
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/detail/descr.h

@@ -0,0 +1,231 @@
+/*
+    pybind11/detail/descr.h: Helper type for concatenating type signatures at compile time
+
+    Copyright (c) 2016 Wenzel Jakob <wenzel.jakob@epfl.ch>
+
+    All rights reserved. Use of this source code is governed by a
+    BSD-style license that can be found in the LICENSE file.
+*/
+
+#pragma once
+
+#include "common.h"
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+#if !defined(_MSC_VER)
+#    define PYBIND11_DESCR_CONSTEXPR static constexpr
+#else
+#    define PYBIND11_DESCR_CONSTEXPR const
+#endif
+
+/* Concatenate type signatures at compile time */
+template <size_t N, typename... Ts>
+struct descr {
+    char text[N + 1]{'\0'};
+
+    constexpr descr() = default;
+    // NOLINTNEXTLINE(google-explicit-constructor)
+    constexpr descr(char const (&s)[N + 1]) : descr(s, make_index_sequence<N>()) {}
+
+    template <size_t... Is>
+    constexpr descr(char const (&s)[N + 1], index_sequence<Is...>) : text{s[Is]..., '\0'} {}
+
+    template <typename... Chars>
+    // NOLINTNEXTLINE(google-explicit-constructor)
+    constexpr descr(char c, Chars... cs) : text{c, static_cast<char>(cs)..., '\0'} {}
+
+    static constexpr std::array<const std::type_info *, sizeof...(Ts) + 1> types() {
+        return {{&typeid(Ts)..., nullptr}};
+    }
+};
+
+template <size_t N1, size_t N2, typename... Ts1, typename... Ts2, size_t... Is1, size_t... Is2>
+constexpr descr<N1 + N2, Ts1..., Ts2...> plus_impl(const descr<N1, Ts1...> &a,
+                                                   const descr<N2, Ts2...> &b,
+                                                   index_sequence<Is1...>,
+                                                   index_sequence<Is2...>) {
+    PYBIND11_WORKAROUND_INCORRECT_MSVC_C4100(b);
+    return {a.text[Is1]..., b.text[Is2]...};
+}
+
+template <size_t N1, size_t N2, typename... Ts1, typename... Ts2>
+constexpr descr<N1 + N2, Ts1..., Ts2...> operator+(const descr<N1, Ts1...> &a,
+                                                   const descr<N2, Ts2...> &b) {
+    return plus_impl(a, b, make_index_sequence<N1>(), make_index_sequence<N2>());
+}
+
+template <size_t N>
+constexpr descr<N - 1> const_name(char const (&text)[N]) {
+    return descr<N - 1>(text);
+}
+constexpr descr<0> const_name(char const (&)[1]) { return {}; }
+
+template <size_t Rem, size_t... Digits>
+struct int_to_str : int_to_str<Rem / 10, Rem % 10, Digits...> {};
+template <size_t... Digits>
+struct int_to_str<0, Digits...> {
+    // WARNING: This only works with C++17 or higher.
+    static constexpr auto digits = descr<sizeof...(Digits)>(('0' + Digits)...);
+};
+
+// Ternary description (like std::conditional)
+template <bool B, size_t N1, size_t N2>
+constexpr enable_if_t<B, descr<N1 - 1>> const_name(char const (&text1)[N1], char const (&)[N2]) {
+    return const_name(text1);
+}
+template <bool B, size_t N1, size_t N2>
+constexpr enable_if_t<!B, descr<N2 - 1>> const_name(char const (&)[N1], char const (&text2)[N2]) {
+    return const_name(text2);
+}
+
+template <bool B, typename T1, typename T2>
+constexpr enable_if_t<B, T1> const_name(const T1 &d, const T2 &) {
+    return d;
+}
+template <bool B, typename T1, typename T2>
+constexpr enable_if_t<!B, T2> const_name(const T1 &, const T2 &d) {
+    return d;
+}
+
+template <size_t Size>
+auto constexpr const_name() -> remove_cv_t<decltype(int_to_str<Size / 10, Size % 10>::digits)> {
+    return int_to_str<Size / 10, Size % 10>::digits;
+}
+
+template <typename Type>
+constexpr descr<1, Type> const_name() {
+    return {'%'};
+}
+
+// Use a different name based on whether the parameter is used as input or output
+template <size_t N1, size_t N2>
+constexpr descr<N1 + N2 + 1> io_name(char const (&text1)[N1], char const (&text2)[N2]) {
+    return const_name("@") + const_name(text1) + const_name("@") + const_name(text2)
+           + const_name("@");
+}
+
+// Ternary description for io_name (like the numeric type_caster)
+template <bool B, size_t N1, size_t N2, size_t N3, size_t N4>
+constexpr enable_if_t<B, descr<N1 + N2 + 1>>
+io_name(char const (&text1)[N1], char const (&text2)[N2], char const (&)[N3], char const (&)[N4]) {
+    return io_name(text1, text2);
+}
+
+template <bool B, size_t N1, size_t N2, size_t N3, size_t N4>
+constexpr enable_if_t<!B, descr<N3 + N4 + 1>>
+io_name(char const (&)[N1], char const (&)[N2], char const (&text3)[N3], char const (&text4)[N4]) {
+    return io_name(text3, text4);
+}
+
+// If "_" is defined as a macro, py::detail::_ cannot be provided.
+// It is therefore best to use py::detail::const_name universally.
+// This block is for backward compatibility only.
+// (The const_name code is repeated to avoid introducing a "_" #define ourselves.)
+#ifndef _
+#    define PYBIND11_DETAIL_UNDERSCORE_BACKWARD_COMPATIBILITY
+template <size_t N>
+constexpr descr<N - 1> _(char const (&text)[N]) {
+    return const_name<N>(text);
+}
+template <bool B, size_t N1, size_t N2>
+constexpr enable_if_t<B, descr<N1 - 1>> _(char const (&text1)[N1], char const (&text2)[N2]) {
+    return const_name<B, N1, N2>(text1, text2);
+}
+template <bool B, size_t N1, size_t N2>
+constexpr enable_if_t<!B, descr<N2 - 1>> _(char const (&text1)[N1], char const (&text2)[N2]) {
+    return const_name<B, N1, N2>(text1, text2);
+}
+template <bool B, typename T1, typename T2>
+constexpr enable_if_t<B, T1> _(const T1 &d1, const T2 &d2) {
+    return const_name<B, T1, T2>(d1, d2);
+}
+template <bool B, typename T1, typename T2>
+constexpr enable_if_t<!B, T2> _(const T1 &d1, const T2 &d2) {
+    return const_name<B, T1, T2>(d1, d2);
+}
+
+template <size_t Size>
+auto constexpr _() -> remove_cv_t<decltype(int_to_str<Size / 10, Size % 10>::digits)> {
+    return const_name<Size>();
+}
+template <typename Type>
+constexpr descr<1, Type> _() {
+    return const_name<Type>();
+}
+#endif // #ifndef _
+
+constexpr descr<0> concat() { return {}; }
+constexpr descr<0> union_concat() { return {}; }
+
+template <size_t N, typename... Ts>
+constexpr descr<N, Ts...> concat(const descr<N, Ts...> &descr) {
+    return descr;
+}
+
+template <size_t N, typename... Ts>
+constexpr descr<N, Ts...> union_concat(const descr<N, Ts...> &descr) {
+    return descr;
+}
+
+template <size_t N1, size_t N2, typename... Ts1, typename... Ts2>
+constexpr descr<N1 + N2 + 3, Ts1..., Ts2...> operator|(const descr<N1, Ts1...> &a,
+                                                       const descr<N2, Ts2...> &b) {
+    return a + const_name(" | ") + b;
+}
+
+#ifdef __cpp_fold_expressions
+template <size_t N1, size_t N2, typename... Ts1, typename... Ts2>
+constexpr descr<N1 + N2 + 2, Ts1..., Ts2...> operator,(const descr<N1, Ts1...> &a,
+                                                       const descr<N2, Ts2...> &b) {
+    return a + const_name(", ") + b;
+}
+
+template <size_t N, typename... Ts, typename... Args>
+constexpr auto concat(const descr<N, Ts...> &d, const Args &...args) {
+    return (d, ..., args);
+}
+
+template <size_t N, typename... Ts, typename... Args>
+constexpr auto union_concat(const descr<N, Ts...> &d, const Args &...args) {
+    return (d | ... | args);
+}
+
+#else
+template <size_t N, typename... Ts, typename... Args>
+constexpr auto concat(const descr<N, Ts...> &d, const Args &...args)
+    -> decltype(std::declval<descr<N + 2, Ts...>>() + concat(args...)) {
+    return d + const_name(", ") + concat(args...);
+}
+
+template <size_t N, typename... Ts, typename... Args>
+constexpr auto union_concat(const descr<N, Ts...> &d, const Args &...args)
+    -> decltype(std::declval<descr<N + 3, Ts...>>() + union_concat(args...)) {
+    return d + const_name(" | ") + union_concat(args...);
+}
+
+#endif
+
+template <size_t N, typename... Ts>
+constexpr descr<N + 2, Ts...> type_descr(const descr<N, Ts...> &descr) {
+    return const_name("{") + descr + const_name("}");
+}
+
+template <size_t N, typename... Ts>
+constexpr descr<N + 4, Ts...> arg_descr(const descr<N, Ts...> &descr) {
+    return const_name("@^") + descr + const_name("@!");
+}
+
+template <size_t N, typename... Ts>
+constexpr descr<N + 4, Ts...> return_descr(const descr<N, Ts...> &descr) {
+    return const_name("@$") + descr + const_name("@!");
+}
+
+template <size_t N, typename... Ts>
+constexpr descr<N + 4, Ts...> inv_descr(const descr<N, Ts...> &descr) {
+    return const_name("@~") + descr + const_name("@!");
+}
+
+PYBIND11_NAMESPACE_END(detail)
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/detail/dynamic_raw_ptr_cast_if_possible.h

@@ -0,0 +1,39 @@
+// Copyright (c) 2021 The Pybind Development Team.
+// All rights reserved. Use of this source code is governed by a
+// BSD-style license that can be found in the LICENSE file.
+
+#pragma once
+
+#include "common.h"
+
+#include <type_traits>
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+template <typename To, typename From, typename SFINAE = void>
+struct dynamic_raw_ptr_cast_is_possible : std::false_type {};
+
+template <typename To, typename From>
+struct dynamic_raw_ptr_cast_is_possible<
+    To,
+    From,
+    detail::enable_if_t<!std::is_same<To, void>::value && std::is_polymorphic<From>::value>>
+    : std::true_type {};
+
+template <typename To,
+          typename From,
+          detail::enable_if_t<!dynamic_raw_ptr_cast_is_possible<To, From>::value, int> = 0>
+To *dynamic_raw_ptr_cast_if_possible(From * /*ptr*/) {
+    return nullptr;
+}
+
+template <typename To,
+          typename From,
+          detail::enable_if_t<dynamic_raw_ptr_cast_is_possible<To, From>::value, int> = 0>
+To *dynamic_raw_ptr_cast_if_possible(From *ptr) {
+    return dynamic_cast<To *>(ptr);
+}
+
+PYBIND11_NAMESPACE_END(detail)
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/detail/exception_translation.h

@@ -0,0 +1,71 @@
+/*
+    pybind11/detail/exception_translation.h: means to translate C++ exceptions to Python exceptions
+
+    Copyright (c) 2024 The Pybind Development Team.
+
+    All rights reserved. Use of this source code is governed by a
+    BSD-style license that can be found in the LICENSE file.
+*/
+
+#pragma once
+
+#include "common.h"
+#include "internals.h"
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+// Apply all the extensions translators from a list
+// Return true if one of the translators completed without raising an exception
+// itself. Return of false indicates that if there are other translators
+// available, they should be tried.
+inline bool apply_exception_translators(std::forward_list<ExceptionTranslator> &translators) {
+    auto last_exception = std::current_exception();
+
+    for (auto &translator : translators) {
+        try {
+            translator(last_exception);
+            return true;
+        } catch (...) {
+            last_exception = std::current_exception();
+        }
+    }
+    return false;
+}
+
+inline void try_translate_exceptions() {
+    /* When an exception is caught, give each registered exception
+        translator a chance to translate it to a Python exception. First
+        all module-local translators will be tried in reverse order of
+        registration. If none of the module-locale translators handle
+        the exception (or there are no module-locale translators) then
+        the global translators will be tried, also in reverse order of
+        registration.
+
+        A translator may choose to do one of the following:
+
+        - catch the exception and call py::set_error()
+            to set a standard (or custom) Python exception, or
+        - do nothing and let the exception fall through to the next translator, or
+        - delegate translation to the next translator by throwing a new type of exception.
+        */
+
+    bool handled = with_exception_translators(
+        [&](std::forward_list<ExceptionTranslator> &exception_translators,
+            std::forward_list<ExceptionTranslator> &local_exception_translators) {
+            if (detail::apply_exception_translators(local_exception_translators)) {
+                return true;
+            }
+            if (detail::apply_exception_translators(exception_translators)) {
+                return true;
+            }
+            return false;
+        });
+
+    if (!handled) {
+        set_error(PyExc_SystemError, "Exception escaped from default exception translator!");
+    }
+}
+
+PYBIND11_NAMESPACE_END(detail)
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/detail/function_record_pyobject.h

@@ -0,0 +1,192 @@
+// Copyright (c) 2024-2025 The Pybind Development Team.
+// All rights reserved. Use of this source code is governed by a
+// BSD-style license that can be found in the LICENSE file.
+
+// For background see the description of PR google/pybind11clif#30099.
+
+#pragma once
+
+#include <pybind11/attr.h>
+#include <pybind11/conduit/pybind11_platform_abi_id.h>
+#include <pybind11/pytypes.h>
+
+#include "common.h"
+
+#include <cstring>
+#include <utility>
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+struct function_record_PyObject {
+    PyObject_HEAD
+    function_record *cpp_func_rec;
+};
+
+PYBIND11_NAMESPACE_BEGIN(function_record_PyTypeObject_methods)
+
+PyObject *tp_new_impl(PyTypeObject *type, PyObject *args, PyObject *kwds);
+PyObject *tp_alloc_impl(PyTypeObject *type, Py_ssize_t nitems);
+int tp_init_impl(PyObject *self, PyObject *args, PyObject *kwds);
+void tp_dealloc_impl(PyObject *self);
+void tp_free_impl(void *self);
+
+static PyObject *reduce_ex_impl(PyObject *self, PyObject *, PyObject *);
+
+static PyMethodDef tp_methods_impl[]
+    = {{"__reduce_ex__",
+        // reduce_ex_impl is a PyCFunctionWithKeywords, but PyMethodDef
+        // requires a PyCFunction. The cast through void* is safe and
+        // idiomatic with METH_KEYWORDS, and it successfully sidesteps
+        // unhelpful compiler warnings.
+        // NOLINTNEXTLINE(bugprone-casting-through-void)
+        reinterpret_cast<PyCFunction>(reinterpret_cast<void *>(reduce_ex_impl)),
+        METH_VARARGS | METH_KEYWORDS,
+        nullptr},
+       {nullptr, nullptr, 0, nullptr}};
+
+// Python 3.12+ emits a DeprecationWarning for heap types whose tp_name does
+// not contain a dot ('.') and that lack a __module__ attribute. For pybind11's
+// internal function_record type, we do not have an actual module object to
+// attach, so we cannot use PyType_FromModuleAndSpec (introduced in Python 3.9)
+// to set __module__ automatically.
+//
+// As a workaround, we define a "qualified" type name that includes a dummy
+// module name (PYBIND11_DUMMY_MODULE_NAME). This is non‑idiomatic but avoids
+// the deprecation warning, and results in reprs like
+//
+//     <class 'pybind11_builtins.pybind11_detail_function_record_...'>
+//
+// even though no real pybind11_builtins module exists. If pybind11 gains an
+// actual module object in the future, this code should switch to
+// PyType_FromModuleAndSpec for Python 3.9+ and drop the dummy module
+// workaround.
+//
+// Note that this name is versioned.
+#define PYBIND11_DETAIL_FUNCTION_RECORD_TP_PLAINNAME                                              \
+    "pybind11_detail_function_record_" PYBIND11_DETAIL_FUNCTION_RECORD_ABI_ID                     \
+    "_" PYBIND11_PLATFORM_ABI_ID
+constexpr char tp_plainname_impl[] = PYBIND11_DETAIL_FUNCTION_RECORD_TP_PLAINNAME;
+constexpr char tp_qualname_impl[]
+    = PYBIND11_DUMMY_MODULE_NAME "." PYBIND11_DETAIL_FUNCTION_RECORD_TP_PLAINNAME;
+
+PYBIND11_NAMESPACE_END(function_record_PyTypeObject_methods)
+
+static PyType_Slot function_record_PyType_Slots[] = {
+    {Py_tp_dealloc,
+     reinterpret_cast<void *>(function_record_PyTypeObject_methods::tp_dealloc_impl)},
+    {Py_tp_methods,
+     reinterpret_cast<void *>(function_record_PyTypeObject_methods::tp_methods_impl)},
+    {Py_tp_init, reinterpret_cast<void *>(function_record_PyTypeObject_methods::tp_init_impl)},
+    {Py_tp_alloc, reinterpret_cast<void *>(function_record_PyTypeObject_methods::tp_alloc_impl)},
+    {Py_tp_new, reinterpret_cast<void *>(function_record_PyTypeObject_methods::tp_new_impl)},
+    {Py_tp_free, reinterpret_cast<void *>(function_record_PyTypeObject_methods::tp_free_impl)},
+    {0, nullptr}};
+
+static PyType_Spec function_record_PyType_Spec
+    = {function_record_PyTypeObject_methods::tp_qualname_impl,
+       sizeof(function_record_PyObject),
+       0,
+       Py_TPFLAGS_DEFAULT | Py_TPFLAGS_HEAPTYPE,
+       function_record_PyType_Slots};
+
+inline PyTypeObject *get_function_record_PyTypeObject() {
+    PYBIND11_LOCK_INTERNALS(get_internals());
+    PyTypeObject *&py_type_obj = detail::get_local_internals().function_record_py_type;
+    if (!py_type_obj) {
+        PyObject *py_obj = PyType_FromSpec(&function_record_PyType_Spec);
+        if (py_obj == nullptr) {
+            throw error_already_set();
+        }
+        py_type_obj = reinterpret_cast<PyTypeObject *>(py_obj);
+    }
+    return py_type_obj;
+}
+
+inline bool is_function_record_PyObject(PyObject *obj) {
+    if (PyType_Check(obj) != 0) {
+        return false;
+    }
+    PyTypeObject *obj_type = Py_TYPE(obj);
+
+    PyTypeObject *frtype = get_function_record_PyTypeObject();
+
+    // Fast path (pointer comparison).
+    if (obj_type == frtype) {
+        return true;
+    }
+    // This works across extension modules. Note that tp_name is versioned.
+    if (strcmp(obj_type->tp_name, function_record_PyTypeObject_methods::tp_qualname_impl) == 0
+        || strcmp(obj_type->tp_name, function_record_PyTypeObject_methods::tp_plainname_impl)
+               == 0) {
+        return true;
+    }
+    return false;
+}
+
+inline function_record *function_record_ptr_from_PyObject(PyObject *obj) {
+    if (is_function_record_PyObject(obj)) {
+        return (reinterpret_cast<detail::function_record_PyObject *>(obj))->cpp_func_rec;
+    }
+    return nullptr;
+}
+
+inline object function_record_PyObject_New() {
+    auto *py_func_rec = PyObject_New(function_record_PyObject, get_function_record_PyTypeObject());
+    if (py_func_rec == nullptr) {
+        throw error_already_set();
+    }
+    py_func_rec->cpp_func_rec = nullptr; // For clarity/purity. Redundant in practice.
+    return reinterpret_steal<object>(reinterpret_cast<PyObject *>(py_func_rec));
+}
+
+PYBIND11_NAMESPACE_BEGIN(function_record_PyTypeObject_methods)
+
+// Guard against accidents & oversights, in particular when porting to future Python versions.
+inline PyObject *tp_new_impl(PyTypeObject *, PyObject *, PyObject *) {
+    pybind11_fail("UNEXPECTED CALL OF function_record_PyTypeObject_methods::tp_new_impl");
+    // return nullptr; // Unreachable.
+}
+
+inline PyObject *tp_alloc_impl(PyTypeObject *, Py_ssize_t) {
+    pybind11_fail("UNEXPECTED CALL OF function_record_PyTypeObject_methods::tp_alloc_impl");
+    // return nullptr; // Unreachable.
+}
+
+inline int tp_init_impl(PyObject *, PyObject *, PyObject *) {
+    pybind11_fail("UNEXPECTED CALL OF function_record_PyTypeObject_methods::tp_init_impl");
+    // return -1; // Unreachable.
+}
+
+inline void tp_free_impl(void *) {
+    pybind11_fail("UNEXPECTED CALL OF function_record_PyTypeObject_methods::tp_free_impl");
+}
+
+inline PyObject *reduce_ex_impl(PyObject *self, PyObject *, PyObject *) {
+    // Deliberately ignoring the arguments for simplicity (expected is `protocol: int`).
+    const function_record *rec = function_record_ptr_from_PyObject(self);
+    if (rec == nullptr) {
+        pybind11_fail(
+            "FATAL: function_record_PyTypeObject reduce_ex_impl(): cannot obtain cpp_func_rec.");
+    }
+    if (rec->name != nullptr && rec->name[0] != '\0' && rec->scope
+        && PyModule_Check(rec->scope.ptr()) != 0) {
+        object scope_module = get_scope_module(rec->scope);
+        if (scope_module) {
+            auto builtins = reinterpret_borrow<dict>(PyEval_GetBuiltins());
+            auto builtins_eval = builtins["eval"];
+            auto reconstruct_args = make_tuple(str("__import__('importlib').import_module('")
+                                               + scope_module + str("')"));
+            return make_tuple(std::move(builtins_eval), std::move(reconstruct_args))
+                .release()
+                .ptr();
+        }
+    }
+    set_error(PyExc_RuntimeError, repr(self) + str(" is not pickleable."));
+    return nullptr;
+}
+
+PYBIND11_NAMESPACE_END(function_record_PyTypeObject_methods)
+
+PYBIND11_NAMESPACE_END(detail)
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/detail/function_ref.h

@@ -0,0 +1,101 @@
+/*
+ * Copyright (c) Meta Platforms, Inc. and affiliates.
+ * All rights reserved.
+ *
+ * This source code is licensed under the BSD-style license found in the
+ * LICENSE file in the root directory of this source tree.
+ */
+
+//===- llvm/ADT/STLFunctionalExtras.h - Extras for <functional> -*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file contains a header-only class template that provides functionality
+// similar to std::function but with non-owning semantics. It is a template-only
+// implementation that requires no additional library linking.
+//
+//===----------------------------------------------------------------------===//
+
+/// An efficient, type-erasing, non-owning reference to a callable. This is
+/// intended for use as the type of a function parameter that is not used
+/// after the function in question returns.
+///
+/// This class does not own the callable, so it is not in general safe to store
+/// a FunctionRef.
+
+// pybind11: modified again from executorch::runtime::FunctionRef
+// - renamed back to function_ref
+// - use pybind11 enable_if_t, remove_cvref_t, and remove_reference_t
+// - lint suppressions
+
+// torch::executor: modified from llvm::function_ref
+// - renamed to FunctionRef
+// - removed LLVM_GSL_POINTER and LLVM_LIFETIME_BOUND macro uses
+// - use namespaced internal::remove_cvref_t
+
+#pragma once
+
+#include <pybind11/detail/common.h>
+
+#include <cstdint>
+#include <type_traits>
+#include <utility>
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+//===----------------------------------------------------------------------===//
+//     Features from C++20
+//===----------------------------------------------------------------------===//
+
+template <typename Fn>
+class function_ref;
+
+template <typename Ret, typename... Params>
+class function_ref<Ret(Params...)> {
+    Ret (*callback)(intptr_t callable, Params... params) = nullptr;
+    intptr_t callable;
+
+    template <typename Callable>
+    // NOLINTNEXTLINE(performance-unnecessary-value-param)
+    static Ret callback_fn(intptr_t callable, Params... params) {
+        // NOLINTNEXTLINE(performance-no-int-to-ptr)
+        return (*reinterpret_cast<Callable *>(callable))(std::forward<Params>(params)...);
+    }
+
+public:
+    function_ref() = default;
+    // NOLINTNEXTLINE(google-explicit-constructor)
+    function_ref(std::nullptr_t) {}
+
+    template <typename Callable>
+    // NOLINTNEXTLINE(google-explicit-constructor)
+    function_ref(
+        Callable &&callable,
+        // This is not the copy-constructor.
+        enable_if_t<!std::is_same<remove_cvref_t<Callable>, function_ref>::value> * = nullptr,
+        // Functor must be callable and return a suitable type.
+        enable_if_t<
+            std::is_void<Ret>::value
+            || std::is_convertible<decltype(std::declval<Callable>()(std::declval<Params>()...)),
+                                   Ret>::value> * = nullptr)
+        : callback(callback_fn<remove_reference_t<Callable>>),
+          callable(reinterpret_cast<intptr_t>(&callable)) {}
+
+    // NOLINTNEXTLINE(performance-unnecessary-value-param)
+    Ret operator()(Params... params) const {
+        return callback(callable, std::forward<Params>(params)...);
+    }
+
+    explicit operator bool() const { return callback; }
+
+    bool operator==(const function_ref<Ret(Params...)> &Other) const {
+        return callable == Other.callable;
+    }
+};
+PYBIND11_NAMESPACE_END(detail)
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/detail/holder_caster_foreign_helpers.h

@@ -0,0 +1,104 @@
+/*
+    pybind11/detail/holder_caster_foreign_helpers.h: Logic to implement
+    set_foreign_holder() in copyable_ and movable_holder_caster.
+
+    Copyright (c) 2025 Hudson River Trading LLC <opensource@hudson-trading.com>
+
+    All rights reserved. Use of this source code is governed by a
+    BSD-style license that can be found in the LICENSE file.
+*/
+
+#pragma once
+
+#include <pybind11/gil.h>
+
+#include "common.h"
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+struct holder_caster_foreign_helpers {
+    struct py_deleter {
+        void operator()(const void *) const noexcept {
+            // Don't run the deleter if the interpreter has been shut down
+            if (Py_IsInitialized() == 0) {
+                return;
+            }
+            gil_scoped_acquire guard;
+            Py_DECREF(o);
+        }
+
+        PyObject *o;
+    };
+
+    // Downcast shared_ptr from the enable_shared_from_this base to the target type.
+    // SFINAE probe: use static_pointer_cast when the static downcast is valid (common case),
+    // fall back to dynamic_pointer_cast when it isn't (virtual inheritance — issue #5989).
+    // We can't use dynamic_pointer_cast unconditionally because it requires polymorphic types;
+    // we can't use is_polymorphic to choose because that's orthogonal to virtual inheritance.
+    // (The implementation uses the "tag dispatch via overload priority" trick.)
+    template <typename type, typename esft_base>
+    static auto esft_downcast(const std::shared_ptr<esft_base> &existing, int /*preferred*/)
+        -> decltype(static_cast<type *>(std::declval<esft_base *>()), std::shared_ptr<type>()) {
+        return std::static_pointer_cast<type>(existing);
+    }
+
+    template <typename type, typename esft_base>
+    static std::shared_ptr<type> esft_downcast(const std::shared_ptr<esft_base> &existing,
+                                               ... /*fallback*/) {
+        return std::dynamic_pointer_cast<type>(existing);
+    }
+
+    template <typename type>
+    static auto set_via_shared_from_this(type *value, std::shared_ptr<type> *holder_out)
+        -> decltype(value->shared_from_this(), bool()) {
+        // object derives from enable_shared_from_this;
+        // try to reuse an existing shared_ptr if one is known
+        if (auto existing = try_get_shared_from_this(value)) {
+            *holder_out = esft_downcast<type>(existing, 0);
+            return true;
+        }
+        return false;
+    }
+
+    template <typename type>
+    static bool set_via_shared_from_this(void *, std::shared_ptr<type> *) {
+        return false;
+    }
+
+    template <typename type>
+    static bool set_foreign_holder(handle src, type *value, std::shared_ptr<type> *holder_out) {
+        // We only support using std::shared_ptr<T> for foreign T, and
+        // it's done by creating a new shared_ptr control block that
+        // owns a reference to the original Python object.
+        if (value == nullptr) {
+            *holder_out = {};
+            return true;
+        }
+        if (set_via_shared_from_this(value, holder_out)) {
+            return true;
+        }
+        *holder_out = std::shared_ptr<type>(value, py_deleter{src.inc_ref().ptr()});
+        return true;
+    }
+
+    template <typename type>
+    static bool
+    set_foreign_holder(handle src, const type *value, std::shared_ptr<const type> *holder_out) {
+        std::shared_ptr<type> holder_mut;
+        if (set_foreign_holder(src, const_cast<type *>(value), &holder_mut)) {
+            *holder_out = holder_mut;
+            return true;
+        }
+        return false;
+    }
+
+    template <typename type>
+    static bool set_foreign_holder(handle, type *, ...) {
+        throw cast_error("Unable to cast foreign type to held instance -- "
+                         "only std::shared_ptr<T> is supported in this case");
+    }
+};
+
+PYBIND11_NAMESPACE_END(detail)
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/detail/init.h

@@ -0,0 +1,548 @@
+/*
+    pybind11/detail/init.h: init factory function implementation and support code.
+
+    Copyright (c) 2017 Jason Rhinelander <jason@imaginary.ca>
+
+    All rights reserved. Use of this source code is governed by a
+    BSD-style license that can be found in the LICENSE file.
+*/
+
+#pragma once
+
+#include "class.h"
+#include "using_smart_holder.h"
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+
+PYBIND11_WARNING_DISABLE_MSVC(4127)
+
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+template <>
+class type_caster<value_and_holder> {
+public:
+    bool load(handle h, bool) {
+        value = reinterpret_cast<value_and_holder *>(h.ptr());
+        return true;
+    }
+
+    template <typename>
+    using cast_op_type = value_and_holder &;
+    explicit operator value_and_holder &() { return *value; }
+    static constexpr auto name = const_name<value_and_holder>();
+
+private:
+    value_and_holder *value = nullptr;
+};
+
+PYBIND11_NAMESPACE_BEGIN(initimpl)
+
+inline void no_nullptr(const void *ptr) {
+    if (!ptr) {
+        throw type_error("pybind11::init(): factory function returned nullptr");
+    }
+}
+
+// Implementing functions for all forms of py::init<...> and py::init(...)
+template <typename Class>
+using Cpp = typename Class::type;
+template <typename Class>
+using Alias = typename Class::type_alias;
+template <typename Class>
+using Holder = typename Class::holder_type;
+
+template <typename Class>
+using is_alias_constructible = std::is_constructible<Alias<Class>, Cpp<Class> &&>;
+
+// Takes a Cpp pointer and returns true if it actually is a polymorphic Alias instance.
+template <typename Class, enable_if_t<Class::has_alias, int> = 0>
+bool is_alias(Cpp<Class> *ptr) {
+    return dynamic_cast<Alias<Class> *>(ptr) != nullptr;
+}
+// Failing fallback version of the above for a no-alias class (always returns false)
+template <typename /*Class*/>
+constexpr bool is_alias(const void *) {
+    return false;
+}
+
+// Constructs and returns a new object; if the given arguments don't map to a constructor, we fall
+// back to brace aggregate initialization so that for aggregate initialization can be used with
+// py::init, e.g.  `py::init<int, int>` to initialize a `struct T { int a; int b; }`.  For
+// non-aggregate types, we need to use an ordinary T(...) constructor (invoking as `T{...}` usually
+// works, but will not do the expected thing when `T` has an `initializer_list<T>` constructor).
+template <typename Class,
+          typename... Args,
+          detail::enable_if_t<std::is_constructible<Class, Args...>::value, int> = 0>
+inline Class *construct_or_initialize(Args &&...args) {
+    return new Class(std::forward<Args>(args)...);
+}
+template <typename Class,
+          typename... Args,
+          detail::enable_if_t<!std::is_constructible<Class, Args...>::value, int> = 0>
+inline Class *construct_or_initialize(Args &&...args) {
+    return new Class{std::forward<Args>(args)...};
+}
+
+// Attempts to constructs an alias using a `Alias(Cpp &&)` constructor.  This allows types with
+// an alias to provide only a single Cpp factory function as long as the Alias can be
+// constructed from an rvalue reference of the base Cpp type.  This means that Alias classes
+// can, when appropriate, simply define a `Alias(Cpp &&)` constructor rather than needing to
+// inherit all the base class constructors.
+template <typename Class>
+void construct_alias_from_cpp(std::true_type /*is_alias_constructible*/,
+                              value_and_holder &v_h,
+                              Cpp<Class> &&base) {
+    v_h.value_ptr() = new Alias<Class>(std::move(base));
+}
+template <typename Class>
+[[noreturn]] void construct_alias_from_cpp(std::false_type /*!is_alias_constructible*/,
+                                           value_and_holder &,
+                                           Cpp<Class> &&) {
+    throw type_error("pybind11::init(): unable to convert returned instance to required "
+                     "alias class: no `Alias<Class>(Class &&)` constructor available");
+}
+
+// Error-generating fallback for factories that don't match one of the below construction
+// mechanisms.
+template <typename Class>
+void construct(...) {
+    static_assert(!std::is_same<Class, Class>::value /* always false */,
+                  "pybind11::init(): init function must return a compatible pointer, "
+                  "holder, or value");
+}
+
+// Pointer return v1: the factory function returns a class pointer for a registered class.
+// If we don't need an alias (because this class doesn't have one, or because the final type is
+// inherited on the Python side) we can simply take over ownership.  Otherwise we need to try to
+// construct an Alias from the returned base instance.
+template <typename Class>
+void construct(value_and_holder &v_h, Cpp<Class> *ptr, bool need_alias) {
+    PYBIND11_WORKAROUND_INCORRECT_MSVC_C4100(need_alias);
+    no_nullptr(ptr);
+    if (Class::has_alias && need_alias && !is_alias<Class>(ptr)) {
+        // We're going to try to construct an alias by moving the cpp type.  Whether or not
+        // that succeeds, we still need to destroy the original cpp pointer (either the
+        // moved away leftover, if the alias construction works, or the value itself if we
+        // throw an error), but we can't just call `delete ptr`: it might have a special
+        // deleter, or might be shared_from_this.  So we construct a holder around it as if
+        // it was a normal instance, then steal the holder away into a local variable; thus
+        // the holder and destruction happens when we leave the C++ scope, and the holder
+        // class gets to handle the destruction however it likes.
+        v_h.value_ptr() = ptr;
+        v_h.set_instance_registered(true); // Trick to prevent init_instance from registering it
+        // DANGER ZONE BEGIN: exceptions will leave v_h in an invalid state.
+        v_h.type->init_instance(v_h.inst, nullptr);                        // Set up the holder
+        Holder<Class> temp_holder(std::move(v_h.holder<Holder<Class>>())); // Steal the holder
+        v_h.type->dealloc(v_h); // Destroys the moved-out holder remains, resets value ptr to null
+        v_h.set_instance_registered(false);
+        // DANGER ZONE END.
+
+        construct_alias_from_cpp<Class>(is_alias_constructible<Class>{}, v_h, std::move(*ptr));
+    } else {
+        // Otherwise the type isn't inherited, so we don't need an Alias
+        v_h.value_ptr() = ptr;
+    }
+}
+
+// Pointer return v2: a factory that always returns an alias instance ptr.  We simply take over
+// ownership of the pointer.
+template <typename Class, enable_if_t<Class::has_alias, int> = 0>
+void construct(value_and_holder &v_h, Alias<Class> *alias_ptr, bool) {
+    no_nullptr(alias_ptr);
+    v_h.value_ptr() = static_cast<Cpp<Class> *>(alias_ptr);
+}
+
+// Holder return: copy its pointer, and move or copy the returned holder into the new instance's
+// holder.  This also handles types like std::shared_ptr<T> and std::unique_ptr<T> where T is a
+// derived type (through those holder's implicit conversion from derived class holder
+// constructors).
+template <typename Class, detail::enable_if_t<!is_smart_holder<Holder<Class>>::value, int> = 0>
+void construct(value_and_holder &v_h, Holder<Class> holder, bool need_alias) {
+    PYBIND11_WORKAROUND_INCORRECT_MSVC_C4100(need_alias);
+    auto *ptr = holder_helper<Holder<Class>>::get(holder);
+    no_nullptr(ptr);
+    // If we need an alias, check that the held pointer is actually an alias instance
+    if (Class::has_alias && need_alias && !is_alias<Class>(ptr)) {
+        throw type_error("pybind11::init(): construction failed: returned holder-wrapped instance "
+                         "is not an alias instance");
+    }
+
+    // Cast away constness to store in void* storage.
+    // The value_and_holder storage is fundamentally untyped (void**), so we lose
+    // const-correctness here by design. The const qualifier will be restored
+    // when the pointer is later retrieved and cast back to the original type.
+    // This explicit const_cast makes the const-removal clearly visible.
+    v_h.value_ptr() = const_cast<void *>(static_cast<const void *>(ptr));
+    v_h.type->init_instance(v_h.inst, &holder);
+}
+
+// return-by-value version 1: returning a cpp class by value.  If the class has an alias and an
+// alias is required the alias must have an `Alias(Cpp &&)` constructor so that we can construct
+// the alias from the base when needed (i.e. because of Python-side inheritance).  When we don't
+// need it, we simply move-construct the cpp value into a new instance.
+template <typename Class>
+void construct(value_and_holder &v_h, Cpp<Class> &&result, bool need_alias) {
+    PYBIND11_WORKAROUND_INCORRECT_MSVC_C4100(need_alias);
+    static_assert(is_move_constructible<Cpp<Class>>::value,
+                  "pybind11::init() return-by-value factory function requires a movable class");
+    if (Class::has_alias && need_alias) {
+        construct_alias_from_cpp<Class>(is_alias_constructible<Class>{}, v_h, std::move(result));
+    } else {
+        v_h.value_ptr() = new Cpp<Class>(std::move(result));
+    }
+}
+
+// return-by-value version 2: returning a value of the alias type itself.  We move-construct an
+// Alias instance (even if no the python-side inheritance is involved).  The is intended for
+// cases where Alias initialization is always desired.
+template <typename Class>
+void construct(value_and_holder &v_h, Alias<Class> &&result, bool) {
+    static_assert(
+        is_move_constructible<Alias<Class>>::value,
+        "pybind11::init() return-by-alias-value factory function requires a movable alias class");
+    v_h.value_ptr() = new Alias<Class>(std::move(result));
+}
+
+template <typename T, typename D>
+smart_holder init_smart_holder_from_unique_ptr(std::unique_ptr<T, D> &&unq_ptr,
+                                               bool void_cast_raw_ptr) {
+    void *void_ptr = void_cast_raw_ptr ? static_cast<void *>(unq_ptr.get()) : nullptr;
+    return smart_holder::from_unique_ptr(std::move(unq_ptr), void_ptr);
+}
+
+template <typename Class,
+          typename D = std::default_delete<Cpp<Class>>,
+          detail::enable_if_t<is_smart_holder<Holder<Class>>::value, int> = 0>
+void construct(value_and_holder &v_h, std::unique_ptr<Cpp<Class>, D> &&unq_ptr, bool need_alias) {
+    PYBIND11_WORKAROUND_INCORRECT_MSVC_C4100(need_alias);
+    auto *ptr = unq_ptr.get();
+    no_nullptr(ptr);
+    if (Class::has_alias && need_alias && !is_alias<Class>(ptr)) {
+        throw type_error("pybind11::init(): construction failed: returned std::unique_ptr pointee "
+                         "is not an alias instance");
+    }
+    // Here and below: if the new object is a trampoline, the shared_from_this mechanism needs
+    // to be prevented from accessing the smart_holder vptr, because it does not keep the
+    // trampoline Python object alive. For types that don't inherit from enable_shared_from_this
+    // it does not matter if void_cast_raw_ptr is true or false, therefore it's not necessary
+    // to also inspect the type.
+    auto smhldr = init_smart_holder_from_unique_ptr(
+        std::move(unq_ptr), /*void_cast_raw_ptr*/ Class::has_alias && is_alias<Class>(ptr));
+    v_h.value_ptr() = ptr;
+    v_h.type->init_instance(v_h.inst, &smhldr);
+}
+
+template <typename Class,
+          typename D = std::default_delete<Alias<Class>>,
+          detail::enable_if_t<is_smart_holder<Holder<Class>>::value, int> = 0>
+void construct(value_and_holder &v_h,
+               std::unique_ptr<Alias<Class>, D> &&unq_ptr,
+               bool /*need_alias*/) {
+    auto *ptr = unq_ptr.get();
+    no_nullptr(ptr);
+    auto smhldr
+        = init_smart_holder_from_unique_ptr(std::move(unq_ptr), /*void_cast_raw_ptr*/ true);
+    v_h.value_ptr() = ptr;
+    v_h.type->init_instance(v_h.inst, &smhldr);
+}
+
+template <typename PtrType, typename Class>
+void construct_from_shared_ptr(value_and_holder &v_h,
+                               std::shared_ptr<PtrType> &&shd_ptr,
+                               bool need_alias) {
+    static_assert(std::is_same<PtrType, Cpp<Class>>::value
+                      || std::is_same<PtrType, const Cpp<Class>>::value,
+                  "Expected (const) Cpp<Class> as shared_ptr pointee");
+    auto *ptr = shd_ptr.get();
+    no_nullptr(ptr);
+    if (Class::has_alias && need_alias && !is_alias<Class>(ptr)) {
+        throw type_error("pybind11::init(): construction failed: returned std::shared_ptr pointee "
+                         "is not an alias instance");
+    }
+    // Cast to non-const if needed, consistent with internal design
+    auto smhldr
+        = smart_holder::from_shared_ptr(std::const_pointer_cast<Cpp<Class>>(std::move(shd_ptr)));
+    v_h.value_ptr() = const_cast<Cpp<Class> *>(ptr);
+    v_h.type->init_instance(v_h.inst, &smhldr);
+}
+
+template <typename Class, detail::enable_if_t<is_smart_holder<Holder<Class>>::value, int> = 0>
+void construct(value_and_holder &v_h, std::shared_ptr<Cpp<Class>> &&shd_ptr, bool need_alias) {
+    construct_from_shared_ptr<Cpp<Class>, Class>(v_h, std::move(shd_ptr), need_alias);
+}
+
+template <typename Class, detail::enable_if_t<is_smart_holder<Holder<Class>>::value, int> = 0>
+void construct(value_and_holder &v_h,
+               std::shared_ptr<const Cpp<Class>> &&shd_ptr,
+               bool need_alias) {
+    construct_from_shared_ptr<const Cpp<Class>, Class>(v_h, std::move(shd_ptr), need_alias);
+}
+
+template <typename Class, detail::enable_if_t<is_smart_holder<Holder<Class>>::value, int> = 0>
+void construct(value_and_holder &v_h,
+               std::shared_ptr<Alias<Class>> &&shd_ptr,
+               bool /*need_alias*/) {
+    auto *ptr = shd_ptr.get();
+    no_nullptr(ptr);
+    auto smhldr = smart_holder::from_shared_ptr(shd_ptr);
+    v_h.value_ptr() = ptr;
+    v_h.type->init_instance(v_h.inst, &smhldr);
+}
+
+// Implementing class for py::init<...>()
+template <typename... Args>
+struct constructor {
+    template <typename Class, typename... Extra, enable_if_t<!Class::has_alias, int> = 0>
+    static void execute(Class &cl, const Extra &...extra) {
+        cl.def(
+            "__init__",
+            [](value_and_holder &v_h,
+               Args... args) { // NOLINT(performance-unnecessary-value-param)
+                v_h.value_ptr() = construct_or_initialize<Cpp<Class>>(std::forward<Args>(args)...);
+            },
+            is_new_style_constructor(),
+            extra...);
+    }
+
+    template <typename Class,
+              typename... Extra,
+              enable_if_t<Class::has_alias && std::is_constructible<Cpp<Class>, Args...>::value,
+                          int> = 0>
+    static void execute(Class &cl, const Extra &...extra) {
+        cl.def(
+            "__init__",
+            [](value_and_holder &v_h, Args... args) {
+                if (Py_TYPE(v_h.inst) == v_h.type->type) {
+                    v_h.value_ptr()
+                        = construct_or_initialize<Cpp<Class>>(std::forward<Args>(args)...);
+                } else {
+                    v_h.value_ptr()
+                        = construct_or_initialize<Alias<Class>>(std::forward<Args>(args)...);
+                }
+            },
+            is_new_style_constructor(),
+            extra...);
+    }
+
+    template <typename Class,
+              typename... Extra,
+              enable_if_t<Class::has_alias && !std::is_constructible<Cpp<Class>, Args...>::value,
+                          int> = 0>
+    static void execute(Class &cl, const Extra &...extra) {
+        cl.def(
+            "__init__",
+            [](value_and_holder &v_h, Args... args) {
+                v_h.value_ptr()
+                    = construct_or_initialize<Alias<Class>>(std::forward<Args>(args)...);
+            },
+            is_new_style_constructor(),
+            extra...);
+    }
+};
+
+// Implementing class for py::init_alias<...>()
+template <typename... Args>
+struct alias_constructor {
+    template <typename Class,
+              typename... Extra,
+              enable_if_t<Class::has_alias && std::is_constructible<Alias<Class>, Args...>::value,
+                          int> = 0>
+    static void execute(Class &cl, const Extra &...extra) {
+        cl.def(
+            "__init__",
+            [](value_and_holder &v_h, Args... args) {
+                v_h.value_ptr()
+                    = construct_or_initialize<Alias<Class>>(std::forward<Args>(args)...);
+            },
+            is_new_style_constructor(),
+            extra...);
+    }
+};
+
+// Implementation class for py::init(Func) and py::init(Func, AliasFunc)
+template <typename CFunc,
+          typename AFunc = void_type (*)(),
+          typename = function_signature_t<CFunc>,
+          typename = function_signature_t<AFunc>>
+struct factory;
+
+// Specialization for py::init(Func)
+// Note: The 4th template parameter `void_type()` is explicitly specified to resolve a
+// template ambiguity with the dual-factory specialization below when compiled with
+// nvcc + GCC (see #5565). Without it, both specializations match equally well for the
+// single-factory case, since the 4th parameter defaults to
+// `function_signature_t<void_type(*)()>` = `void_type()`, which the dual-factory
+// specialization can also decompose as `AReturn(AArgs...)` with `AReturn=void_type`
+// and `AArgs={}`.
+template <typename Func, typename Return, typename... Args>
+struct factory<Func, void_type (*)(), Return(Args...), void_type()> {
+    remove_reference_t<Func> class_factory;
+
+    // NOLINTNEXTLINE(google-explicit-constructor)
+    factory(Func &&f) : class_factory(std::forward<Func>(f)) {}
+
+    // The given class either has no alias or has no separate alias factory;
+    // this always constructs the class itself.  If the class is registered with an alias
+    // type and an alias instance is needed (i.e. because the final type is a Python class
+    // inheriting from the C++ type) the returned value needs to either already be an alias
+    // instance, or the alias needs to be constructible from a `Class &&` argument.
+    template <typename Class, typename... Extra>
+    void execute(Class &cl, const Extra &...extra) && {
+#if defined(PYBIND11_CPP14)
+        cl.def(
+            "__init__",
+            [func = std::move(class_factory)]
+#else
+        auto &func = class_factory;
+        cl.def(
+            "__init__",
+            [func]
+#endif
+            (value_and_holder &v_h, Args... args) {
+                construct<Class>(
+                    v_h, func(std::forward<Args>(args)...), Py_TYPE(v_h.inst) != v_h.type->type);
+            },
+            is_new_style_constructor(),
+            extra...);
+    }
+};
+
+// Specialization for py::init(Func, AliasFunc)
+template <typename CFunc,
+          typename AFunc,
+          typename CReturn,
+          typename... CArgs,
+          typename AReturn,
+          typename... AArgs>
+struct factory<CFunc, AFunc, CReturn(CArgs...), AReturn(AArgs...)> {
+    static_assert(sizeof...(CArgs) == sizeof...(AArgs),
+                  "pybind11::init(class_factory, alias_factory): class and alias factories "
+                  "must have identical argument signatures");
+    static_assert(all_of<std::is_same<CArgs, AArgs>...>::value,
+                  "pybind11::init(class_factory, alias_factory): class and alias factories "
+                  "must have identical argument signatures");
+
+    remove_reference_t<CFunc> class_factory;
+    remove_reference_t<AFunc> alias_factory;
+
+    factory(CFunc &&c, AFunc &&a)
+        : class_factory(std::forward<CFunc>(c)), alias_factory(std::forward<AFunc>(a)) {}
+
+    // The class factory is called when the `self` type passed to `__init__` is the direct
+    // class (i.e. not inherited), the alias factory when `self` is a Python-side subtype.
+    template <typename Class, typename... Extra>
+    void execute(Class &cl, const Extra &...extra) && {
+        static_assert(Class::has_alias,
+                      "The two-argument version of `py::init()` can "
+                      "only be used if the class has an alias");
+#if defined(PYBIND11_CPP14)
+        cl.def(
+            "__init__",
+            [class_func = std::move(class_factory), alias_func = std::move(alias_factory)]
+#else
+        auto &class_func = class_factory;
+        auto &alias_func = alias_factory;
+        cl.def(
+            "__init__",
+            [class_func, alias_func]
+#endif
+            (value_and_holder &v_h, CArgs... args) {
+                if (Py_TYPE(v_h.inst) == v_h.type->type) {
+                    // If the instance type equals the registered type we don't have inheritance,
+                    // so don't need the alias and can construct using the class function:
+                    construct<Class>(v_h, class_func(std::forward<CArgs>(args)...), false);
+                } else {
+                    construct<Class>(v_h, alias_func(std::forward<CArgs>(args)...), true);
+                }
+            },
+            is_new_style_constructor(),
+            extra...);
+    }
+};
+
+/// Set just the C++ state. Same as `__init__`.
+template <typename Class, typename T>
+void setstate(value_and_holder &v_h, T &&result, bool need_alias) {
+    construct<Class>(v_h, std::forward<T>(result), need_alias);
+}
+
+/// Set both the C++ and Python states
+template <typename Class,
+          typename T,
+          typename O,
+          enable_if_t<std::is_convertible<O, handle>::value, int> = 0>
+void setstate(value_and_holder &v_h, std::pair<T, O> &&result, bool need_alias) {
+    construct<Class>(v_h, std::move(result.first), need_alias);
+    auto d = handle(result.second);
+    if (PyDict_Check(d.ptr()) && PyDict_Size(d.ptr()) == 0) {
+        // Skipping setattr below, to not force use of py::dynamic_attr() for Class unnecessarily.
+        // See PR #2972 for details.
+        return;
+    }
+    // Our tests never run into an unset dict, but being careful here for now (see #5658)
+    auto dict = getattr(reinterpret_cast<PyObject *>(v_h.inst), "__dict__", none());
+    if (dict.is_none()) {
+        setattr(reinterpret_cast<PyObject *>(v_h.inst), "__dict__", d);
+    } else {
+        // Keep the original object dict and just update it
+        if (PyDict_Update(dict.ptr(), d.ptr()) < 0) {
+            throw error_already_set();
+        }
+    }
+}
+
+/// Implementation for py::pickle(GetState, SetState)
+template <typename Get,
+          typename Set,
+          typename = function_signature_t<Get>,
+          typename = function_signature_t<Set>>
+struct pickle_factory;
+
+template <typename Get,
+          typename Set,
+          typename RetState,
+          typename Self,
+          typename NewInstance,
+          typename ArgState>
+struct pickle_factory<Get, Set, RetState(Self), NewInstance(ArgState)> {
+    using Ret = intrinsic_t<RetState>;
+    using Arg = intrinsic_t<ArgState>;
+
+    // Subclasses are now allowed for support between type hint and generic versions of types
+    // (e.g.) typing::List <--> list
+    static_assert(std::is_same<Ret, Arg>::value || std::is_base_of<Ret, Arg>::value
+                      || std::is_base_of<Arg, Ret>::value,
+                  "The type returned by `__getstate__` must be the same or subclass of the "
+                  "argument accepted by `__setstate__`");
+
+    remove_reference_t<Get> get;
+    remove_reference_t<Set> set;
+
+    pickle_factory(Get get, Set set) : get(std::forward<Get>(get)), set(std::forward<Set>(set)) {}
+
+    template <typename Class, typename... Extra>
+    void execute(Class &cl, const Extra &...extra) && {
+        cl.def("__getstate__", std::move(get), pos_only());
+
+#if defined(PYBIND11_CPP14)
+        cl.def(
+            "__setstate__",
+            [func = std::move(set)]
+#else
+        auto &func = set;
+        cl.def(
+            "__setstate__",
+            [func]
+#endif
+            (value_and_holder &v_h, ArgState state) {
+                setstate<Class>(
+                    v_h, func(std::forward<ArgState>(state)), Py_TYPE(v_h.inst) != v_h.type->type);
+            },
+            is_new_style_constructor(),
+            extra...);
+    }
+};
+
+PYBIND11_NAMESPACE_END(initimpl)
+PYBIND11_NAMESPACE_END(detail)
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/detail/native_enum_data.h

@@ -0,0 +1,227 @@
+// Copyright (c) 2022-2025 The pybind Community.
+// All rights reserved. Use of this source code is governed by a
+// BSD-style license that can be found in the LICENSE file.
+
+#pragma once
+
+#include "../pytypes.h"
+#include "common.h"
+#include "internals.h"
+
+#include <cassert>
+#include <sstream>
+#include <string>
+#include <typeindex>
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+// This is a separate function only to enable easy unit testing.
+inline std::string
+native_enum_missing_finalize_error_message(const std::string &enum_name_encoded) {
+    return "pybind11::native_enum<...>(\"" + enum_name_encoded + "\", ...): MISSING .finalize()";
+}
+
+// Internals for pybind11::native_enum; one native_enum_data object exists
+// inside each pybind11::native_enum and lives only for the duration of the
+// native_enum binding statement.
+class native_enum_data {
+public:
+    native_enum_data(handle parent_scope_,
+                     const char *enum_name,
+                     const char *native_type_name,
+                     const char *class_doc,
+                     const native_enum_record &enum_record_)
+        : enum_name_encoded{enum_name}, native_type_name_encoded{native_type_name},
+          enum_type_index{*enum_record_.cpptype},
+          parent_scope(reinterpret_borrow<object>(parent_scope_)), enum_name{enum_name},
+          native_type_name{native_type_name}, class_doc(class_doc), export_values_flag{false},
+          finalize_needed{false} {
+        // Create the enum record capsule. It will be installed on the enum
+        // type object during finalize(). Its destructor removes the enum
+        // mapping from our internals, so that we won't try to convert to an
+        // enum type that's been destroyed.
+        enum_record = capsule(
+            new native_enum_record{enum_record_},
+            native_enum_record::attribute_name(),
+            +[](void *record_) {
+                auto *record = static_cast<native_enum_record *>(record_);
+                with_internals([&](internals &internals) {
+                    internals.native_enum_type_map.erase(*record->cpptype);
+                });
+                delete record;
+            });
+    }
+
+    void finalize();
+
+    native_enum_data(const native_enum_data &) = delete;
+    native_enum_data &operator=(const native_enum_data &) = delete;
+
+#if !defined(NDEBUG)
+    // This dtor cannot easily be unit tested because it terminates the process.
+    ~native_enum_data() {
+        if (finalize_needed) {
+            pybind11_fail(native_enum_missing_finalize_error_message(enum_name_encoded));
+        }
+    }
+#endif
+
+protected:
+    void disarm_finalize_check(const char *error_context) {
+        if (!finalize_needed) {
+            pybind11_fail("pybind11::native_enum<...>(\"" + enum_name_encoded
+                          + "\"): " + error_context);
+        }
+        finalize_needed = false;
+    }
+
+    void arm_finalize_check() {
+        assert(!finalize_needed); // Catch redundant calls.
+        finalize_needed = true;
+    }
+
+    std::string enum_name_encoded;
+    std::string native_type_name_encoded;
+    std::type_index enum_type_index;
+
+private:
+    object parent_scope;
+    str enum_name;
+    str native_type_name;
+    std::string class_doc;
+    capsule enum_record;
+
+protected:
+    list members;
+    list member_docs;
+    bool export_values_flag : 1; // Attention: It is best to keep the bools together.
+
+private:
+    bool finalize_needed : 1;
+};
+
+inline handle
+global_internals_native_enum_type_map_get_item(const std::type_index &enum_type_index) {
+    return with_internals([&](internals &internals) {
+        auto found = internals.native_enum_type_map.find(enum_type_index);
+        if (found != internals.native_enum_type_map.end()) {
+            return handle(found->second);
+        }
+        return handle();
+    });
+}
+
+inline bool
+global_internals_native_enum_type_map_contains(const std::type_index &enum_type_index) {
+    return with_internals([&](internals &internals) {
+        return internals.native_enum_type_map.count(enum_type_index) != 0;
+    });
+}
+
+inline object import_or_getattr(const std::string &fully_qualified_name,
+                                const std::string &append_to_exception_message) {
+    std::istringstream stream(fully_qualified_name);
+    std::string part;
+
+    if (!std::getline(stream, part, '.') || part.empty()) {
+        std::string msg = "Invalid fully-qualified name `";
+        msg += fully_qualified_name;
+        msg += "`";
+        msg += append_to_exception_message;
+        throw value_error(msg);
+    }
+
+    auto curr_scope = reinterpret_steal<object>(PyImport_ImportModule(part.c_str()));
+    if (!curr_scope) {
+        std::string msg = "Failed to import top-level module `";
+        msg += part;
+        msg += "`";
+        msg += append_to_exception_message;
+        raise_from(PyExc_ImportError, msg.c_str());
+        throw error_already_set();
+    }
+
+    // Now recursively getattr or import remaining parts
+    std::string curr_path = part;
+    while (std::getline(stream, part, '.')) {
+        if (part.empty()) {
+            std::string msg = "Invalid fully-qualified name `";
+            msg += fully_qualified_name;
+            msg += "`";
+            msg += append_to_exception_message;
+            throw value_error(msg);
+        }
+        std::string next_path = curr_path;
+        next_path += ".";
+        next_path += part;
+        auto next_scope
+            = reinterpret_steal<object>(PyObject_GetAttrString(curr_scope.ptr(), part.c_str()));
+        if (!next_scope) {
+            error_fetch_and_normalize stored_getattr_error("getattr");
+            // Try importing the next level
+            next_scope = reinterpret_steal<object>(PyImport_ImportModule(next_path.c_str()));
+            if (!next_scope) {
+                error_fetch_and_normalize stored_import_error("import");
+                std::string msg = "Failed to import or getattr `";
+                msg += part;
+                msg += "` from `";
+                msg += curr_path;
+                msg += "`";
+                msg += append_to_exception_message;
+                msg += "\n-------- getattr exception --------\n";
+                msg += stored_getattr_error.error_string();
+                msg += "\n-------- import exception --------\n";
+                msg += stored_import_error.error_string();
+                throw import_error(msg.c_str());
+            }
+        }
+        curr_scope = next_scope;
+        curr_path = next_path;
+    }
+    return curr_scope;
+}
+
+inline void native_enum_data::finalize() {
+    disarm_finalize_check("DOUBLE finalize");
+    if (hasattr(parent_scope, enum_name)) {
+        pybind11_fail("pybind11::native_enum<...>(\"" + enum_name_encoded
+                      + "\"): an object with that name is already defined");
+    }
+    auto py_enum_type = import_or_getattr(native_type_name, " (native_type_name)");
+    auto py_enum = py_enum_type(enum_name, members);
+    object module_name = get_module_name_if_available(parent_scope);
+    if (module_name) {
+        py_enum.attr("__module__") = module_name;
+    }
+    if (hasattr(parent_scope, "__qualname__")) {
+        const auto parent_qualname = parent_scope.attr("__qualname__").cast<std::string>();
+        py_enum.attr("__qualname__") = str(parent_qualname + "." + enum_name.cast<std::string>());
+    }
+    parent_scope.attr(enum_name) = py_enum;
+    if (export_values_flag) {
+        for (auto member : members) {
+            auto member_name = member[int_(0)];
+            if (hasattr(parent_scope, member_name)) {
+                pybind11_fail("pybind11::native_enum<...>(\"" + enum_name_encoded + "\").value(\""
+                              + member_name.cast<std::string>()
+                              + "\"): an object with that name is already defined");
+            }
+            parent_scope.attr(member_name) = py_enum[member_name];
+        }
+    }
+    if (!class_doc.empty()) {
+        py_enum.attr("__doc__") = class_doc.c_str();
+    }
+    for (auto doc : member_docs) {
+        py_enum[doc[int_(0)]].attr("__doc__") = doc[int_(1)];
+    }
+
+    py_enum.attr(native_enum_record::attribute_name()) = enum_record;
+    with_internals([&](internals &internals) {
+        internals.native_enum_type_map[enum_type_index] = py_enum.ptr();
+    });
+}
+
+PYBIND11_NAMESPACE_END(detail)
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/detail/pybind11_namespace_macros.h

@@ -0,0 +1,82 @@
+// Copyright (c) 2016-2025 The Pybind Development Team.
+// All rights reserved. Use of this source code is governed by a
+// BSD-style license that can be found in the LICENSE file.
+
+#pragma once
+
+// PLEASE DO NOT ADD ANY INCLUDES HERE
+
+// Define some generic pybind11 helper macros for warning management.
+//
+// Note that compiler-specific push/pop pairs are baked into the
+// PYBIND11_NAMESPACE_BEGIN/PYBIND11_NAMESPACE_END pair of macros. Therefore manual
+// PYBIND11_WARNING_PUSH/PYBIND11_WARNING_POP are usually only needed in `#include` sections.
+//
+// If you find you need to suppress a warning, please try to make the suppression as local as
+// possible using these macros. Please also be sure to push/pop with the pybind11 macros. Please
+// only use compiler specifics if you need to check specific versions, e.g. Apple Clang vs. vanilla
+// Clang.
+#if defined(__INTEL_COMPILER)
+#    define PYBIND11_COMPILER_INTEL
+#    define PYBIND11_PRAGMA(...) _Pragma(#__VA_ARGS__)
+#    define PYBIND11_WARNING_PUSH PYBIND11_PRAGMA(warning push)
+#    define PYBIND11_WARNING_POP PYBIND11_PRAGMA(warning pop)
+#elif defined(__clang__)
+#    define PYBIND11_COMPILER_CLANG
+#    define PYBIND11_PRAGMA(...) _Pragma(#__VA_ARGS__)
+#    define PYBIND11_WARNING_PUSH PYBIND11_PRAGMA(clang diagnostic push)
+#    define PYBIND11_WARNING_POP PYBIND11_PRAGMA(clang diagnostic pop)
+#elif defined(__GNUC__)
+#    define PYBIND11_COMPILER_GCC
+#    define PYBIND11_PRAGMA(...) _Pragma(#__VA_ARGS__)
+#    define PYBIND11_WARNING_PUSH PYBIND11_PRAGMA(GCC diagnostic push)
+#    define PYBIND11_WARNING_POP PYBIND11_PRAGMA(GCC diagnostic pop)
+#elif defined(_MSC_VER) // Must be after the clang branch because clang-cl also defines _MSC_VER
+#    define PYBIND11_COMPILER_MSVC
+#    define PYBIND11_PRAGMA(...) __pragma(__VA_ARGS__)
+#    define PYBIND11_WARNING_PUSH PYBIND11_PRAGMA(warning(push))
+#    define PYBIND11_WARNING_POP PYBIND11_PRAGMA(warning(pop))
+#endif
+
+#ifdef PYBIND11_COMPILER_MSVC
+#    define PYBIND11_WARNING_DISABLE_MSVC(name) PYBIND11_PRAGMA(warning(disable : name))
+#else
+#    define PYBIND11_WARNING_DISABLE_MSVC(name)
+#endif
+
+#ifdef PYBIND11_COMPILER_CLANG
+#    define PYBIND11_WARNING_DISABLE_CLANG(name) PYBIND11_PRAGMA(clang diagnostic ignored name)
+#else
+#    define PYBIND11_WARNING_DISABLE_CLANG(name)
+#endif
+
+#ifdef PYBIND11_COMPILER_GCC
+#    define PYBIND11_WARNING_DISABLE_GCC(name) PYBIND11_PRAGMA(GCC diagnostic ignored name)
+#else
+#    define PYBIND11_WARNING_DISABLE_GCC(name)
+#endif
+
+#ifdef PYBIND11_COMPILER_INTEL
+#    define PYBIND11_WARNING_DISABLE_INTEL(name) PYBIND11_PRAGMA(warning disable name)
+#else
+#    define PYBIND11_WARNING_DISABLE_INTEL(name)
+#endif
+
+#define PYBIND11_NAMESPACE_BEGIN(name)                                                            \
+    namespace name {                                                                              \
+    PYBIND11_WARNING_PUSH
+
+#define PYBIND11_NAMESPACE_END(name)                                                              \
+    PYBIND11_WARNING_POP                                                                          \
+    }
+
+// Robust support for some features and loading modules compiled against different pybind versions
+// requires forcing hidden visibility on pybind code, so we enforce this by setting the attribute
+// on the main `pybind11` namespace.
+#if !defined(PYBIND11_NAMESPACE)
+#    if defined(__GNUG__) && !defined(_WIN32)
+#        define PYBIND11_NAMESPACE pybind11 __attribute__((visibility("hidden")))
+#    else
+#        define PYBIND11_NAMESPACE pybind11
+#    endif
+#endif

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/detail/struct_smart_holder.h

@@ -0,0 +1,398 @@
+// Copyright (c) 2020-2024 The Pybind Development Team.
+// All rights reserved. Use of this source code is governed by a
+// BSD-style license that can be found in the LICENSE file.
+
+/* Proof-of-Concept for smart pointer interoperability.
+
+High-level aspects:
+
+* Support all `unique_ptr`, `shared_ptr` interops that are feasible.
+
+* Cleanly and clearly report all interops that are infeasible.
+
+* Meant to fit into a `PyObject`, as a holder for C++ objects.
+
+* Support a system design that makes it impossible to trigger
+  C++ Undefined Behavior, especially from Python.
+
+* Support a system design with clean runtime inheritance casting. From this
+  it follows that the `smart_holder` needs to be type-erased (`void*`).
+
+* Handling of RTTI for the type-erased held pointer is NOT implemented here.
+  It is the responsibility of the caller to ensure that `static_cast<T *>`
+  is well-formed when calling `as_*` member functions. Inheritance casting
+  needs to be handled in a different layer (similar to the code organization
+  in boost/python/object/inheritance.hpp).
+
+Details:
+
+* The "root holder" chosen here is a `shared_ptr<void>` (named `vptr` in this
+  implementation). This choice is practically inevitable because `shared_ptr`
+  has only very limited support for inspecting and accessing its deleter.
+
+* If created from a raw pointer, or a `unique_ptr` without a custom deleter,
+  `vptr` always uses a custom deleter, to support `unique_ptr`-like disowning.
+  The custom deleters could be extended to included life-time management for
+  external objects (e.g. `PyObject`).
+
+* If created from an external `shared_ptr`, or a `unique_ptr` with a custom
+  deleter, including life-time management for external objects is infeasible.
+
+* By choice, the smart_holder is movable but not copyable, to keep the design
+  simple, and to guard against accidental copying overhead.
+
+* The `void_cast_raw_ptr` option is needed to make the `smart_holder` `vptr`
+  member invisible to the `shared_from_this` mechanism, in case the lifetime
+  of a `PyObject` is tied to the pointee.
+*/
+
+#pragma once
+
+#include "pybind11_namespace_macros.h"
+
+#include <cstring>
+#include <functional>
+#include <memory>
+#include <stdexcept>
+#include <string>
+#include <type_traits>
+#include <typeinfo>
+#include <utility>
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+PYBIND11_NAMESPACE_BEGIN(memory)
+
+// Default fallback.
+static constexpr bool type_has_shared_from_this(...) { return false; }
+
+// This overload uses SFINAE to skip enable_shared_from_this checks when the
+// base is inaccessible (e.g. private inheritance).
+template <typename T>
+static auto type_has_shared_from_this(const T *ptr)
+    -> decltype(static_cast<const std::enable_shared_from_this<T> *>(ptr), true) {
+    return true;
+}
+
+// Inaccessible base → substitution failure → fallback overload selected
+template <typename T>
+static constexpr bool type_has_shared_from_this(const void *) {
+    return false;
+}
+
+struct guarded_delete {
+    // NOTE: PYBIND11_INTERNALS_VERSION needs to be bumped if changes are made to this struct.
+    std::weak_ptr<void> released_ptr;    // Trick to keep the smart_holder memory footprint small.
+    std::function<void(void *)> del_fun; // Rare case.
+    void (*del_ptr)(void *);             // Common case.
+    bool use_del_fun;
+    bool armed_flag;
+    guarded_delete(std::function<void(void *)> &&del_fun, bool armed_flag)
+        : del_fun{std::move(del_fun)}, del_ptr{nullptr}, use_del_fun{true},
+          armed_flag{armed_flag} {}
+    guarded_delete(void (*del_ptr)(void *), bool armed_flag)
+        : del_ptr{del_ptr}, use_del_fun{false}, armed_flag{armed_flag} {}
+    void operator()(void *raw_ptr) const {
+        if (armed_flag) {
+            if (use_del_fun) {
+                del_fun(raw_ptr);
+            } else {
+                del_ptr(raw_ptr);
+            }
+        }
+    }
+};
+
+inline guarded_delete *get_guarded_delete(const std::shared_ptr<void> &ptr) {
+    return std::get_deleter<guarded_delete>(ptr);
+}
+
+using get_guarded_delete_fn = guarded_delete *(*) (const std::shared_ptr<void> &);
+
+template <typename T, typename std::enable_if<std::is_destructible<T>::value, int>::type = 0>
+inline void std_default_delete_if_destructible(void *raw_ptr) {
+    std::default_delete<T>{}(static_cast<T *>(raw_ptr));
+}
+
+template <typename T, typename std::enable_if<!std::is_destructible<T>::value, int>::type = 0>
+inline void std_default_delete_if_destructible(void *) {
+    // This noop operator is needed to avoid a compilation error (for `delete raw_ptr;`), but
+    // throwing an exception from a destructor will std::terminate the process. Therefore the
+    // runtime check for lifetime-management correctness is implemented elsewhere (in
+    // ensure_pointee_is_destructible()).
+}
+
+template <typename T>
+guarded_delete make_guarded_std_default_delete(bool armed_flag) {
+    return guarded_delete(std_default_delete_if_destructible<T>, armed_flag);
+}
+
+template <typename T, typename D>
+struct custom_deleter {
+    // NOTE: PYBIND11_INTERNALS_VERSION needs to be bumped if changes are made to this struct.
+    D deleter;
+    explicit custom_deleter(D &&deleter) : deleter{std::forward<D>(deleter)} {}
+    void operator()(void *raw_ptr) { deleter(static_cast<T *>(raw_ptr)); }
+};
+
+template <typename T, typename D>
+guarded_delete make_guarded_custom_deleter(D &&uqp_del, bool armed_flag) {
+    return guarded_delete(
+        std::function<void(void *)>(custom_deleter<T, D>(std::forward<D>(uqp_del))), armed_flag);
+}
+
+template <typename T, typename D>
+constexpr bool uqp_del_is_std_default_delete() {
+    return std::is_same<D, std::default_delete<T>>::value
+           || std::is_same<D, std::default_delete<T const>>::value;
+}
+
+inline bool type_info_equal_across_dso_boundaries(const std::type_info &a,
+                                                  const std::type_info &b) {
+    // RTTI pointer comparison may fail across DSOs (e.g., macOS libc++).
+    // Fallback to name comparison, which is generally safe and ABI-stable enough for our use.
+    return a == b || std::strcmp(a.name(), b.name()) == 0;
+}
+
+struct smart_holder {
+    // NOTE: PYBIND11_INTERNALS_VERSION needs to be bumped if changes are made to this struct.
+    const std::type_info *rtti_uqp_del = nullptr;
+    std::shared_ptr<void> vptr;
+    bool vptr_is_using_noop_deleter : 1;
+    bool vptr_is_using_std_default_delete : 1;
+    bool vptr_is_external_shared_ptr : 1;
+    bool is_populated : 1;
+    bool is_disowned : 1;
+
+    // Design choice: smart_holder is movable but not copyable.
+    smart_holder(smart_holder &&) = default;
+    smart_holder(const smart_holder &) = delete;
+    smart_holder &operator=(smart_holder &&) = delete;
+    smart_holder &operator=(const smart_holder &) = delete;
+
+    smart_holder()
+        : vptr_is_using_noop_deleter{false}, vptr_is_using_std_default_delete{false},
+          vptr_is_external_shared_ptr{false}, is_populated{false}, is_disowned{false} {}
+
+    bool has_pointee() const { return vptr != nullptr; }
+
+    template <typename T>
+    static void ensure_pointee_is_destructible(const char *context) {
+        if (!std::is_destructible<T>::value) {
+            throw std::invalid_argument(std::string("Pointee is not destructible (") + context
+                                        + ").");
+        }
+    }
+
+    void ensure_is_populated(const char *context) const {
+        if (!is_populated) {
+            throw std::runtime_error(std::string("Unpopulated holder (") + context + ").");
+        }
+    }
+    void ensure_is_not_disowned(const char *context) const {
+        if (is_disowned) {
+            throw std::runtime_error(std::string("Holder was disowned already (") + context
+                                     + ").");
+        }
+    }
+
+    void ensure_vptr_is_using_std_default_delete(const char *context) const {
+        if (vptr_is_external_shared_ptr) {
+            throw std::invalid_argument(std::string("Cannot disown external shared_ptr (")
+                                        + context + ").");
+        }
+        if (vptr_is_using_noop_deleter) {
+            throw std::invalid_argument(std::string("Cannot disown non-owning holder (") + context
+                                        + ").");
+        }
+        if (!vptr_is_using_std_default_delete) {
+            throw std::invalid_argument(std::string("Cannot disown custom deleter (") + context
+                                        + ").");
+        }
+    }
+
+    template <typename T, typename D>
+    void ensure_compatible_uqp_del(const char *context) const {
+        if (!rtti_uqp_del) {
+            if (!uqp_del_is_std_default_delete<T, D>()) {
+                throw std::invalid_argument(std::string("Missing unique_ptr deleter (") + context
+                                            + ").");
+            }
+            ensure_vptr_is_using_std_default_delete(context);
+            return;
+        }
+        if (uqp_del_is_std_default_delete<T, D>() && vptr_is_using_std_default_delete) {
+            return;
+        }
+        if (!type_info_equal_across_dso_boundaries(typeid(D), *rtti_uqp_del)) {
+            throw std::invalid_argument(std::string("Incompatible unique_ptr deleter (") + context
+                                        + ").");
+        }
+    }
+
+    void ensure_has_pointee(const char *context) const {
+        if (!has_pointee()) {
+            throw std::invalid_argument(std::string("Disowned holder (") + context + ").");
+        }
+    }
+
+    void ensure_use_count_1(const char *context) const {
+        if (vptr == nullptr) {
+            throw std::invalid_argument(std::string("Cannot disown nullptr (") + context + ").");
+        }
+        // In multithreaded environments accessing use_count can lead to
+        // race conditions, but in the context of Python it is a bug (elsewhere)
+        // if the Global Interpreter Lock (GIL) is not being held when this code
+        // is reached.
+        // PYBIND11:REMINDER: This may need to be protected by a mutex in free-threaded Python.
+        if (vptr.use_count() != 1) {
+            throw std::invalid_argument(std::string("Cannot disown use_count != 1 (") + context
+                                        + ").");
+        }
+    }
+
+    void reset_vptr_deleter_armed_flag(const get_guarded_delete_fn ggd_fn, bool armed_flag) const {
+        auto *gd = ggd_fn(vptr);
+        if (gd == nullptr) {
+            throw std::runtime_error(
+                "smart_holder::reset_vptr_deleter_armed_flag() called in an invalid context.");
+        }
+        gd->armed_flag = armed_flag;
+    }
+
+    // Caller is responsible for precondition: ensure_compatible_uqp_del<T, D>() must succeed.
+    template <typename T, typename D>
+    std::unique_ptr<D> extract_deleter(const char *context,
+                                       const get_guarded_delete_fn ggd_fn) const {
+        auto *gd = ggd_fn(vptr);
+        if (gd && gd->use_del_fun) {
+            const auto &custom_deleter_ptr = gd->del_fun.template target<custom_deleter<T, D>>();
+            if (custom_deleter_ptr == nullptr) {
+                throw std::runtime_error(
+                    std::string("smart_holder::extract_deleter() precondition failure (") + context
+                    + ").");
+            }
+            static_assert(std::is_copy_constructible<D>::value,
+                          "Required for compatibility with smart_holder functionality.");
+            return std::unique_ptr<D>(new D(custom_deleter_ptr->deleter));
+        }
+        return nullptr;
+    }
+
+    static smart_holder from_raw_ptr_unowned(void *raw_ptr) {
+        smart_holder hld;
+        hld.vptr.reset(raw_ptr, [](void *) {});
+        hld.vptr_is_using_noop_deleter = true;
+        hld.is_populated = true;
+        return hld;
+    }
+
+    template <typename T>
+    T *as_raw_ptr_unowned() const {
+        return static_cast<T *>(vptr.get());
+    }
+
+    template <typename T>
+    static smart_holder from_raw_ptr_take_ownership(T *raw_ptr, bool void_cast_raw_ptr = false) {
+        ensure_pointee_is_destructible<T>("from_raw_ptr_take_ownership");
+        smart_holder hld;
+        auto gd = make_guarded_std_default_delete<T>(true);
+        if (void_cast_raw_ptr) {
+            hld.vptr.reset(static_cast<void *>(raw_ptr), std::move(gd));
+        } else {
+            hld.vptr.reset(raw_ptr, std::move(gd));
+        }
+        hld.vptr_is_using_std_default_delete = true;
+        hld.is_populated = true;
+        return hld;
+    }
+
+    // Caller is responsible for ensuring the complex preconditions
+    // (see `smart_holder_type_caster_support::load_helper`).
+    void disown(const get_guarded_delete_fn ggd_fn) {
+        reset_vptr_deleter_armed_flag(ggd_fn, false);
+        is_disowned = true;
+    }
+
+    // Caller is responsible for ensuring the complex preconditions
+    // (see `smart_holder_type_caster_support::load_helper`).
+    void reclaim_disowned(const get_guarded_delete_fn ggd_fn) {
+        reset_vptr_deleter_armed_flag(ggd_fn, true);
+        is_disowned = false;
+    }
+
+    // Caller is responsible for ensuring the complex preconditions
+    // (see `smart_holder_type_caster_support::load_helper`).
+    void release_disowned() { vptr.reset(); }
+
+    void ensure_can_release_ownership(const char *context = "ensure_can_release_ownership") const {
+        ensure_is_not_disowned(context);
+        ensure_vptr_is_using_std_default_delete(context);
+        ensure_use_count_1(context);
+    }
+
+    // Caller is responsible for ensuring the complex preconditions
+    // (see `smart_holder_type_caster_support::load_helper`).
+    void release_ownership(const get_guarded_delete_fn ggd_fn) {
+        reset_vptr_deleter_armed_flag(ggd_fn, false);
+        release_disowned();
+    }
+
+    template <typename T, typename D>
+    static smart_holder from_unique_ptr(std::unique_ptr<T, D> &&unq_ptr,
+                                        void *mi_subobject_ptr = nullptr) {
+        smart_holder hld;
+        hld.rtti_uqp_del = &typeid(D);
+        hld.vptr_is_using_std_default_delete = uqp_del_is_std_default_delete<T, D>();
+
+        // Build the owning control block on the *real object start* (T*).
+        guarded_delete gd
+            = hld.vptr_is_using_std_default_delete
+                  ? make_guarded_std_default_delete<T>(true)
+                  : make_guarded_custom_deleter<T, D>(std::move(unq_ptr.get_deleter()), true);
+        // Critical: construct owner with pointer we intend to delete
+        std::shared_ptr<T> owner(unq_ptr.get(), std::move(gd));
+        // Relinquish ownership only after successful construction of owner
+        (void) unq_ptr.release();
+
+        // Publish either the MI/VI subobject pointer (if provided) or the full object.
+        // Why this is needed:
+        //   * The `owner` shared_ptr must always manage the true object start (T*).
+        //     That ensures the deleter is invoked on a valid object header, so the
+        //     virtual destructor can dispatch safely (critical on MSVC with virtual
+        //     inheritance, where base subobjects are not at offset 0).
+        //   * However, pybind11 needs to *register* and expose the subobject pointer
+        //     appropriate for the type being bound.
+        //     This pointer may differ from the T* object start under multiple/virtual
+        //     inheritance.
+        // This is achieved by using an aliasing shared_ptr<void>:
+        //   - `owner` retains lifetime of the actual T* object start for deletion.
+        //   - `vptr` points at the adjusted subobject (mi_subobject_ptr), giving
+        //     Python the correct identity/registration address.
+        // If no subobject pointer is passed, we simply publish the full object.
+        if (mi_subobject_ptr) {
+            hld.vptr = std::shared_ptr<void>(owner, mi_subobject_ptr);
+        } else {
+            hld.vptr = std::static_pointer_cast<void>(owner);
+        }
+
+        hld.is_populated = true;
+        return hld;
+    }
+
+    template <typename T>
+    static smart_holder from_shared_ptr(const std::shared_ptr<T> &shd_ptr) {
+        smart_holder hld;
+        hld.vptr = std::static_pointer_cast<void>(shd_ptr);
+        hld.vptr_is_external_shared_ptr = true;
+        hld.is_populated = true;
+        return hld;
+    }
+
+    template <typename T>
+    std::shared_ptr<T> as_shared_ptr() const {
+        return std::static_pointer_cast<T>(vptr);
+    }
+};
+
+PYBIND11_NAMESPACE_END(memory)
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/detail/typeid.h

@@ -0,0 +1,65 @@
+/*
+    pybind11/detail/typeid.h: Compiler-independent access to type identifiers
+
+    Copyright (c) 2016 Wenzel Jakob <wenzel.jakob@epfl.ch>
+
+    All rights reserved. Use of this source code is governed by a
+    BSD-style license that can be found in the LICENSE file.
+*/
+
+#pragma once
+
+#include <cstdio>
+#include <cstdlib>
+
+#if defined(__GNUG__)
+#    include <cxxabi.h>
+#endif
+
+#include "common.h"
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+/// Erase all occurrences of a substring
+inline void erase_all(std::string &string, const std::string &search) {
+    for (size_t pos = 0;;) {
+        pos = string.find(search, pos);
+        if (pos == std::string::npos) {
+            break;
+        }
+        string.erase(pos, search.length());
+    }
+}
+
+PYBIND11_NOINLINE void clean_type_id(std::string &name) {
+#if defined(__GNUG__)
+    int status = 0;
+    std::unique_ptr<char, void (*)(void *)> res{
+        abi::__cxa_demangle(name.c_str(), nullptr, nullptr, &status), std::free};
+    if (status == 0) {
+        name = res.get();
+    }
+#else
+    detail::erase_all(name, "class ");
+    detail::erase_all(name, "struct ");
+    detail::erase_all(name, "enum ");
+#endif
+    detail::erase_all(name, "pybind11::");
+}
+
+inline std::string clean_type_id(const char *typeid_name) {
+    std::string name(typeid_name);
+    detail::clean_type_id(name);
+    return name;
+}
+
+PYBIND11_NAMESPACE_END(detail)
+
+/// Return a string representation of a C++ type
+template <typename T>
+static std::string type_id() {
+    return detail::clean_type_id(typeid(T).name());
+}
+
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/detail/using_smart_holder.h

@@ -0,0 +1,22 @@
+// Copyright (c) 2024 The Pybind Development Team.
+// All rights reserved. Use of this source code is governed by a
+// BSD-style license that can be found in the LICENSE file.
+
+#pragma once
+
+#include "common.h"
+#include "struct_smart_holder.h"
+
+#include <type_traits>
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+
+using pybind11::memory::smart_holder;
+
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+template <typename H>
+using is_smart_holder = std::is_same<H, smart_holder>;
+
+PYBIND11_NAMESPACE_END(detail)
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/detail/value_and_holder.h

@@ -0,0 +1,92 @@
+// Copyright (c) 2016-2024 The Pybind Development Team.
+// All rights reserved. Use of this source code is governed by a
+// BSD-style license that can be found in the LICENSE file.
+
+#pragma once
+
+#include "common.h"
+
+#include <cstddef>
+#include <cstdint>
+#include <typeinfo>
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+struct value_and_holder {
+    instance *inst = nullptr;
+    size_t index = 0u;
+    const detail::type_info *type = nullptr;
+    void **vh = nullptr;
+
+    // Main constructor for a found value/holder:
+    value_and_holder(instance *i, const detail::type_info *type, size_t vpos, size_t index)
+        : inst{i}, index{index}, type{type},
+          vh{inst->simple_layout ? inst->simple_value_holder
+                                 : &inst->nonsimple.values_and_holders[vpos]} {}
+
+    // Default constructor (used to signal a value-and-holder not found by get_value_and_holder())
+    value_and_holder() = default;
+
+    // Used for past-the-end iterator
+    explicit value_and_holder(size_t index) : index{index} {}
+
+    template <typename V = void>
+    V *&value_ptr() const {
+        return reinterpret_cast<V *&>(vh[0]);
+    }
+    // True if this `value_and_holder` has a non-null value pointer
+    explicit operator bool() const { return value_ptr() != nullptr; }
+
+    template <typename H>
+    H &holder() const {
+        return reinterpret_cast<H &>(vh[1]);
+    }
+    bool holder_constructed() const {
+        return inst->simple_layout
+                   ? inst->simple_holder_constructed
+                   : (inst->nonsimple.status[index] & instance::status_holder_constructed) != 0u;
+    }
+    // NOLINTNEXTLINE(readability-make-member-function-const)
+    void set_holder_constructed(bool v = true) {
+        if (inst->simple_layout) {
+            inst->simple_holder_constructed = v;
+        } else if (v) {
+            inst->nonsimple.status[index] |= instance::status_holder_constructed;
+        } else {
+            inst->nonsimple.status[index]
+                &= static_cast<std::uint8_t>(~instance::status_holder_constructed);
+        }
+    }
+    bool instance_registered() const {
+        return inst->simple_layout
+                   ? inst->simple_instance_registered
+                   : ((inst->nonsimple.status[index] & instance::status_instance_registered) != 0);
+    }
+    // NOLINTNEXTLINE(readability-make-member-function-const)
+    void set_instance_registered(bool v = true) {
+        if (inst->simple_layout) {
+            inst->simple_instance_registered = v;
+        } else if (v) {
+            inst->nonsimple.status[index] |= instance::status_instance_registered;
+        } else {
+            inst->nonsimple.status[index]
+                &= static_cast<std::uint8_t>(~instance::status_instance_registered);
+        }
+    }
+};
+
+// This is a semi-public API to check if the corresponding instance has been constructed with a
+// holder. That is, if the instance has been constructed with a holder, the `__init__` method is
+// called and the C++ object is valid. Otherwise, the C++ object might only be allocated, but not
+// initialized. This will lead to **SEGMENTATION FAULTS** if the C++ object is used in any way.
+// Example usage: https://pybind11.readthedocs.io/en/stable/advanced/classes.html#custom-type-setup
+//                for `tp_traverse` and `tp_clear` implementations.
+// WARNING: The caller is responsible for ensuring that the `reinterpret_cast` is valid.
+inline bool is_holder_constructed(PyObject *obj) {
+    auto *const instance = reinterpret_cast<pybind11::detail::instance *>(obj);
+    return instance->get_value_and_holder().holder_constructed();
+}
+
+PYBIND11_NAMESPACE_END(detail)
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/eigen.h

@@ -0,0 +1,12 @@
+/*
+    pybind11/eigen.h: Transparent conversion for dense and sparse Eigen matrices
+
+    Copyright (c) 2016 Wenzel Jakob <wenzel.jakob@epfl.ch>
+
+    All rights reserved. Use of this source code is governed by a
+    BSD-style license that can be found in the LICENSE file.
+*/
+
+#pragma once
+
+#include "eigen/matrix.h"

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/eigen/common.h

@@ -0,0 +1,9 @@
+// Copyright (c) 2023 The pybind Community.
+
+#pragma once
+
+// Common message for `static_assert()`s, which are useful to easily
+// preempt much less obvious errors.
+#define PYBIND11_EIGEN_MESSAGE_POINTER_TYPES_ARE_NOT_SUPPORTED                                    \
+    "Pointer types (in particular `PyObject *`) are not supported as scalar types for Eigen "     \
+    "types."

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/eigen/matrix.h

@@ -0,0 +1,723 @@
+/*
+    pybind11/eigen/matrix.h: Transparent conversion for dense and sparse Eigen matrices
+
+    Copyright (c) 2016 Wenzel Jakob <wenzel.jakob@epfl.ch>
+
+    All rights reserved. Use of this source code is governed by a
+    BSD-style license that can be found in the LICENSE file.
+*/
+
+#pragma once
+
+#include <pybind11/numpy.h>
+
+#include "common.h"
+
+/* HINT: To suppress warnings originating from the Eigen headers, use -isystem.
+   See also:
+       https://stackoverflow.com/questions/2579576/i-dir-vs-isystem-dir
+       https://stackoverflow.com/questions/1741816/isystem-for-ms-visual-studio-c-compiler
+*/
+PYBIND11_WARNING_PUSH
+PYBIND11_WARNING_DISABLE_MSVC(5054) // https://github.com/pybind/pybind11/pull/3741
+//       C5054: operator '&': deprecated between enumerations of different types
+#if defined(__MINGW32__)
+PYBIND11_WARNING_DISABLE_GCC("-Wmaybe-uninitialized")
+#endif
+
+#include <Eigen/Core>
+#include <Eigen/SparseCore>
+
+PYBIND11_WARNING_POP
+
+// Eigen prior to 3.2.7 doesn't have proper move constructors--but worse, some classes get implicit
+// move constructors that break things.  We could detect this an explicitly copy, but an extra copy
+// of matrices seems highly undesirable.
+static_assert(EIGEN_VERSION_AT_LEAST(3, 2, 7),
+              "Eigen matrix support in pybind11 requires Eigen >= 3.2.7");
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+
+PYBIND11_WARNING_DISABLE_MSVC(4127)
+
+// Provide a convenience alias for easier pass-by-ref usage with fully dynamic strides:
+using EigenDStride = Eigen::Stride<Eigen::Dynamic, Eigen::Dynamic>;
+template <typename MatrixType>
+using EigenDRef = Eigen::Ref<MatrixType, 0, EigenDStride>;
+template <typename MatrixType>
+using EigenDMap = Eigen::Map<MatrixType, 0, EigenDStride>;
+
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+#if EIGEN_VERSION_AT_LEAST(3, 3, 0)
+using EigenIndex = Eigen::Index;
+template <typename Scalar, int Flags, typename StorageIndex>
+using EigenMapSparseMatrix = Eigen::Map<Eigen::SparseMatrix<Scalar, Flags, StorageIndex>>;
+#else
+using EigenIndex = EIGEN_DEFAULT_DENSE_INDEX_TYPE;
+template <typename Scalar, int Flags, typename StorageIndex>
+using EigenMapSparseMatrix = Eigen::MappedSparseMatrix<Scalar, Flags, StorageIndex>;
+#endif
+
+// Matches Eigen::Map, Eigen::Ref, blocks, etc:
+template <typename T>
+using is_eigen_dense_map = all_of<is_template_base_of<Eigen::DenseBase, T>,
+                                  std::is_base_of<Eigen::MapBase<T, Eigen::ReadOnlyAccessors>, T>>;
+template <typename T>
+using is_eigen_mutable_map = std::is_base_of<Eigen::MapBase<T, Eigen::WriteAccessors>, T>;
+template <typename T>
+using is_eigen_dense_plain
+    = all_of<negation<is_eigen_dense_map<T>>, is_template_base_of<Eigen::PlainObjectBase, T>>;
+template <typename T>
+using is_eigen_sparse = is_template_base_of<Eigen::SparseMatrixBase, T>;
+// Test for objects inheriting from EigenBase<Derived> that aren't captured by the above.  This
+// basically covers anything that can be assigned to a dense matrix but that don't have a typical
+// matrix data layout that can be copied from their .data().  For example, DiagonalMatrix and
+// SelfAdjointView fall into this category.
+template <typename T>
+using is_eigen_other
+    = all_of<is_template_base_of<Eigen::EigenBase, T>,
+             negation<any_of<is_eigen_dense_map<T>, is_eigen_dense_plain<T>, is_eigen_sparse<T>>>>;
+
+// Captures numpy/eigen conformability status (returned by EigenProps::conformable()):
+template <bool EigenRowMajor>
+struct EigenConformable {
+    bool conformable = false;
+    EigenIndex rows = 0, cols = 0;
+    EigenDStride stride{0, 0};    // Only valid if negativestrides is false!
+    bool negativestrides = false; // If true, do not use stride!
+
+    // NOLINTNEXTLINE(google-explicit-constructor)
+    EigenConformable(bool fits = false) : conformable{fits} {}
+    // Matrix type:
+    EigenConformable(EigenIndex r, EigenIndex c, EigenIndex rstride, EigenIndex cstride)
+        : conformable{true}, rows{r}, cols{c},
+          // TODO: when Eigen bug #747 is fixed, remove the tests for non-negativity.
+          // http://eigen.tuxfamily.org/bz/show_bug.cgi?id=747
+          stride{EigenRowMajor ? (rstride > 0 ? rstride : 0)
+                               : (cstride > 0 ? cstride : 0) /* outer stride */,
+                 EigenRowMajor ? (cstride > 0 ? cstride : 0)
+                               : (rstride > 0 ? rstride : 0) /* inner stride */},
+          negativestrides{rstride < 0 || cstride < 0} {}
+    // Vector type:
+    EigenConformable(EigenIndex r, EigenIndex c, EigenIndex stride)
+        : EigenConformable(r, c, r == 1 ? c * stride : stride, c == 1 ? r : r * stride) {}
+
+    template <typename props>
+    bool stride_compatible() const {
+        // To have compatible strides, we need (on both dimensions) one of fully dynamic strides,
+        // matching strides, or a dimension size of 1 (in which case the stride value is
+        // irrelevant). Alternatively, if any dimension size is 0, the strides are not relevant
+        // (and numpy ≥ 1.23 sets the strides to 0 in that case, so we need to check explicitly).
+        if (negativestrides) {
+            return false;
+        }
+        if (rows == 0 || cols == 0) {
+            return true;
+        }
+        return (props::inner_stride == Eigen::Dynamic || props::inner_stride == stride.inner()
+                || (EigenRowMajor ? cols : rows) == 1)
+               && (props::outer_stride == Eigen::Dynamic || props::outer_stride == stride.outer()
+                   || (EigenRowMajor ? rows : cols) == 1);
+    }
+    // NOLINTNEXTLINE(google-explicit-constructor)
+    operator bool() const { return conformable; }
+};
+
+template <typename Type>
+struct eigen_extract_stride {
+    using type = Type;
+};
+template <typename PlainObjectType, int MapOptions, typename StrideType>
+struct eigen_extract_stride<Eigen::Map<PlainObjectType, MapOptions, StrideType>> {
+    using type = StrideType;
+};
+template <typename PlainObjectType, int Options, typename StrideType>
+struct eigen_extract_stride<Eigen::Ref<PlainObjectType, Options, StrideType>> {
+    using type = StrideType;
+};
+
+// Helper struct for extracting information from an Eigen type
+template <typename Type_>
+struct EigenProps {
+    using Type = Type_;
+    using Scalar = typename Type::Scalar;
+    using StrideType = typename eigen_extract_stride<Type>::type;
+    static constexpr EigenIndex rows = Type::RowsAtCompileTime, cols = Type::ColsAtCompileTime,
+                                size = Type::SizeAtCompileTime;
+    static constexpr bool row_major = Type::IsRowMajor,
+                          vector
+                          = Type::IsVectorAtCompileTime, // At least one dimension has fixed size 1
+        fixed_rows = rows != Eigen::Dynamic, fixed_cols = cols != Eigen::Dynamic,
+                          fixed = size != Eigen::Dynamic, // Fully-fixed size
+        dynamic = !fixed_rows && !fixed_cols;             // Fully-dynamic size
+
+    template <EigenIndex i, EigenIndex ifzero>
+    using if_zero = std::integral_constant<EigenIndex, i == 0 ? ifzero : i>;
+    static constexpr EigenIndex inner_stride
+        = if_zero<StrideType::InnerStrideAtCompileTime, 1>::value,
+        outer_stride = if_zero < StrideType::OuterStrideAtCompileTime,
+        vector      ? size
+        : row_major ? cols
+                    : rows > ::value;
+    static constexpr bool dynamic_stride
+        = inner_stride == Eigen::Dynamic && outer_stride == Eigen::Dynamic;
+    static constexpr bool requires_row_major
+        = !dynamic_stride && !vector && (row_major ? inner_stride : outer_stride) == 1;
+    static constexpr bool requires_col_major
+        = !dynamic_stride && !vector && (row_major ? outer_stride : inner_stride) == 1;
+
+    // Takes an input array and determines whether we can make it fit into the Eigen type.  If
+    // the array is a vector, we attempt to fit it into either an Eigen 1xN or Nx1 vector
+    // (preferring the latter if it will fit in either, i.e. for a fully dynamic matrix type).
+    static EigenConformable<row_major> conformable(const array &a) {
+        const auto dims = a.ndim();
+        if (dims < 1 || dims > 2) {
+            return false;
+        }
+
+        if (dims == 2) { // Matrix type: require exact match (or dynamic)
+
+            EigenIndex np_rows = a.shape(0), np_cols = a.shape(1),
+                       np_rstride = a.strides(0) / static_cast<ssize_t>(sizeof(Scalar)),
+                       np_cstride = a.strides(1) / static_cast<ssize_t>(sizeof(Scalar));
+            if ((fixed_rows && np_rows != rows) || (fixed_cols && np_cols != cols)) {
+                return false;
+            }
+
+            return {np_rows, np_cols, np_rstride, np_cstride};
+        }
+
+        // Otherwise we're storing an n-vector.  Only one of the strides will be used, but
+        // whichever is used, we want the (single) numpy stride value.
+        const EigenIndex n = a.shape(0),
+                         stride = a.strides(0) / static_cast<ssize_t>(sizeof(Scalar));
+
+        if (vector) { // Eigen type is a compile-time vector
+            if (fixed && size != n) {
+                return false; // Vector size mismatch
+            }
+            return {rows == 1 ? 1 : n, cols == 1 ? 1 : n, stride};
+        }
+        if (fixed) {
+            // The type has a fixed size, but is not a vector: abort
+            return false;
+        }
+        if (fixed_cols) {
+            // Since this isn't a vector, cols must be != 1.  We allow this only if it exactly
+            // equals the number of elements (rows is Dynamic, and so 1 row is allowed).
+            if (cols != n) {
+                return false;
+            }
+            return {1, n, stride};
+        } // Otherwise it's either fully dynamic, or column dynamic; both become a column vector
+        if (fixed_rows && rows != n) {
+            return false;
+        }
+        return {n, 1, stride};
+    }
+
+    static constexpr bool show_writeable
+        = is_eigen_dense_map<Type>::value && is_eigen_mutable_map<Type>::value;
+    static constexpr bool show_order = is_eigen_dense_map<Type>::value;
+    static constexpr bool show_c_contiguous = show_order && requires_row_major;
+    static constexpr bool show_f_contiguous
+        = !show_c_contiguous && show_order && requires_col_major;
+
+    static constexpr auto descriptor
+        = const_name("typing.Annotated[")
+          + io_name("numpy.typing.ArrayLike, ", "numpy.typing.NDArray[")
+          + npy_format_descriptor<Scalar>::name + io_name("", "]") + const_name(", \"[")
+          + const_name<fixed_rows>(const_name<(size_t) rows>(), const_name("m")) + const_name(", ")
+          + const_name<fixed_cols>(const_name<(size_t) cols>(), const_name("n"))
+          + const_name("]\"")
+          // For a reference type (e.g. Ref<MatrixXd>) we have other constraints that might need to
+          // be satisfied: writeable=True (for a mutable reference), and, depending on the map's
+          // stride options, possibly f_contiguous or c_contiguous.  We include them in the
+          // descriptor output to provide some hint as to why a TypeError is occurring (otherwise
+          // it can be confusing to see that a function accepts a
+          // 'typing.Annotated[numpy.typing.NDArray[numpy.float64], "[3,2]"]' and an error message
+          // that you *gave* a numpy.ndarray of the right type and dimensions.
+          + const_name<show_writeable>(", \"flags.writeable\"", "")
+          + const_name<show_c_contiguous>(", \"flags.c_contiguous\"", "")
+          + const_name<show_f_contiguous>(", \"flags.f_contiguous\"", "") + const_name("]");
+};
+
+// Casts an Eigen type to numpy array.  If given a base, the numpy array references the src data,
+// otherwise it'll make a copy.  writeable lets you turn off the writeable flag for the array.
+template <typename props>
+handle
+eigen_array_cast(typename props::Type const &src, handle base = handle(), bool writeable = true) {
+    constexpr ssize_t elem_size = sizeof(typename props::Scalar);
+    array a;
+    if (props::vector) {
+        a = array({src.size()}, {elem_size * src.innerStride()}, src.data(), base);
+    } else {
+        a = array({src.rows(), src.cols()},
+                  {elem_size * src.rowStride(), elem_size * src.colStride()},
+                  src.data(),
+                  base);
+    }
+
+    if (!writeable) {
+        array_proxy(a.ptr())->flags &= ~detail::npy_api::NPY_ARRAY_WRITEABLE_;
+    }
+
+    return a.release();
+}
+
+// Takes an lvalue ref to some Eigen type and a (python) base object, creating a numpy array that
+// reference the Eigen object's data with `base` as the python-registered base class (if omitted,
+// the base will be set to None, and lifetime management is up to the caller).  The numpy array is
+// non-writeable if the given type is const.
+template <typename props, typename Type>
+handle eigen_ref_array(Type &src, handle parent = none()) {
+    // none here is to get past array's should-we-copy detection, which currently always
+    // copies when there is no base.  Setting the base to None should be harmless.
+    return eigen_array_cast<props>(src, parent, !std::is_const<Type>::value);
+}
+
+// Takes a pointer to some dense, plain Eigen type, builds a capsule around it, then returns a
+// numpy array that references the encapsulated data with a python-side reference to the capsule to
+// tie its destruction to that of any dependent python objects.  Const-ness is determined by
+// whether or not the Type of the pointer given is const.
+template <typename props, typename Type, typename = enable_if_t<is_eigen_dense_plain<Type>::value>>
+handle eigen_encapsulate(Type *src) {
+    capsule base(src, [](void *o) { delete static_cast<Type *>(o); });
+    return eigen_ref_array<props>(*src, base);
+}
+
+// Type caster for regular, dense matrix types (e.g. MatrixXd), but not maps/refs/etc. of dense
+// types.
+template <typename Type>
+struct type_caster<Type, enable_if_t<is_eigen_dense_plain<Type>::value>> {
+    using Scalar = typename Type::Scalar;
+    static_assert(!std::is_pointer<Scalar>::value,
+                  PYBIND11_EIGEN_MESSAGE_POINTER_TYPES_ARE_NOT_SUPPORTED);
+    using props = EigenProps<Type>;
+
+    bool load(handle src, bool convert) {
+        // If we're in no-convert mode, only load if given an array of the correct type
+        if (!convert && !isinstance<array_t<Scalar>>(src)) {
+            return false;
+        }
+
+        // Coerce into an array, but don't do type conversion yet; the copy below handles it.
+        auto buf = array::ensure(src);
+
+        if (!buf) {
+            return false;
+        }
+
+        auto dims = buf.ndim();
+        if (dims < 1 || dims > 2) {
+            return false;
+        }
+
+        auto fits = props::conformable(buf);
+        if (!fits) {
+            return false;
+        }
+
+        PYBIND11_WARNING_PUSH
+        PYBIND11_WARNING_DISABLE_GCC("-Wmaybe-uninitialized") // See PR #5516
+        // Allocate the new type, then build a numpy reference into it
+        value = Type(fits.rows, fits.cols);
+        PYBIND11_WARNING_POP
+        auto ref = reinterpret_steal<array>(eigen_ref_array<props>(value));
+        if (dims == 1) {
+            ref = ref.squeeze();
+        } else if (ref.ndim() == 1) {
+            buf = buf.squeeze();
+        }
+
+        int result = detail::npy_api::get().PyArray_CopyInto_(ref.ptr(), buf.ptr());
+
+        if (result < 0) { // Copy failed!
+            PyErr_Clear();
+            return false;
+        }
+
+        return true;
+    }
+
+private:
+    // Cast implementation
+    template <typename CType>
+    static handle cast_impl(CType *src, return_value_policy policy, handle parent) {
+        switch (policy) {
+            case return_value_policy::take_ownership:
+            case return_value_policy::automatic:
+                return eigen_encapsulate<props>(src);
+            case return_value_policy::move:
+                return eigen_encapsulate<props>(new CType(std::move(*src)));
+            case return_value_policy::copy:
+                return eigen_array_cast<props>(*src);
+            case return_value_policy::reference:
+            case return_value_policy::automatic_reference:
+                return eigen_ref_array<props>(*src);
+            case return_value_policy::reference_internal:
+                return eigen_ref_array<props>(*src, parent);
+            default:
+                throw cast_error("unhandled return_value_policy: should not happen!");
+        };
+    }
+
+public:
+    // Normal returned non-reference, non-const value:
+    static handle cast(Type &&src, return_value_policy /* policy */, handle parent) {
+        return cast_impl(&src, return_value_policy::move, parent);
+    }
+    // If you return a non-reference const, we mark the numpy array readonly:
+    static handle cast(const Type &&src, return_value_policy /* policy */, handle parent) {
+        return cast_impl(&src, return_value_policy::move, parent);
+    }
+    // lvalue reference return; default (automatic) becomes copy
+    static handle cast(Type &src, return_value_policy policy, handle parent) {
+        if (policy == return_value_policy::automatic
+            || policy == return_value_policy::automatic_reference) {
+            policy = return_value_policy::copy;
+        }
+        return cast_impl(&src, policy, parent);
+    }
+    // const lvalue reference return; default (automatic) becomes copy
+    static handle cast(const Type &src, return_value_policy policy, handle parent) {
+        if (policy == return_value_policy::automatic
+            || policy == return_value_policy::automatic_reference) {
+            policy = return_value_policy::copy;
+        }
+        return cast(&src, policy, parent);
+    }
+    // non-const pointer return
+    static handle cast(Type *src, return_value_policy policy, handle parent) {
+        return cast_impl(src, policy, parent);
+    }
+    // const pointer return
+    static handle cast(const Type *src, return_value_policy policy, handle parent) {
+        return cast_impl(src, policy, parent);
+    }
+
+    static constexpr auto name = props::descriptor;
+
+    // NOLINTNEXTLINE(google-explicit-constructor)
+    operator Type *() { return &value; }
+    // NOLINTNEXTLINE(google-explicit-constructor)
+    operator Type &() { return value; }
+    // NOLINTNEXTLINE(google-explicit-constructor)
+    operator Type &&() && { return std::move(value); }
+    template <typename T>
+    using cast_op_type = movable_cast_op_type<T>;
+
+private:
+    Type value;
+};
+
+// Base class for casting reference/map/block/etc. objects back to python.
+template <typename MapType>
+struct eigen_map_caster {
+    static_assert(!std::is_pointer<typename MapType::Scalar>::value,
+                  PYBIND11_EIGEN_MESSAGE_POINTER_TYPES_ARE_NOT_SUPPORTED);
+
+private:
+    using props = EigenProps<MapType>;
+
+public:
+    // Directly referencing a ref/map's data is a bit dangerous (whatever the map/ref points to has
+    // to stay around), but we'll allow it under the assumption that you know what you're doing
+    // (and have an appropriate keep_alive in place).  We return a numpy array pointing directly at
+    // the ref's data (The numpy array ends up read-only if the ref was to a const matrix type.)
+    // Note that this means you need to ensure you don't destroy the object in some other way (e.g.
+    // with an appropriate keep_alive, or with a reference to a statically allocated matrix).
+    static handle cast(const MapType &src, return_value_policy policy, handle parent) {
+        switch (policy) {
+            case return_value_policy::copy:
+                return eigen_array_cast<props>(src);
+            case return_value_policy::reference_internal:
+                return eigen_array_cast<props>(src, parent, is_eigen_mutable_map<MapType>::value);
+            case return_value_policy::reference:
+            case return_value_policy::automatic:
+            case return_value_policy::automatic_reference:
+                return eigen_array_cast<props>(src, none(), is_eigen_mutable_map<MapType>::value);
+            default:
+                // move, take_ownership don't make any sense for a ref/map:
+                pybind11_fail("Invalid return_value_policy for Eigen Map/Ref/Block type");
+        }
+    }
+
+    // return_descr forces the use of NDArray instead of ArrayLike in args
+    // since Ref<...> args can only accept arrays.
+    static constexpr auto name = return_descr(props::descriptor);
+
+    // Explicitly delete these: support python -> C++ conversion on these (i.e. these can be return
+    // types but not bound arguments).  We still provide them (with an explicitly delete) so that
+    // you end up here if you try anyway.
+    bool load(handle, bool) = delete;
+    operator MapType() = delete;
+    template <typename>
+    using cast_op_type = MapType;
+};
+
+// We can return any map-like object (but can only load Refs, specialized next):
+template <typename Type>
+struct type_caster<Type, enable_if_t<is_eigen_dense_map<Type>::value>> : eigen_map_caster<Type> {};
+
+// Loader for Ref<...> arguments.  See the documentation for info on how to make this work without
+// copying (it requires some extra effort in many cases).
+template <typename PlainObjectType, typename StrideType>
+struct type_caster<
+    Eigen::Ref<PlainObjectType, 0, StrideType>,
+    enable_if_t<is_eigen_dense_map<Eigen::Ref<PlainObjectType, 0, StrideType>>::value>>
+    : public eigen_map_caster<Eigen::Ref<PlainObjectType, 0, StrideType>> {
+private:
+    using Type = Eigen::Ref<PlainObjectType, 0, StrideType>;
+    using props = EigenProps<Type>;
+    using Scalar = typename props::Scalar;
+    static_assert(!std::is_pointer<Scalar>::value,
+                  PYBIND11_EIGEN_MESSAGE_POINTER_TYPES_ARE_NOT_SUPPORTED);
+    using MapType = Eigen::Map<PlainObjectType, 0, StrideType>;
+    using Array
+        = array_t<Scalar,
+                  array::forcecast
+                      | ((props::row_major ? props::inner_stride : props::outer_stride) == 1
+                             ? array::c_style
+                         : (props::row_major ? props::outer_stride : props::inner_stride) == 1
+                             ? array::f_style
+                             : 0)>;
+    static constexpr bool need_writeable = is_eigen_mutable_map<Type>::value;
+    // Delay construction (these have no default constructor)
+    std::unique_ptr<MapType> map;
+    std::unique_ptr<Type> ref;
+    // Our array.  When possible, this is just a numpy array pointing to the source data, but
+    // sometimes we can't avoid copying (e.g. input is not a numpy array at all, has an
+    // incompatible layout, or is an array of a type that needs to be converted).  Using a numpy
+    // temporary (rather than an Eigen temporary) saves an extra copy when we need both type
+    // conversion and storage order conversion.  (Note that we refuse to use this temporary copy
+    // when loading an argument for a Ref<M> with M non-const, i.e. a read-write reference).
+    Array copy_or_ref;
+
+public:
+    bool load(handle src, bool convert) {
+        // First check whether what we have is already an array of the right type.  If not, we
+        // can't avoid a copy (because the copy is also going to do type conversion).
+        bool need_copy = !isinstance<Array>(src);
+
+        EigenConformable<props::row_major> fits;
+        if (!need_copy) {
+            // We don't need a converting copy, but we also need to check whether the strides are
+            // compatible with the Ref's stride requirements
+            auto aref = reinterpret_borrow<Array>(src);
+
+            if (aref && (!need_writeable || aref.writeable())) {
+                fits = props::conformable(aref);
+                if (!fits) {
+                    return false; // Incompatible dimensions
+                }
+                if (!fits.template stride_compatible<props>()) {
+                    need_copy = true;
+                } else {
+                    copy_or_ref = std::move(aref);
+                }
+            } else {
+                need_copy = true;
+            }
+        }
+
+        if (need_copy) {
+            // We need to copy: If we need a mutable reference, or we're not supposed to convert
+            // (either because we're in the no-convert overload pass, or because we're explicitly
+            // instructed not to copy (via `py::arg().noconvert()`) we have to fail loading.
+            if (!convert || need_writeable) {
+                return false;
+            }
+
+            Array copy = Array::ensure(src);
+            if (!copy) {
+                return false;
+            }
+            fits = props::conformable(copy);
+            if (!fits || !fits.template stride_compatible<props>()) {
+                return false;
+            }
+            copy_or_ref = std::move(copy);
+            loader_life_support::add_patient(copy_or_ref);
+        }
+
+        ref.reset();
+        map.reset(new MapType(data(copy_or_ref),
+                              fits.rows,
+                              fits.cols,
+                              make_stride(fits.stride.outer(), fits.stride.inner())));
+        ref.reset(new Type(*map));
+
+        return true;
+    }
+
+    // NOLINTNEXTLINE(google-explicit-constructor)
+    operator Type *() { return ref.get(); }
+    // NOLINTNEXTLINE(google-explicit-constructor)
+    operator Type &() { return *ref; }
+    template <typename _T>
+    using cast_op_type = pybind11::detail::cast_op_type<_T>;
+
+private:
+    template <typename T = Type, enable_if_t<is_eigen_mutable_map<T>::value, int> = 0>
+    Scalar *data(Array &a) {
+        return a.mutable_data();
+    }
+
+    template <typename T = Type, enable_if_t<!is_eigen_mutable_map<T>::value, int> = 0>
+    const Scalar *data(Array &a) {
+        return a.data();
+    }
+
+    // Attempt to figure out a constructor of `Stride` that will work.
+    // If both strides are fixed, use a default constructor:
+    template <typename S>
+    using stride_ctor_default = bool_constant<S::InnerStrideAtCompileTime != Eigen::Dynamic
+                                              && S::OuterStrideAtCompileTime != Eigen::Dynamic
+                                              && std::is_default_constructible<S>::value>;
+    // Otherwise, if there is a two-index constructor, assume it is (outer,inner) like
+    // Eigen::Stride, and use it:
+    template <typename S>
+    using stride_ctor_dual
+        = bool_constant<!stride_ctor_default<S>::value
+                        && std::is_constructible<S, EigenIndex, EigenIndex>::value>;
+    // Otherwise, if there is a one-index constructor, and just one of the strides is dynamic, use
+    // it (passing whichever stride is dynamic).
+    template <typename S>
+    using stride_ctor_outer
+        = bool_constant<!any_of<stride_ctor_default<S>, stride_ctor_dual<S>>::value
+                        && S::OuterStrideAtCompileTime == Eigen::Dynamic
+                        && S::InnerStrideAtCompileTime != Eigen::Dynamic
+                        && std::is_constructible<S, EigenIndex>::value>;
+    template <typename S>
+    using stride_ctor_inner
+        = bool_constant<!any_of<stride_ctor_default<S>, stride_ctor_dual<S>>::value
+                        && S::InnerStrideAtCompileTime == Eigen::Dynamic
+                        && S::OuterStrideAtCompileTime != Eigen::Dynamic
+                        && std::is_constructible<S, EigenIndex>::value>;
+
+    template <typename S = StrideType, enable_if_t<stride_ctor_default<S>::value, int> = 0>
+    static S make_stride(EigenIndex, EigenIndex) {
+        return S();
+    }
+    template <typename S = StrideType, enable_if_t<stride_ctor_dual<S>::value, int> = 0>
+    static S make_stride(EigenIndex outer, EigenIndex inner) {
+        return S(outer, inner);
+    }
+    template <typename S = StrideType, enable_if_t<stride_ctor_outer<S>::value, int> = 0>
+    static S make_stride(EigenIndex outer, EigenIndex) {
+        return S(outer);
+    }
+    template <typename S = StrideType, enable_if_t<stride_ctor_inner<S>::value, int> = 0>
+    static S make_stride(EigenIndex, EigenIndex inner) {
+        return S(inner);
+    }
+};
+
+// type_caster for special matrix types (e.g. DiagonalMatrix), which are EigenBase, but not
+// EigenDense (i.e. they don't have a data(), at least not with the usual matrix layout).
+// load() is not supported, but we can cast them into the python domain by first copying to a
+// regular Eigen::Matrix, then casting that.
+template <typename Type>
+struct type_caster<Type, enable_if_t<is_eigen_other<Type>::value>> {
+    static_assert(!std::is_pointer<typename Type::Scalar>::value,
+                  PYBIND11_EIGEN_MESSAGE_POINTER_TYPES_ARE_NOT_SUPPORTED);
+
+protected:
+    using Matrix
+        = Eigen::Matrix<typename Type::Scalar, Type::RowsAtCompileTime, Type::ColsAtCompileTime>;
+    using props = EigenProps<Matrix>;
+
+public:
+    static handle cast(const Type &src, return_value_policy /* policy */, handle /* parent */) {
+        handle h = eigen_encapsulate<props>(new Matrix(src));
+        return h;
+    }
+    static handle cast(const Type *src, return_value_policy policy, handle parent) {
+        return cast(*src, policy, parent);
+    }
+
+    static constexpr auto name = props::descriptor;
+
+    // Explicitly delete these: support python -> C++ conversion on these (i.e. these can be return
+    // types but not bound arguments).  We still provide them (with an explicitly delete) so that
+    // you end up here if you try anyway.
+    bool load(handle, bool) = delete;
+    operator Type() = delete;
+    template <typename>
+    using cast_op_type = Type;
+};
+
+template <typename Type>
+struct type_caster<Type, enable_if_t<is_eigen_sparse<Type>::value>> {
+    using Scalar = typename Type::Scalar;
+    static_assert(!std::is_pointer<Scalar>::value,
+                  PYBIND11_EIGEN_MESSAGE_POINTER_TYPES_ARE_NOT_SUPPORTED);
+    using StorageIndex = remove_reference_t<decltype(*std::declval<Type>().outerIndexPtr())>;
+    using Index = typename Type::Index;
+    static constexpr bool rowMajor = Type::IsRowMajor;
+
+    bool load(handle src, bool) {
+        if (!src) {
+            return false;
+        }
+
+        auto obj = reinterpret_borrow<object>(src);
+        object sparse_module = module_::import("scipy.sparse");
+        object matrix_type = sparse_module.attr(rowMajor ? "csr_matrix" : "csc_matrix");
+
+        if (!type::handle_of(obj).is(matrix_type)) {
+            try {
+                obj = matrix_type(obj);
+            } catch (const error_already_set &) {
+                return false;
+            }
+        }
+
+        auto values = array_t<Scalar>((object) obj.attr("data"));
+        auto innerIndices = array_t<StorageIndex>((object) obj.attr("indices"));
+        auto outerIndices = array_t<StorageIndex>((object) obj.attr("indptr"));
+        auto shape = pybind11::tuple((pybind11::object) obj.attr("shape"));
+        auto nnz = obj.attr("nnz").cast<Index>();
+
+        if (!values || !innerIndices || !outerIndices) {
+            return false;
+        }
+
+        value = EigenMapSparseMatrix<Scalar,
+                                     Type::Flags &(Eigen::RowMajor | Eigen::ColMajor),
+                                     StorageIndex>(shape[0].cast<Index>(),
+                                                   shape[1].cast<Index>(),
+                                                   std::move(nnz),
+                                                   outerIndices.mutable_data(),
+                                                   innerIndices.mutable_data(),
+                                                   values.mutable_data());
+
+        return true;
+    }
+
+    static handle cast(const Type &src, return_value_policy /* policy */, handle /* parent */) {
+        const_cast<Type &>(src).makeCompressed();
+
+        object matrix_type
+            = module_::import("scipy.sparse").attr(rowMajor ? "csr_matrix" : "csc_matrix");
+
+        array data(src.nonZeros(), src.valuePtr());
+        array outerIndices((rowMajor ? src.rows() : src.cols()) + 1, src.outerIndexPtr());
+        array innerIndices(src.nonZeros(), src.innerIndexPtr());
+
+        return matrix_type(pybind11::make_tuple(
+                               std::move(data), std::move(innerIndices), std::move(outerIndices)),
+                           pybind11::make_tuple(src.rows(), src.cols()))
+            .release();
+    }
+
+    PYBIND11_TYPE_CASTER(Type,
+                         const_name<(Type::IsRowMajor) != 0>("scipy.sparse.csr_matrix[",
+                                                             "scipy.sparse.csc_matrix[")
+                             + npy_format_descriptor<Scalar>::name + const_name("]"));
+};
+
+PYBIND11_NAMESPACE_END(detail)
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/eigen/tensor.h

@@ -0,0 +1,520 @@
+/*
+    pybind11/eigen/tensor.h: Transparent conversion for Eigen tensors
+
+    All rights reserved. Use of this source code is governed by a
+    BSD-style license that can be found in the LICENSE file.
+*/
+
+#pragma once
+
+#include <pybind11/numpy.h>
+
+#include "common.h"
+
+#if defined(__GNUC__) && !defined(__clang__) && !defined(__INTEL_COMPILER)
+static_assert(__GNUC__ > 5, "Eigen Tensor support in pybind11 requires GCC > 5.0");
+#endif
+
+// Disable warnings for Eigen
+PYBIND11_WARNING_PUSH
+PYBIND11_WARNING_DISABLE_MSVC(4554)
+PYBIND11_WARNING_DISABLE_MSVC(4127)
+#if defined(__MINGW32__)
+PYBIND11_WARNING_DISABLE_GCC("-Wmaybe-uninitialized")
+#endif
+
+#include <unsupported/Eigen/CXX11/Tensor>
+
+PYBIND11_WARNING_POP
+
+static_assert(EIGEN_VERSION_AT_LEAST(3, 3, 0),
+              "Eigen Tensor support in pybind11 requires Eigen >= 3.3.0");
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+
+PYBIND11_WARNING_DISABLE_MSVC(4127)
+
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+inline bool is_tensor_aligned(const void *data) {
+    return (reinterpret_cast<std::size_t>(data) % EIGEN_DEFAULT_ALIGN_BYTES) == 0;
+}
+
+template <typename T>
+constexpr int compute_array_flag_from_tensor() {
+    static_assert((static_cast<int>(T::Layout) == static_cast<int>(Eigen::RowMajor))
+                      || (static_cast<int>(T::Layout) == static_cast<int>(Eigen::ColMajor)),
+                  "Layout must be row or column major");
+    return (static_cast<int>(T::Layout) == static_cast<int>(Eigen::RowMajor)) ? array::c_style
+                                                                              : array::f_style;
+}
+
+template <typename T>
+struct eigen_tensor_helper {};
+
+template <typename Scalar_, int NumIndices_, int Options_, typename IndexType>
+struct eigen_tensor_helper<Eigen::Tensor<Scalar_, NumIndices_, Options_, IndexType>> {
+    using Type = Eigen::Tensor<Scalar_, NumIndices_, Options_, IndexType>;
+    using ValidType = void;
+
+    static Eigen::DSizes<typename Type::Index, Type::NumIndices> get_shape(const Type &f) {
+        return f.dimensions();
+    }
+
+    static constexpr bool
+    is_correct_shape(const Eigen::DSizes<typename Type::Index, Type::NumIndices> & /*shape*/) {
+        return true;
+    }
+
+    template <typename T>
+    struct helper {};
+
+    template <size_t... Is>
+    struct helper<index_sequence<Is...>> {
+        static constexpr auto value = ::pybind11::detail::concat(const_name(((void) Is, "?"))...);
+    };
+
+    static constexpr auto dimensions_descriptor
+        = helper<decltype(make_index_sequence<Type::NumIndices>())>::value;
+
+    template <typename... Args>
+    static Type *alloc(Args &&...args) {
+        return new Type(std::forward<Args>(args)...);
+    }
+
+    static void free(Type *tensor) { delete tensor; }
+};
+
+template <typename Scalar_, typename std::ptrdiff_t... Indices, int Options_, typename IndexType>
+struct eigen_tensor_helper<
+    Eigen::TensorFixedSize<Scalar_, Eigen::Sizes<Indices...>, Options_, IndexType>> {
+    using Type = Eigen::TensorFixedSize<Scalar_, Eigen::Sizes<Indices...>, Options_, IndexType>;
+    using ValidType = void;
+
+    static constexpr Eigen::DSizes<typename Type::Index, Type::NumIndices>
+    get_shape(const Type & /*f*/) {
+        return get_shape();
+    }
+
+    static constexpr Eigen::DSizes<typename Type::Index, Type::NumIndices> get_shape() {
+        return Eigen::DSizes<typename Type::Index, Type::NumIndices>(Indices...);
+    }
+
+    static bool
+    is_correct_shape(const Eigen::DSizes<typename Type::Index, Type::NumIndices> &shape) {
+        return get_shape() == shape;
+    }
+
+    static constexpr auto dimensions_descriptor
+        = ::pybind11::detail::concat(const_name<Indices>()...);
+
+    template <typename... Args>
+    static Type *alloc(Args &&...args) {
+        Eigen::aligned_allocator<Type> allocator;
+        return ::new (allocator.allocate(1)) Type(std::forward<Args>(args)...);
+    }
+
+    static void free(Type *tensor) {
+        Eigen::aligned_allocator<Type> allocator;
+        tensor->~Type();
+        allocator.deallocate(tensor, 1);
+    }
+};
+
+template <typename Type, bool ShowDetails, bool NeedsWriteable = false>
+struct get_tensor_descriptor {
+    static constexpr auto details
+        = const_name<NeedsWriteable>(", \"flags.writeable\"", "")
+          + const_name<static_cast<int>(Type::Layout) == static_cast<int>(Eigen::RowMajor)>(
+              ", \"flags.c_contiguous\"", ", \"flags.f_contiguous\"");
+    static constexpr auto value
+        = const_name("typing.Annotated[")
+          + io_name("numpy.typing.ArrayLike, ", "numpy.typing.NDArray[")
+          + npy_format_descriptor<typename Type::Scalar>::name + io_name("", "]")
+          + const_name(", \"[") + eigen_tensor_helper<remove_cv_t<Type>>::dimensions_descriptor
+          + const_name("]\"") + const_name<ShowDetails>(details, const_name("")) + const_name("]");
+};
+
+// When EIGEN_AVOID_STL_ARRAY is defined, Eigen::DSizes<T, 0> does not have the begin() member
+// function. Falling back to a simple loop works around this issue.
+//
+// We need to disable the type-limits warning for the inner loop when size = 0.
+
+PYBIND11_WARNING_PUSH
+PYBIND11_WARNING_DISABLE_GCC("-Wtype-limits")
+
+template <typename T, int size>
+std::vector<T> convert_dsizes_to_vector(const Eigen::DSizes<T, size> &arr) {
+    std::vector<T> result(size);
+
+    for (size_t i = 0; i < size; i++) {
+        result[i] = arr[i];
+    }
+
+    return result;
+}
+
+template <typename T, int size>
+Eigen::DSizes<T, size> get_shape_for_array(const array &arr) {
+    Eigen::DSizes<T, size> result;
+    const T *shape = arr.shape();
+    for (size_t i = 0; i < size; i++) {
+        result[i] = shape[i];
+    }
+
+    return result;
+}
+
+PYBIND11_WARNING_POP
+
+template <typename Type>
+struct type_caster<Type, typename eigen_tensor_helper<Type>::ValidType> {
+    static_assert(!std::is_pointer<typename Type::Scalar>::value,
+                  PYBIND11_EIGEN_MESSAGE_POINTER_TYPES_ARE_NOT_SUPPORTED);
+    using Helper = eigen_tensor_helper<Type>;
+    static constexpr auto temp_name = get_tensor_descriptor<Type, false>::value;
+    PYBIND11_TYPE_CASTER(Type, temp_name);
+
+    bool load(handle src, bool convert) {
+        if (!convert) {
+            if (!isinstance<array>(src)) {
+                return false;
+            }
+            array temp = array::ensure(src);
+            if (!temp) {
+                return false;
+            }
+
+            if (!temp.dtype().is(dtype::of<typename Type::Scalar>())) {
+                return false;
+            }
+        }
+
+        array_t<typename Type::Scalar, compute_array_flag_from_tensor<Type>()> arr(
+            reinterpret_borrow<object>(src));
+
+        if (arr.ndim() != Type::NumIndices) {
+            return false;
+        }
+        auto shape = get_shape_for_array<typename Type::Index, Type::NumIndices>(arr);
+
+        if (!Helper::is_correct_shape(shape)) {
+            return false;
+        }
+
+#if EIGEN_VERSION_AT_LEAST(3, 4, 0)
+        auto data_pointer = arr.data();
+#else
+        // Handle Eigen bug
+        auto data_pointer = const_cast<typename Type::Scalar *>(arr.data());
+#endif
+
+        if (is_tensor_aligned(arr.data())) {
+            value = Eigen::TensorMap<const Type, Eigen::Aligned>(data_pointer, shape);
+        } else {
+            value = Eigen::TensorMap<const Type>(data_pointer, shape);
+        }
+
+        return true;
+    }
+
+    static handle cast(Type &&src, return_value_policy policy, handle parent) {
+        if (policy == return_value_policy::reference
+            || policy == return_value_policy::reference_internal) {
+            pybind11_fail("Cannot use a reference return value policy for an rvalue");
+        }
+        return cast_impl(&src, return_value_policy::move, parent);
+    }
+
+    static handle cast(const Type &&src, return_value_policy policy, handle parent) {
+        if (policy == return_value_policy::reference
+            || policy == return_value_policy::reference_internal) {
+            pybind11_fail("Cannot use a reference return value policy for an rvalue");
+        }
+        return cast_impl(&src, return_value_policy::move, parent);
+    }
+
+    static handle cast(Type &src, return_value_policy policy, handle parent) {
+        if (policy == return_value_policy::automatic
+            || policy == return_value_policy::automatic_reference) {
+            policy = return_value_policy::copy;
+        }
+        return cast_impl(&src, policy, parent);
+    }
+
+    static handle cast(const Type &src, return_value_policy policy, handle parent) {
+        if (policy == return_value_policy::automatic
+            || policy == return_value_policy::automatic_reference) {
+            policy = return_value_policy::copy;
+        }
+        return cast(&src, policy, parent);
+    }
+
+    static handle cast(Type *src, return_value_policy policy, handle parent) {
+        if (policy == return_value_policy::automatic) {
+            policy = return_value_policy::take_ownership;
+        } else if (policy == return_value_policy::automatic_reference) {
+            policy = return_value_policy::reference;
+        }
+        return cast_impl(src, policy, parent);
+    }
+
+    static handle cast(const Type *src, return_value_policy policy, handle parent) {
+        if (policy == return_value_policy::automatic) {
+            policy = return_value_policy::take_ownership;
+        } else if (policy == return_value_policy::automatic_reference) {
+            policy = return_value_policy::reference;
+        }
+        return cast_impl(src, policy, parent);
+    }
+
+    template <typename C>
+    static handle cast_impl(C *src, return_value_policy policy, handle parent) {
+        object parent_object;
+        bool writeable = false;
+        switch (policy) {
+            case return_value_policy::move:
+                if (std::is_const<C>::value) {
+                    pybind11_fail("Cannot move from a constant reference");
+                }
+
+                src = Helper::alloc(std::move(*src));
+
+                parent_object
+                    = capsule(src, [](void *ptr) { Helper::free(reinterpret_cast<Type *>(ptr)); });
+                writeable = true;
+                break;
+
+            case return_value_policy::take_ownership:
+                if (std::is_const<C>::value) {
+                    // This cast is ugly, and might be UB in some cases, but we don't have an
+                    // alternative here as we must free that memory
+                    Helper::free(const_cast<Type *>(src));
+                    pybind11_fail("Cannot take ownership of a const reference");
+                }
+
+                parent_object
+                    = capsule(src, [](void *ptr) { Helper::free(reinterpret_cast<Type *>(ptr)); });
+                writeable = true;
+                break;
+
+            case return_value_policy::copy:
+                writeable = true;
+                break;
+
+            case return_value_policy::reference:
+                parent_object = none();
+                writeable = !std::is_const<C>::value;
+                break;
+
+            case return_value_policy::reference_internal:
+                // Default should do the right thing
+                if (!parent) {
+                    pybind11_fail("Cannot use reference internal when there is no parent");
+                }
+                parent_object = reinterpret_borrow<object>(parent);
+                writeable = !std::is_const<C>::value;
+                break;
+
+            default:
+                pybind11_fail("pybind11 bug in eigen.h, please file a bug report");
+        }
+
+        auto result = array_t<typename Type::Scalar, compute_array_flag_from_tensor<Type>()>(
+            convert_dsizes_to_vector(Helper::get_shape(*src)), src->data(), parent_object);
+
+        if (!writeable) {
+            array_proxy(result.ptr())->flags &= ~detail::npy_api::NPY_ARRAY_WRITEABLE_;
+        }
+
+        return result.release();
+    }
+};
+
+template <typename StoragePointerType,
+          bool needs_writeable,
+          enable_if_t<!needs_writeable, bool> = true>
+StoragePointerType get_array_data_for_type(array &arr) {
+#if EIGEN_VERSION_AT_LEAST(3, 4, 0)
+    return reinterpret_cast<StoragePointerType>(arr.data());
+#else
+    // Handle Eigen bug
+    return reinterpret_cast<StoragePointerType>(const_cast<void *>(arr.data()));
+#endif
+}
+
+template <typename StoragePointerType,
+          bool needs_writeable,
+          enable_if_t<needs_writeable, bool> = true>
+StoragePointerType get_array_data_for_type(array &arr) {
+    return reinterpret_cast<StoragePointerType>(arr.mutable_data());
+}
+
+template <typename T, typename = void>
+struct get_storage_pointer_type;
+
+template <typename MapType>
+struct get_storage_pointer_type<MapType, void_t<typename MapType::StoragePointerType>> {
+    using SPT = typename MapType::StoragePointerType;
+};
+
+template <typename MapType>
+struct get_storage_pointer_type<MapType, void_t<typename MapType::PointerArgType>> {
+    using SPT = typename MapType::PointerArgType;
+};
+
+template <typename Type, int Options>
+struct type_caster<Eigen::TensorMap<Type, Options>,
+                   typename eigen_tensor_helper<remove_cv_t<Type>>::ValidType> {
+    static_assert(!std::is_pointer<typename Type::Scalar>::value,
+                  PYBIND11_EIGEN_MESSAGE_POINTER_TYPES_ARE_NOT_SUPPORTED);
+    using MapType = Eigen::TensorMap<Type, Options>;
+    using Helper = eigen_tensor_helper<remove_cv_t<Type>>;
+
+    bool load(handle src, bool /*convert*/) {
+        // Note that we have a lot more checks here as we want to make sure to avoid copies
+        if (!isinstance<array>(src)) {
+            return false;
+        }
+        auto arr = reinterpret_borrow<array>(src);
+        if ((arr.flags() & compute_array_flag_from_tensor<Type>()) == 0) {
+            return false;
+        }
+
+        if (!arr.dtype().is(dtype::of<typename Type::Scalar>())) {
+            return false;
+        }
+
+        if (arr.ndim() != Type::NumIndices) {
+            return false;
+        }
+
+        constexpr bool is_aligned = (Options & Eigen::Aligned) != 0;
+
+        if (is_aligned && !is_tensor_aligned(arr.data())) {
+            return false;
+        }
+
+        auto shape = get_shape_for_array<typename Type::Index, Type::NumIndices>(arr);
+
+        if (!Helper::is_correct_shape(shape)) {
+            return false;
+        }
+
+        if (needs_writeable && !arr.writeable()) {
+            return false;
+        }
+
+        auto result = get_array_data_for_type<typename get_storage_pointer_type<MapType>::SPT,
+                                              needs_writeable>(arr);
+
+        value.reset(new MapType(std::move(result), std::move(shape)));
+
+        return true;
+    }
+
+    static handle cast(MapType &&src, return_value_policy policy, handle parent) {
+        return cast_impl(&src, policy, parent);
+    }
+
+    static handle cast(const MapType &&src, return_value_policy policy, handle parent) {
+        return cast_impl(&src, policy, parent);
+    }
+
+    static handle cast(MapType &src, return_value_policy policy, handle parent) {
+        if (policy == return_value_policy::automatic
+            || policy == return_value_policy::automatic_reference) {
+            policy = return_value_policy::copy;
+        }
+        return cast_impl(&src, policy, parent);
+    }
+
+    static handle cast(const MapType &src, return_value_policy policy, handle parent) {
+        if (policy == return_value_policy::automatic
+            || policy == return_value_policy::automatic_reference) {
+            policy = return_value_policy::copy;
+        }
+        return cast(&src, policy, parent);
+    }
+
+    static handle cast(MapType *src, return_value_policy policy, handle parent) {
+        if (policy == return_value_policy::automatic) {
+            policy = return_value_policy::take_ownership;
+        } else if (policy == return_value_policy::automatic_reference) {
+            policy = return_value_policy::reference;
+        }
+        return cast_impl(src, policy, parent);
+    }
+
+    static handle cast(const MapType *src, return_value_policy policy, handle parent) {
+        if (policy == return_value_policy::automatic) {
+            policy = return_value_policy::take_ownership;
+        } else if (policy == return_value_policy::automatic_reference) {
+            policy = return_value_policy::reference;
+        }
+        return cast_impl(src, policy, parent);
+    }
+
+    template <typename C>
+    static handle cast_impl(C *src, return_value_policy policy, handle parent) {
+        object parent_object;
+        constexpr bool writeable = !std::is_const<C>::value;
+        switch (policy) {
+            case return_value_policy::reference:
+                parent_object = none();
+                break;
+
+            case return_value_policy::reference_internal:
+                // Default should do the right thing
+                if (!parent) {
+                    pybind11_fail("Cannot use reference internal when there is no parent");
+                }
+                parent_object = reinterpret_borrow<object>(parent);
+                break;
+
+            default:
+                // move, take_ownership don't make any sense for a ref/map:
+                pybind11_fail("Invalid return_value_policy for Eigen Map type, must be either "
+                              "reference or reference_internal");
+        }
+
+        auto result = array_t<typename Type::Scalar, compute_array_flag_from_tensor<Type>()>(
+            convert_dsizes_to_vector(Helper::get_shape(*src)),
+            src->data(),
+            std::move(parent_object));
+
+        if (!writeable) {
+            array_proxy(result.ptr())->flags &= ~detail::npy_api::NPY_ARRAY_WRITEABLE_;
+        }
+
+        return result.release();
+    }
+
+#if EIGEN_VERSION_AT_LEAST(3, 4, 0)
+
+    static constexpr bool needs_writeable = !std::is_const<typename std::remove_pointer<
+        typename get_storage_pointer_type<MapType>::SPT>::type>::value;
+#else
+    // Handle Eigen bug
+    static constexpr bool needs_writeable = !std::is_const<Type>::value;
+#endif
+
+protected:
+    // TODO: Move to std::optional once std::optional has more support
+    std::unique_ptr<MapType> value;
+
+public:
+    // return_descr forces the use of NDArray instead of ArrayLike since refs can only reference
+    // arrays
+    static constexpr auto name
+        = return_descr(get_tensor_descriptor<Type, true, needs_writeable>::value);
+    explicit operator MapType *() { return value.get(); }
+    explicit operator MapType &() { return *value; }
+    explicit operator MapType &&() && { return std::move(*value); }
+
+    template <typename T_>
+    using cast_op_type = ::pybind11::detail::movable_cast_op_type<T_>;
+};
+
+PYBIND11_NAMESPACE_END(detail)
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/embed.h

@@ -0,0 +1,320 @@
+/*
+    pybind11/embed.h: Support for embedding the interpreter
+
+    Copyright (c) 2017 Wenzel Jakob <wenzel.jakob@epfl.ch>
+
+    All rights reserved. Use of this source code is governed by a
+    BSD-style license that can be found in the LICENSE file.
+*/
+
+#pragma once
+
+#include "pybind11.h"
+#include "eval.h"
+
+#include <memory>
+#include <vector>
+
+#if defined(PYPY_VERSION)
+#    error Embedding the interpreter is not supported with PyPy
+#endif
+
+#define PYBIND11_EMBEDDED_MODULE_IMPL(name)                                                       \
+    extern "C" PyObject *pybind11_init_impl_##name();                                             \
+    extern "C" PyObject *pybind11_init_impl_##name() { return pybind11_init_wrapper_##name(); }
+
+/** \rst
+    Add a new module to the table of builtins for the interpreter. Must be
+    defined in global scope. The first macro parameter is the name of the
+    module (without quotes). The second parameter is the variable which will
+    be used as the interface to add functions and classes to the module.
+
+    .. code-block:: cpp
+
+        PYBIND11_EMBEDDED_MODULE(example, m) {
+            // ... initialize functions and classes here
+            m.def("foo", []() {
+                return "Hello, World!";
+            });
+        }
+
+    The third and subsequent macro arguments are optional, and can be used to
+    mark the module as supporting various Python features.
+
+    - ``mod_gil_not_used()``
+    - ``multiple_interpreters::per_interpreter_gil()``
+    - ``multiple_interpreters::shared_gil()``
+    - ``multiple_interpreters::not_supported()``
+
+    .. code-block:: cpp
+
+        PYBIND11_EMBEDDED_MODULE(example, m, py::mod_gil_not_used()) {
+            m.def("foo", []() {
+                return "Hello, Free-threaded World!";
+            });
+        }
+
+ \endrst */
+PYBIND11_WARNING_PUSH
+PYBIND11_WARNING_DISABLE_CLANG("-Wgnu-zero-variadic-macro-arguments")
+#define PYBIND11_EMBEDDED_MODULE(name, variable, ...)                                             \
+    PYBIND11_MODULE_PYINIT(name, ##__VA_ARGS__)                                                   \
+    ::pybind11::detail::embedded_module PYBIND11_CONCAT(pybind11_module_, name)(                  \
+        PYBIND11_TOSTRING(name), PYBIND11_CONCAT(PyInit_, name));                                 \
+    PYBIND11_MODULE_EXEC(name, variable)
+PYBIND11_WARNING_POP
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+/// Python 2.7/3.x compatible version of `PyImport_AppendInittab` and error checks.
+struct embedded_module {
+    using init_t = PyObject *(*) ();
+    embedded_module(const char *name, init_t init) {
+        if (Py_IsInitialized() != 0) {
+            pybind11_fail("Can't add new modules after the interpreter has been initialized");
+        }
+
+        auto result = PyImport_AppendInittab(name, init);
+        if (result == -1) {
+            pybind11_fail("Insufficient memory to add a new module");
+        }
+    }
+};
+
+struct wide_char_arg_deleter {
+    void operator()(wchar_t *ptr) const {
+        // API docs: https://docs.python.org/3/c-api/sys.html#c.Py_DecodeLocale
+        PyMem_RawFree(ptr);
+    }
+};
+
+inline wchar_t *widen_chars(const char *safe_arg) {
+    wchar_t *widened_arg = Py_DecodeLocale(safe_arg, nullptr);
+    return widened_arg;
+}
+
+inline void precheck_interpreter() {
+    if (Py_IsInitialized() != 0) {
+        pybind11_fail("The interpreter is already running");
+    }
+}
+
+#if !defined(PYBIND11_PYCONFIG_SUPPORT_PY_VERSION_HEX)
+#    define PYBIND11_PYCONFIG_SUPPORT_PY_VERSION_HEX (0x03080000)
+#endif
+
+#if PY_VERSION_HEX < PYBIND11_PYCONFIG_SUPPORT_PY_VERSION_HEX
+inline void initialize_interpreter_pre_pyconfig(bool init_signal_handlers,
+                                                int argc,
+                                                const char *const *argv,
+                                                bool add_program_dir_to_path) {
+    detail::precheck_interpreter();
+    Py_InitializeEx(init_signal_handlers ? 1 : 0);
+
+    auto argv_size = static_cast<size_t>(argc);
+    // SetArgv* on python 3 takes wchar_t, so we have to convert.
+    std::unique_ptr<wchar_t *[]> widened_argv(new wchar_t *[argv_size]);
+    std::vector<std::unique_ptr<wchar_t[], detail::wide_char_arg_deleter>> widened_argv_entries;
+    widened_argv_entries.reserve(argv_size);
+    for (size_t ii = 0; ii < argv_size; ++ii) {
+        widened_argv_entries.emplace_back(detail::widen_chars(argv[ii]));
+        if (!widened_argv_entries.back()) {
+            // A null here indicates a character-encoding failure or the python
+            // interpreter out of memory. Give up.
+            return;
+        }
+        widened_argv[ii] = widened_argv_entries.back().get();
+    }
+
+    auto *pysys_argv = widened_argv.get();
+
+    PySys_SetArgvEx(argc, pysys_argv, static_cast<int>(add_program_dir_to_path));
+}
+#endif
+
+PYBIND11_NAMESPACE_END(detail)
+
+#if PY_VERSION_HEX >= PYBIND11_PYCONFIG_SUPPORT_PY_VERSION_HEX
+inline void initialize_interpreter(PyConfig *config,
+                                   int argc = 0,
+                                   const char *const *argv = nullptr,
+                                   bool add_program_dir_to_path = true) {
+    detail::precheck_interpreter();
+    PyStatus status = PyConfig_SetBytesArgv(config, argc, const_cast<char *const *>(argv));
+    if (PyStatus_Exception(status) != 0) {
+        // A failure here indicates a character-encoding failure or the python
+        // interpreter out of memory. Give up.
+        PyConfig_Clear(config);
+        throw std::runtime_error(PyStatus_IsError(status) != 0 ? status.err_msg
+                                                               : "Failed to prepare CPython");
+    }
+    status = Py_InitializeFromConfig(config);
+    if (PyStatus_Exception(status) != 0) {
+        PyConfig_Clear(config);
+        throw std::runtime_error(PyStatus_IsError(status) != 0 ? status.err_msg
+                                                               : "Failed to init CPython");
+    }
+    if (add_program_dir_to_path) {
+        PyRun_SimpleString("import sys, os.path; "
+                           "sys.path.insert(0, "
+                           "os.path.abspath(os.path.dirname(sys.argv[0])) "
+                           "if sys.argv and os.path.exists(sys.argv[0]) else '')");
+    }
+    PyConfig_Clear(config);
+}
+#endif
+
+/** \rst
+    Initialize the Python interpreter. No other pybind11 or CPython API functions can be
+    called before this is done; with the exception of `PYBIND11_EMBEDDED_MODULE`. The
+    optional `init_signal_handlers` parameter can be used to skip the registration of
+    signal handlers (see the `Python documentation`_ for details). Calling this function
+    again after the interpreter has already been initialized is a fatal error.
+
+    If initializing the Python interpreter fails, then the program is terminated.  (This
+    is controlled by the CPython runtime and is an exception to pybind11's normal behavior
+    of throwing exceptions on errors.)
+
+    The remaining optional parameters, `argc`, `argv`, and `add_program_dir_to_path` are
+    used to populate ``sys.argv`` and ``sys.path``.
+    See the |PySys_SetArgvEx documentation|_ for details.
+
+    .. _Python documentation: https://docs.python.org/3/c-api/init.html#c.Py_InitializeEx
+    .. |PySys_SetArgvEx documentation| replace:: ``PySys_SetArgvEx`` documentation
+    .. _PySys_SetArgvEx documentation: https://docs.python.org/3/c-api/init.html#c.PySys_SetArgvEx
+ \endrst */
+inline void initialize_interpreter(bool init_signal_handlers = true,
+                                   int argc = 0,
+                                   const char *const *argv = nullptr,
+                                   bool add_program_dir_to_path = true) {
+#if PY_VERSION_HEX < PYBIND11_PYCONFIG_SUPPORT_PY_VERSION_HEX
+    detail::initialize_interpreter_pre_pyconfig(
+        init_signal_handlers, argc, argv, add_program_dir_to_path);
+#else
+    PyConfig config;
+    PyConfig_InitPythonConfig(&config);
+    // See PR #4473 for background
+    config.parse_argv = 0;
+
+    config.install_signal_handlers = init_signal_handlers ? 1 : 0;
+    initialize_interpreter(&config, argc, argv, add_program_dir_to_path);
+#endif
+
+    // There is exactly one interpreter alive currently.
+    detail::has_seen_non_main_interpreter() = false;
+}
+
+/** \rst
+    Shut down the Python interpreter. No pybind11 or CPython API functions can be called
+    after this. In addition, pybind11 objects must not outlive the interpreter:
+
+    .. code-block:: cpp
+
+        { // BAD
+            py::initialize_interpreter();
+            auto hello = py::str("Hello, World!");
+            py::finalize_interpreter();
+        } // <-- BOOM, hello's destructor is called after interpreter shutdown
+
+        { // GOOD
+            py::initialize_interpreter();
+            { // scoped
+                auto hello = py::str("Hello, World!");
+            } // <-- OK, hello is cleaned up properly
+            py::finalize_interpreter();
+        }
+
+        { // BETTER
+            py::scoped_interpreter guard{};
+            auto hello = py::str("Hello, World!");
+        }
+
+    .. warning::
+
+        The interpreter can be restarted by calling `initialize_interpreter` again.
+        Modules created using pybind11 can be safely re-initialized. However, Python
+        itself cannot completely unload binary extension modules and there are several
+        caveats with regard to interpreter restarting. All the details can be found
+        in the CPython documentation. In short, not all interpreter memory may be
+        freed, either due to reference cycles or user-created global data.
+
+ \endrst */
+inline void finalize_interpreter() {
+    // get rid of any thread-local interpreter cache that currently exists
+    if (detail::has_seen_non_main_interpreter()) {
+        detail::get_internals_pp_manager().unref();
+        detail::get_local_internals_pp_manager().unref();
+
+        // We know there can be no other interpreter alive now
+        detail::has_seen_non_main_interpreter() = false;
+    }
+
+    // Re-fetch the internals pointer-to-pointer (but not the internals itself, which might not
+    // exist). It's possible for the  internals to be created during Py_Finalize() (e.g. if a
+    // py::capsule calls `get_internals()` during destruction), so we get the pointer-pointer here
+    // and check it after Py_Finalize().
+    detail::get_internals_pp_manager().get_pp();
+    detail::get_local_internals_pp_manager().get_pp();
+
+    Py_Finalize();
+
+    detail::get_internals_pp_manager().destroy();
+
+    // Local internals contains data managed by the current interpreter, so we must clear them to
+    // avoid undefined behaviors when initializing another interpreter
+    detail::get_local_internals_pp_manager().destroy();
+
+    // We know there is no interpreter alive now, so we can reset the multi-flag
+    detail::has_seen_non_main_interpreter() = false;
+}
+
+/** \rst
+    Scope guard version of `initialize_interpreter` and `finalize_interpreter`.
+    This a move-only guard and only a single instance can exist.
+
+    See `initialize_interpreter` for a discussion of its constructor arguments.
+
+    .. code-block:: cpp
+
+        #include <pybind11/embed.h>
+
+        int main() {
+            py::scoped_interpreter guard{};
+            py::print(Hello, World!);
+        } // <-- interpreter shutdown
+ \endrst */
+class scoped_interpreter {
+public:
+    explicit scoped_interpreter(bool init_signal_handlers = true,
+                                int argc = 0,
+                                const char *const *argv = nullptr,
+                                bool add_program_dir_to_path = true) {
+        initialize_interpreter(init_signal_handlers, argc, argv, add_program_dir_to_path);
+    }
+
+#if PY_VERSION_HEX >= PYBIND11_PYCONFIG_SUPPORT_PY_VERSION_HEX
+    explicit scoped_interpreter(PyConfig *config,
+                                int argc = 0,
+                                const char *const *argv = nullptr,
+                                bool add_program_dir_to_path = true) {
+        initialize_interpreter(config, argc, argv, add_program_dir_to_path);
+    }
+#endif
+
+    scoped_interpreter(const scoped_interpreter &) = delete;
+    scoped_interpreter(scoped_interpreter &&other) noexcept { other.is_valid = false; }
+    scoped_interpreter &operator=(const scoped_interpreter &) = delete;
+    scoped_interpreter &operator=(scoped_interpreter &&) = delete;
+
+    ~scoped_interpreter() {
+        if (is_valid) {
+            finalize_interpreter();
+        }
+    }
+
+private:
+    bool is_valid = true;
+};
+
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/eval.h

@@ -0,0 +1,161 @@
+/*
+    pybind11/eval.h: Support for evaluating Python expressions and statements
+    from strings and files
+
+    Copyright (c) 2016 Klemens Morgenstern <klemens.morgenstern@ed-chemnitz.de> and
+                       Wenzel Jakob <wenzel.jakob@epfl.ch>
+
+    All rights reserved. Use of this source code is governed by a
+    BSD-style license that can be found in the LICENSE file.
+*/
+
+#pragma once
+
+#include "pybind11.h"
+
+#include <utility>
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+inline void ensure_builtins_in_globals(object &global) {
+#if defined(PYPY_VERSION)
+    // Running exec and eval adds `builtins` module under `__builtins__` key to
+    // globals if not yet present.  Python 3.8 made PyRun_String behave
+    // similarly. Let's also do that for older versions, for consistency. This
+    // was missing from PyPy3.8 7.3.7.
+    if (!global.contains("__builtins__"))
+        global["__builtins__"] = module_::import(PYBIND11_BUILTINS_MODULE);
+#else
+    (void) global;
+#endif
+}
+
+PYBIND11_NAMESPACE_END(detail)
+
+enum eval_mode {
+    /// Evaluate a string containing an isolated expression
+    eval_expr,
+
+    /// Evaluate a string containing a single statement. Returns \c none
+    eval_single_statement,
+
+    /// Evaluate a string containing a sequence of statement. Returns \c none
+    eval_statements
+};
+
+template <eval_mode mode = eval_expr>
+object eval(const str &expr, object global = globals(), object local = object()) {
+    if (!local) {
+        local = global;
+    }
+
+    detail::ensure_builtins_in_globals(global);
+
+    /* PyRun_String does not accept a PyObject / encoding specifier,
+       this seems to be the only alternative */
+    std::string buffer = "# -*- coding: utf-8 -*-\n" + (std::string) expr;
+
+    int start = 0;
+    switch (mode) {
+        case eval_expr:
+            start = Py_eval_input;
+            break;
+        case eval_single_statement:
+            start = Py_single_input;
+            break;
+        case eval_statements:
+            start = Py_file_input;
+            break;
+        default:
+            pybind11_fail("invalid evaluation mode");
+    }
+
+    PyObject *result = PyRun_String(buffer.c_str(), start, global.ptr(), local.ptr());
+    if (!result) {
+        throw error_already_set();
+    }
+    return reinterpret_steal<object>(result);
+}
+
+template <eval_mode mode = eval_expr, size_t N>
+object eval(const char (&s)[N], object global = globals(), object local = object()) {
+    /* Support raw string literals by removing common leading whitespace */
+    auto expr = (s[0] == '\n') ? str(module_::import("textwrap").attr("dedent")(s)) : str(s);
+    return eval<mode>(expr, std::move(global), std::move(local));
+}
+
+inline void exec(const str &expr, object global = globals(), object local = object()) {
+    eval<eval_statements>(expr, std::move(global), std::move(local));
+}
+
+template <size_t N>
+void exec(const char (&s)[N], object global = globals(), object local = object()) {
+    eval<eval_statements>(s, std::move(global), std::move(local));
+}
+
+#if defined(PYPY_VERSION) || defined(GRAALVM_PYTHON)
+template <eval_mode mode = eval_statements>
+object eval_file(str, object, object) {
+    pybind11_fail("eval_file not supported in this interpreter. Use eval");
+}
+template <eval_mode mode = eval_statements>
+object eval_file(str, object) {
+    pybind11_fail("eval_file not supported in this interpreter. Use eval");
+}
+template <eval_mode mode = eval_statements>
+object eval_file(str) {
+    pybind11_fail("eval_file not supported in this interpreter. Use eval");
+}
+#else
+template <eval_mode mode = eval_statements>
+object eval_file(str fname, object global = globals(), object local = object()) {
+    if (!local) {
+        local = global;
+    }
+
+    detail::ensure_builtins_in_globals(global);
+
+    int start = 0;
+    switch (mode) {
+        case eval_expr:
+            start = Py_eval_input;
+            break;
+        case eval_single_statement:
+            start = Py_single_input;
+            break;
+        case eval_statements:
+            start = Py_file_input;
+            break;
+        default:
+            pybind11_fail("invalid evaluation mode");
+    }
+
+    int closeFile = 1;
+    std::string fname_str = (std::string) fname;
+    FILE *f =
+#    if PY_VERSION_HEX >= 0x030E0000
+        Py_fopen(fname.ptr(), "r");
+#    else
+        _Py_fopen_obj(fname.ptr(), "r");
+#    endif
+    if (!f) {
+        PyErr_Clear();
+        pybind11_fail("File \"" + fname_str + "\" could not be opened!");
+    }
+
+    if (!global.contains("__file__")) {
+        global["__file__"] = std::move(fname);
+    }
+
+    PyObject *result
+        = PyRun_FileEx(f, fname_str.c_str(), start, global.ptr(), local.ptr(), closeFile);
+
+    if (!result) {
+        throw error_already_set();
+    }
+    return reinterpret_steal<object>(result);
+}
+#endif
+
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/functional.h

@@ -0,0 +1,146 @@
+/*
+    pybind11/functional.h: std::function<> support
+
+    Copyright (c) 2016 Wenzel Jakob <wenzel.jakob@epfl.ch>
+
+    All rights reserved. Use of this source code is governed by a
+    BSD-style license that can be found in the LICENSE file.
+*/
+
+#pragma once
+
+#include "pybind11.h"
+
+#include <functional>
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+PYBIND11_NAMESPACE_BEGIN(detail)
+PYBIND11_NAMESPACE_BEGIN(type_caster_std_function_specializations)
+
+// ensure GIL is held during functor destruction
+struct func_handle {
+    function f;
+#if !(defined(_MSC_VER) && _MSC_VER == 1916 && defined(PYBIND11_CPP17))
+    // This triggers a syntax error under very special conditions (very weird indeed).
+    explicit
+#endif
+        func_handle(function &&f_) noexcept
+        : f(std::move(f_)) {
+    }
+    func_handle(const func_handle &f_) { operator=(f_); }
+    func_handle &operator=(const func_handle &f_) {
+        gil_scoped_acquire acq;
+        f = f_.f;
+        return *this;
+    }
+    ~func_handle() {
+        gil_scoped_acquire acq;
+        function kill_f(std::move(f));
+    }
+};
+
+// to emulate 'move initialization capture' in C++11
+struct func_wrapper_base {
+    func_handle hfunc;
+    explicit func_wrapper_base(func_handle &&hf) noexcept : hfunc(hf) {}
+};
+
+template <typename Return, typename... Args>
+struct func_wrapper : func_wrapper_base {
+    using func_wrapper_base::func_wrapper_base;
+    Return operator()(Args... args) const { // NOLINT(performance-unnecessary-value-param)
+        gil_scoped_acquire acq;
+        // casts the returned object as a rvalue to the return type
+        return hfunc.f(std::forward<Args>(args)...).template cast<Return>();
+    }
+};
+
+PYBIND11_NAMESPACE_END(type_caster_std_function_specializations)
+
+template <typename Return, typename... Args>
+struct type_caster<std::function<Return(Args...)>> {
+    using type = std::function<Return(Args...)>;
+    using retval_type = conditional_t<std::is_same<Return, void>::value, void_type, Return>;
+    using function_type = Return (*)(Args...);
+
+public:
+    bool load(handle src, bool convert) {
+        if (src.is_none()) {
+            // Defer accepting None to other overloads (if we aren't in convert mode):
+            if (!convert) {
+                return false;
+            }
+            return true;
+        }
+
+        if (!isinstance<function>(src)) {
+            return false;
+        }
+
+        auto func = reinterpret_borrow<function>(src);
+
+        /*
+           When passing a C++ function as an argument to another C++
+           function via Python, every function call would normally involve
+           a full C++ -> Python -> C++ roundtrip, which can be prohibitive.
+           Here, we try to at least detect the case where the function is
+           stateless (i.e. function pointer or lambda function without
+           captured variables), in which case the roundtrip can be avoided.
+         */
+        if (auto cfunc = func.cpp_function()) {
+            auto *cfunc_self = PyCFunction_GET_SELF(cfunc.ptr());
+            if (cfunc_self == nullptr) {
+                PyErr_Clear();
+            } else {
+                function_record *rec = function_record_ptr_from_PyObject(cfunc_self);
+                while (rec != nullptr) {
+                    if (rec->is_stateless
+                        && same_type(typeid(function_type),
+                                     *reinterpret_cast<const std::type_info *>(rec->data[1]))) {
+                        struct capture {
+                            function_type f;
+
+                            static capture *from_data(void **data) {
+                                return PYBIND11_STD_LAUNDER(reinterpret_cast<capture *>(data));
+                            }
+                        };
+                        PYBIND11_ENSURE_PRECONDITION_FOR_FUNCTIONAL_H_PERFORMANCE_OPTIMIZATIONS(
+                            std::is_standard_layout<capture>::value);
+                        value = capture::from_data(rec->data)->f;
+                        return true;
+                    }
+                    rec = rec->next;
+                }
+            }
+            // PYPY segfaults here when passing builtin function like sum.
+            // Raising an fail exception here works to prevent the segfault, but only on gcc.
+            // See PR #1413 for full details
+        }
+
+        value = type_caster_std_function_specializations::func_wrapper<Return, Args...>(
+            type_caster_std_function_specializations::func_handle(std::move(func)));
+        return true;
+    }
+
+    template <typename Func>
+    static handle cast(Func &&f_, return_value_policy policy, handle /* parent */) {
+        if (!f_) {
+            return none().release();
+        }
+
+        auto result = f_.template target<function_type>();
+        if (result) {
+            return cpp_function(*result, policy).release();
+        }
+        return cpp_function(std::forward<Func>(f_), policy).release();
+    }
+
+    PYBIND11_TYPE_CASTER(
+        type,
+        const_name("collections.abc.Callable[[")
+            + ::pybind11::detail::concat(::pybind11::detail::inv_descr(make_caster<Args>::name)...)
+            + const_name("], ") + make_caster<retval_type>::name + const_name("]"));
+};
+
+PYBIND11_NAMESPACE_END(detail)
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/gil.h

@@ -0,0 +1,203 @@
+/*
+    pybind11/gil.h: RAII helpers for managing the GIL
+
+    Copyright (c) 2016 Wenzel Jakob <wenzel.jakob@epfl.ch>
+
+    All rights reserved. Use of this source code is governed by a
+    BSD-style license that can be found in the LICENSE file.
+*/
+
+#pragma once
+
+#if defined(PYBIND11_SIMPLE_GIL_MANAGEMENT)
+
+#    include "detail/common.h"
+#    include "gil_simple.h"
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+
+using gil_scoped_acquire = gil_scoped_acquire_simple;
+using gil_scoped_release = gil_scoped_release_simple;
+
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)
+
+#else
+
+#    include "detail/common.h"
+#    include "detail/internals.h"
+
+#    include <cassert>
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+PYBIND11_WARNING_PUSH
+PYBIND11_WARNING_DISABLE_GCC("-Wredundant-decls")
+
+// forward declarations
+PyThreadState *get_thread_state_unchecked();
+
+PYBIND11_WARNING_POP
+
+PYBIND11_NAMESPACE_END(detail)
+
+/* The functions below essentially reproduce the PyGILState_* API using a RAII
+ * pattern, but there are a few important differences:
+ *
+ * 1. When acquiring the GIL from an non-main thread during the finalization
+ *    phase, the GILState API blindly terminates the calling thread, which
+ *    is often not what is wanted. This API does not do this.
+ *
+ * 2. The gil_scoped_release function can optionally cut the relationship
+ *    of a PyThreadState and its associated thread, which allows moving it to
+ *    another thread (this is a fairly rare/advanced use case).
+ *
+ * 3. The reference count of an acquired thread state can be controlled. This
+ *    can be handy to prevent cases where callbacks issued from an external
+ *    thread would otherwise constantly construct and destroy thread state data
+ *    structures.
+ *
+ * See the Python bindings of NanoGUI (http://github.com/wjakob/nanogui) for an
+ * example which uses features 2 and 3 to migrate the Python thread of
+ * execution to another thread (to run the event loop on the original thread,
+ * in this case).
+ */
+
+class gil_scoped_acquire {
+public:
+    PYBIND11_NOINLINE gil_scoped_acquire() {
+        auto &internals = detail::get_internals();
+        tstate = internals.tstate.get();
+
+        if (!tstate) {
+            /* Check if the GIL was acquired using the PyGILState_* API instead (e.g. if
+               calling from a Python thread). Since we use a different key, this ensures
+               we don't create a new thread state and deadlock in PyEval_AcquireThread
+               below. Note we don't save this state with internals.tstate, since we don't
+               create it we would fail to clear it (its reference count should be > 0). */
+            tstate = PyGILState_GetThisThreadState();
+        }
+
+        if (!tstate) {
+            tstate = PyThreadState_New(internals.istate);
+#    if defined(PYBIND11_DETAILED_ERROR_MESSAGES)
+            if (!tstate) {
+                pybind11_fail("scoped_acquire: could not create thread state!");
+            }
+#    endif
+            tstate->gilstate_counter = 0;
+            internals.tstate = tstate;
+        } else {
+            release = detail::get_thread_state_unchecked() != tstate;
+        }
+
+        if (release) {
+            PyEval_AcquireThread(tstate);
+        }
+
+        inc_ref();
+    }
+
+    gil_scoped_acquire(const gil_scoped_acquire &) = delete;
+    gil_scoped_acquire &operator=(const gil_scoped_acquire &) = delete;
+
+    void inc_ref() { ++tstate->gilstate_counter; }
+
+    PYBIND11_NOINLINE void dec_ref() {
+        --tstate->gilstate_counter;
+#    if defined(PYBIND11_DETAILED_ERROR_MESSAGES)
+        if (detail::get_thread_state_unchecked() != tstate) {
+            pybind11_fail("scoped_acquire::dec_ref(): thread state must be current!");
+        }
+        if (tstate->gilstate_counter < 0) {
+            pybind11_fail("scoped_acquire::dec_ref(): reference count underflow!");
+        }
+#    endif
+        if (tstate->gilstate_counter == 0) {
+#    if defined(PYBIND11_DETAILED_ERROR_MESSAGES)
+            if (!release) {
+                pybind11_fail("scoped_acquire::dec_ref(): internal error!");
+            }
+#    endif
+            // Make sure that PyThreadState_Clear is not recursively called by finalizers.
+            // See issue #5827
+            ++tstate->gilstate_counter;
+            PyThreadState_Clear(tstate);
+            --tstate->gilstate_counter;
+            if (active) {
+                PyThreadState_DeleteCurrent();
+            }
+            detail::get_internals().tstate.reset();
+            release = false;
+        }
+    }
+
+    /// This method will disable the PyThreadState_DeleteCurrent call and the
+    /// GIL won't be released. This method should be used if the interpreter
+    /// could be shutting down when this is called, as thread deletion is not
+    /// allowed during shutdown. Check _Py_IsFinalizing() on Python 3.7+, and
+    /// protect subsequent code.
+    PYBIND11_NOINLINE void disarm() { active = false; }
+
+    PYBIND11_NOINLINE ~gil_scoped_acquire() {
+        dec_ref();
+        if (release) {
+            PyEval_SaveThread();
+        }
+    }
+
+private:
+    PyThreadState *tstate = nullptr;
+    bool release = true;
+    bool active = true;
+};
+
+class gil_scoped_release {
+public:
+    // PRECONDITION: The GIL must be held when this constructor is called.
+    explicit gil_scoped_release(bool disassoc = false) : disassoc(disassoc) {
+        assert(PyGILState_Check());
+        // `get_internals()` must be called here unconditionally in order to initialize
+        // `internals.tstate` for subsequent `gil_scoped_acquire` calls. Otherwise, an
+        // initialization race could occur as multiple threads try `gil_scoped_acquire`.
+        auto &internals = detail::get_internals();
+        // NOLINTNEXTLINE(cppcoreguidelines-prefer-member-initializer)
+        tstate = PyEval_SaveThread();
+        if (disassoc) {
+            internals.tstate.reset();
+        }
+    }
+
+    gil_scoped_release(const gil_scoped_release &) = delete;
+    gil_scoped_release &operator=(const gil_scoped_release &) = delete;
+
+    /// This method will disable the PyThreadState_DeleteCurrent call and the
+    /// GIL won't be acquired. This method should be used if the interpreter
+    /// could be shutting down when this is called, as thread deletion is not
+    /// allowed during shutdown. Check _Py_IsFinalizing() on Python 3.7+, and
+    /// protect subsequent code.
+    PYBIND11_NOINLINE void disarm() { active = false; }
+
+    ~gil_scoped_release() {
+        if (!tstate) {
+            return;
+        }
+        // `PyEval_RestoreThread()` should not be called if runtime is finalizing
+        if (active) {
+            PyEval_RestoreThread(tstate);
+        }
+        if (disassoc) {
+            detail::get_internals().tstate = tstate;
+        }
+    }
+
+private:
+    PyThreadState *tstate;
+    bool disassoc;
+    bool active = true;
+};
+
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)
+
+#endif // !PYBIND11_SIMPLE_GIL_MANAGEMENT

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/gil_safe_call_once.h

@@ -0,0 +1,273 @@
+// Copyright (c) 2023 The pybind Community.
+
+#pragma once
+
+#include "detail/common.h"
+#include "detail/internals.h"
+#include "gil.h"
+
+#include <cassert>
+#include <mutex>
+
+#if defined(Py_GIL_DISABLED) || defined(PYBIND11_HAS_SUBINTERPRETER_SUPPORT)
+#    include <atomic>
+#endif
+#ifdef PYBIND11_HAS_SUBINTERPRETER_SUPPORT
+#    include <cstdint>
+#    include <memory>
+#    include <string>
+#endif
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+
+PYBIND11_NAMESPACE_BEGIN(detail)
+#if defined(Py_GIL_DISABLED) || defined(PYBIND11_HAS_SUBINTERPRETER_SUPPORT)
+using atomic_bool = std::atomic_bool;
+#else
+using atomic_bool = bool;
+#endif
+PYBIND11_NAMESPACE_END(detail)
+
+// Use the `gil_safe_call_once_and_store` class below instead of the naive
+//
+//   static auto imported_obj = py::module_::import("module_name"); // BAD, DO NOT USE!
+//
+// which has two serious issues:
+//
+//     1. Py_DECREF() calls potentially after the Python interpreter was finalized already, and
+//     2. deadlocks in multi-threaded processes (because of missing lock ordering).
+//
+// The following alternative avoids both problems:
+//
+//   PYBIND11_CONSTINIT static py::gil_safe_call_once_and_store<py::object> storage;
+//   auto &imported_obj = storage // Do NOT make this `static`!
+//       .call_once_and_store_result([]() {
+//           return py::module_::import("module_name");
+//       })
+//       .get_stored();
+//
+// The parameter of `call_once_and_store_result()` must be callable. It can make
+// CPython API calls, and in particular, it can temporarily release the GIL.
+//
+// `T` can be any C++ type, it does not have to involve CPython API types.
+//
+// The behavior with regard to signals, e.g. `SIGINT` (`KeyboardInterrupt`),
+// is not ideal. If the main thread is the one to actually run the `Callable`,
+// then a `KeyboardInterrupt` will interrupt it if it is running normal Python
+// code. The situation is different if a non-main thread runs the
+// `Callable`, and then the main thread starts waiting for it to complete:
+// a `KeyboardInterrupt` will not interrupt the non-main thread, but it will
+// get processed only when it is the main thread's turn again and it is running
+// normal Python code. However, this will be unnoticeable for quick call-once
+// functions, which is usually the case.
+//
+// For in-depth background, see docs/advanced/deadlock.md
+#ifndef PYBIND11_HAS_SUBINTERPRETER_SUPPORT
+// Subinterpreter support is disabled.
+// In this case, we can store the result globally, because there is only a single interpreter.
+//
+// The life span of the stored result is the entire process lifetime. It is leaked on process
+// termination to avoid destructor calls after the Python interpreter was finalized.
+template <typename T>
+class gil_safe_call_once_and_store {
+public:
+    // PRECONDITION: The GIL must be held when `call_once_and_store_result()` is called.
+    //
+    // NOTE: The second parameter (finalize callback) is intentionally unused when subinterpreter
+    // support is disabled. In that case, storage is process-global and intentionally leaked to
+    // avoid calling destructors after the Python interpreter has been finalized.
+    template <typename Callable>
+    gil_safe_call_once_and_store &call_once_and_store_result(Callable &&fn,
+                                                             void (*)(T &) /*unused*/ = nullptr) {
+        if (!is_initialized_) { // This read is guarded by the GIL.
+            // Multiple threads may enter here, because the GIL is released in the next line and
+            // CPython API calls in the `fn()` call below may release and reacquire the GIL.
+            gil_scoped_release gil_rel; // Needed to establish lock ordering.
+            std::call_once(once_flag_, [&] {
+                // Only one thread will ever enter here.
+                gil_scoped_acquire gil_acq;
+                ::new (storage_) T(fn()); // fn may release, but will reacquire, the GIL.
+                is_initialized_ = true;   // This write is guarded by the GIL.
+            });
+            // All threads will observe `is_initialized_` as true here.
+        }
+        // Intentionally not returning `T &` to ensure the calling code is self-documenting.
+        return *this;
+    }
+
+    // This must only be called after `call_once_and_store_result()` was called.
+    T &get_stored() {
+        assert(is_initialized_);
+        PYBIND11_WARNING_PUSH
+#    if !defined(__clang__) && defined(__GNUC__) && __GNUC__ < 5
+        // Needed for gcc 4.8.5
+        PYBIND11_WARNING_DISABLE_GCC("-Wstrict-aliasing")
+#    endif
+        return *reinterpret_cast<T *>(storage_);
+        PYBIND11_WARNING_POP
+    }
+
+    constexpr gil_safe_call_once_and_store() = default;
+    // The instance is a global static, so its destructor runs when the process
+    // is terminating. Therefore, do nothing here because the Python interpreter
+    // may have been finalized already.
+    PYBIND11_DTOR_CONSTEXPR ~gil_safe_call_once_and_store() = default;
+
+    // Disable copy and move operations.
+    gil_safe_call_once_and_store(const gil_safe_call_once_and_store &) = delete;
+    gil_safe_call_once_and_store(gil_safe_call_once_and_store &&) = delete;
+    gil_safe_call_once_and_store &operator=(const gil_safe_call_once_and_store &) = delete;
+    gil_safe_call_once_and_store &operator=(gil_safe_call_once_and_store &&) = delete;
+
+private:
+    // The global static storage (per-process) when subinterpreter support is disabled.
+    alignas(T) char storage_[sizeof(T)] = {};
+    std::once_flag once_flag_;
+
+    // The `is_initialized_`-`storage_` pair is very similar to `std::optional`,
+    // but the latter does not have the triviality properties of former,
+    // therefore `std::optional` is not a viable alternative here.
+    detail::atomic_bool is_initialized_{false};
+};
+#else
+// Subinterpreter support is enabled.
+// In this case, we should store the result per-interpreter instead of globally, because each
+// subinterpreter has its own separate state. The cached result may not shareable across
+// interpreters (e.g., imported modules and their members).
+
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+template <typename T>
+struct call_once_storage {
+    alignas(T) char storage[sizeof(T)] = {};
+    std::once_flag once_flag;
+    void (*finalize)(T &) = nullptr;
+    std::atomic_bool is_initialized{false};
+
+    call_once_storage() = default;
+    ~call_once_storage() {
+        if (is_initialized) {
+            if (finalize != nullptr) {
+                finalize(*reinterpret_cast<T *>(storage));
+            } else {
+                reinterpret_cast<T *>(storage)->~T();
+            }
+        }
+    }
+    call_once_storage(const call_once_storage &) = delete;
+    call_once_storage(call_once_storage &&) = delete;
+    call_once_storage &operator=(const call_once_storage &) = delete;
+    call_once_storage &operator=(call_once_storage &&) = delete;
+};
+
+PYBIND11_NAMESPACE_END(detail)
+
+// Prefix for storage keys in the interpreter state dict.
+#    define PYBIND11_CALL_ONCE_STORAGE_KEY_PREFIX PYBIND11_INTERNALS_ID "_call_once_storage__"
+
+// The life span of the stored result is the entire interpreter lifetime. An additional
+// `finalize_fn` can be provided to clean up the stored result when the interpreter is destroyed.
+template <typename T>
+class gil_safe_call_once_and_store {
+public:
+    // PRECONDITION: The GIL must be held when `call_once_and_store_result()` is called.
+    template <typename Callable>
+    gil_safe_call_once_and_store &call_once_and_store_result(Callable &&fn,
+                                                             void (*finalize_fn)(T &) = nullptr) {
+        if (!is_last_storage_valid()) {
+            // Multiple threads may enter here, because the GIL is released in the next line and
+            // CPython API calls in the `fn()` call below may release and reacquire the GIL.
+            gil_scoped_release gil_rel; // Needed to establish lock ordering.
+            // There can be multiple threads going through here.
+            storage_type *value = nullptr;
+            {
+                gil_scoped_acquire gil_acq; // Restore lock ordering.
+                // This function is thread-safe under free-threading.
+                value = get_or_create_storage_in_state_dict();
+            }
+            assert(value != nullptr);
+            std::call_once(value->once_flag, [&] {
+                // Only one thread will ever enter here.
+                gil_scoped_acquire gil_acq;
+                // fn may release, but will reacquire, the GIL.
+                ::new (value->storage) T(fn());
+                value->finalize = finalize_fn;
+                value->is_initialized = true;
+                last_storage_ptr_ = reinterpret_cast<T *>(value->storage);
+                is_initialized_by_at_least_one_interpreter_ = true;
+            });
+            // All threads will observe `is_initialized_by_at_least_one_interpreter_` as true here.
+        }
+        // Intentionally not returning `T &` to ensure the calling code is self-documenting.
+        return *this;
+    }
+
+    // This must only be called after `call_once_and_store_result()` was called.
+    T &get_stored() {
+        T *result = last_storage_ptr_;
+        if (!is_last_storage_valid()) {
+            gil_scoped_acquire gil_acq;
+            auto *value = get_or_create_storage_in_state_dict();
+            result = last_storage_ptr_ = reinterpret_cast<T *>(value->storage);
+        }
+        assert(result != nullptr);
+        return *result;
+    }
+
+    constexpr gil_safe_call_once_and_store() = default;
+    // The instance is a global static, so its destructor runs when the process
+    // is terminating. Therefore, do nothing here because the Python interpreter
+    // may have been finalized already.
+    PYBIND11_DTOR_CONSTEXPR ~gil_safe_call_once_and_store() = default;
+
+    // Disable copy and move operations because the memory address is used as key.
+    gil_safe_call_once_and_store(const gil_safe_call_once_and_store &) = delete;
+    gil_safe_call_once_and_store(gil_safe_call_once_and_store &&) = delete;
+    gil_safe_call_once_and_store &operator=(const gil_safe_call_once_and_store &) = delete;
+    gil_safe_call_once_and_store &operator=(gil_safe_call_once_and_store &&) = delete;
+
+private:
+    using storage_type = detail::call_once_storage<T>;
+
+    // Indicator of fast path for single-interpreter case.
+    bool is_last_storage_valid() const {
+        return is_initialized_by_at_least_one_interpreter_
+               && !detail::has_seen_non_main_interpreter();
+    }
+
+    // Get the unique key for this storage instance in the interpreter's state dict.
+    // The return type should not be `py::str` because PyObject is interpreter-dependent.
+    std::string get_storage_key() const {
+        // The instance is expected to be global static, so using its address as unique identifier.
+        // The typical usage is like:
+        //
+        //   PYBIND11_CONSTINIT static gil_safe_call_once_and_store<T> storage;
+        //
+        return PYBIND11_CALL_ONCE_STORAGE_KEY_PREFIX
+               + std::to_string(reinterpret_cast<std::uintptr_t>(this));
+    }
+
+    // Get or create per-storage capsule in the current interpreter's state dict.
+    // The storage is interpreter-dependent and will not be shared across interpreters.
+    storage_type *get_or_create_storage_in_state_dict() {
+        return detail::atomic_get_or_create_in_state_dict<storage_type>(get_storage_key().c_str())
+            .first;
+    }
+
+    // No storage needed when subinterpreter support is enabled.
+    // The actual storage is stored in the per-interpreter state dict via
+    // `get_or_create_storage_in_state_dict()`.
+
+    // Fast local cache to avoid repeated lookups when there are no multiple interpreters.
+    // This is only valid if there is a single interpreter. Otherwise, it is not used.
+    // WARNING: We cannot use thread local cache similar to `internals_pp_manager::internals_p_tls`
+    //          because the thread local storage cannot be explicitly invalidated when interpreters
+    //          are destroyed (unlike `internals_pp_manager` which has explicit hooks for that).
+    T *last_storage_ptr_ = nullptr;
+    // This flag is true if the value has been initialized by any interpreter (may not be the
+    // current one).
+    detail::atomic_bool is_initialized_by_at_least_one_interpreter_{false};
+};
+#endif
+
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/gil_simple.h

@@ -0,0 +1,37 @@
+// Copyright (c) 2016-2025 The Pybind Development Team.
+// All rights reserved. Use of this source code is governed by a
+// BSD-style license that can be found in the LICENSE file.
+
+#pragma once
+
+#include "detail/common.h"
+
+#include <cassert>
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+
+class gil_scoped_acquire_simple {
+    PyGILState_STATE state;
+
+public:
+    gil_scoped_acquire_simple() : state{PyGILState_Ensure()} {}
+    gil_scoped_acquire_simple(const gil_scoped_acquire_simple &) = delete;
+    gil_scoped_acquire_simple &operator=(const gil_scoped_acquire_simple &) = delete;
+    ~gil_scoped_acquire_simple() { PyGILState_Release(state); }
+};
+
+class gil_scoped_release_simple {
+    PyThreadState *state;
+
+public:
+    // PRECONDITION: The GIL must be held when this constructor is called.
+    gil_scoped_release_simple() {
+        assert(PyGILState_Check());
+        state = PyEval_SaveThread();
+    }
+    gil_scoped_release_simple(const gil_scoped_release_simple &) = delete;
+    gil_scoped_release_simple &operator=(const gil_scoped_release_simple &) = delete;
+    ~gil_scoped_release_simple() { PyEval_RestoreThread(state); }
+};
+
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/iostream.h

@@ -0,0 +1,301 @@
+/*
+    pybind11/iostream.h -- Tools to assist with redirecting cout and cerr to Python
+
+    Copyright (c) 2017 Henry F. Schreiner
+
+    All rights reserved. Use of this source code is governed by a
+    BSD-style license that can be found in the LICENSE file.
+
+    WARNING: The implementation in this file is NOT thread safe. Multiple
+    threads writing to a redirected ostream concurrently cause data races
+    and potentially buffer overflows. Therefore it is currently a requirement
+    that all (possibly) concurrent redirected ostream writes are protected by
+    a mutex.
+    #HelpAppreciated: Work on iostream.h thread safety.
+    For more background see the discussions under
+    https://github.com/pybind/pybind11/pull/2982 and
+    https://github.com/pybind/pybind11/pull/2995.
+*/
+
+#pragma once
+
+#include "pybind11.h"
+
+#include <algorithm>
+#include <cstring>
+#include <iostream>
+#include <iterator>
+#include <memory>
+#include <ostream>
+#include <streambuf>
+#include <string>
+#include <utility>
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+// Buffer that writes to Python instead of C++
+class pythonbuf : public std::streambuf {
+private:
+    using traits_type = std::streambuf::traits_type;
+
+    const size_t buf_size;
+    std::unique_ptr<char[]> d_buffer;
+    object pywrite;
+    object pyflush;
+
+    int overflow(int c) override {
+        if (!traits_type::eq_int_type(c, traits_type::eof())) {
+            *pptr() = traits_type::to_char_type(c);
+            pbump(1);
+        }
+        return sync() == 0 ? traits_type::not_eof(c) : traits_type::eof();
+    }
+
+    // Computes how many bytes at the end of the buffer are part of an
+    // incomplete sequence of UTF-8 bytes.
+    // Precondition: pbase() < pptr()
+    size_t utf8_remainder() const {
+        const auto rbase = std::reverse_iterator<char *>(pbase());
+        const auto rpptr = std::reverse_iterator<char *>(pptr());
+        auto is_ascii = [](char c) { return (static_cast<unsigned char>(c) & 0x80) == 0x00; };
+        auto is_leading = [](char c) { return (static_cast<unsigned char>(c) & 0xC0) == 0xC0; };
+        auto is_leading_2b = [](char c) { return static_cast<unsigned char>(c) <= 0xDF; };
+        auto is_leading_3b = [](char c) { return static_cast<unsigned char>(c) <= 0xEF; };
+        // If the last character is ASCII, there are no incomplete code points
+        if (is_ascii(*rpptr)) {
+            return 0;
+        }
+        // Otherwise, work back from the end of the buffer and find the first
+        // UTF-8 leading byte
+        const auto rpend = rbase - rpptr >= 3 ? rpptr + 3 : rbase;
+        const auto leading = std::find_if(rpptr, rpend, is_leading);
+        if (leading == rbase) {
+            return 0;
+        }
+        const auto dist = static_cast<size_t>(leading - rpptr);
+        size_t remainder = 0;
+
+        if (dist == 0) {
+            remainder = 1; // 1-byte code point is impossible
+        } else if (dist == 1) {
+            remainder = is_leading_2b(*leading) ? 0 : dist + 1;
+        } else if (dist == 2) {
+            remainder = is_leading_3b(*leading) ? 0 : dist + 1;
+        }
+        // else if (dist >= 3), at least 4 bytes before encountering an UTF-8
+        // leading byte, either no remainder or invalid UTF-8.
+        // Invalid UTF-8 will cause an exception later when converting
+        // to a Python string, so that's not handled here.
+        return remainder;
+    }
+
+    // This function must be non-virtual to be called in a destructor.
+    int _sync() {
+        if (pbase() != pptr()) { // If buffer is not empty
+            gil_scoped_acquire tmp;
+            // This subtraction cannot be negative, so dropping the sign.
+            auto size = static_cast<size_t>(pptr() - pbase());
+            size_t remainder = utf8_remainder();
+
+            if (size > remainder) {
+                str line(pbase(), size - remainder);
+                pywrite(std::move(line));
+                pyflush();
+            }
+
+            // Copy the remainder at the end of the buffer to the beginning:
+            if (remainder > 0) {
+                std::memmove(pbase(), pptr() - remainder, remainder);
+            }
+            setp(pbase(), epptr());
+            pbump(static_cast<int>(remainder));
+        }
+        return 0;
+    }
+
+    int sync() override { return _sync(); }
+
+public:
+    // Minimum buffer size must accommodate the largest incomplete UTF-8 sequence
+    // (3 bytes) plus one position reserved for overflow(), i.e. 4 bytes total.
+    static constexpr size_t minimum_buffer_size = 4;
+
+    explicit pythonbuf(const object &pyostream, size_t buffer_size = 1024)
+        : buf_size(buffer_size < minimum_buffer_size // ternary avoids C++14 std::max ODR-use of
+                                                     // static constexpr
+                       ? minimum_buffer_size
+                       : buffer_size),
+          d_buffer(new char[buf_size]), pywrite(pyostream.attr("write")),
+          pyflush(pyostream.attr("flush")) {
+        setp(d_buffer.get(), d_buffer.get() + buf_size - 1);
+    }
+
+    pythonbuf(pythonbuf &&other) noexcept
+        : buf_size(other.buf_size), d_buffer(std::move(other.d_buffer)),
+          pywrite(std::move(other.pywrite)), pyflush(std::move(other.pyflush)) {
+        const auto pending = (other.pbase() != nullptr && other.pptr() != nullptr)
+                                 ? static_cast<int>(other.pptr() - other.pbase())
+                                 : 0;
+        if (d_buffer != nullptr) {
+            // Rebuild the put area from the transferred storage.
+            setp(d_buffer.get(), d_buffer.get() + buf_size - 1);
+            pbump(pending);
+        } else {
+            setp(nullptr, nullptr);
+        }
+        // Prevent the moved-from destructor from flushing through moved-out handles.
+        other.setp(nullptr, nullptr);
+    }
+
+    /// Sync before destroy
+    ~pythonbuf() override { _sync(); }
+};
+
+PYBIND11_NAMESPACE_END(detail)
+
+/** \rst
+    This a move-only guard that redirects output.
+
+    .. code-block:: cpp
+
+        #include <pybind11/iostream.h>
+
+        ...
+
+        {
+            py::scoped_ostream_redirect output;
+            std::cout << "Hello, World!"; // Python stdout
+        } // <-- return std::cout to normal
+
+    You can explicitly pass the c++ stream and the python object,
+    for example to guard stderr instead.
+
+    .. code-block:: cpp
+
+        {
+            py::scoped_ostream_redirect output{
+                std::cerr, py::module::import("sys").attr("stderr")};
+            std::cout << "Hello, World!";
+        }
+ \endrst */
+class scoped_ostream_redirect {
+protected:
+    std::streambuf *old;
+    std::ostream &costream;
+    detail::pythonbuf buffer;
+    bool active = true;
+
+public:
+    explicit scoped_ostream_redirect(std::ostream &costream = std::cout,
+                                     const object &pyostream
+                                     = module_::import("sys").attr("stdout"))
+        : costream(costream), buffer(pyostream) {
+        old = costream.rdbuf(&buffer);
+    }
+
+    ~scoped_ostream_redirect() {
+        if (active) {
+            costream.rdbuf(old);
+        }
+    }
+
+    scoped_ostream_redirect(const scoped_ostream_redirect &) = delete;
+    // NOLINTNEXTLINE(performance-noexcept-move-constructor)
+    scoped_ostream_redirect(scoped_ostream_redirect &&other)
+        : old(other.old), costream(other.costream), buffer(std::move(other.buffer)),
+          active(other.active) {
+        if (active) {
+            costream.rdbuf(&buffer); // Re-point stream to our buffer
+            other.active = false;
+        }
+    }
+    scoped_ostream_redirect &operator=(const scoped_ostream_redirect &) = delete;
+    scoped_ostream_redirect &operator=(scoped_ostream_redirect &&) = delete;
+};
+
+/** \rst
+    Like `scoped_ostream_redirect`, but redirects cerr by default. This class
+    is provided primary to make ``py::call_guard`` easier to make.
+
+    .. code-block:: cpp
+
+     m.def("noisy_func", &noisy_func,
+           py::call_guard<scoped_ostream_redirect,
+                          scoped_estream_redirect>());
+
+\endrst */
+class scoped_estream_redirect : public scoped_ostream_redirect {
+public:
+    explicit scoped_estream_redirect(std::ostream &costream = std::cerr,
+                                     const object &pyostream
+                                     = module_::import("sys").attr("stderr"))
+        : scoped_ostream_redirect(costream, pyostream) {}
+};
+
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+// Class to redirect output as a context manager. C++ backend.
+class OstreamRedirect {
+    bool do_stdout_;
+    bool do_stderr_;
+    std::unique_ptr<scoped_ostream_redirect> redirect_stdout;
+    std::unique_ptr<scoped_estream_redirect> redirect_stderr;
+
+public:
+    explicit OstreamRedirect(bool do_stdout = true, bool do_stderr = true)
+        : do_stdout_(do_stdout), do_stderr_(do_stderr) {}
+
+    void enter() {
+        if (do_stdout_) {
+            redirect_stdout.reset(new scoped_ostream_redirect());
+        }
+        if (do_stderr_) {
+            redirect_stderr.reset(new scoped_estream_redirect());
+        }
+    }
+
+    void exit() {
+        redirect_stdout.reset();
+        redirect_stderr.reset();
+    }
+};
+
+PYBIND11_NAMESPACE_END(detail)
+
+/** \rst
+    This is a helper function to add a C++ redirect context manager to Python
+    instead of using a C++ guard. To use it, add the following to your binding code:
+
+    .. code-block:: cpp
+
+        #include <pybind11/iostream.h>
+
+        ...
+
+        py::add_ostream_redirect(m, "ostream_redirect");
+
+    You now have a Python context manager that redirects your output:
+
+    .. code-block:: python
+
+        with m.ostream_redirect():
+            m.print_to_cout_function()
+
+    This manager can optionally be told which streams to operate on:
+
+    .. code-block:: python
+
+        with m.ostream_redirect(stdout=true, stderr=true):
+            m.noisy_function_with_error_printing()
+
+ \endrst */
+inline class_<detail::OstreamRedirect>
+add_ostream_redirect(module_ m, const std::string &name = "ostream_redirect") {
+    return class_<detail::OstreamRedirect>(std::move(m), name.c_str(), module_local())
+        .def(init<bool, bool>(), arg("stdout") = true, arg("stderr") = true)
+        .def("__enter__", &detail::OstreamRedirect::enter)
+        .def("__exit__", [](detail::OstreamRedirect &self_, const args &) { self_.exit(); });
+}
+
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/native_enum.h

@@ -0,0 +1,76 @@
+// Copyright (c) 2022-2025 The pybind Community.
+// All rights reserved. Use of this source code is governed by a
+// BSD-style license that can be found in the LICENSE file.
+
+#pragma once
+
+#include "detail/common.h"
+#include "detail/native_enum_data.h"
+#include "detail/type_caster_base.h"
+#include "cast.h"
+
+#include <cassert>
+#include <limits>
+#include <type_traits>
+#include <typeindex>
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+
+/// Conversions between Python's native (stdlib) enum types and C++ enums.
+template <typename EnumType>
+class native_enum : public detail::native_enum_data {
+public:
+    using Underlying = typename std::underlying_type<EnumType>::type;
+
+    native_enum(handle parent_scope,
+                const char *name,
+                const char *native_type_name,
+                const char *class_doc = "")
+        : detail::native_enum_data(
+              parent_scope, name, native_type_name, class_doc, make_record()) {
+        if (detail::get_local_type_info(typeid(EnumType)) != nullptr
+            || detail::get_global_type_info(typeid(EnumType)) != nullptr) {
+            pybind11_fail(
+                "pybind11::native_enum<...>(\"" + enum_name_encoded
+                + "\") is already registered as a `pybind11::enum_` or `pybind11::class_`!");
+        }
+        if (detail::global_internals_native_enum_type_map_contains(enum_type_index)) {
+            pybind11_fail("pybind11::native_enum<...>(\"" + enum_name_encoded
+                          + "\") is already registered!");
+        }
+        arm_finalize_check();
+    }
+
+    /// Export enumeration entries into the parent scope
+    native_enum &export_values() {
+        assert(!export_values_flag); // Catch redundant calls.
+        export_values_flag = true;
+        return *this;
+    }
+
+    /// Add an enumeration entry
+    native_enum &value(char const *name, EnumType value, const char *doc = nullptr) {
+        // Disarm for the case that the native_enum_data dtor runs during exception unwinding.
+        disarm_finalize_check("value after finalize");
+        members.append(make_tuple(name, static_cast<Underlying>(value)));
+        if (doc) {
+            member_docs.append(make_tuple(name, doc));
+        }
+        arm_finalize_check(); // There was no exception.
+        return *this;
+    }
+
+    native_enum(const native_enum &) = delete;
+    native_enum &operator=(const native_enum &) = delete;
+
+private:
+    static detail::native_enum_record make_record() {
+        detail::native_enum_record ret;
+        ret.cpptype = &typeid(EnumType);
+        ret.size_bytes = sizeof(EnumType);
+        ret.is_signed = std::is_signed<Underlying>::value;
+        return ret;
+    }
+};
+
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/operators.h

@@ -0,0 +1,202 @@
+/*
+    pybind11/operator.h: Metatemplates for operator overloading
+
+    Copyright (c) 2016 Wenzel Jakob <wenzel.jakob@epfl.ch>
+
+    All rights reserved. Use of this source code is governed by a
+    BSD-style license that can be found in the LICENSE file.
+*/
+
+#pragma once
+
+#include "pybind11.h"
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+/// Enumeration with all supported operator types
+enum op_id : int {
+    op_add,
+    op_sub,
+    op_mul,
+    op_div,
+    op_mod,
+    op_divmod,
+    op_pow,
+    op_lshift,
+    op_rshift,
+    op_and,
+    op_xor,
+    op_or,
+    op_neg,
+    op_pos,
+    op_abs,
+    op_invert,
+    op_int,
+    op_long,
+    op_float,
+    op_str,
+    op_cmp,
+    op_gt,
+    op_ge,
+    op_lt,
+    op_le,
+    op_eq,
+    op_ne,
+    op_iadd,
+    op_isub,
+    op_imul,
+    op_idiv,
+    op_imod,
+    op_ilshift,
+    op_irshift,
+    op_iand,
+    op_ixor,
+    op_ior,
+    op_complex,
+    op_bool,
+    op_nonzero,
+    op_repr,
+    op_truediv,
+    op_itruediv,
+    op_hash
+};
+
+enum op_type : int {
+    op_l, /* base type on left */
+    op_r, /* base type on right */
+    op_u  /* unary operator */
+};
+
+struct self_t {};
+static const self_t self = self_t();
+
+/// Type for an unused type slot
+struct undefined_t {};
+
+/// Don't warn about an unused variable
+inline self_t __self() { return self; }
+
+/// base template of operator implementations
+template <op_id, op_type, typename B, typename L, typename R>
+struct op_impl {};
+
+/// Operator implementation generator
+template <op_id id, op_type ot, typename L, typename R>
+struct op_ {
+    static constexpr bool op_enable_if_hook = true;
+    template <typename Class, typename... Extra>
+    void execute(Class &cl, const Extra &...extra) const {
+        using Base = typename Class::type;
+        using L_type = conditional_t<std::is_same<L, self_t>::value, Base, L>;
+        using R_type = conditional_t<std::is_same<R, self_t>::value, Base, R>;
+        using op = op_impl<id, ot, Base, L_type, R_type>;
+        cl.def(op::name(), &op::execute, is_operator(), extra...);
+    }
+    template <typename Class, typename... Extra>
+    void execute_cast(Class &cl, const Extra &...extra) const {
+        using Base = typename Class::type;
+        using L_type = conditional_t<std::is_same<L, self_t>::value, Base, L>;
+        using R_type = conditional_t<std::is_same<R, self_t>::value, Base, R>;
+        using op = op_impl<id, ot, Base, L_type, R_type>;
+        cl.def(op::name(), &op::execute_cast, is_operator(), extra...);
+    }
+};
+
+#define PYBIND11_BINARY_OPERATOR(id, rid, op, expr)                                               \
+    template <typename B, typename L, typename R>                                                 \
+    struct op_impl<op_##id, op_l, B, L, R> {                                                      \
+        static char const *name() { return "__" #id "__"; }                                       \
+        static auto execute(const L &l, const R &r) -> decltype(expr) { return (expr); }          \
+        static B execute_cast(const L &l, const R &r) { return B(expr); }                         \
+    };                                                                                            \
+    template <typename B, typename L, typename R>                                                 \
+    struct op_impl<op_##id, op_r, B, L, R> {                                                      \
+        static char const *name() { return "__" #rid "__"; }                                      \
+        static auto execute(const R &r, const L &l) -> decltype(expr) { return (expr); }          \
+        static B execute_cast(const R &r, const L &l) { return B(expr); }                         \
+    };                                                                                            \
+    inline op_<op_##id, op_l, self_t, self_t> op(const self_t &, const self_t &) {                \
+        return op_<op_##id, op_l, self_t, self_t>();                                              \
+    }                                                                                             \
+    template <typename T>                                                                         \
+    op_<op_##id, op_l, self_t, T> op(const self_t &, const T &) {                                 \
+        return op_<op_##id, op_l, self_t, T>();                                                   \
+    }                                                                                             \
+    template <typename T>                                                                         \
+    op_<op_##id, op_r, T, self_t> op(const T &, const self_t &) {                                 \
+        return op_<op_##id, op_r, T, self_t>();                                                   \
+    }
+
+#define PYBIND11_INPLACE_OPERATOR(id, op, expr)                                                   \
+    template <typename B, typename L, typename R>                                                 \
+    struct op_impl<op_##id, op_l, B, L, R> {                                                      \
+        static char const *name() { return "__" #id "__"; }                                       \
+        static auto execute(L &l, const R &r) -> decltype(expr) { return expr; }                  \
+        static B execute_cast(L &l, const R &r) { return B(expr); }                               \
+    };                                                                                            \
+    template <typename T>                                                                         \
+    op_<op_##id, op_l, self_t, T> op(const self_t &, const T &) {                                 \
+        return op_<op_##id, op_l, self_t, T>();                                                   \
+    }
+
+#define PYBIND11_UNARY_OPERATOR(id, op, expr)                                                     \
+    template <typename B, typename L>                                                             \
+    struct op_impl<op_##id, op_u, B, L, undefined_t> {                                            \
+        static char const *name() { return "__" #id "__"; }                                       \
+        static auto execute(const L &l) -> decltype(expr) { return expr; }                        \
+        static B execute_cast(const L &l) { return B(expr); }                                     \
+    };                                                                                            \
+    inline op_<op_##id, op_u, self_t, undefined_t> op(const self_t &) {                           \
+        return op_<op_##id, op_u, self_t, undefined_t>();                                         \
+    }
+
+PYBIND11_BINARY_OPERATOR(sub, rsub, operator-, l - r)
+PYBIND11_BINARY_OPERATOR(add, radd, operator+, l + r)
+PYBIND11_BINARY_OPERATOR(mul, rmul, operator*, l *r)
+PYBIND11_BINARY_OPERATOR(truediv, rtruediv, operator/, l / r)
+PYBIND11_BINARY_OPERATOR(mod, rmod, operator%, l % r)
+PYBIND11_BINARY_OPERATOR(lshift, rlshift, operator<<, l << r)
+PYBIND11_BINARY_OPERATOR(rshift, rrshift, operator>>, l >> r)
+PYBIND11_BINARY_OPERATOR(and, rand, operator&, l &r)
+PYBIND11_BINARY_OPERATOR(xor, rxor, operator^, l ^ r)
+PYBIND11_BINARY_OPERATOR(eq, eq, operator==, l == r)
+PYBIND11_BINARY_OPERATOR(ne, ne, operator!=, l != r)
+PYBIND11_BINARY_OPERATOR(or, ror, operator|, l | r)
+PYBIND11_BINARY_OPERATOR(gt, lt, operator>, l > r)
+PYBIND11_BINARY_OPERATOR(ge, le, operator>=, l >= r)
+PYBIND11_BINARY_OPERATOR(lt, gt, operator<, l < r)
+PYBIND11_BINARY_OPERATOR(le, ge, operator<=, l <= r)
+// PYBIND11_BINARY_OPERATOR(pow,       rpow,         pow,          std::pow(l,  r))
+PYBIND11_INPLACE_OPERATOR(iadd, operator+=, l += r)
+PYBIND11_INPLACE_OPERATOR(isub, operator-=, l -= r)
+PYBIND11_INPLACE_OPERATOR(imul, operator*=, l *= r)
+PYBIND11_INPLACE_OPERATOR(itruediv, operator/=, l /= r)
+PYBIND11_INPLACE_OPERATOR(imod, operator%=, l %= r)
+PYBIND11_INPLACE_OPERATOR(ilshift, operator<<=, l <<= r)
+PYBIND11_INPLACE_OPERATOR(irshift, operator>>=, l >>= r)
+PYBIND11_INPLACE_OPERATOR(iand, operator&=, l &= r)
+PYBIND11_INPLACE_OPERATOR(ixor, operator^=, l ^= r)
+PYBIND11_INPLACE_OPERATOR(ior, operator|=, l |= r)
+PYBIND11_UNARY_OPERATOR(neg, operator-, -l)
+PYBIND11_UNARY_OPERATOR(pos, operator+, +l)
+// WARNING: This usage of `abs` should only be done for existing STL overloads.
+// Adding overloads directly in to the `std::` namespace is advised against:
+// https://en.cppreference.com/w/cpp/language/extending_std
+PYBIND11_UNARY_OPERATOR(abs, abs, std::abs(l))
+PYBIND11_UNARY_OPERATOR(hash, hash, std::hash<L>()(l))
+PYBIND11_UNARY_OPERATOR(invert, operator~, (~l))
+PYBIND11_UNARY_OPERATOR(bool, operator!, !!l)
+PYBIND11_UNARY_OPERATOR(int, int_, (int) l)
+PYBIND11_UNARY_OPERATOR(float, float_, (double) l)
+
+#undef PYBIND11_BINARY_OPERATOR
+#undef PYBIND11_INPLACE_OPERATOR
+#undef PYBIND11_UNARY_OPERATOR
+PYBIND11_NAMESPACE_END(detail)
+
+using detail::self;
+// Add named operators so that they are accessible via `py::`.
+using detail::hash;
+
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/options.h

@@ -0,0 +1,92 @@
+/*
+    pybind11/options.h: global settings that are configurable at runtime.
+
+    Copyright (c) 2016 Wenzel Jakob <wenzel.jakob@epfl.ch>
+
+    All rights reserved. Use of this source code is governed by a
+    BSD-style license that can be found in the LICENSE file.
+*/
+
+#pragma once
+
+#include "detail/common.h"
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+
+class options {
+public:
+    // Default RAII constructor, which leaves settings as they currently are.
+    options() : previous_state(global_state()) {}
+
+    // Class is non-copyable.
+    options(const options &) = delete;
+    options &operator=(const options &) = delete;
+
+    // Destructor, which restores settings that were in effect before.
+    ~options() { global_state() = previous_state; }
+
+    // Setter methods (affect the global state):
+
+    options &disable_user_defined_docstrings() & {
+        global_state().show_user_defined_docstrings = false;
+        return *this;
+    }
+
+    options &enable_user_defined_docstrings() & {
+        global_state().show_user_defined_docstrings = true;
+        return *this;
+    }
+
+    options &disable_function_signatures() & {
+        global_state().show_function_signatures = false;
+        return *this;
+    }
+
+    options &enable_function_signatures() & {
+        global_state().show_function_signatures = true;
+        return *this;
+    }
+
+    options &disable_enum_members_docstring() & {
+        global_state().show_enum_members_docstring = false;
+        return *this;
+    }
+
+    options &enable_enum_members_docstring() & {
+        global_state().show_enum_members_docstring = true;
+        return *this;
+    }
+
+    // Getter methods (return the global state):
+
+    static bool show_user_defined_docstrings() {
+        return global_state().show_user_defined_docstrings;
+    }
+
+    static bool show_function_signatures() { return global_state().show_function_signatures; }
+
+    static bool show_enum_members_docstring() {
+        return global_state().show_enum_members_docstring;
+    }
+
+    // This type is not meant to be allocated on the heap.
+    void *operator new(size_t) = delete;
+
+private:
+    struct state {
+        bool show_user_defined_docstrings = true; //< Include user-supplied texts in docstrings.
+        bool show_function_signatures = true;     //< Include auto-generated function signatures
+                                                  //  in docstrings.
+        bool show_enum_members_docstring = true;  //< Include auto-generated member list in enum
+                                                  //  docstrings.
+    };
+
+    static state &global_state() {
+        static state instance;
+        return instance;
+    }
+
+    state previous_state;
+};
+
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/stl.h

@@ -0,0 +1,666 @@
+/*
+    pybind11/stl.h: Transparent conversion for STL data types
+
+    Copyright (c) 2016 Wenzel Jakob <wenzel.jakob@epfl.ch>
+
+    All rights reserved. Use of this source code is governed by a
+    BSD-style license that can be found in the LICENSE file.
+*/
+
+#pragma once
+
+#include "pybind11.h"
+#include "detail/common.h"
+#include "detail/descr.h"
+#include "detail/type_caster_base.h"
+
+#include <deque>
+#include <initializer_list>
+#include <list>
+#include <map>
+#include <memory>
+#include <ostream>
+#include <set>
+#include <unordered_map>
+#include <unordered_set>
+#include <valarray>
+
+// See `detail/common.h` for implementation of these guards.
+#if defined(PYBIND11_HAS_OPTIONAL)
+#    include <optional>
+#elif defined(PYBIND11_HAS_EXP_OPTIONAL)
+#    include <experimental/optional>
+#endif
+
+#if defined(PYBIND11_HAS_VARIANT)
+#    include <variant>
+#endif
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+//
+// Begin: Equivalent of
+//        https://github.com/google/clif/blob/ae4eee1de07cdf115c0c9bf9fec9ff28efce6f6c/clif/python/runtime.cc#L388-L438
+/*
+The three `object_is_convertible_to_*()` functions below are
+the result of converging the behaviors of pybind11 and PyCLIF
+(http://github.com/google/clif).
+
+Originally PyCLIF was extremely far on the permissive side of the spectrum,
+while pybind11 was very far on the strict side. Originally PyCLIF accepted any
+Python iterable as input for a C++ `vector`/`set`/`map` argument, as long as
+the elements were convertible. The obvious (in hindsight) problem was that
+any empty Python iterable could be passed to any of these C++ types, e.g. `{}`
+was accepted for C++ `vector`/`set` arguments, or `[]` for C++ `map` arguments.
+
+The functions below strike a practical permissive-vs-strict compromise,
+informed by tens of thousands of use cases in the wild. A main objective is
+to prevent accidents and improve readability:
+
+- Python literals must match the C++ types.
+
+- For C++ `set`: The potentially reducing conversion from a Python sequence
+  (e.g. Python `list` or `tuple`) to a C++ `set` must be explicit, by going
+  through a Python `set`.
+
+- However, a Python `set` can still be passed to a C++ `vector`. The rationale
+  is that this conversion is not reducing. Implicit conversions of this kind
+  are also fairly commonly used, therefore enforcing explicit conversions
+  would have an unfavorable cost : benefit ratio; more sloppily speaking,
+  such an enforcement would be more annoying than helpful.
+
+Additional checks have been added to allow types derived from `collections.abc.Set` and
+`collections.abc.Mapping` (`collections.abc.Sequence` is already allowed by `PySequence_Check`).
+*/
+
+inline bool object_is_instance_with_one_of_tp_names(PyObject *obj,
+                                                    std::initializer_list<const char *> tp_names) {
+    if (PyType_Check(obj)) {
+        return false;
+    }
+    const char *obj_tp_name = Py_TYPE(obj)->tp_name;
+    for (const auto *tp_name : tp_names) {
+        if (std::strcmp(obj_tp_name, tp_name) == 0) {
+            return true;
+        }
+    }
+    return false;
+}
+
+inline bool object_is_convertible_to_std_vector(const handle &src) {
+    // Allow sequence-like objects, but not (byte-)string-like objects.
+    if (PySequence_Check(src.ptr()) != 0) {
+        return !PyUnicode_Check(src.ptr()) && !PyBytes_Check(src.ptr());
+    }
+    // Allow generators, set/frozenset and several common iterable types.
+    return (PyGen_Check(src.ptr()) != 0) || (PyAnySet_Check(src.ptr()) != 0)
+           || object_is_instance_with_one_of_tp_names(
+               src.ptr(), {"dict_keys", "dict_values", "dict_items", "map", "zip"});
+}
+
+inline bool object_is_convertible_to_std_set(const handle &src, bool convert) {
+    // Allow set/frozenset and dict keys.
+    // In convert mode: also allow types derived from collections.abc.Set.
+    return ((PyAnySet_Check(src.ptr()) != 0)
+            || object_is_instance_with_one_of_tp_names(src.ptr(), {"dict_keys"}))
+           || (convert && isinstance(src, module_::import("collections.abc").attr("Set")));
+}
+
+inline bool object_is_convertible_to_std_map(const handle &src, bool convert) {
+    // Allow dict.
+    if (PyDict_Check(src.ptr())) {
+        return true;
+    }
+    // Allow types conforming to Mapping Protocol.
+    // According to https://docs.python.org/3/c-api/mapping.html, `PyMappingCheck()` checks for
+    // `__getitem__()` without checking the type of keys. In order to restrict the allowed types
+    // closer to actual Mapping-like types, we also check for the `items()` method.
+    if (PyMapping_Check(src.ptr()) != 0) {
+        PyObject *items = PyObject_GetAttrString(src.ptr(), "items");
+        if (items != nullptr) {
+            bool is_convertible = (PyCallable_Check(items) != 0);
+            Py_DECREF(items);
+            if (is_convertible) {
+                return true;
+            }
+        } else {
+            PyErr_Clear();
+        }
+    }
+    // In convert mode: Allow types derived from collections.abc.Mapping
+    return convert && isinstance(src, module_::import("collections.abc").attr("Mapping"));
+}
+
+//
+// End: Equivalent of clif/python/runtime.cc
+//
+
+/// Extracts an const lvalue reference or rvalue reference for U based on the type of T (e.g. for
+/// forwarding a container element).  Typically used indirect via forwarded_type(), below.
+template <typename T, typename U>
+using forwarded_type = conditional_t<std::is_lvalue_reference<T>::value,
+                                     remove_reference_t<U> &,
+                                     remove_reference_t<U> &&>;
+
+/// Forwards a value U as rvalue or lvalue according to whether T is rvalue or lvalue; typically
+/// used for forwarding a container's elements.
+template <typename T, typename U>
+constexpr forwarded_type<T, U> forward_like(U &&u) {
+    return std::forward<detail::forwarded_type<T, U>>(std::forward<U>(u));
+}
+
+// Checks if a container has a STL style reserve method.
+// This will only return true for a `reserve()` with a `void` return.
+template <typename C>
+using has_reserve_method = std::is_same<decltype(std::declval<C>().reserve(0)), void>;
+
+template <typename Type, typename Key>
+struct set_caster {
+    using type = Type;
+    using key_conv = make_caster<Key>;
+
+private:
+    template <typename T = Type, enable_if_t<has_reserve_method<T>::value, int> = 0>
+    void reserve_maybe(const anyset &s, Type *) {
+        value.reserve(s.size());
+    }
+    void reserve_maybe(const anyset &, void *) {}
+
+    bool convert_iterable(const iterable &itbl, bool convert) {
+        for (const auto &it : itbl) {
+            key_conv conv;
+            if (!conv.load(it, convert)) {
+                return false;
+            }
+            value.insert(cast_op<Key &&>(std::move(conv)));
+        }
+        return true;
+    }
+
+    bool convert_anyset(const anyset &s, bool convert) {
+        value.clear();
+        reserve_maybe(s, &value);
+        return convert_iterable(s, convert);
+    }
+
+public:
+    bool load(handle src, bool convert) {
+        if (!object_is_convertible_to_std_set(src, convert)) {
+            return false;
+        }
+        if (isinstance<anyset>(src)) {
+            value.clear();
+            return convert_anyset(reinterpret_borrow<anyset>(src), convert);
+        }
+        if (!convert) {
+            return false;
+        }
+        assert(isinstance<iterable>(src));
+        value.clear();
+        return convert_iterable(reinterpret_borrow<iterable>(src), convert);
+    }
+
+    template <typename T>
+    static handle cast(T &&src, return_value_policy policy, handle parent) {
+        if (!std::is_lvalue_reference<T>::value) {
+            policy = return_value_policy_override<Key>::policy(policy);
+        }
+        pybind11::set s;
+        for (auto &&value : src) {
+            auto value_ = reinterpret_steal<object>(
+                key_conv::cast(detail::forward_like<T>(value), policy, parent));
+            if (!value_ || !s.add(std::move(value_))) {
+                return handle();
+            }
+        }
+        return s.release();
+    }
+
+    PYBIND11_TYPE_CASTER(type,
+                         io_name("collections.abc.Set", "set") + const_name("[") + key_conv::name
+                             + const_name("]"));
+};
+
+template <typename Type, typename Key, typename Value>
+struct map_caster {
+    using key_conv = make_caster<Key>;
+    using value_conv = make_caster<Value>;
+
+private:
+    template <typename T = Type, enable_if_t<has_reserve_method<T>::value, int> = 0>
+    void reserve_maybe(const dict &d, Type *) {
+        value.reserve(d.size());
+    }
+    void reserve_maybe(const dict &, void *) {}
+
+    bool convert_elements(const dict &d, bool convert) {
+        value.clear();
+        reserve_maybe(d, &value);
+        for (const auto &it : d) {
+            key_conv kconv;
+            value_conv vconv;
+            if (!kconv.load(it.first.ptr(), convert) || !vconv.load(it.second.ptr(), convert)) {
+                return false;
+            }
+            value.emplace(cast_op<Key &&>(std::move(kconv)), cast_op<Value &&>(std::move(vconv)));
+        }
+        return true;
+    }
+
+public:
+    bool load(handle src, bool convert) {
+        if (!object_is_convertible_to_std_map(src, convert)) {
+            return false;
+        }
+        if (isinstance<dict>(src)) {
+            return convert_elements(reinterpret_borrow<dict>(src), convert);
+        }
+        if (!convert) {
+            return false;
+        }
+        auto items = reinterpret_steal<object>(PyMapping_Items(src.ptr()));
+        if (!items) {
+            throw error_already_set();
+        }
+        assert(isinstance<iterable>(items));
+        return convert_elements(dict(reinterpret_borrow<iterable>(items)), convert);
+    }
+
+    template <typename T>
+    static handle cast(T &&src, return_value_policy policy, handle parent) {
+        dict d;
+        return_value_policy policy_key = policy;
+        return_value_policy policy_value = policy;
+        if (!std::is_lvalue_reference<T>::value) {
+            policy_key = return_value_policy_override<Key>::policy(policy_key);
+            policy_value = return_value_policy_override<Value>::policy(policy_value);
+        }
+        for (auto &&kv : src) {
+            auto key = reinterpret_steal<object>(
+                key_conv::cast(detail::forward_like<T>(kv.first), policy_key, parent));
+            auto value = reinterpret_steal<object>(
+                value_conv::cast(detail::forward_like<T>(kv.second), policy_value, parent));
+            if (!key || !value) {
+                return handle();
+            }
+            d[std::move(key)] = std::move(value);
+        }
+        return d.release();
+    }
+
+    PYBIND11_TYPE_CASTER(Type,
+                         io_name("collections.abc.Mapping", "dict") + const_name("[")
+                             + key_conv::name + const_name(", ") + value_conv::name
+                             + const_name("]"));
+};
+
+template <typename Type, typename Value>
+struct list_caster {
+    using value_conv = make_caster<Value>;
+
+    bool load(handle src, bool convert) {
+        if (!object_is_convertible_to_std_vector(src)) {
+            return false;
+        }
+        if (isinstance<sequence>(src)) {
+            return convert_elements(src, convert);
+        }
+        if (!convert) {
+            return false;
+        }
+        // Designed to be behavior-equivalent to passing tuple(src) from Python:
+        // The conversion to a tuple will first exhaust the generator object, to ensure that
+        // the generator is not left in an unpredictable (to the caller) partially-consumed
+        // state.
+        assert(isinstance<iterable>(src));
+        return convert_elements(tuple(reinterpret_borrow<iterable>(src)), convert);
+    }
+
+private:
+    template <typename T = Type, enable_if_t<has_reserve_method<T>::value, int> = 0>
+    void reserve_maybe(const sequence &s, Type *) {
+        value.reserve(s.size());
+    }
+    void reserve_maybe(const sequence &, void *) {}
+
+    bool convert_elements(handle seq, bool convert) {
+        auto s = reinterpret_borrow<sequence>(seq);
+        value.clear();
+        reserve_maybe(s, &value);
+        for (const auto &it : seq) {
+            value_conv conv;
+            if (!conv.load(it, convert)) {
+                return false;
+            }
+            value.push_back(cast_op<Value &&>(std::move(conv)));
+        }
+        return true;
+    }
+
+public:
+    template <typename T>
+    static handle cast(T &&src, return_value_policy policy, handle parent) {
+        if (!std::is_lvalue_reference<T>::value) {
+            policy = return_value_policy_override<Value>::policy(policy);
+        }
+        list l(src.size());
+        ssize_t index = 0;
+        for (auto &&value : src) {
+            auto value_ = reinterpret_steal<object>(
+                value_conv::cast(detail::forward_like<T>(value), policy, parent));
+            if (!value_) {
+                return handle();
+            }
+            PyList_SET_ITEM(l.ptr(), index++, value_.release().ptr()); // steals a reference
+        }
+        return l.release();
+    }
+
+    PYBIND11_TYPE_CASTER(Type,
+                         io_name("collections.abc.Sequence", "list") + const_name("[")
+                             + value_conv::name + const_name("]"));
+};
+
+template <typename Type, typename Alloc>
+struct type_caster<std::vector<Type, Alloc>> : list_caster<std::vector<Type, Alloc>, Type> {};
+
+template <typename Type, typename Alloc>
+struct type_caster<std::deque<Type, Alloc>> : list_caster<std::deque<Type, Alloc>, Type> {};
+
+template <typename Type, typename Alloc>
+struct type_caster<std::list<Type, Alloc>> : list_caster<std::list<Type, Alloc>, Type> {};
+
+template <typename ArrayType, typename V, size_t... I>
+ArrayType vector_to_array_impl(V &&v, index_sequence<I...>) {
+    return {{std::move(v[I])...}};
+}
+
+// Based on https://en.cppreference.com/w/cpp/container/array/to_array
+template <typename ArrayType, size_t N, typename V>
+ArrayType vector_to_array(V &&v) {
+    return vector_to_array_impl<ArrayType, V>(std::forward<V>(v), make_index_sequence<N>{});
+}
+
+template <typename ArrayType, typename Value, bool Resizable, size_t Size = 0>
+struct array_caster {
+    using value_conv = make_caster<Value>;
+
+private:
+    std::unique_ptr<ArrayType> value;
+
+    template <bool R = Resizable, enable_if_t<R, int> = 0>
+    bool convert_elements(handle seq, bool convert) {
+        auto l = reinterpret_borrow<sequence>(seq);
+        value.reset(new ArrayType{});
+        // Using `resize` to preserve the behavior exactly as it was before PR #5305
+        // For the `resize` to work, `Value` must be default constructible.
+        // For `std::valarray`, this is a requirement:
+        // https://en.cppreference.com/w/cpp/named_req/NumericType
+        value->resize(l.size());
+        size_t ctr = 0;
+        for (const auto &it : l) {
+            value_conv conv;
+            if (!conv.load(it, convert)) {
+                return false;
+            }
+            (*value)[ctr++] = cast_op<Value &&>(std::move(conv));
+        }
+        return true;
+    }
+
+    template <bool R = Resizable, enable_if_t<!R, int> = 0>
+    bool convert_elements(handle seq, bool convert) {
+        auto l = reinterpret_borrow<sequence>(seq);
+        if (l.size() != Size) {
+            return false;
+        }
+        // The `temp` storage is needed to support `Value` types that are not
+        // default-constructible.
+        // Deliberate choice: no template specializations, for simplicity, and
+        // because the compile time overhead for the specializations is deemed
+        // more significant than the runtime overhead for the `temp` storage.
+        std::vector<Value> temp;
+        temp.reserve(l.size());
+        for (auto it : l) {
+            value_conv conv;
+            if (!conv.load(it, convert)) {
+                return false;
+            }
+            temp.emplace_back(cast_op<Value &&>(std::move(conv)));
+        }
+        value.reset(new ArrayType(vector_to_array<ArrayType, Size>(std::move(temp))));
+        return true;
+    }
+
+public:
+    bool load(handle src, bool convert) {
+        if (!object_is_convertible_to_std_vector(src)) {
+            return false;
+        }
+        if (isinstance<sequence>(src)) {
+            return convert_elements(src, convert);
+        }
+        if (!convert) {
+            return false;
+        }
+        // Designed to be behavior-equivalent to passing tuple(src) from Python:
+        // The conversion to a tuple will first exhaust the generator object, to ensure that
+        // the generator is not left in an unpredictable (to the caller) partially-consumed
+        // state.
+        assert(isinstance<iterable>(src));
+        return convert_elements(tuple(reinterpret_borrow<iterable>(src)), convert);
+    }
+
+    template <typename T>
+    static handle cast(T &&src, return_value_policy policy, handle parent) {
+        list l(src.size());
+        ssize_t index = 0;
+        for (auto &&value : src) {
+            auto value_ = reinterpret_steal<object>(
+                value_conv::cast(detail::forward_like<T>(value), policy, parent));
+            if (!value_) {
+                return handle();
+            }
+            PyList_SET_ITEM(l.ptr(), index++, value_.release().ptr()); // steals a reference
+        }
+        return l.release();
+    }
+
+    // Code copied from PYBIND11_TYPE_CASTER macro.
+    // Intentionally preserving the behavior exactly as it was before PR #5305
+    template <typename T_, enable_if_t<std::is_same<ArrayType, remove_cv_t<T_>>::value, int> = 0>
+    static handle cast(T_ *src, return_value_policy policy, handle parent) {
+        if (!src) {
+            return none().release();
+        }
+        if (policy == return_value_policy::take_ownership) {
+            auto h = cast(std::move(*src), policy, parent);
+            delete src; // WARNING: Assumes `src` was allocated with `new`.
+            return h;
+        }
+        return cast(*src, policy, parent);
+    }
+
+    // NOLINTNEXTLINE(google-explicit-constructor)
+    operator ArrayType *() { return &(*value); }
+    // NOLINTNEXTLINE(google-explicit-constructor)
+    operator ArrayType &() { return *value; }
+    // NOLINTNEXTLINE(google-explicit-constructor)
+    operator ArrayType &&() && { return std::move(*value); }
+
+    template <typename T_>
+    using cast_op_type = movable_cast_op_type<T_>;
+
+    static constexpr auto name
+        = const_name<Resizable>(const_name(""), const_name("typing.Annotated["))
+          + io_name("collections.abc.Sequence", "list") + const_name("[") + value_conv::name
+          + const_name("]")
+          + const_name<Resizable>(const_name(""),
+                                  const_name(", \"FixedSize(") + const_name<Size>()
+                                      + const_name(")\"]"));
+};
+
+template <typename Type, size_t Size>
+struct type_caster<std::array<Type, Size>>
+    : array_caster<std::array<Type, Size>, Type, false, Size> {};
+
+template <typename Type>
+struct type_caster<std::valarray<Type>> : array_caster<std::valarray<Type>, Type, true> {};
+
+template <typename Key, typename Compare, typename Alloc>
+struct type_caster<std::set<Key, Compare, Alloc>>
+    : set_caster<std::set<Key, Compare, Alloc>, Key> {};
+
+template <typename Key, typename Hash, typename Equal, typename Alloc>
+struct type_caster<std::unordered_set<Key, Hash, Equal, Alloc>>
+    : set_caster<std::unordered_set<Key, Hash, Equal, Alloc>, Key> {};
+
+template <typename Key, typename Value, typename Compare, typename Alloc>
+struct type_caster<std::map<Key, Value, Compare, Alloc>>
+    : map_caster<std::map<Key, Value, Compare, Alloc>, Key, Value> {};
+
+template <typename Key, typename Value, typename Hash, typename Equal, typename Alloc>
+struct type_caster<std::unordered_map<Key, Value, Hash, Equal, Alloc>>
+    : map_caster<std::unordered_map<Key, Value, Hash, Equal, Alloc>, Key, Value> {};
+
+// This type caster is intended to be used for std::optional and std::experimental::optional
+template <typename Type, typename Value = typename Type::value_type>
+struct optional_caster {
+    using value_conv = make_caster<Value>;
+
+    template <typename T>
+    static handle cast(T &&src, return_value_policy policy, handle parent) {
+        if (!src) {
+            return none().release();
+        }
+        if (!std::is_lvalue_reference<T>::value) {
+            policy = return_value_policy_override<Value>::policy(policy);
+        }
+        // NOLINTNEXTLINE(bugprone-unchecked-optional-access)
+        return value_conv::cast(*std::forward<T>(src), policy, parent);
+    }
+
+    bool load(handle src, bool convert) {
+        if (!src) {
+            return false;
+        }
+        if (src.is_none()) {
+            return true; // default-constructed value is already empty
+        }
+        value_conv inner_caster;
+        if (!inner_caster.load(src, convert)) {
+            return false;
+        }
+
+        value.emplace(cast_op<Value &&>(std::move(inner_caster)));
+        return true;
+    }
+
+    PYBIND11_TYPE_CASTER(Type, value_conv::name | make_caster<none>::name);
+};
+
+#if defined(PYBIND11_HAS_OPTIONAL)
+template <typename T>
+struct type_caster<std::optional<T>> : public optional_caster<std::optional<T>> {};
+
+template <>
+struct type_caster<std::nullopt_t> : public void_caster<std::nullopt_t> {};
+#endif
+
+#if defined(PYBIND11_HAS_EXP_OPTIONAL)
+template <typename T>
+struct type_caster<std::experimental::optional<T>>
+    : public optional_caster<std::experimental::optional<T>> {};
+
+template <>
+struct type_caster<std::experimental::nullopt_t>
+    : public void_caster<std::experimental::nullopt_t> {};
+#endif
+
+/// Visit a variant and cast any found type to Python
+struct variant_caster_visitor {
+    return_value_policy policy;
+    handle parent;
+
+    using result_type = handle; // required by boost::variant in C++11
+
+    template <typename T>
+    result_type operator()(T &&src) const {
+        return make_caster<T>::cast(std::forward<T>(src), policy, parent);
+    }
+};
+
+/// Helper class which abstracts away variant's `visit` function. `std::variant` and similar
+/// `namespace::variant` types which provide a `namespace::visit()` function are handled here
+/// automatically using argument-dependent lookup. Users can provide specializations for other
+/// variant-like classes, e.g. `boost::variant` and `boost::apply_visitor`.
+template <template <typename...> class Variant>
+struct visit_helper {
+    template <typename... Args>
+    static auto call(Args &&...args) -> decltype(visit(std::forward<Args>(args)...)) {
+        return visit(std::forward<Args>(args)...);
+    }
+};
+
+/// Generic variant caster
+template <typename Variant>
+struct variant_caster;
+
+template <template <typename...> class V, typename... Ts>
+struct variant_caster<V<Ts...>> {
+    static_assert(sizeof...(Ts) > 0, "Variant must consist of at least one alternative.");
+
+    template <typename U, typename... Us>
+    bool load_alternative(handle src, bool convert, type_list<U, Us...>) {
+        auto caster = make_caster<U>();
+        if (caster.load(src, convert)) {
+            value = cast_op<U>(std::move(caster));
+            return true;
+        }
+        return load_alternative(src, convert, type_list<Us...>{});
+    }
+
+    bool load_alternative(handle, bool, type_list<>) { return false; }
+
+    bool load(handle src, bool convert) {
+        // Do a first pass without conversions to improve constructor resolution.
+        // E.g. `py::int_(1).cast<variant<double, int>>()` needs to fill the `int`
+        // slot of the variant. Without two-pass loading `double` would be filled
+        // because it appears first and a conversion is possible.
+        if (convert && load_alternative(src, false, type_list<Ts...>{})) {
+            return true;
+        }
+        return load_alternative(src, convert, type_list<Ts...>{});
+    }
+
+    template <typename Variant>
+    static handle cast(Variant &&src, return_value_policy policy, handle parent) {
+        return visit_helper<V>::call(variant_caster_visitor{policy, parent},
+                                     std::forward<Variant>(src));
+    }
+
+    using Type = V<Ts...>;
+    PYBIND11_TYPE_CASTER(Type, ::pybind11::detail::union_concat(make_caster<Ts>::name...));
+};
+
+#if defined(PYBIND11_HAS_VARIANT)
+template <typename... Ts>
+struct type_caster<std::variant<Ts...>> : variant_caster<std::variant<Ts...>> {};
+
+template <>
+struct type_caster<std::monostate> : public void_caster<std::monostate> {};
+#endif
+
+PYBIND11_NAMESPACE_END(detail)
+
+inline std::ostream &operator<<(std::ostream &os, const handle &obj) {
+#ifdef PYBIND11_HAS_STRING_VIEW
+    os << str(obj).cast<std::string_view>();
+#else
+    os << (std::string) str(obj);
+#endif
+    return os;
+}
+
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/stl/filesystem.h

@@ -0,0 +1,114 @@
+// Copyright (c) 2021 The Pybind Development Team.
+// All rights reserved. Use of this source code is governed by a
+// BSD-style license that can be found in the LICENSE file.
+
+#pragma once
+
+#include <pybind11/cast.h>
+#include <pybind11/detail/common.h>
+#include <pybind11/detail/descr.h>
+#include <pybind11/pybind11.h>
+#include <pybind11/pytypes.h>
+
+#include <string>
+
+#if defined(PYBIND11_HAS_FILESYSTEM)
+#    include <filesystem>
+#elif defined(PYBIND11_HAS_EXPERIMENTAL_FILESYSTEM)
+#    include <experimental/filesystem>
+#else
+#    error "Neither #include <filesystem> nor #include <experimental/filesystem> is available."
+#endif
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+#ifdef PYPY_VERSION
+#    define PYBIND11_REINTERPRET_CAST_VOID_PTR_IF_NOT_PYPY(...) (__VA_ARGS__)
+#else
+#    define PYBIND11_REINTERPRET_CAST_VOID_PTR_IF_NOT_PYPY(...)                                   \
+        (reinterpret_cast<void *>(__VA_ARGS__))
+#endif
+
+#if defined(PYBIND11_HAS_FILESYSTEM) || defined(PYBIND11_HAS_EXPERIMENTAL_FILESYSTEM)
+template <typename T>
+struct path_caster {
+
+private:
+    static PyObject *unicode_from_fs_native(const std::string &w) {
+#    if !defined(PYPY_VERSION)
+        return PyUnicode_DecodeFSDefaultAndSize(w.c_str(), ssize_t(w.size()));
+#    else
+        // PyPy mistakenly declares the first parameter as non-const.
+        return PyUnicode_DecodeFSDefaultAndSize(const_cast<char *>(w.c_str()), ssize_t(w.size()));
+#    endif
+    }
+
+    static PyObject *unicode_from_fs_native(const std::wstring &w) {
+        return PyUnicode_FromWideChar(w.c_str(), ssize_t(w.size()));
+    }
+
+public:
+    static handle cast(const T &path, return_value_policy, handle) {
+        if (auto py_str = unicode_from_fs_native(path.native())) {
+            return module_::import("pathlib")
+                .attr("Path")(reinterpret_steal<object>(py_str))
+                .release();
+        }
+        return nullptr;
+    }
+
+    bool load(handle handle, bool) {
+        // PyUnicode_FSConverter and PyUnicode_FSDecoder normally take care of
+        // calling PyOS_FSPath themselves, but that's broken on PyPy (PyPy
+        // issue #3168) so we do it ourselves instead.
+        PyObject *buf = PyOS_FSPath(handle.ptr());
+        if (!buf) {
+            PyErr_Clear();
+            return false;
+        }
+        PyObject *native = nullptr;
+        if constexpr (std::is_same_v<typename T::value_type, char>) {
+            if (PyUnicode_FSConverter(buf, PYBIND11_REINTERPRET_CAST_VOID_PTR_IF_NOT_PYPY(&native))
+                != 0) {
+                if (auto *c_str = PyBytes_AsString(native)) {
+                    // AsString returns a pointer to the internal buffer, which
+                    // must not be free'd.
+                    value = c_str;
+                }
+            }
+        } else if constexpr (std::is_same_v<typename T::value_type, wchar_t>) {
+            if (PyUnicode_FSDecoder(buf, PYBIND11_REINTERPRET_CAST_VOID_PTR_IF_NOT_PYPY(&native))
+                != 0) {
+                if (auto *c_str = PyUnicode_AsWideCharString(native, nullptr)) {
+                    // AsWideCharString returns a new string that must be free'd.
+                    value = c_str; // Copies the string.
+                    PyMem_Free(c_str);
+                }
+            }
+        }
+        Py_XDECREF(native);
+        Py_DECREF(buf);
+        if (PyErr_Occurred()) {
+            PyErr_Clear();
+            return false;
+        }
+        return true;
+    }
+
+    PYBIND11_TYPE_CASTER(T, io_name("os.PathLike | str | bytes", "pathlib.Path"));
+};
+
+#endif // PYBIND11_HAS_FILESYSTEM || defined(PYBIND11_HAS_EXPERIMENTAL_FILESYSTEM)
+
+#if defined(PYBIND11_HAS_FILESYSTEM)
+template <>
+struct type_caster<std::filesystem::path> : public path_caster<std::filesystem::path> {};
+#elif defined(PYBIND11_HAS_EXPERIMENTAL_FILESYSTEM)
+template <>
+struct type_caster<std::experimental::filesystem::path>
+    : public path_caster<std::experimental::filesystem::path> {};
+#endif
+
+PYBIND11_NAMESPACE_END(detail)
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/stl_bind.h

@@ -0,0 +1,858 @@
+/*
+    pybind11/std_bind.h: Binding generators for STL data types
+
+    Copyright (c) 2016 Sergey Lyskov and Wenzel Jakob
+
+    All rights reserved. Use of this source code is governed by a
+    BSD-style license that can be found in the LICENSE file.
+*/
+
+#pragma once
+
+#include "detail/common.h"
+#include "detail/type_caster_base.h"
+#include "cast.h"
+#include "operators.h"
+
+#include <algorithm>
+#include <sstream>
+#include <type_traits>
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+/* SFINAE helper class used by 'is_comparable */
+template <typename T>
+struct container_traits {
+    template <typename T2>
+    static std::true_type
+    test_comparable(decltype(std::declval<const T2 &>() == std::declval<const T2 &>()) *);
+    template <typename T2>
+    static std::false_type test_comparable(...);
+    template <typename T2>
+    static std::true_type test_value(typename T2::value_type *);
+    template <typename T2>
+    static std::false_type test_value(...);
+    template <typename T2>
+    static std::true_type test_pair(typename T2::first_type *, typename T2::second_type *);
+    template <typename T2>
+    static std::false_type test_pair(...);
+
+    static constexpr const bool is_comparable
+        = std::is_same<std::true_type, decltype(test_comparable<T>(nullptr))>::value;
+    static constexpr const bool is_pair
+        = std::is_same<std::true_type, decltype(test_pair<T>(nullptr, nullptr))>::value;
+    static constexpr const bool is_vector
+        = std::is_same<std::true_type, decltype(test_value<T>(nullptr))>::value;
+    static constexpr const bool is_element = !is_pair && !is_vector;
+};
+
+/* Default: is_comparable -> std::false_type */
+template <typename T, typename SFINAE = void>
+struct is_comparable : std::false_type {};
+
+/* For non-map data structures, check whether operator== can be instantiated */
+template <typename T>
+struct is_comparable<
+    T,
+    enable_if_t<container_traits<T>::is_element && container_traits<T>::is_comparable>>
+    : std::true_type {};
+
+/* For a vector/map data structure, recursively check the value type
+   (which is std::pair for maps) */
+template <typename T>
+struct is_comparable<T, enable_if_t<container_traits<T>::is_vector>>
+    : is_comparable<typename recursive_container_traits<T>::type_to_check_recursively> {};
+
+template <>
+struct is_comparable<recursive_bottom> : std::true_type {};
+
+/* For pairs, recursively check the two data types */
+template <typename T>
+struct is_comparable<T, enable_if_t<container_traits<T>::is_pair>> {
+    static constexpr const bool value = is_comparable<typename T::first_type>::value
+                                        && is_comparable<typename T::second_type>::value;
+};
+
+/* Fallback functions */
+template <typename, typename, typename... Args>
+void vector_if_copy_constructible(const Args &...) {}
+template <typename, typename, typename... Args>
+void vector_if_equal_operator(const Args &...) {}
+template <typename, typename, typename... Args>
+void vector_if_insertion_operator(const Args &...) {}
+template <typename, typename, typename... Args>
+void vector_modifiers(const Args &...) {}
+
+template <typename Vector, typename Class_>
+void vector_if_copy_constructible(enable_if_t<is_copy_constructible<Vector>::value, Class_> &cl) {
+    cl.def(init<const Vector &>(), "Copy constructor");
+}
+
+template <typename Vector, typename Class_>
+void vector_if_equal_operator(enable_if_t<is_comparable<Vector>::value, Class_> &cl) {
+    using T = typename Vector::value_type;
+
+    cl.def(self == self);
+    cl.def(self != self);
+
+    cl.def(
+        "count",
+        [](const Vector &v, const T &x) { return std::count(v.begin(), v.end(), x); },
+        arg("x"),
+        "Return the number of times ``x`` appears in the list");
+
+    cl.def(
+        "remove",
+        [](Vector &v, const T &x) {
+            auto p = std::find(v.begin(), v.end(), x);
+            if (p != v.end()) {
+                v.erase(p);
+            } else {
+                throw value_error();
+            }
+        },
+        arg("x"),
+        "Remove the first item from the list whose value is x. "
+        "It is an error if there is no such item.");
+
+    cl.def(
+        "__contains__",
+        [](const Vector &v, const T &x) { return std::find(v.begin(), v.end(), x) != v.end(); },
+        arg("x"),
+        "Return true the container contains ``x``");
+}
+
+// Vector modifiers -- requires a copyable vector_type:
+// (Technically, some of these (pop and __delitem__) don't actually require copyability, but it
+// seems silly to allow deletion but not insertion, so include them here too.)
+template <typename Vector, typename Class_>
+void vector_modifiers(
+    enable_if_t<is_copy_constructible<typename Vector::value_type>::value, Class_> &cl) {
+    using T = typename Vector::value_type;
+    using SizeType = typename Vector::size_type;
+    using DiffType = typename Vector::difference_type;
+
+    auto wrap_i = [](DiffType i, SizeType n) {
+        if (i < 0) {
+            i += n;
+        }
+        if (i < 0 || (SizeType) i >= n) {
+            throw index_error();
+        }
+        return i;
+    };
+
+    cl.def(
+        "append",
+        [](Vector &v, const T &value) { v.push_back(value); },
+        arg("x"),
+        "Add an item to the end of the list");
+
+    cl.def(init([](const iterable &it) {
+        auto v = std::unique_ptr<Vector>(new Vector());
+        v->reserve(len_hint(it));
+        for (handle h : it) {
+            v->push_back(h.cast<T>());
+        }
+        return v.release();
+    }));
+
+    cl.def("clear", [](Vector &v) { v.clear(); }, "Clear the contents");
+
+    cl.def(
+        "extend",
+        [](Vector &v, const Vector &src) { v.insert(v.end(), src.begin(), src.end()); },
+        arg("L"),
+        "Extend the list by appending all the items in the given list");
+
+    cl.def(
+        "extend",
+        [](Vector &v, const iterable &it) {
+            const size_t old_size = v.size();
+            v.reserve(old_size + len_hint(it));
+            try {
+                for (handle h : it) {
+                    v.push_back(h.cast<T>());
+                }
+            } catch (const cast_error &) {
+                v.erase(v.begin() + static_cast<typename Vector::difference_type>(old_size),
+                        v.end());
+                try {
+                    v.shrink_to_fit();
+                } catch (const std::exception &) { // NOLINT(bugprone-empty-catch)
+                    // Do nothing
+                }
+                throw;
+            }
+        },
+        arg("L"),
+        "Extend the list by appending all the items in the given list");
+
+    cl.def(
+        "insert",
+        [](Vector &v, DiffType i, const T &x) {
+            // Can't use wrap_i; i == v.size() is OK
+            if (i < 0) {
+                i += v.size();
+            }
+            if (i < 0 || (SizeType) i > v.size()) {
+                throw index_error();
+            }
+            v.insert(v.begin() + i, x);
+        },
+        arg("i"),
+        arg("x"),
+        "Insert an item at a given position.");
+
+    cl.def(
+        "pop",
+        [](Vector &v) {
+            if (v.empty()) {
+                throw index_error();
+            }
+            T t = std::move(v.back());
+            v.pop_back();
+            return t;
+        },
+        "Remove and return the last item");
+
+    cl.def(
+        "pop",
+        [wrap_i](Vector &v, DiffType i) {
+            i = wrap_i(i, v.size());
+            T t = std::move(v[(SizeType) i]);
+            v.erase(std::next(v.begin(), i));
+            return t;
+        },
+        arg("i"),
+        "Remove and return the item at index ``i``");
+
+    cl.def("__setitem__", [wrap_i](Vector &v, DiffType i, const T &t) {
+        i = wrap_i(i, v.size());
+        v[(SizeType) i] = t;
+    });
+
+    /// Slicing protocol
+    cl.def(
+        "__getitem__",
+        [](const Vector &v, const slice &slice) -> Vector * {
+            size_t start = 0, stop = 0, step = 0, slicelength = 0;
+
+            if (!slice.compute(v.size(), &start, &stop, &step, &slicelength)) {
+                throw error_already_set();
+            }
+
+            auto *seq = new Vector();
+            seq->reserve(slicelength);
+
+            for (size_t i = 0; i < slicelength; ++i) {
+                seq->push_back(v[start]);
+                start += step;
+            }
+            return seq;
+        },
+        arg("s"),
+        "Retrieve list elements using a slice object");
+
+    cl.def(
+        "__setitem__",
+        [](Vector &v, const slice &slice, const Vector &value) {
+            size_t start = 0, stop = 0, step = 0, slicelength = 0;
+            if (!slice.compute(v.size(), &start, &stop, &step, &slicelength)) {
+                throw error_already_set();
+            }
+
+            if (slicelength != value.size()) {
+                throw std::runtime_error(
+                    "Left and right hand size of slice assignment have different sizes!");
+            }
+
+            for (size_t i = 0; i < slicelength; ++i) {
+                v[start] = value[i];
+                start += step;
+            }
+        },
+        "Assign list elements using a slice object");
+
+    cl.def(
+        "__delitem__",
+        [wrap_i](Vector &v, DiffType i) {
+            i = wrap_i(i, v.size());
+            v.erase(v.begin() + i);
+        },
+        "Delete the list elements at index ``i``");
+
+    cl.def(
+        "__delitem__",
+        [](Vector &v, const slice &slice) {
+            size_t start = 0, stop = 0, step = 0, slicelength = 0;
+
+            if (!slice.compute(v.size(), &start, &stop, &step, &slicelength)) {
+                throw error_already_set();
+            }
+
+            if (step == 1 && false) {
+                v.erase(v.begin() + (DiffType) start, v.begin() + DiffType(start + slicelength));
+            } else {
+                for (size_t i = 0; i < slicelength; ++i) {
+                    v.erase(v.begin() + DiffType(start));
+                    start += step - 1;
+                }
+            }
+        },
+        "Delete list elements using a slice object");
+}
+
+// If the type has an operator[] that doesn't return a reference (most notably std::vector<bool>),
+// we have to access by copying; otherwise we return by reference.
+template <typename Vector>
+using vector_needs_copy
+    = negation<std::is_same<decltype(std::declval<Vector>()[typename Vector::size_type()]),
+                            typename Vector::value_type &>>;
+
+// The usual case: access and iterate by reference
+template <typename Vector, typename Class_>
+void vector_accessor(enable_if_t<!vector_needs_copy<Vector>::value, Class_> &cl) {
+    using T = typename Vector::value_type;
+    using SizeType = typename Vector::size_type;
+    using DiffType = typename Vector::difference_type;
+    using ItType = typename Vector::iterator;
+
+    auto wrap_i = [](DiffType i, SizeType n) {
+        if (i < 0) {
+            i += n;
+        }
+        if (i < 0 || (SizeType) i >= n) {
+            throw index_error();
+        }
+        return i;
+    };
+
+    cl.def(
+        "__getitem__",
+        [wrap_i](Vector &v, DiffType i) -> T & {
+            i = wrap_i(i, v.size());
+            return v[(SizeType) i];
+        },
+        return_value_policy::reference_internal // ref + keepalive
+    );
+
+    cl.def(
+        "__iter__",
+        [](Vector &v) {
+            return make_iterator<return_value_policy::reference_internal, ItType, ItType, T &>(
+                v.begin(), v.end());
+        },
+        keep_alive<0, 1>() /* Essential: keep list alive while iterator exists */
+    );
+}
+
+// The case for special objects, like std::vector<bool>, that have to be returned-by-copy:
+template <typename Vector, typename Class_>
+void vector_accessor(enable_if_t<vector_needs_copy<Vector>::value, Class_> &cl) {
+    using T = typename Vector::value_type;
+    using SizeType = typename Vector::size_type;
+    using DiffType = typename Vector::difference_type;
+    using ItType = typename Vector::iterator;
+    cl.def("__getitem__", [](const Vector &v, DiffType i) -> T {
+        if (i < 0) {
+            i += v.size();
+            if (i < 0) {
+                throw index_error();
+            }
+        }
+        auto i_st = static_cast<SizeType>(i);
+        if (i_st >= v.size()) {
+            throw index_error();
+        }
+        return v[i_st];
+    });
+
+    cl.def(
+        "__iter__",
+        [](Vector &v) {
+            return make_iterator<return_value_policy::copy, ItType, ItType, T>(v.begin(), v.end());
+        },
+        keep_alive<0, 1>() /* Essential: keep list alive while iterator exists */
+    );
+}
+
+template <typename Vector, typename Class_>
+auto vector_if_insertion_operator(Class_ &cl, std::string const &name)
+    -> decltype(std::declval<std::ostream &>() << std::declval<typename Vector::value_type>(),
+                void()) {
+    using size_type = typename Vector::size_type;
+
+    cl.def(
+        "__repr__",
+        [name](Vector &v) {
+            std::ostringstream s;
+            s << name << '[';
+            for (size_type i = 0; i < v.size(); ++i) {
+                s << v[i];
+                if (i != v.size() - 1) {
+                    s << ", ";
+                }
+            }
+            s << ']';
+            return s.str();
+        },
+        "Return the canonical string representation of this list.");
+}
+
+// Provide the buffer interface for vectors if we have data() and we have a format for it
+// GCC seems to have "void std::vector<bool>::data()" - doing SFINAE on the existence of data()
+// is insufficient, we need to check it returns an appropriate pointer
+template <typename Vector, typename = void>
+struct vector_has_data_and_format : std::false_type {};
+template <typename Vector>
+struct vector_has_data_and_format<
+    Vector,
+    enable_if_t<std::is_same<decltype(format_descriptor<typename Vector::value_type>::format(),
+                                      std::declval<Vector>().data()),
+                             typename Vector::value_type *>::value>> : std::true_type {};
+
+// [workaround(intel)] Separate function required here
+// Workaround as the Intel compiler does not compile the enable_if_t part below
+// (tested with icc (ICC) 2021.1 Beta 20200827)
+template <typename... Args>
+constexpr bool args_any_are_buffer() {
+    return detail::any_of<std::is_same<Args, buffer_protocol>...>::value;
+}
+
+// [workaround(intel)] Separate function required here
+// [workaround(msvc)] Can't use constexpr bool in return type
+
+// Add the buffer interface to a vector
+template <typename Vector, typename Class_, typename... Args>
+void vector_buffer_impl(Class_ &cl, std::true_type) {
+    using T = typename Vector::value_type;
+
+    static_assert(vector_has_data_and_format<Vector>::value,
+                  "There is not an appropriate format descriptor for this vector");
+
+    // numpy.h declares this for arbitrary types, but it may raise an exception and crash hard
+    // at runtime if PYBIND11_NUMPY_DTYPE hasn't been called, so check here
+    format_descriptor<T>::format();
+
+    cl.def_buffer([](Vector &v) -> buffer_info {
+        return buffer_info(v.data(),
+                           static_cast<ssize_t>(sizeof(T)),
+                           format_descriptor<T>::format(),
+                           1,
+                           {v.size()},
+                           {sizeof(T)});
+    });
+
+    cl.def(init([](const buffer &buf) {
+        auto info = buf.request();
+        if (info.ndim != 1 || info.strides[0] % static_cast<ssize_t>(sizeof(T))) {
+            throw type_error("Only valid 1D buffers can be copied to a vector");
+        }
+        if (!detail::compare_buffer_info<T>::compare(info)
+            || (ssize_t) sizeof(T) != info.itemsize) {
+            throw type_error("Format mismatch (Python: " + info.format
+                             + " C++: " + format_descriptor<T>::format() + ")");
+        }
+
+        T *p = static_cast<T *>(info.ptr);
+        ssize_t step = info.strides[0] / static_cast<ssize_t>(sizeof(T));
+        T *end = p + info.shape[0] * step;
+        if (step == 1) {
+            return Vector(p, end);
+        }
+        Vector vec;
+        vec.reserve((size_t) info.shape[0]);
+        for (; p != end; p += step) {
+            vec.push_back(*p);
+        }
+        return vec;
+    }));
+
+    return;
+}
+
+template <typename Vector, typename Class_, typename... Args>
+void vector_buffer_impl(Class_ &, std::false_type) {}
+
+template <typename Vector, typename Class_, typename... Args>
+void vector_buffer(Class_ &cl) {
+    vector_buffer_impl<Vector, Class_, Args...>(
+        cl, detail::any_of<std::is_same<Args, buffer_protocol>...>{});
+}
+
+PYBIND11_NAMESPACE_END(detail)
+
+//
+// std::vector
+//
+template <typename Vector, typename holder_type = default_holder_type<Vector>, typename... Args>
+class_<Vector, holder_type> bind_vector(handle scope, std::string const &name, Args &&...args) {
+    using Class_ = class_<Vector, holder_type>;
+
+    // If the value_type is unregistered (e.g. a converting type) or is itself registered
+    // module-local then make the vector binding module-local as well:
+    using vtype = typename Vector::value_type;
+    auto *vtype_info = detail::get_type_info(typeid(vtype));
+    bool local = !vtype_info || vtype_info->module_local;
+
+    Class_ cl(scope, name.c_str(), pybind11::module_local(local), std::forward<Args>(args)...);
+
+    // Declare the buffer interface if a buffer_protocol() is passed in
+    detail::vector_buffer<Vector, Class_, Args...>(cl);
+
+    cl.def(init<>());
+
+    // Register copy constructor (if possible)
+    detail::vector_if_copy_constructible<Vector, Class_>(cl);
+
+    // Register comparison-related operators and functions (if possible)
+    detail::vector_if_equal_operator<Vector, Class_>(cl);
+
+    // Register stream insertion operator (if possible)
+    detail::vector_if_insertion_operator<Vector, Class_>(cl, name);
+
+    // Modifiers require copyable vector value type
+    detail::vector_modifiers<Vector, Class_>(cl);
+
+    // Accessor and iterator; return by value if copyable, otherwise we return by ref + keep-alive
+    detail::vector_accessor<Vector, Class_>(cl);
+
+    cl.def(
+        "__bool__",
+        [](const Vector &v) -> bool { return !v.empty(); },
+        "Check whether the list is nonempty");
+
+    cl.def("__len__", [](const Vector &vec) { return vec.size(); });
+
+#if 0
+    // C++ style functions deprecated, leaving it here as an example
+    cl.def(init<size_type>());
+
+    cl.def("resize",
+         (void (Vector::*) (size_type count)) & Vector::resize,
+         "changes the number of elements stored");
+
+    cl.def("erase",
+        [](Vector &v, SizeType i) {
+        if (i >= v.size())
+            throw index_error();
+        v.erase(v.begin() + i);
+    }, "erases element at index ``i``");
+
+    cl.def("empty",         &Vector::empty,         "checks whether the container is empty");
+    cl.def("size",          &Vector::size,          "returns the number of elements");
+    cl.def("push_back", (void (Vector::*)(const T&)) &Vector::push_back, "adds an element to the end");
+    cl.def("pop_back",                               &Vector::pop_back, "removes the last element");
+
+    cl.def("max_size",      &Vector::max_size,      "returns the maximum possible number of elements");
+    cl.def("reserve",       &Vector::reserve,       "reserves storage");
+    cl.def("capacity",      &Vector::capacity,      "returns the number of elements that can be held in currently allocated storage");
+    cl.def("shrink_to_fit", &Vector::shrink_to_fit, "reduces memory usage by freeing unused memory");
+
+    cl.def("clear", &Vector::clear, "clears the contents");
+    cl.def("swap",   &Vector::swap, "swaps the contents");
+
+    cl.def("front", [](Vector &v) {
+        if (v.size()) return v.front();
+        else throw index_error();
+    }, "access the first element");
+
+    cl.def("back", [](Vector &v) {
+        if (v.size()) return v.back();
+        else throw index_error();
+    }, "access the last element ");
+
+#endif
+
+    return cl;
+}
+
+//
+// std::map, std::unordered_map
+//
+
+PYBIND11_NAMESPACE_BEGIN(detail)
+
+/* Fallback functions */
+template <typename, typename, typename... Args>
+void map_if_insertion_operator(const Args &...) {}
+template <typename, typename, typename... Args>
+void map_assignment(const Args &...) {}
+
+// Map assignment when copy-assignable: just copy the value
+template <typename Map, typename Class_>
+void map_assignment(
+    enable_if_t<is_copy_assignable<typename Map::mapped_type>::value, Class_> &cl) {
+    using KeyType = typename Map::key_type;
+    using MappedType = typename Map::mapped_type;
+
+    cl.def("__setitem__", [](Map &m, const KeyType &k, const MappedType &v) {
+        auto it = m.find(k);
+        if (it != m.end()) {
+            it->second = v;
+        } else {
+            m.emplace(k, v);
+        }
+    });
+}
+
+// Not copy-assignable, but still copy-constructible: we can update the value by erasing and
+// reinserting
+template <typename Map, typename Class_>
+void map_assignment(enable_if_t<!is_copy_assignable<typename Map::mapped_type>::value
+                                    && is_copy_constructible<typename Map::mapped_type>::value,
+                                Class_> &cl) {
+    using KeyType = typename Map::key_type;
+    using MappedType = typename Map::mapped_type;
+
+    cl.def("__setitem__", [](Map &m, const KeyType &k, const MappedType &v) {
+        // We can't use m[k] = v; because value type might not be default constructable
+        auto r = m.emplace(k, v);
+        if (!r.second) {
+            // value type is not copy assignable so the only way to insert it is to erase it
+            // first...
+            m.erase(r.first);
+            m.emplace(k, v);
+        }
+    });
+}
+
+template <typename Map, typename Class_>
+auto map_if_insertion_operator(Class_ &cl, std::string const &name)
+    -> decltype(std::declval<std::ostream &>() << std::declval<typename Map::key_type>()
+                                               << std::declval<typename Map::mapped_type>(),
+                void()) {
+
+    cl.def(
+        "__repr__",
+        [name](Map &m) {
+            std::ostringstream s;
+            s << name << '{';
+            bool f = false;
+            for (auto const &kv : m) {
+                if (f) {
+                    s << ", ";
+                }
+                s << kv.first << ": " << kv.second;
+                f = true;
+            }
+            s << '}';
+            return s.str();
+        },
+        "Return the canonical string representation of this map.");
+}
+
+struct keys_view {
+    virtual size_t len() = 0;
+    virtual iterator iter() = 0;
+    virtual bool contains(const handle &k) = 0;
+    virtual ~keys_view() = default;
+};
+
+struct values_view {
+    virtual size_t len() = 0;
+    virtual iterator iter() = 0;
+    virtual ~values_view() = default;
+};
+
+struct items_view {
+    virtual size_t len() = 0;
+    virtual iterator iter() = 0;
+    virtual ~items_view() = default;
+};
+
+template <typename Map>
+struct KeysViewImpl : public detail::keys_view {
+    explicit KeysViewImpl(Map &map) : map(map) {}
+    size_t len() override { return map.size(); }
+    iterator iter() override { return make_key_iterator(map.begin(), map.end()); }
+    bool contains(const handle &k) override {
+        try {
+            return map.find(k.template cast<typename Map::key_type>()) != map.end();
+        } catch (const cast_error &) {
+            return false;
+        }
+    }
+    Map &map;
+};
+
+template <typename Map>
+struct ValuesViewImpl : public detail::values_view {
+    explicit ValuesViewImpl(Map &map) : map(map) {}
+    size_t len() override { return map.size(); }
+    iterator iter() override { return make_value_iterator(map.begin(), map.end()); }
+    Map &map;
+};
+
+template <typename Map>
+struct ItemsViewImpl : public detail::items_view {
+    explicit ItemsViewImpl(Map &map) : map(map) {}
+    size_t len() override { return map.size(); }
+    iterator iter() override { return make_iterator(map.begin(), map.end()); }
+    Map &map;
+};
+
+inline str format_message_key_error_key_object(handle py_key) {
+    str message = "pybind11::bind_map key";
+    if (!py_key) {
+        return message;
+    }
+    try {
+        message = str(py_key);
+    } catch (const std::exception &) {
+        try {
+            message = repr(py_key);
+        } catch (const std::exception &) {
+            return message;
+        }
+    }
+    const ssize_t cut_length = 100;
+    if (len(message) > 2 * cut_length + 3) {
+        return str(message[slice(0, cut_length, 1)]) + str("✄✄✄")
+               + str(message[slice(-cut_length, static_cast<ssize_t>(len(message)), 1)]);
+    }
+    return message;
+}
+
+template <typename KeyType>
+str format_message_key_error(const KeyType &key) {
+    object py_key;
+    try {
+        py_key = cast(key);
+    } catch (const std::exception &) {
+        do { // Trick to avoid "empty catch" warning/error.
+        } while (false);
+    }
+    return format_message_key_error_key_object(py_key);
+}
+
+PYBIND11_NAMESPACE_END(detail)
+
+template <typename Map, typename holder_type = default_holder_type<Map>, typename... Args>
+class_<Map, holder_type> bind_map(handle scope, const std::string &name, Args &&...args) {
+    using KeyType = typename Map::key_type;
+    using MappedType = typename Map::mapped_type;
+    using KeysView = detail::keys_view;
+    using ValuesView = detail::values_view;
+    using ItemsView = detail::items_view;
+    using Class_ = class_<Map, holder_type>;
+
+    // If either type is a non-module-local bound type then make the map binding non-local as well;
+    // otherwise (e.g. both types are either module-local or converting) the map will be
+    // module-local.
+    auto *tinfo = detail::get_type_info(typeid(MappedType));
+    bool local = !tinfo || tinfo->module_local;
+    if (local) {
+        tinfo = detail::get_type_info(typeid(KeyType));
+        local = !tinfo || tinfo->module_local;
+    }
+
+    Class_ cl(scope, name.c_str(), pybind11::module_local(local), std::forward<Args>(args)...);
+
+    // Wrap KeysView if it wasn't already wrapped
+    if (!detail::get_type_info(typeid(KeysView))) {
+        class_<KeysView> keys_view(scope, "KeysView", pybind11::module_local(local));
+        keys_view.def("__len__", &KeysView::len);
+        keys_view.def("__iter__",
+                      &KeysView::iter,
+                      keep_alive<0, 1>() /* Essential: keep view alive while iterator exists */
+        );
+        keys_view.def("__contains__", &KeysView::contains);
+    }
+    // Similarly for ValuesView:
+    if (!detail::get_type_info(typeid(ValuesView))) {
+        class_<ValuesView> values_view(scope, "ValuesView", pybind11::module_local(local));
+        values_view.def("__len__", &ValuesView::len);
+        values_view.def("__iter__",
+                        &ValuesView::iter,
+                        keep_alive<0, 1>() /* Essential: keep view alive while iterator exists */
+        );
+    }
+    // Similarly for ItemsView:
+    if (!detail::get_type_info(typeid(ItemsView))) {
+        class_<ItemsView> items_view(scope, "ItemsView", pybind11::module_local(local));
+        items_view.def("__len__", &ItemsView::len);
+        items_view.def("__iter__",
+                       &ItemsView::iter,
+                       keep_alive<0, 1>() /* Essential: keep view alive while iterator exists */
+        );
+    }
+
+    cl.def(init<>());
+
+    // Register stream insertion operator (if possible)
+    detail::map_if_insertion_operator<Map, Class_>(cl, name);
+
+    cl.def(
+        "__bool__",
+        [](const Map &m) -> bool { return !m.empty(); },
+        "Check whether the map is nonempty");
+
+    cl.def(
+        "__iter__",
+        [](Map &m) { return make_key_iterator(m.begin(), m.end()); },
+        keep_alive<0, 1>() /* Essential: keep map alive while iterator exists */
+    );
+
+    cl.def(
+        "keys",
+        [](Map &m) { return std::unique_ptr<KeysView>(new detail::KeysViewImpl<Map>(m)); },
+        keep_alive<0, 1>() /* Essential: keep map alive while view exists */
+    );
+
+    cl.def(
+        "values",
+        [](Map &m) { return std::unique_ptr<ValuesView>(new detail::ValuesViewImpl<Map>(m)); },
+        keep_alive<0, 1>() /* Essential: keep map alive while view exists */
+    );
+
+    cl.def(
+        "items",
+        [](Map &m) { return std::unique_ptr<ItemsView>(new detail::ItemsViewImpl<Map>(m)); },
+        keep_alive<0, 1>() /* Essential: keep map alive while view exists */
+    );
+
+    cl.def(
+        "__getitem__",
+        [](Map &m, const KeyType &k) -> MappedType & {
+            auto it = m.find(k);
+            if (it == m.end()) {
+                set_error(PyExc_KeyError, detail::format_message_key_error(k));
+                throw error_already_set();
+            }
+            return it->second;
+        },
+        return_value_policy::reference_internal // ref + keepalive
+    );
+
+    cl.def("__contains__", [](Map &m, const KeyType &k) -> bool {
+        auto it = m.find(k);
+        if (it == m.end()) {
+            return false;
+        }
+        return true;
+    });
+    // Fallback for when the object is not of the key type
+    cl.def("__contains__", [](Map &, const object &) -> bool { return false; });
+
+    // Assignment provided only if the type is copyable
+    detail::map_assignment<Map, Class_>(cl);
+
+    cl.def("__delitem__", [](Map &m, const KeyType &k) {
+        auto it = m.find(k);
+        if (it == m.end()) {
+            set_error(PyExc_KeyError, detail::format_message_key_error(k));
+            throw error_already_set();
+        }
+        m.erase(it);
+    });
+
+    // Always use a lambda in case of `using` declaration
+    cl.def("__len__", [](const Map &m) { return m.size(); });
+
+    return cl;
+}
+
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)

基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/subinterpreter.h

@@ -0,0 +1,462 @@
+/*
+    pybind11/subinterpreter.h: Support for creating and using subinterpreters
+
+    Copyright (c) 2025 The Pybind Development Team.
+
+    All rights reserved. Use of this source code is governed by a
+    BSD-style license that can be found in the LICENSE file.
+*/
+
+#pragma once
+
+#include "detail/common.h"
+#include "detail/internals.h"
+#include "gil.h"
+
+#include <stdexcept>
+
+#ifndef PYBIND11_HAS_SUBINTERPRETER_SUPPORT
+#    error "This platform does not support subinterpreters, do not include this file."
+#endif
+
+PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE)
+
+class subinterpreter;
+class subinterpreter_thread_state;
+
+/// Activate the subinterpreter and acquire its GIL, while also releasing any GIL and interpreter
+/// currently held. Upon exiting the scope, the previous subinterpreter (if any) and its
+/// associated GIL are restored to their state as they were before the scope was entered.
+///
+/// Two construction modes are supported:
+///
+/// 1. `subinterpreter_scoped_activate(subinterpreter const &)`:
+///    Transient mode (the default).  A fresh PyThreadState is created on entry and destroyed on
+///    exit.  This is the established behavior; existing code is unaffected.
+///
+/// 2. `subinterpreter_scoped_activate(subinterpreter_thread_state &)`:
+///    Reuse mode.  The PyThreadState owned by the given subinterpreter_thread_state is swapped
+///    in on entry and swapped out (but NOT destroyed) on exit, so repeated activations on the
+///    same OS thread reuse the same PyThreadState and preserve its per-thread interpreter state.
+///    Use this when a single OS thread re-enters one or more subinterpreters many times.
+class subinterpreter_scoped_activate {
+public:
+    explicit subinterpreter_scoped_activate(subinterpreter const &si);
+    explicit subinterpreter_scoped_activate(subinterpreter_thread_state &ts);
+    ~subinterpreter_scoped_activate();
+
+    subinterpreter_scoped_activate(subinterpreter_scoped_activate &&) = delete;
+    subinterpreter_scoped_activate(subinterpreter_scoped_activate const &) = delete;
+    subinterpreter_scoped_activate &operator=(subinterpreter_scoped_activate &) = delete;
+    subinterpreter_scoped_activate &operator=(subinterpreter_scoped_activate const &) = delete;
+
+private:
+    PyThreadState *old_tstate_ = nullptr;
+    PyThreadState *tstate_ = nullptr;
+    PyGILState_STATE gil_state_;
+    bool simple_gil_ = false;
+    // When true, tstate_ is owned by a subinterpreter_thread_state and must NOT be destroyed
+    // when this scope exits (only swapped out).
+    bool borrowed_ = false;
+};
+
+/// Holds a Python subinterpreter instance
+class subinterpreter {
+public:
+    /// empty/unusable, but move-assignable.  use create() to create a subinterpreter.
+    subinterpreter() = default;
+
+    subinterpreter(subinterpreter const &copy) = delete;
+    subinterpreter &operator=(subinterpreter const &copy) = delete;
+
+    subinterpreter(subinterpreter &&old) noexcept
+        : istate_(old.istate_), creation_tstate_(old.creation_tstate_) {
+        old.istate_ = nullptr;
+        old.creation_tstate_ = nullptr;
+    }
+
+    subinterpreter &operator=(subinterpreter &&old) noexcept {
+        std::swap(old.istate_, istate_);
+        std::swap(old.creation_tstate_, creation_tstate_);
+        return *this;
+    }
+
+    /// Create a new subinterpreter with the specified configuration
+    /// @note This function acquires (and then releases) the main interpreter GIL, but the main
+    /// interpreter and its GIL are not required to be held prior to calling this function.
+    static subinterpreter create(PyInterpreterConfig const &cfg) {
+
+        error_scope err_scope;
+        subinterpreter result;
+        {
+            // we must hold the main GIL in order to create a subinterpreter
+            subinterpreter_scoped_activate main_guard(main());
+
+            auto *prev_tstate = PyThreadState_Get();
+
+            PyStatus status;
+
+            {
+                /*
+                Several internal CPython modules are lacking proper subinterpreter support in 3.12
+                even though it is "stable" in that version.  This most commonly seems to cause
+                crashes when two interpreters concurrently initialize, which imports several things
+                (like builtins, unicode, codecs).
+                */
+#if PY_VERSION_HEX < 0x030D0000 && defined(Py_MOD_PER_INTERPRETER_GIL_SUPPORTED)
+                static std::mutex one_at_a_time;
+                std::lock_guard<std::mutex> guard(one_at_a_time);
+#endif
+                status = Py_NewInterpreterFromConfig(&result.creation_tstate_, &cfg);
+            }
+
+            // this doesn't raise a normal Python exception, it provides an exit() status code.
+            if (PyStatus_Exception(status) != 0) {
+                pybind11_fail("failed to create new sub-interpreter");
+            }
+
+            // upon success, the new interpreter is activated in this thread
+            result.istate_ = result.creation_tstate_->interp;
+            detail::has_seen_non_main_interpreter() = true;
+            detail::get_internals(); // initialize internals.tstate, amongst other things...
+
+            // In 3.13+ this state should be deleted right away, and the memory will be reused for
+            // the next threadstate on this interpreter. However, on 3.12 we cannot do that, we
+            // must keep it around (but not use it) ... see destructor.
+#if PY_VERSION_HEX >= 0x030D0000
+            PyThreadState_Clear(result.creation_tstate_);
+            PyThreadState_DeleteCurrent();
+#endif
+
+            // we have to switch back to main, and then the scopes will handle cleanup
+            PyThreadState_Swap(prev_tstate);
+        }
+        return result;
+    }
+
+    /// Calls create() with a default configuration of an isolated interpreter that disallows fork,
+    /// exec, and Python threads.
+    static subinterpreter create() {
+        // same as the default config in the python docs
+        PyInterpreterConfig cfg;
+        std::memset(&cfg, 0, sizeof(cfg));
+        cfg.allow_threads = 1;
+        cfg.check_multi_interp_extensions = 1;
+        cfg.gil = PyInterpreterConfig_OWN_GIL;
+        return create(cfg);
+    }
+
+    ~subinterpreter() {
+        if (!creation_tstate_) {
+            // non-owning wrapper, do nothing.
+            return;
+        }
+
+        PyThreadState *destroy_tstate = nullptr;
+        PyThreadState *old_tstate = nullptr;
+
+        // Python 3.12 requires us to keep the original PyThreadState alive until we are ready to
+        // destroy the interpreter.  We prefer to use that to destroy the interpreter.
+#if PY_VERSION_HEX < 0x030D0000
+        // The tstate passed to Py_EndInterpreter MUST have been created on the current OS thread.
+        bool same_thread = false;
+#    ifdef PY_HAVE_THREAD_NATIVE_ID
+        same_thread = PyThread_get_thread_native_id() == creation_tstate_->native_thread_id;
+#    endif
+        if (same_thread) {
+            // OK it is safe to use the creation state here
+            destroy_tstate = creation_tstate_;
+            old_tstate = PyThreadState_Swap(destroy_tstate);
+        } else {
+            // We have to make a new tstate on this thread and use that.
+            destroy_tstate = PyThreadState_New(istate_);
+            old_tstate = PyThreadState_Swap(destroy_tstate);
+
+            // We can use the one we just created, so we must delete the creation state.
+            PyThreadState_Clear(creation_tstate_);
+            PyThreadState_Delete(creation_tstate_);
+        }
+#else
+        destroy_tstate = PyThreadState_New(istate_);
+        old_tstate = PyThreadState_Swap(destroy_tstate);
+#endif
+
+        bool switch_back = (old_tstate != nullptr) && old_tstate->interp != istate_;
+
+        // Internals always exists in the subinterpreter, this class enforces it when it creates
+        // the subinterpreter. Even if it didn't, this only creates the pointer-to-pointer, not the
+        // internals themselves.
+        detail::get_internals_pp_manager().get_pp();
+        detail::get_local_internals_pp_manager().get_pp();
+
+        // End it
+        Py_EndInterpreter(destroy_tstate);
+
+        // It's possible for the  internals to be created during endinterpreter (e.g. if a
+        // py::capsule calls `get_internals()` during destruction), so we destroy afterward.
+        detail::get_internals_pp_manager().destroy();
+        detail::get_local_internals_pp_manager().destroy();
+
+        // switch back to the old tstate and old GIL (if there was one)
+        if (switch_back) {
+            PyThreadState_Swap(old_tstate);
+        }
+    }
+
+    /// Get a handle to the main interpreter that can be used with subinterpreter_scoped_activate
+    /// Note that destructing the handle is a noop, the main interpreter can only be ended by
+    /// py::finalize_interpreter()
+    static subinterpreter main() {
+        subinterpreter m;
+        m.istate_ = PyInterpreterState_Main();
+        m.disarm(); // make destruct a noop
+        return m;
+    }
+
+    /// Get a non-owning wrapper of the currently active interpreter (if any)
+    static subinterpreter current() {
+        subinterpreter c;
+        c.istate_ = detail::get_interpreter_state_unchecked();
+        c.disarm(); // make destruct a noop, we don't own this...
+        return c;
+    }
+
+    /// Get the numerical identifier for the sub-interpreter
+    int64_t id() const {
+        if (istate_ != nullptr) {
+            return PyInterpreterState_GetID(istate_);
+        }
+        return -1; // CPython uses one-up numbers from 0, so negative should be safe to return
+                   // here.
+    }
+
+    /// Get the interpreter's state dict.  This interpreter's GIL must be held before calling!
+    dict state_dict() { return reinterpret_borrow<dict>(PyInterpreterState_GetDict(istate_)); }
+
+    /// abandon cleanup of this subinterpreter (leak it). this might be needed during
+    /// finalization...
+    void disarm() { creation_tstate_ = nullptr; }
+
+    /// An empty wrapper cannot be activated
+    bool empty() const { return istate_ == nullptr; }
+
+    /// Is this wrapper non-empty
+    explicit operator bool() const { return !empty(); }
+
+private:
+    friend class subinterpreter_scoped_activate;
+    friend class subinterpreter_thread_state;
+    PyInterpreterState *istate_ = nullptr;
+    PyThreadState *creation_tstate_ = nullptr;
+};
+
+/// RAII wrapper that owns a PyThreadState bound to a specific subinterpreter on the OS thread
+/// that constructed it.  Intended to be held long-lived (e.g. as a `thread_local`, or inside a
+/// per-thread struct) so that many subinterpreter_scoped_activate scopes on the same OS thread
+/// can reuse a single PyThreadState instead of creating and destroying one each time.
+///
+/// The PyThreadState is created on construction in a *released* state: it is NOT made current,
+/// and no GIL is acquired.  Activation is the job of subinterpreter_scoped_activate.
+///
+/// A single OS thread can hold one of these per subinterpreter and alternate between them via
+/// subinterpreter_scoped_activate without churning PyThreadState objects.
+///
+/// Lifetime / threading requirements:
+///
+/// - Construction and destruction must happen on the SAME OS thread (a PyThreadState is bound
+///   to the OS thread that created it; deleting it on a different thread is undefined behavior).
+/// - The owning subinterpreter must still be alive when this object is destroyed.
+/// - This object must NOT be destroyed while a subinterpreter_scoped_activate referring to it is
+///   still alive (the activator holds a reference into it).
+///
+/// Typical usage:
+///
+/// @code
+///   thread_local py::subinterpreter_thread_state ts(sub);
+///   {
+///       py::subinterpreter_scoped_activate guard(ts);   // swap-in only
+///       // ... use the subinterpreter ...
+///   }                                                    // swap-out, tstate kept alive
+///   {
+///       py::subinterpreter_scoped_activate guard(ts);   // reuses the same PyThreadState
+///       // ...
+///   }
+/// @endcode
+class subinterpreter_thread_state {
+public:
+    /// Create a PyThreadState for `si` on the calling OS thread.  The new state is left in a
+    /// released state (not current, no GIL acquired).
+    explicit subinterpreter_thread_state(subinterpreter const &si);
+
+    /// Destroy the owned PyThreadState.  Must run on the same OS thread that constructed this
+    /// object, while the owning subinterpreter is still alive, and while no
+    /// subinterpreter_scoped_activate referring to this object is alive.
+    ~subinterpreter_thread_state();
+
+    subinterpreter_thread_state(subinterpreter_thread_state const &) = delete;
+    subinterpreter_thread_state(subinterpreter_thread_state &&) = delete;
+    subinterpreter_thread_state &operator=(subinterpreter_thread_state const &) = delete;
+    subinterpreter_thread_state &operator=(subinterpreter_thread_state &&) = delete;
+
+    /// The interpreter this thread state belongs to.
+    PyInterpreterState *interpreter_state() const { return istate_; }
+
+    /// The owned PyThreadState pointer; valid for the lifetime of this object.
+    PyThreadState *raw_thread_state() const { return tstate_; }
+
+private:
+    friend class subinterpreter_scoped_activate;
+    PyThreadState *tstate_ = nullptr;
+    PyInterpreterState *istate_ = nullptr;
+};
+
+class scoped_subinterpreter {
+public:
+    scoped_subinterpreter() : si_(subinterpreter::create()), scope_(si_) {}
+
+    explicit scoped_subinterpreter(PyInterpreterConfig const &cfg)
+        : si_(subinterpreter::create(cfg)), scope_(si_) {}
+
+private:
+    subinterpreter si_;
+    subinterpreter_scoped_activate scope_;
+};
+
+// --- subinterpreter_scoped_activate -----------------------------------------------------------
+
+inline subinterpreter_scoped_activate::subinterpreter_scoped_activate(subinterpreter const &si) {
+    if (!si.istate_) {
+        pybind11_fail("null subinterpreter");
+    }
+
+    if (detail::get_interpreter_state_unchecked() == si.istate_) {
+        // we are already on this interpreter, make sure we hold the GIL
+        simple_gil_ = true;
+        gil_state_ = PyGILState_Ensure();
+        return;
+    }
+
+    // we can't really interact with the interpreter at all until we switch to it
+    // not even to, for example, look in its state dict or touch its internals
+    tstate_ = PyThreadState_New(si.istate_);
+
+    // make the interpreter active and acquire the GIL
+    old_tstate_ = PyThreadState_Swap(tstate_);
+
+    // save this in internals for scoped_gil calls (see also: PR #5870)
+    detail::get_internals().tstate = tstate_;
+}
+
+inline subinterpreter_scoped_activate::subinterpreter_scoped_activate(
+    subinterpreter_thread_state &ts) {
+    if (ts.tstate_ == nullptr) {
+        pybind11_fail("subinterpreter_scoped_activate: empty subinterpreter_thread_state");
+    }
+
+    if (detail::get_interpreter_state_unchecked() == ts.istate_) {
+        // We are already on this interpreter -- e.g. nested activation, or a different
+        // PyThreadState for the same interpreter is already current on this thread.  Match the
+        // fast path of the (subinterpreter const&) overload: just ensure the GIL is held.  The
+        // `ts` argument's PyThreadState is intentionally NOT swapped to here; the already-current
+        // tstate keeps being used until the outer scope exits.
+        simple_gil_ = true;
+        gil_state_ = PyGILState_Ensure();
+        return;
+    }
+
+#if defined(PYBIND11_DETAILED_ERROR_MESSAGES)
+    {
+        // A PyThreadState is bound to its creating OS thread; it may only be activated there.
+        bool same_thread = true;
+#    ifdef PY_HAVE_THREAD_NATIVE_ID
+        same_thread = PyThread_get_thread_native_id() == ts.tstate_->native_thread_id;
+#    endif
+        if (!same_thread) {
+            pybind11_fail("subinterpreter_scoped_activate: a subinterpreter_thread_state must be "
+                          "activated on the same OS thread that constructed it!");
+        }
+    }
+#endif
+
+    tstate_ = ts.tstate_;
+    borrowed_ = true;
+
+    // make the interpreter active and acquire the GIL
+    old_tstate_ = PyThreadState_Swap(tstate_);
+
+    // save this in internals for scoped_gil calls (see also: PR #5870)
+    detail::get_internals().tstate = tstate_;
+}
+
+inline subinterpreter_scoped_activate::~subinterpreter_scoped_activate() {
+    if (simple_gil_) {
+        // We were on this interpreter already, so just make sure the GIL goes back as it was
+        PyGILState_Release(gil_state_);
+    } else {
+        if (tstate_) {
+#if defined(PYBIND11_DETAILED_ERROR_MESSAGES)
+            if (detail::get_thread_state_unchecked() != tstate_) {
+                pybind11_fail("~subinterpreter_scoped_activate: thread state must be current!");
+            }
+#endif
+            detail::get_internals().tstate.reset();
+            if (!borrowed_) {
+                PyThreadState_Clear(tstate_);
+                PyThreadState_DeleteCurrent();
+            }
+            // When borrowed_, tstate_ stays alive in its owning subinterpreter_thread_state for
+            // reuse; the PyThreadState_Swap below merely detaches it from this thread.
+        }
+
+        // Go back the previous interpreter (if any) and acquire THAT gil
+        PyThreadState_Swap(old_tstate_);
+    }
+}
+
+// --- subinterpreter_thread_state --------------------------------------------------------------
+
+inline subinterpreter_thread_state::subinterpreter_thread_state(subinterpreter const &si) {
+    if (!si.istate_) {
+        pybind11_fail("subinterpreter_thread_state: null subinterpreter");
+    }
+    istate_ = si.istate_;
+    // PyThreadState_New does not require holding any GIL and does not make the new state current.
+    tstate_ = PyThreadState_New(istate_);
+    if (tstate_ == nullptr) {
+        pybind11_fail("subinterpreter_thread_state: PyThreadState_New returned null");
+    }
+}
+
+inline subinterpreter_thread_state::~subinterpreter_thread_state() {
+    if (tstate_ == nullptr) {
+        return;
+    }
+#if defined(PYBIND11_DETAILED_ERROR_MESSAGES)
+    {
+        // A PyThreadState must be cleared and deleted on the OS thread that created it.
+        bool same_thread = true;
+#    ifdef PY_HAVE_THREAD_NATIVE_ID
+        same_thread = PyThread_get_thread_native_id() == tstate_->native_thread_id;
+#    endif
+        if (!same_thread) {
+            pybind11_fail("~subinterpreter_thread_state: must be destroyed on the same OS thread "
+                          "that constructed it!");
+        }
+    }
+#endif
+    // The PyThreadState must be made current to be cleared and deleted on the owning OS thread.
+    // Swap it in (which acquires the subinterpreter's GIL), clear+delete, then restore whatever
+    // was active before.
+    PyThreadState *prev = PyThreadState_Swap(tstate_);
+    PyThreadState_Clear(tstate_);
+    PyThreadState_DeleteCurrent();
+    // If `prev` is tstate_ itself, the user destroyed this object while it was active via a
+    // subinterpreter_scoped_activate -- a contract violation, but be defensive: do NOT swap back
+    // to a now-deleted pointer.  Leaving the thread with no current interpreter is consistent
+    // with the cached state having just been destroyed.
+    if (prev != nullptr && prev != tstate_) {
+        PyThreadState_Swap(prev);
+    }
+}
+
+PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE)