diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/Fused MoE Baseline入门：快速跑通最小闭环教程.md b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/Fused MoE Baseline入门：快速跑通最小闭环教程.md index 6d3b3ee..abd11e2 100644 --- a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/Fused MoE Baseline入门：快速跑通最小闭环教程.md +++ b/基于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 **操作：**切换到指定项目路径。 diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/.vscode/settings.json b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/.vscode/settings.json new file mode 100644 index 0000000..c497274 --- /dev/null +++ b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/.vscode/settings.json @@ -0,0 +1,3 @@ +{ + "cmake.sourceDirectory": "/root/Project/fusedmoe_v2/standalone/fused_moe_i8_tn" +} \ No newline at end of file diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/query/client-vscode/query.json b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/query/client-vscode/query.json new file mode 100644 index 0000000..82bb964 --- /dev/null +++ b/基于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}]} \ No newline at end of file diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/cache-v2-ea2ef11d05674d96d761.json b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/cache-v2-ea2ef11d05674d96d761.json new file mode 100644 index 0000000..3a5515b --- /dev/null +++ b/基于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 + } +} diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/cmakeFiles-v1-7899829d23c1c1ae3e98.json b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/cmakeFiles-v1-7899829d23c1c1ae3e98.json new file mode 100644 index 0000000..99000ff --- /dev/null +++ b/基于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 + } +} diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/codemodel-v2-7cd0b1b00876e71f364f.json b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/codemodel-v2-7cd0b1b00876e71f364f.json new file mode 100644 index 0000000..507753c --- /dev/null +++ b/基于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 + } +} diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/directory-.-Debug-f5ebdc15457944623624.json b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/directory-.-Debug-f5ebdc15457944623624.json new file mode 100644 index 0000000..3a67af9 --- /dev/null +++ b/基于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" : "." + } +} diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/index-2026-05-29T05-34-43-0302.json b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/index-2026-05-29T05-34-43-0302.json new file mode 100644 index 0000000..109ba15 --- /dev/null +++ b/基于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 + } + } + ] + } + } + } +} diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/target-build_fused_moe_i8_tn_example-Debug-a97a299baa6c6c6d83d0.json b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/target-build_fused_moe_i8_tn_example-Debug-a97a299baa6c6c6d83d0.json new file mode 100644 index 0000000..b2bb402 --- /dev/null +++ b/基于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" +} diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/target-run-Debug-0d66e135afa1376e0f20.json b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/target-run-Debug-0d66e135afa1376e0f20.json new file mode 100644 index 0000000..3a3883e --- /dev/null +++ b/基于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" +} diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/toolchains-v1-8ae3cf416ede58af34e6.json b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/toolchains-v1-8ae3cf416ede58af34e6.json new file mode 100644 index 0000000..2fdb147 --- /dev/null +++ b/基于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 + } +} diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeCache.txt b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeCache.txt new file mode 100644 index 0000000..d8b74ab --- /dev/null +++ b/基于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 + diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeFiles/3.28.3/CMakeSystem.cmake b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeFiles/3.28.3/CMakeSystem.cmake new file mode 100644 index 0000000..2e26099 --- /dev/null +++ b/基于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) diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeFiles/CMakeConfigureLog.yaml b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeFiles/CMakeConfigureLog.yaml new file mode 100644 index 0000000..17b4ac2 --- /dev/null +++ b/基于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 +... diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeFiles/TargetDirectories.txt b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeFiles/TargetDirectories.txt new file mode 100644 index 0000000..8336d29 --- /dev/null +++ b/基于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 diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeFiles/cmake.check_cache b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeFiles/cmake.check_cache new file mode 100644 index 0000000..3dccd73 --- /dev/null +++ b/基于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 diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeFiles/rules.ninja b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeFiles/rules.ninja new file mode 100644 index 0000000..ede407b --- /dev/null +++ b/基于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: + diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/build.ninja b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/build.ninja new file mode 100644 index 0000000..76ee09c --- /dev/null +++ b/基于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 diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/cmake_install.cmake b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/cmake_install.cmake new file mode 100644 index 0000000..9a85ad9 --- /dev/null +++ b/基于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}") diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/scripts/build_fused_moe_i8_tn_pybind.sh b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/scripts/build_fused_moe_i8_tn_pybind.sh new file mode 100644 index 0000000..785ed79 --- /dev/null +++ b/基于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" diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/scripts/run_fused_moe_i8_tn_benchmark.sh b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/scripts/run_fused_moe_i8_tn_benchmark.sh new file mode 100644 index 0000000..8469e64 --- /dev/null +++ b/基于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" "$@" diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/scripts/run_fused_moe_i8_tn_pybind_test.sh b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/scripts/run_fused_moe_i8_tn_pybind_test.sh new file mode 100644 index 0000000..72108bb --- /dev/null +++ b/基于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" "$@" diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_bf16_tn/CMakeLists.txt b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_bf16_tn/CMakeLists.txt new file mode 100644 index 0000000..ba191a2 --- /dev/null +++ b/基于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) diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_bf16_tn/Makefile b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_bf16_tn/Makefile new file mode 100644 index 0000000..1a343b6 --- /dev/null +++ b/基于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) diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_bf16_tn/src/fused_moe_bf16_tn_example.cpp b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_bf16_tn/src/fused_moe_bf16_tn_example.cpp new file mode 100644 index 0000000..bdce6ef --- /dev/null +++ b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_bf16_tn/src/fused_moe_bf16_tn_example.cpp @@ -0,0 +1,246 @@ +#include + +#include +#include +#include +#include +#include +#include + +#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; + +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(mctlass::bfloat16_t(value)); +} + +void fill_inputs(std::vector &a, + std::vector &b_col_major, + std::vector &moe_weights, + std::vector &token_ids, + std::vector &expert_ids) { + a.resize(static_cast(kNumTokens) * kK); + b_col_major.resize(static_cast(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(row) * kK + kk] = float_to_bf16(value); + } + token_ids[row] = row; + moe_weights[row] = 0.5f + 0.03125f * static_cast(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(expert) * kN + col) * kK + kk] = float_to_bf16(value); + } + } + } +} + +std::vector reference_fused_moe(const std::vector &a, + const std::vector &b_col_major, + const std::vector &moe_weights, + const std::vector &token_ids, + const std::vector &expert_ids) { + std::vector out(static_cast(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(token) * kK + kk]) * + bf16_to_float(b_col_major[(static_cast(expert) * kN + col) * kK + kk]); + } + out[static_cast(routed_row) * kN + col] = float_to_bf16(acc * moe_weight); + } + } + + return out; +} + +bool validate_result(const std::vector &got, const std::vector &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(first_bad / kN); + const int col = static_cast(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 host_a; + std::vector host_b; + std::vector host_moe_weights; + std::vector host_token_ids; + std::vector host_expert_ids; + std::vector host_num_tokens_post_padded(1, kEM); + std::vector host_output(static_cast(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 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(&dev_a), host_a.size() * sizeof(Bf16)), "mcMalloc(dev_a)"); + check_mc(mcMalloc(reinterpret_cast(&dev_b), host_b.size() * sizeof(Bf16)), "mcMalloc(dev_b)"); + check_mc(mcMalloc(reinterpret_cast(&dev_moe_weights), host_moe_weights.size() * sizeof(float)), + "mcMalloc(dev_moe_weights)"); + check_mc(mcMalloc(reinterpret_cast(&dev_token_ids), host_token_ids.size() * sizeof(int)), + "mcMalloc(dev_token_ids)"); + check_mc(mcMalloc(reinterpret_cast(&dev_expert_ids), host_expert_ids.size() * sizeof(int)), + "mcMalloc(dev_expert_ids)"); + check_mc(mcMalloc(reinterpret_cast(&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(&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; +} diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/CMakeLists.txt b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/CMakeLists.txt new file mode 100644 index 0000000..897693f --- /dev/null +++ b/基于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) diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/Makefile b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/Makefile new file mode 100644 index 0000000..869704f --- /dev/null +++ b/基于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) diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/benchmark_results.md b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/benchmark_results.md new file mode 100644 index 0000000..4fe76e0 --- /dev/null +++ b/基于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 diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_pybind.o b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_pybind.o new file mode 100644 index 0000000..ac296d8 Binary files /dev/null and b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_pybind.o differ diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_pybind.so b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_pybind.so new file mode 100644 index 0000000..e55d11c Binary files /dev/null and b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_pybind.so differ diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/fused-moe-i8-tn-pybind-triton-guide.md b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/fused-moe-i8-tn-pybind-triton-guide.md new file mode 100644 index 0000000..aabe813 --- /dev/null +++ b/基于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` 的性能差距 diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/__pycache__/benchmark_fused_moe_i8_tn.cpython-312.pyc b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/__pycache__/benchmark_fused_moe_i8_tn.cpython-312.pyc new file mode 100644 index 0000000..d8d9f6f Binary files /dev/null and b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/__pycache__/benchmark_fused_moe_i8_tn.cpython-312.pyc differ diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/__pycache__/fused_moe_i8_tn_triton.cpython-310.pyc b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/__pycache__/fused_moe_i8_tn_triton.cpython-310.pyc new file mode 100644 index 0000000..8f49683 Binary files /dev/null and b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/__pycache__/fused_moe_i8_tn_triton.cpython-310.pyc differ diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/__pycache__/fused_moe_i8_tn_triton.cpython-312.pyc b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/__pycache__/fused_moe_i8_tn_triton.cpython-312.pyc new file mode 100644 index 0000000..3092674 Binary files /dev/null and b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/__pycache__/fused_moe_i8_tn_triton.cpython-312.pyc differ diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/__pycache__/test_fused_moe_i8_tn_pybind.cpython-310.pyc b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/__pycache__/test_fused_moe_i8_tn_pybind.cpython-310.pyc new file mode 100644 index 0000000..14ca558 Binary files /dev/null and b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/__pycache__/test_fused_moe_i8_tn_pybind.cpython-310.pyc differ diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/__pycache__/test_fused_moe_i8_tn_pybind.cpython-312.pyc b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/__pycache__/test_fused_moe_i8_tn_pybind.cpython-312.pyc new file mode 100644 index 0000000..073dbb6 Binary files /dev/null and b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/__pycache__/test_fused_moe_i8_tn_pybind.cpython-312.pyc differ diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/benchmark_fused_moe_i8_tn.py b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/benchmark_fused_moe_i8_tn.py new file mode 100644 index 0000000..587c4f2 --- /dev/null +++ b/基于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() diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/fused_moe_i8_tn_triton.py b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/fused_moe_i8_tn_triton.py new file mode 100644 index 0000000..de905cd --- /dev/null +++ b/基于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() diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/test_fused_moe_i8_tn_pybind.py b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/test_fused_moe_i8_tn_pybind.py new file mode 100644 index 0000000..516cdcd --- /dev/null +++ b/基于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() diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_example.cpp b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_example.cpp new file mode 100644 index 0000000..b58c7b5 --- /dev/null +++ b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_example.cpp @@ -0,0 +1,324 @@ +#include + +#include +#include +#include +#include +#include +#include + +#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 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(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(m) * static_cast(n) * static_cast(k); + return operations / (static_cast(avg_ms) * 1.0e9); +} + +void fill_inputs(const CaseConfig &cfg, + std::vector &a, + std::vector &b_col_major, + std::vector &scale_a, + std::vector &scale_b, + std::vector &moe_weights, + std::vector &token_ids, + std::vector &expert_ids) { + const int total_rows = cfg.num_tokens * cfg.topk; + + a.resize(static_cast(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(row) * kK + kk] = + static_cast(((row * 13 + kk * 7 + cfg.topk * 5 + 3) % 11) - 5); + } + scale_a[row] = 0.125f + 0.015625f * static_cast((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((routed_row + cfg.topk) % 5); + } + + b_col_major.resize(static_cast(kNumExperts) * kN * kK); + scale_b.resize(static_cast(kNumExperts) * kN); + for (int expert = 0; expert < kNumExperts; ++expert) { + for (int col = 0; col < kN; ++col) { + scale_b[static_cast(expert) * kN + col] = + 0.25f + 0.03125f * static_cast((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(expert) * kN + col) * kK + kk] = + static_cast(value); + } + } + } +} + +std::vector reference_fused_moe(const CaseConfig &cfg, + const std::vector &a, + const std::vector &b_col_major, + const std::vector &scale_a, + const std::vector &scale_b, + const std::vector &moe_weights, + const std::vector &token_ids, + const std::vector &expert_ids) { + const int total_rows = cfg.num_tokens * cfg.topk; + std::vector out(static_cast(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(a[static_cast(token) * kK + kk]); + const int32_t rhs = + static_cast(b_col_major[(static_cast(expert) * kN + col) * kK + kk]); + acc += lhs * rhs; + } + const float scaled = static_cast(acc) * row_scale * + scale_b[static_cast(expert) * kN + col]; + out[static_cast(routed_row) * kN + col] = fused_moe_i8_tn::float_to_bf16(scaled); + } + } + + return out; +} + +bool validate_result(const std::vector &got, + const std::vector &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(first_bad / kN); + const int col = static_cast(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 +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(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 host_a; + std::vector host_b; + std::vector host_scale_a; + std::vector host_scale_b; + std::vector host_moe_weights; + std::vector host_token_ids; + std::vector 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 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 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; +} diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_kernel.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_kernel.h new file mode 100644 index 0000000..8e09e6b --- /dev/null +++ b/基于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 +#include + +#include + +#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 +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 +__global__ void direct_moe_kernel(typename DirectMoeKernel::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{}, Int{}), 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(&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{}, Int{}), + make_stride(Int{}, Int<1>{})); + Tensor sB = make_tensor(make_smem_ptr((ElementB *)smem_B), + make_shape(Int{}, Int{}), + make_stride(Int{}, 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(&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(&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(&weights[i]) + j) = + __builtin_mxc_ldg_b32_predicator(const_cast(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(&a_scale[i]) + j) = + __builtin_mxc_ldg_b32_predicator(const_cast(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(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(&b_scale[0])[0], a_scale_f2, zero2); + scale[1] = __builtin_mxc_pk_fma_f32(reinterpret_cast(&b_scale[0])[1], a_scale_f2, zero2); + scale[2] = __builtin_mxc_pk_fma_f32(reinterpret_cast(&b_scale[1])[0], a_scale_f2, zero2); + scale[3] = __builtin_mxc_pk_fma_f32(reinterpret_cast(&b_scale[1])[1], a_scale_f2, zero2); + *reinterpret_cast(&out[0]) = + __builtin_mxc_pk_fma_f32(*reinterpret_cast(&out[0]), scale[0], zero2); + *reinterpret_cast(&out[2]) = + __builtin_mxc_pk_fma_f32(*reinterpret_cast(&out[2]), scale[1], zero2); + *reinterpret_cast(&out[4]) = + __builtin_mxc_pk_fma_f32(*reinterpret_cast(&out[4]), scale[2], zero2); + *reinterpret_cast(&out[6]) = + __builtin_mxc_pk_fma_f32(*reinterpret_cast(&out[6]), scale[3], zero2); + + CVT_F32_TO_BF16(tempC[0], reinterpret_cast(&out)[0], reinterpret_cast(&out)[1]); + CVT_F32_TO_BF16(tempC[1], reinterpret_cast(&out)[2], reinterpret_cast(&out)[3]); + __builtin_mxc_stg_b64_predicator(Caddr + rowC[i * 4 + j] * args.problem_size.n() + colC[0], + 0, + *(reinterpret_cast(&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(&out)[4], reinterpret_cast(&out)[5]); + CVT_F32_TO_BF16(tempC[1], reinterpret_cast(&out)[6], reinterpret_cast(&out)[7]); + __builtin_mxc_stg_b64_predicator(Caddr + rowC[i * 4 + j] * args.problem_size.n() + colC[1], + 0, + *(reinterpret_cast(&tempC)), + true, + false, + false, + (rowC[i * 4 + j] < args.problem_size.m()) && colC_mask[1], + 1, + MACA_ICMP_EQ); + } + } +} + +template +using DirectMoeGemmKernel = DirectMoeKernel; + +template +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 +inline Status launch(typename Kernel::Arguments const &args, mcStream_t stream = nullptr) { + dim3 const block(kThreadCount, 1, 1); + dim3 const grid = get_grid_shape(args); + direct_moe_kernel<<>>(args); + return Status::kSuccess; +} + +} // namespace fused_moe_i8_tn diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_macros.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_macros.h new file mode 100644 index 0000000..4977f6a --- /dev/null +++ b/基于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(&(dst)) = *reinterpret_cast(&(src)); \ + FUSED_MOE_CP_ASYNC_FENC() + +#define FUSED_MOE_STS(dst, src, type_) \ + FUSED_MOE_CP_ASYNC_FENC(); \ + *reinterpret_cast(&(dst)) = *reinterpret_cast(&(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 diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_pybind.cpp b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_pybind.cpp new file mode 100644 index 0000000..692bb3e --- /dev/null +++ b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_pybind.cpp @@ -0,0 +1,101 @@ +#include + +#include +#include + +#include +#include +#include +#include + +#include "fused_moe_i8_tn_runner.h" + +namespace py = pybind11; + +namespace { + +template +std::vector copy_1d_array(const py::array &array, const char *name) { + auto view = py::array_t::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(view.data()); + return std::vector(ptr, ptr + view.size()); +} + +std::vector copy_2d_int8_array(const py::array &array, const char *name) { + auto view = py::array_t::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(view.data()); + return std::vector(ptr, ptr + view.size()); +} + +py::array_t 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::ensure(a); + auto b_view = py::array_t::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(a_view.shape(0)); + inputs.k = static_cast(a_view.shape(1)); + inputs.num_experts = num_experts; + inputs.n = static_cast(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(b_view.data()), + static_cast(b_view.data()) + b_view.size()); + inputs.scale_a = copy_1d_array(scale_a, "scale_a"); + inputs.scale_b = copy_1d_array(scale_b, "scale_b"); + inputs.moe_weights = copy_1d_array(moe_weights, "moe_weights"); + inputs.token_ids = copy_1d_array(token_ids, "token_ids"); + inputs.expert_ids = copy_1d_array(expert_ids, "expert_ids"); + + fused_moe_i8_tn::RunResult result = fused_moe_i8_tn::run_fused_moe_i8_tn(inputs, device_id); + std::vector output = fused_moe_i8_tn::bf16_vector_to_float(result.output); + + py::array_t 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); +} diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_runner.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_runner.h new file mode 100644 index 0000000..fb55783 --- /dev/null +++ b/基于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 + +#include +#include +#include +#include +#include +#include + +#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 a; + std::vector b_col_major; + std::vector scale_a; + std::vector scale_b; + std::vector moe_weights; + std::vector token_ids; + std::vector expert_ids; +}; + +struct RunResult { + int rows = 0; + int cols = 0; + std::vector output; +}; + +inline float bf16_to_float(BFloat16 value) { + return static_cast(value); +} + +inline BFloat16 float_to_bf16(float value) { + return BFloat16(value); +} + +inline std::vector bf16_vector_to_float(const std::vector &input) { + std::vector 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(inputs.em); + const size_t expected_a = static_cast(inputs.num_tokens) * inputs.k; + const size_t expected_b = static_cast(inputs.num_experts) * inputs.n * inputs.k; + const size_t expected_scale_a = static_cast(inputs.num_tokens); + const size_t expected_scale_b = static_cast(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(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 +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 host_num_tokens_post_padded(1, inputs.em); + std::vector host_output(static_cast(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(&dev_a), inputs.a.size() * sizeof(int8_t)), "mcMalloc(dev_a)"); + check_mc_or_throw(mcMalloc(reinterpret_cast(&dev_b), inputs.b_col_major.size() * sizeof(int8_t)), "mcMalloc(dev_b)"); + check_mc_or_throw(mcMalloc(reinterpret_cast(&dev_scale_a), inputs.scale_a.size() * sizeof(float)), + "mcMalloc(dev_scale_a)"); + check_mc_or_throw(mcMalloc(reinterpret_cast(&dev_scale_b), inputs.scale_b.size() * sizeof(float)), + "mcMalloc(dev_scale_b)"); + check_mc_or_throw(mcMalloc(reinterpret_cast(&dev_moe_weights), inputs.moe_weights.size() * sizeof(float)), + "mcMalloc(dev_moe_weights)"); + check_mc_or_throw(mcMalloc(reinterpret_cast(&dev_token_ids), inputs.token_ids.size() * sizeof(int)), + "mcMalloc(dev_token_ids)"); + check_mc_or_throw(mcMalloc(reinterpret_cast(&dev_expert_ids), inputs.expert_ids.size() * sizeof(int)), + "mcMalloc(dev_expert_ids)"); + check_mc_or_throw(mcMalloc(reinterpret_cast(&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(&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; + auto const args = make_direct_arguments(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(args), "direct_moe_kernel"); + } else { + using GemmKernel = DirectMoeKernel; + auto const args = make_direct_arguments(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(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 diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_types.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_types.h new file mode 100644 index 0000000..b4d0f32 --- /dev/null +++ b/基于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 +#include + +#include +#include +#include + +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(tmp); +#else + uint32_t bits; + std::memcpy(&bits, &x, sizeof(bits)); + bits += ((bits >> 16) & 1) + 0x7fff; + storage = static_cast(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(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 ¶ms) + : scale_a_(params.scale_a), scale_b_(params.scale_b), moe_weights_(params.moe_weights) {} +}; + +} // namespace fused_moe_i8_tn diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/i8_tn_256x256x128_raw_arrays/CMakeLists.txt b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/i8_tn_256x256x128_raw_arrays/CMakeLists.txt new file mode 100644 index 0000000..ab62ec2 --- /dev/null +++ b/基于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) diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/i8_tn_256x256x128_raw_arrays/include/gemm_i8_tn_256x256x128_raw_arrays.hpp b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/i8_tn_256x256x128_raw_arrays/include/gemm_i8_tn_256x256x128_raw_arrays.hpp new file mode 100644 index 0000000..9983e9b --- /dev/null +++ b/基于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 +__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(reinterpret_cast( + a_ptr + static_cast(row) * K + k_tile_base + ctx.a_load_k[Pass])), + 0, + -1, + true, + true, + false, + true); +} + +template +__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(reinterpret_cast( + b_ptr + static_cast(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(&a_pack); + AbTypeI8Mma const *b0_frag = reinterpret_cast(&b0_pack); + AbTypeI8Mma const *b1_frag = reinterpret_cast(&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(smem_b + ctx.b_lds_offset[column_group][col_block][0][0]), + LdsTypeI8Mma); + STANDALONE_LDS( + b1_pack_low[column_group][col_block], + *const_cast(smem_b + ctx.b_lds_offset[column_group][col_block][0][1]), + LdsTypeI8Mma); + STANDALONE_LDS( + b0_pack_high[column_group][col_block], + *const_cast(smem_b + ctx.b_lds_offset[column_group][col_block][1][0]), + LdsTypeI8Mma); + STANDALONE_LDS( + b1_pack_high[column_group][col_block], + *const_cast(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(smem_a + ctx.a_lds_offset[row_block][0]), LdsTypeI8Mma); + STANDALONE_LDS(a_pack_high, *const_cast(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(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(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(bidz) * M * K; + const int8_t *b_ptr = B + static_cast(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 diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/i8_tn_256x256x128_raw_arrays/include/standalone_maca_kernel_utils.hpp b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/i8_tn_256x256x128_raw_arrays/include/standalone_maca_kernel_utils.hpp new file mode 100644 index 0000000..87be72e --- /dev/null +++ b/基于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 + +#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(&(dst)) = *reinterpret_cast(&(src)); \ + standalone_cp_async_fenc() diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/i8_tn_256x256x128_raw_arrays/src/test.cpp b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/i8_tn_256x256x128_raw_arrays/src/test.cpp new file mode 100644 index 0000000..fe67aa0 --- /dev/null +++ b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/i8_tn_256x256x128_raw_arrays/src/test.cpp @@ -0,0 +1,229 @@ +#include + +#include +#include +#include +#include +#include + +#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 &a, int m, int k) { + a.resize(static_cast(m) * k); + for (int row = 0; row < m; ++row) { + for (int kk = 0; kk < k; ++kk) { + a[static_cast(row) * k + kk] = + static_cast(((row * 13 + kk * 7 + 5) % 9) - 4); + } + } +} + +void fill_col_major_b(std::vector &b, int n, int k) { + b.resize(static_cast(n) * k); + for (int col = 0; col < n; ++col) { + for (int kk = 0; kk < k; ++kk) { + b[static_cast(col) * k + kk] = + static_cast(((kk * 11 + col * 5 + 3) % 7) - 3); + } + } +} + +std::vector reference_gemm_tn(const std::vector &a, + const std::vector &b_col_major, + int m, + int n, + int k) { + std::vector out(static_cast(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(a[static_cast(row) * k + kk]) * + static_cast(b_col_major[static_cast(col) * k + kk]); + } + out[static_cast(row) * n + col] = acc; + } + } + return out; +} + +bool run_case(int m, int n, int k, const char *tag) { + std::vector host_a; + std::vector host_b; + std::vector host_d(static_cast(m) * n, -1); + fill_row_major_a(host_a, m, k); + fill_col_major_b(host_b, n, k); + const std::vector 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(&dev_a), host_a.size() * sizeof(int8_t)), "mcMalloc(dev_a)"); + check_mc(mcMalloc(reinterpret_cast(&dev_b), host_b.size() * sizeof(int8_t)), "mcMalloc(dev_b)"); + check_mc(mcMalloc(reinterpret_cast(&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<<>>(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(first_bad / n); + const int col = static_cast(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(m) * n * k; + return operations / (static_cast(avg_ms) * 1.0e9); +} + +BenchmarkResult run_benchmark(int m, int n, int k, int warmup_iterations, int measured_iterations) { + std::vector host_a; + std::vector host_b; + std::vector host_d(static_cast(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(&dev_a), host_a.size() * sizeof(int8_t)), "mcMalloc(dev_a)"); + check_mc(mcMalloc(reinterpret_cast(&dev_b), host_b.size() * sizeof(int8_t)), "mcMalloc(dev_b)"); + check_mc(mcMalloc(reinterpret_cast(&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<<>>(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<<>>(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(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; +} diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/attr.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/attr.h new file mode 100644 index 0000000..d337595 --- /dev/null +++ b/基于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 + + 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 + +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 +struct base { + + PYBIND11_DEPRECATED( + "base() was deprecated in favor of specifying 'T' as a template argument to class_") + base() = default; +}; + +/// Keep patient alive while nurse lives +template +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; + + 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()); + + is equivalent to the following pseudocode: + + .. code-block:: cpp + + m.def("foo", [](args...) { + T scope_guard; + return foo(args...); // forwarded arguments + }); + \endrst */ +template +struct call_guard; + +template <> +struct call_guard<> { + using type = detail::void_type; +}; + +template +struct call_guard { + static_assert(std::is_default_constructible::value, + "The guard type must be default constructible"); + + using type = T; +}; + +template +struct call_guard { + struct type { + T guard{}; // Compose multiple guard types with left-to-right default-constructor order + typename call_guard::type next{}; + }; +}; + +/// @} annotations + +PYBIND11_NAMESPACE_BEGIN(detail) +/* Forward declarations */ +enum op_id : int; +enum op_type : int; +struct undefined_t; +template +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 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(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 +struct process_attribute; + +template +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 : process_attribute_default { + static void init(const name &n, function_record *r) { r->name = const_cast(n.value); } +}; + +/// Process an attribute specifying the function's docstring +template <> +struct process_attribute : process_attribute_default { + static void init(const doc &n, function_record *r) { r->doc = const_cast(n.value); } +}; + +/// Process an attribute specifying the function's docstring (provided as a C-style string) +template <> +struct process_attribute : process_attribute_default { + static void init(const char *d, function_record *r) { r->doc = const_cast(d); } + static void init(const char *d, type_record *r) { r->doc = d; } +}; +template <> +struct process_attribute : process_attribute {}; + +/// Process an attribute indicating the function's return value policy +template <> +struct process_attribute : process_attribute_default { + 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 : process_attribute_default { + 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 : process_attribute_default { + 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 : process_attribute_default { + 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 : process_attribute_default { + 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 : process_attribute_default { + static void init(const is_operator &, function_record *r) { r->is_operator = true; } +}; + +template <> +struct process_attribute + : process_attribute_default { + 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 : process_attribute_default { + 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 : process_attribute_default { + 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 : process_attribute_default { + static void init(const kw_only &, function_record *r) { + append_self_arg_if_needed(r); + if (r->has_args && r->nargs_pos != static_cast(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(r->args.size()); + } +}; + +/// Process a positional-only-argument maker +template <> +struct process_attribute : process_attribute_default { + static void init(const pos_only &, function_record *r) { + append_self_arg_if_needed(r); + r->nargs_pos_only = static_cast(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 +struct process_attribute::value>> + : process_attribute_default { + 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 +struct process_attribute> : process_attribute_default> { + static void init(const base &, type_record *r) { r->add_base(typeid(T), nullptr); } +}; + +/// Process a multiple inheritance attribute +template <> +struct process_attribute : process_attribute_default { + static void init(const multiple_inheritance &, type_record *r) { + r->multiple_inheritance = true; + } +}; + +template <> +struct process_attribute : process_attribute_default { + static void init(const dynamic_attr &, type_record *r) { r->dynamic_attr = true; } +}; + +template <> +struct process_attribute { + static void init(const custom_type_setup &value, type_record *r) { + r->custom_type_setup_callback = value.value; + } +}; + +template <> +struct process_attribute : process_attribute_default { + static void init(const is_final &, type_record *r) { r->is_final = true; } +}; + +template <> +struct process_attribute : process_attribute_default { + static void init(const buffer_protocol &, type_record *r) { r->buffer_protocol = true; } +}; + +template <> +struct process_attribute : process_attribute_default { + static void init(const metaclass &m, type_record *r) { r->metaclass = m.value; } +}; + +template <> +struct process_attribute : process_attribute_default { + static void init(const module_local &l, type_record *r) { r->module_local = l.value; } +}; + +template <> +struct process_attribute + : process_attribute_default { + 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 : process_attribute_default { + static void init(const prepend &, function_record *r) { r->prepend = true; } +}; + +/// Process an 'arithmetic' attribute for enums (does nothing here) +template <> +struct process_attribute : process_attribute_default {}; + +template +struct process_attribute> : process_attribute_default> {}; + +/** + * 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 +struct process_attribute> + : public process_attribute_default> { + template = 0> + static void precall(function_call &call) { + keep_alive_impl(Nurse, Patient, call, handle()); + } + template = 0> + static void postcall(function_call &, handle) {} + template = 0> + static void precall(function_call &) {} + template = 0> + static void postcall(function_call &call, handle ret) { + keep_alive_impl(Nurse, Patient, call, ret); + } +}; + +/// Recursively iterate over variadic template arguments +template +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::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::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::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::type>::postcall(call, fn_ret), 0)...}; + } +}; + +template +struct is_keep_alive : std::false_type {}; + +template +struct is_keep_alive> : std::true_type {}; + +template +using is_call_guard = is_instantiation; + +/// Extract the ``type`` from the first `call_guard` in `Extras...` (or `void_type` if none found) +template +using extract_guard_t = typename exactly_one_t, Extra...>::type; + +/// Check the number of named arguments at compile time +template ::value...), + size_t self = constexpr_sum(std::is_same::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(has_args) + static_cast(has_kwargs)) + == nargs; +} + +PYBIND11_NAMESPACE_END(detail) +PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE) diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/buffer_info.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/buffer_info.h new file mode 100644 index 0000000..10fa825 --- /dev/null +++ b/基于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 + + 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 c_strides(const std::vector &shape, ssize_t itemsize) { + auto ndim = shape.size(); + std::vector 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 f_strides(const std::vector &shape, ssize_t itemsize) { + auto ndim = shape.size(); + std::vector strides(ndim, itemsize); + for (size_t i = 1; i < ndim; ++i) { + strides[i] = strides[i - 1] * shape[i - 1]; + } + return strides; +} + +template +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::format() + ssize_t ndim = 0; // Number of dimensions + std::vector shape; // Shape of the tensor (1 entry per dimension) + std::vector 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 shape_in, + detail::any_container 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(shape.size()) + || ndim != static_cast(strides.size())) { + pybind11_fail("buffer_info: ndim doesn't match shape and/or strides length"); + } + for (size_t i = 0; i < static_cast(ndim); ++i) { + size *= shape[i]; + } + } + + template + buffer_info(T *ptr, + detail::any_container shape_in, + detail::any_container strides_in, + bool readonly = false) + : buffer_info(private_ctr_tag(), + ptr, + sizeof(T), + format_descriptor::format(), + static_cast(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 + buffer_info(T *ptr, ssize_t size, bool readonly = false) + : buffer_info(ptr, sizeof(T), format_descriptor::format(), size, readonly) {} + + template + buffer_info(const T *ptr, ssize_t size, bool readonly = true) + : buffer_info( + const_cast(ptr), sizeof(T), format_descriptor::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(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(b)` and + // `buffer_info::item_type_is_equivalent_to(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 + bool item_type_is_equivalent_to() const { + return detail::compare_buffer_info::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 &&shape_in, + detail::any_container &&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 +struct compare_buffer_info { + static bool compare(const buffer_info &b) { + // NOLINTNEXTLINE(bugprone-sizeof-expression) Needed for `PyObject *` + return b.format == format_descriptor::format() && b.itemsize == (ssize_t) sizeof(T); + } +}; + +template +struct compare_buffer_info::value>> { + static bool compare(const buffer_info &b) { + return static_cast(b.itemsize) == sizeof(T) + && (b.format == format_descriptor::value + || ((sizeof(T) == sizeof(long)) + && b.format == (std::is_unsigned::value ? "L" : "l")) + || ((sizeof(T) == sizeof(size_t)) + && b.format == (std::is_unsigned::value ? "N" : "n"))); + } +}; + +PYBIND11_NAMESPACE_END(detail) +PYBIND11_NAMESPACE_END(PYBIND11_NAMESPACE) diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/cast.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/cast.h new file mode 100644 index 0000000..b7a4c2b --- /dev/null +++ b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/third_party/pybind11/include/pybind11/cast.h @@ -0,0 +1,2447 @@ +/* + pybind11/cast.h: Partial template specializations to cast between + C++ and Python types + + Copyright (c) 2016 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/argument_vector.h" +#include "detail/common.h" +#include "detail/descr.h" +#include "detail/holder_caster_foreign_helpers.h" +#include "detail/native_enum_data.h" +#include "detail/type_caster_base.h" +#include "detail/typeid.h" +#include "pytypes.h" + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +PYBIND11_NAMESPACE_BEGIN(PYBIND11_NAMESPACE) + +PYBIND11_WARNING_DISABLE_MSVC(4127) + +PYBIND11_NAMESPACE_BEGIN(detail) + +template +class type_caster : public type_caster_base {}; +template +using make_caster = type_caster>; + +// Shortcut for calling a caster's `cast_op_type` cast operator for casting a type_caster to a T +template +typename make_caster::template cast_op_type cast_op(make_caster &caster) { + using result_t = typename make_caster::template cast_op_type; // See PR #4893 + return caster.operator result_t(); +} +template +typename make_caster::template cast_op_type::type> +cast_op(make_caster &&caster) { + using result_t = typename make_caster::template cast_op_type< + typename std::add_rvalue_reference::type>; // See PR #4893 + return std::move(caster).operator result_t(); +} + +template +class type_caster_enum_type { +private: + using Underlying = typename std::underlying_type::type; + +public: + static constexpr auto name = const_name(); + + template + static handle cast(SrcType &&src, return_value_policy, handle parent) { + handle native_enum + = global_internals_native_enum_type_map_get_item(std::type_index(typeid(EnumType))); + if (native_enum) { + return native_enum(static_cast(src)).release(); + } + return type_caster_base::cast( + std::forward(src), + // Fixes https://github.com/pybind/pybind11/pull/3643#issuecomment-1022987818: + return_value_policy::copy, + parent); + } + + template + static handle cast(SrcType *src, return_value_policy policy, handle parent) { + return cast(*src, policy, parent); + } + + bool load(handle src, bool convert) { + handle native_enum + = global_internals_native_enum_type_map_get_item(std::type_index(typeid(EnumType))); + if (native_enum) { + if (!isinstance(src, native_enum)) { + return false; + } + type_caster underlying_caster; + if (!underlying_caster.load(src.attr("value"), convert)) { + pybind11_fail("native_enum internal consistency failure."); + } + native_value = static_cast(static_cast(underlying_caster)); + native_loaded = true; + return true; + } + + type_caster_base legacy_caster; + if (legacy_caster.load(src, convert)) { + legacy_ptr = static_cast(legacy_caster); + return true; + } + return false; + } + + template + using cast_op_type = detail::cast_op_type; + + // NOLINTNEXTLINE(google-explicit-constructor) + operator EnumType *() { return native_loaded ? &native_value : legacy_ptr; } + + // NOLINTNEXTLINE(google-explicit-constructor) + operator EnumType &() { + if (!native_loaded && !legacy_ptr) { + throw reference_cast_error(); + } + return native_loaded ? native_value : *legacy_ptr; + } + +private: + EnumType native_value; // if loading a py::native_enum + bool native_loaded = false; + EnumType *legacy_ptr = nullptr; // if loading a py::enum_ +}; + +template +struct type_caster_enum_type_enabled : std::true_type {}; + +template +struct type_uses_type_caster_enum_type { + static constexpr bool value + = std::is_enum::value && type_caster_enum_type_enabled::value; +}; + +template +class type_caster::value>> + : public type_caster_enum_type {}; + +template ::value, int> = 0> +bool isinstance_native_enum_impl(handle obj, const std::type_info &tp) { + handle native_enum = global_internals_native_enum_type_map_get_item(tp); + if (!native_enum) { + return false; + } + return isinstance(obj, native_enum); +} + +template ::value, int> = 0> +bool isinstance_native_enum_impl(handle, const std::type_info &) { + return false; +} + +template +bool isinstance_native_enum(handle obj, const std::type_info &tp) { + return isinstance_native_enum_impl>(obj, tp); +} + +template +class type_caster> { +private: + using caster_t = make_caster; + caster_t subcaster; + using reference_t = type &; + using subcaster_cast_op_type = typename caster_t::template cast_op_type; + + static_assert( + std::is_same::type &, subcaster_cast_op_type>::value + || std::is_same::value, + "std::reference_wrapper caster requires T to have a caster with an " + "`operator T &()` or `operator const T &()`"); + +public: + bool load(handle src, bool convert) { return subcaster.load(src, convert); } + static constexpr auto name = caster_t::name; + static handle + cast(const std::reference_wrapper &src, return_value_policy policy, handle parent) { + // It is definitely wrong to take ownership of this pointer, so mask that rvp + if (policy == return_value_policy::take_ownership + || policy == return_value_policy::automatic) { + policy = return_value_policy::automatic_reference; + } + return caster_t::cast(&src.get(), policy, parent); + } + template + using cast_op_type = std::reference_wrapper; + explicit operator std::reference_wrapper() { return cast_op(subcaster); } +}; + +#define PYBIND11_TYPE_CASTER(type, py_name) \ +protected: \ + type value; \ + \ +public: \ + static constexpr auto name = py_name; \ + template >::value, \ + int> = 0> \ + static ::pybind11::handle cast( \ + T_ *src, ::pybind11::return_value_policy policy, ::pybind11::handle parent) { \ + if (!src) \ + return ::pybind11::none().release(); \ + if (policy == ::pybind11::return_value_policy::take_ownership) { \ + auto h = cast(std::move(*src), policy, parent); \ + delete src; \ + return h; \ + } \ + return cast(*src, policy, parent); \ + } \ + operator type *() { return &value; } /* NOLINT(bugprone-macro-parentheses) */ \ + operator type &() { return value; } /* NOLINT(bugprone-macro-parentheses) */ \ + operator type &&() && { return std::move(value); } /* NOLINT(bugprone-macro-parentheses) */ \ + template \ + using cast_op_type = ::pybind11::detail::movable_cast_op_type + +template +using is_std_char_type = any_of, /* std::string */ +#if defined(PYBIND11_HAS_U8STRING) + std::is_same, /* std::u8string */ +#endif + std::is_same, /* std::u16string */ + std::is_same, /* std::u32string */ + std::is_same /* std::wstring */ + >; + +template +struct type_caster::value && !is_std_char_type::value>> { + using _py_type_0 = conditional_t; + using _py_type_1 = conditional_t::value, + _py_type_0, + typename std::make_unsigned<_py_type_0>::type>; + using py_type = conditional_t::value, double, _py_type_1>; + +public: + bool load(handle src, bool convert) { + py_type py_value; + + if (!src) { + return false; + } + + if (std::is_floating_point::value) { + if (convert || PyFloat_Check(src.ptr()) || PYBIND11_LONG_CHECK(src.ptr())) { + py_value = (py_type) PyFloat_AsDouble(src.ptr()); + } else { + return false; + } + } else if (PyFloat_Check(src.ptr()) + || !(convert || PYBIND11_LONG_CHECK(src.ptr()) + || PYBIND11_INDEX_CHECK(src.ptr()))) { + // Explicitly reject float → int conversion even in convert mode. + // This prevents silent truncation (e.g., 1.9 → 1). + // Only int → float conversion is allowed (widening, no precision loss). + // Also reject if none of the conversion conditions are met. + return false; + } else { + handle src_or_index = src; + // PyPy: 7.3.7's 3.8 does not implement PyLong_*'s __index__ calls. +#if defined(PYPY_VERSION) + object index; + // If not a PyLong, we need to call PyNumber_Index explicitly on PyPy. + // When convert is false, we only reach here if PYBIND11_INDEX_CHECK passed above. + if (!PYBIND11_LONG_CHECK(src.ptr())) { + index = reinterpret_steal(PyNumber_Index(src.ptr())); + if (!index) { + PyErr_Clear(); + if (!convert) + return false; + } else { + src_or_index = index; + } + } +#endif + if (std::is_unsigned::value) { + py_value = as_unsigned(src_or_index.ptr()); + } else { // signed integer: + py_value = sizeof(T) <= sizeof(long) + ? (py_type) PyLong_AsLong(src_or_index.ptr()) + : (py_type) PYBIND11_LONG_AS_LONGLONG(src_or_index.ptr()); + } + } + + bool py_err = (PyErr_Occurred() != nullptr); + if (py_err) { + assert(py_value == static_cast(-1)); + } + + // Check to see if the conversion is valid (integers should match exactly) + // Signed/unsigned checks happen elsewhere + if (py_err + || (std::is_integral::value && sizeof(py_type) != sizeof(T) + && py_value != (py_type) (T) py_value)) { + PyErr_Clear(); + if (py_err && convert && (PyNumber_Check(src.ptr()) != 0)) { + auto tmp = reinterpret_steal(std::is_floating_point::value + ? PyNumber_Float(src.ptr()) + : PyNumber_Long(src.ptr())); + PyErr_Clear(); + return load(tmp, false); + } + return false; + } + + value = (T) py_value; + return true; + } + + template + static typename std::enable_if::value, handle>::type + cast(U src, return_value_policy /* policy */, handle /* parent */) { + return PyFloat_FromDouble((double) src); + } + + template + static typename std::enable_if::value && std::is_signed::value + && (sizeof(U) <= sizeof(long)), + handle>::type + cast(U src, return_value_policy /* policy */, handle /* parent */) { + return PYBIND11_LONG_FROM_SIGNED((long) src); + } + + template + static typename std::enable_if::value && std::is_unsigned::value + && (sizeof(U) <= sizeof(unsigned long)), + handle>::type + cast(U src, return_value_policy /* policy */, handle /* parent */) { + return PYBIND11_LONG_FROM_UNSIGNED((unsigned long) src); + } + + template + static typename std::enable_if::value && std::is_signed::value + && (sizeof(U) > sizeof(long)), + handle>::type + cast(U src, return_value_policy /* policy */, handle /* parent */) { + return PyLong_FromLongLong((long long) src); + } + + template + static typename std::enable_if::value && std::is_unsigned::value + && (sizeof(U) > sizeof(unsigned long)), + handle>::type + cast(U src, return_value_policy /* policy */, handle /* parent */) { + return PyLong_FromUnsignedLongLong((unsigned long long) src); + } + + PYBIND11_TYPE_CASTER( + T, + io_name::value>("typing.SupportsInt | typing.SupportsIndex", + "int", + "typing.SupportsFloat | typing.SupportsIndex", + "float")); +}; + +template +struct void_caster { +public: + bool load(handle src, bool) { + if (src && src.is_none()) { + return true; + } + return false; + } + static handle cast(T, return_value_policy /* policy */, handle /* parent */) { + return none().release(); + } + PYBIND11_TYPE_CASTER(T, const_name("None")); +}; + +template <> +class type_caster : public void_caster {}; + +template <> +class type_caster : public type_caster { +public: + using type_caster::cast; + + bool load(handle h, bool) { + if (!h) { + return false; + } + if (h.is_none()) { + value = nullptr; + return true; + } + + /* Check if this is a capsule */ + if (isinstance(h)) { + value = reinterpret_borrow(h); + return true; + } + + /* Check if this is a C++ type */ + const auto &bases + = all_type_info(reinterpret_cast(type::handle_of(h).ptr())); + if (bases.size() == 1) { // Only allowing loading from a single-value type + value = values_and_holders(reinterpret_cast(h.ptr())).begin()->value_ptr(); + return true; + } + + /* Fail */ + return false; + } + + static handle cast(const void *ptr, return_value_policy /* policy */, handle /* parent */) { + if (ptr) { + return capsule(ptr).release(); + } + return none().release(); + } + + template + using cast_op_type = void *&; + explicit operator void *&() { return value; } + static constexpr auto name = const_name(PYBIND11_CAPSULE_TYPE_TYPE_HINT); + +private: + void *value = nullptr; +}; + +template <> +class type_caster : public void_caster {}; + +template <> +class type_caster { +public: + bool load(handle src, bool convert) { + if (!src) { + return false; + } + if (src.ptr() == Py_True) { + value = true; + return true; + } + if (src.ptr() == Py_False) { + value = false; + return true; + } + if (convert || is_numpy_bool(src)) { + // (allow non-implicit conversion for numpy booleans), use strncmp + // since NumPy 1.x had an additional trailing underscore. + + Py_ssize_t res = -1; + if (src.is_none()) { + res = 0; // None is implicitly converted to False + } +#if defined(PYPY_VERSION) + // On PyPy, check that "__bool__" attr exists + else if (hasattr(src, PYBIND11_BOOL_ATTR)) { + res = PyObject_IsTrue(src.ptr()); + } +#else + // Alternate approach for CPython: this does the same as the above, but optimized + // using the CPython API so as to avoid an unneeded attribute lookup. + else if (auto *tp_as_number = Py_TYPE(src.ptr())->tp_as_number) { + if (PYBIND11_NB_BOOL(tp_as_number)) { + res = (*PYBIND11_NB_BOOL(tp_as_number))(src.ptr()); + } + } +#endif + if (res == 0 || res == 1) { + value = (res != 0); + return true; + } + PyErr_Clear(); + } + return false; + } + static handle cast(bool src, return_value_policy /* policy */, handle /* parent */) { + return handle(src ? Py_True : Py_False).inc_ref(); + } + PYBIND11_TYPE_CASTER(bool, const_name("bool")); + +private: + // Test if an object is a NumPy boolean (without fetching the type). + static bool is_numpy_bool(handle object) { + const char *type_name = Py_TYPE(object.ptr())->tp_name; + // Name changed to `numpy.bool` in NumPy 2, `numpy.bool_` is needed for 1.x support + return std::strcmp("numpy.bool", type_name) == 0 + || std::strcmp("numpy.bool_", type_name) == 0; + } +}; + +// Helper class for UTF-{8,16,32} C++ stl strings: +template +struct string_caster { + using CharT = typename StringType::value_type; + + // Simplify life by being able to assume standard char sizes (the standard only guarantees + // minimums, but Python requires exact sizes) + static_assert(!std::is_same::value || sizeof(CharT) == 1, + "Unsupported char size != 1"); +#if defined(PYBIND11_HAS_U8STRING) + static_assert(!std::is_same::value || sizeof(CharT) == 1, + "Unsupported char8_t size != 1"); +#endif + static_assert(!std::is_same::value || sizeof(CharT) == 2, + "Unsupported char16_t size != 2"); + static_assert(!std::is_same::value || sizeof(CharT) == 4, + "Unsupported char32_t size != 4"); + // wchar_t can be either 16 bits (Windows) or 32 (everywhere else) + static_assert(!std::is_same::value || sizeof(CharT) == 2 || sizeof(CharT) == 4, + "Unsupported wchar_t size != 2/4"); + static constexpr size_t UTF_N = 8 * sizeof(CharT); + + bool load(handle src, bool) { + handle load_src = src; + if (!src) { + return false; + } + if (!PyUnicode_Check(load_src.ptr())) { + return load_raw(load_src); + } + + // For UTF-8 we avoid the need for a temporary `bytes` object by using + // `PyUnicode_AsUTF8AndSize`. + if (UTF_N == 8) { + Py_ssize_t size = -1; + const auto *buffer + = reinterpret_cast(PyUnicode_AsUTF8AndSize(load_src.ptr(), &size)); + if (!buffer) { + PyErr_Clear(); + return false; + } + value = StringType(buffer, static_cast(size)); + return true; + } + + auto utfNbytes + = reinterpret_steal(PyUnicode_AsEncodedString(load_src.ptr(), + UTF_N == 8 ? "utf-8" + : UTF_N == 16 ? "utf-16" + : "utf-32", + nullptr)); + if (!utfNbytes) { + PyErr_Clear(); + return false; + } + + const auto *buffer + = reinterpret_cast(PYBIND11_BYTES_AS_STRING(utfNbytes.ptr())); + size_t length = static_cast(PYBIND11_BYTES_SIZE(utfNbytes.ptr())) / sizeof(CharT); + // Skip BOM for UTF-16/32 + if (UTF_N > 8) { + buffer++; + length--; + } + value = StringType(buffer, length); + + // If we're loading a string_view we need to keep the encoded Python object alive: + if (IsView) { + loader_life_support::add_patient(utfNbytes); + } + + return true; + } + + static handle + cast(const StringType &src, return_value_policy /* policy */, handle /* parent */) { + const char *buffer = reinterpret_cast(src.data()); + auto nbytes = ssize_t(src.size() * sizeof(CharT)); + handle s = decode_utfN(buffer, nbytes); + if (!s) { + throw error_already_set(); + } + return s; + } + + PYBIND11_TYPE_CASTER(StringType, const_name(PYBIND11_STRING_NAME)); + +private: + static handle decode_utfN(const char *buffer, ssize_t nbytes) { +#if !defined(PYPY_VERSION) + return UTF_N == 8 ? PyUnicode_DecodeUTF8(buffer, nbytes, nullptr) + : UTF_N == 16 ? PyUnicode_DecodeUTF16(buffer, nbytes, nullptr, nullptr) + : PyUnicode_DecodeUTF32(buffer, nbytes, nullptr, nullptr); +#else + // PyPy segfaults when on PyUnicode_DecodeUTF16 (and possibly on PyUnicode_DecodeUTF32 as + // well), so bypass the whole thing by just passing the encoding as a string value, which + // works properly: + return PyUnicode_Decode(buffer, + nbytes, + UTF_N == 8 ? "utf-8" + : UTF_N == 16 ? "utf-16" + : "utf-32", + nullptr); +#endif + } + + // When loading into a std::string or char*, accept a bytes/bytearray object as-is (i.e. + // without any encoding/decoding attempt). For other C++ char sizes this is a no-op. + // which supports loading a unicode from a str, doesn't take this path. + template + bool load_raw(enable_if_t::value, handle> src) { + if (PYBIND11_BYTES_CHECK(src.ptr())) { + // We were passed raw bytes; accept it into a std::string or char* + // without any encoding attempt. + const char *bytes = PYBIND11_BYTES_AS_STRING(src.ptr()); + if (!bytes) { + pybind11_fail("Unexpected PYBIND11_BYTES_AS_STRING() failure."); + } + value = StringType(bytes, (size_t) PYBIND11_BYTES_SIZE(src.ptr())); + return true; + } + if (PyByteArray_Check(src.ptr())) { + // We were passed a bytearray; accept it into a std::string or char* + // without any encoding attempt. + const char *bytearray = PyByteArray_AsString(src.ptr()); + if (!bytearray) { + pybind11_fail("Unexpected PyByteArray_AsString() failure."); + } + value = StringType(bytearray, (size_t) PyByteArray_Size(src.ptr())); + return true; + } + + return false; + } + + template + bool load_raw(enable_if_t::value, handle>) { + return false; + } +}; + +template +struct type_caster, + enable_if_t::value>> + : string_caster> {}; + +#ifdef PYBIND11_HAS_STRING_VIEW +template +struct type_caster, + enable_if_t::value>> + : string_caster, true> {}; +#endif + +// Type caster for C-style strings. We basically use a std::string type caster, but also add the +// ability to use None as a nullptr char* (which the string caster doesn't allow). +template +struct type_caster::value>> { + using StringType = std::basic_string; + using StringCaster = make_caster; + StringCaster str_caster; + bool none = false; + CharT one_char = 0; + +public: + bool load(handle src, bool convert) { + if (!src) { + return false; + } + if (src.is_none()) { + // Defer accepting None to other overloads (if we aren't in convert mode): + if (!convert) { + return false; + } + none = true; + return true; + } + return str_caster.load(src, convert); + } + + static handle cast(const CharT *src, return_value_policy policy, handle parent) { + if (src == nullptr) { + return pybind11::none().release(); + } + return StringCaster::cast(StringType(src), policy, parent); + } + + static handle cast(CharT src, return_value_policy policy, handle parent) { + if (std::is_same::value) { + handle s = PyUnicode_DecodeLatin1((const char *) &src, 1, nullptr); + if (!s) { + throw error_already_set(); + } + return s; + } + return StringCaster::cast(StringType(1, src), policy, parent); + } + + explicit operator CharT *() { + return none ? nullptr : const_cast(static_cast(str_caster).c_str()); + } + explicit operator CharT &() { + if (none) { + throw value_error("Cannot convert None to a character"); + } + + auto &value = static_cast(str_caster); + size_t str_len = value.size(); + if (str_len == 0) { + throw value_error("Cannot convert empty string to a character"); + } + + // If we're in UTF-8 mode, we have two possible failures: one for a unicode character that + // is too high, and one for multiple unicode characters (caught later), so we need to + // figure out how long the first encoded character is in bytes to distinguish between these + // two errors. We also allow want to allow unicode characters U+0080 through U+00FF, as + // those can fit into a single char value. + if (StringCaster::UTF_N == 8 && str_len > 1 && str_len <= 4) { + auto v0 = static_cast(value[0]); + // low bits only: 0-127 + // 0b110xxxxx - start of 2-byte sequence + // 0b1110xxxx - start of 3-byte sequence + // 0b11110xxx - start of 4-byte sequence + size_t char0_bytes = (v0 & 0x80) == 0 ? 1 + : (v0 & 0xE0) == 0xC0 ? 2 + : (v0 & 0xF0) == 0xE0 ? 3 + : 4; + + if (char0_bytes == str_len) { + // If we have a 128-255 value, we can decode it into a single char: + if (char0_bytes == 2 && (v0 & 0xFC) == 0xC0) { // 0x110000xx 0x10xxxxxx + one_char = static_cast(((v0 & 3) << 6) + + (static_cast(value[1]) & 0x3F)); + return one_char; + } + // Otherwise we have a single character, but it's > U+00FF + throw value_error("Character code point not in range(0x100)"); + } + } + + // UTF-16 is much easier: we can only have a surrogate pair for values above U+FFFF, thus a + // surrogate pair with total length 2 instantly indicates a range error (but not a "your + // string was too long" error). + else if (StringCaster::UTF_N == 16 && str_len == 2) { + one_char = static_cast(value[0]); + if (one_char >= 0xD800 && one_char < 0xE000) { + throw value_error("Character code point not in range(0x10000)"); + } + } + + if (str_len != 1) { + throw value_error("Expected a character, but multi-character string found"); + } + + one_char = value[0]; + return one_char; + } + + static constexpr auto name = const_name(PYBIND11_STRING_NAME); + template + using cast_op_type = pybind11::detail::cast_op_type<_T>; +}; + +// Base implementation for std::tuple and std::pair +template