feat(ascendc): add Huawei Ascend C toolchain support

Add toolchain, language, rules, and compiler modules for the
bisheng compiler driver from CANN (Compute Architecture for
Neural Networks) toolkit. Supports .asc (Ascend C kernel) and
.aicpu (AI-CPU operator) source kinds with binary/static/shared
target types.

- Toolchain: SDK auto-detection via ASCEND_HOME_PATH env or --sdk
- Compiler: bisheng module inherits from clang, adds add_sourceflags
  for explicit sourcekind overrides
- Rules: ascendc.env (default C++17), npuarchs (--npu-arch), build
- Language: check_main detection, add_ascnpuarchs API
- Tests: unit tests for check_main, project tests for libs/mixed

Fixes #7506

21 files changed, 910 insertions, 3 deletions

diff --git a/tests/modules/language/ascendc/test.lua b/tests/modules/language/ascendc/test.lua
new file mode 100644
index 000000000..507bff9ec
--- /dev/null
+++ b/tests/modules/language/ascendc/test.lua
@@ -0,0 +1,34 @@
+import("core.language.language")
+
+function _write_asc_file(name, source)
+    local sourcefile = os.tmpfile(name) .. ".asc"
+    io.writefile(sourcefile, source)
+    return sourcefile
+end
+
+function test_check_main(t)
+    local instance = language.load("ascendc")
+    t:require(instance)
+
+    local check_main = instance:get("check_main")
+    t:require(check_main)
+
+    local mainfile = _write_asc_file("ascendc_check_main", [[
+#include "acl/acl.h"
+
+int32_t main(int argc, char const *argv[]) {
+    return 0;
+}
+]])
+    t:are_equal(check_main(mainfile), true)
+
+    local commentfile = _write_asc_file("ascendc_check_main_comment", [[
+// int main() { return 0; }
+/*
+int main() { return 1; }
+*/
+__global__ __vector__ void kernel() {
+}
+]])
+    t:are_equal(check_main(commentfile), false)
+end
diff --git a/tests/projects/ascendc/libs/src/lib.asc b/tests/projects/ascendc/libs/src/lib.asc
new file mode 100644
index 000000000..c0e0dad5d
--- /dev/null
+++ b/tests/projects/ascendc/libs/src/lib.asc
@@ -0,0 +1,5 @@
+#include "kernel_operator.h"
+
+__global__ __vector__ void hello_world() {
+    AscendC::printf("[Block (%lu/%lu)]: Hello World!!!\n", AscendC::GetBlockIdx(), AscendC::GetBlockNum());
+}
diff --git a/tests/projects/ascendc/libs/src/main.asc b/tests/projects/ascendc/libs/src/main.asc
new file mode 100644
index 000000000..a8390825c
--- /dev/null
+++ b/tests/projects/ascendc/libs/src/main.asc
@@ -0,0 +1,21 @@
+#include "kernel_operator.h"
+#include "acl/acl.h"
+
+extern __global__ __vector__ void hello_world();
+
+int32_t main(int argc, char const *argv[]) {
+    aclInit(nullptr);
+    int32_t deviceId = 0;
+    aclrtSetDevice(deviceId);
+    aclrtStream stream = nullptr;
+    aclrtCreateStream(&stream);
+
+    constexpr uint32_t numBlocks = 8;
+    hello_world<<<numBlocks, nullptr, stream>>>();
+    aclrtSynchronizeStream(stream);
+
+    aclrtDestroyStream(stream);
+    aclrtResetDevice(deviceId);
+    aclFinalize();
+    return 0;
+}
diff --git a/tests/projects/ascendc/libs/test.lua b/tests/projects/ascendc/libs/test.lua
new file mode 100644
index 000000000..0a36e62e4
--- /dev/null
+++ b/tests/projects/ascendc/libs/test.lua
@@ -0,0 +1,9 @@
+function main(t)
+    if not is_host("linux") then
+        return t:skip("wrong host platform")
+    end
+    if not os.getenv("ASCEND_HOME_PATH") then
+        return t:skip("ASCEND_HOME_PATH not set")
+    end
+    t:build()
+end
diff --git a/tests/projects/ascendc/libs/xmake.lua b/tests/projects/ascendc/libs/xmake.lua
new file mode 100644
index 000000000..90ca89b1b
--- /dev/null
+++ b/tests/projects/ascendc/libs/xmake.lua
@@ -0,0 +1,17 @@
+add_rules("mode.debug", "mode.release")
+
+target("ascendc_static")
+    set_kind("static")
+    add_files("src/lib.asc")
+    add_ascnpuarchs("dav-2201")
+
+target("ascendc_shared")
+    set_kind("shared")
+    add_files("src/lib.asc")
+    add_ascnpuarchs("dav-2201")
+
+target("ascendc_libs_bin")
+    add_deps("ascendc_static", "ascendc_shared")
+    set_kind("binary")
+    add_files("src/main.asc")
+    add_ascnpuarchs("dav-2201")
diff --git a/tests/projects/ascendc/mixed_binary/src/helper.aicpu b/tests/projects/ascendc/mixed_binary/src/helper.aicpu
new file mode 100644
index 000000000..a11afdd90
--- /dev/null
+++ b/tests/projects/ascendc/mixed_binary/src/helper.aicpu
@@ -0,0 +1,6 @@
+#include "aicpu_api.h"
+
+__global__ __aicpu__ uint32_t hello_world(void *args) {
+    AscendC::printf("Hello World!!!\n");
+    return 0;
+}
diff --git a/tests/projects/ascendc/mixed_binary/src/main.asc b/tests/projects/ascendc/mixed_binary/src/main.asc
new file mode 100644
index 000000000..9bc0999bd
--- /dev/null
+++ b/tests/projects/ascendc/mixed_binary/src/main.asc
@@ -0,0 +1,25 @@
+#include "acl/acl.h"
+
+struct KernelArgs {
+    int mode;
+};
+
+extern __global__ __aicpu__ uint32_t hello_world(void *args);
+
+int32_t main(int argc, char const *argv[]) {
+    aclInit(nullptr);
+    int32_t deviceId = 0;
+    aclrtSetDevice(deviceId);
+    aclrtStream stream = nullptr;
+    aclrtCreateStream(&stream);
+
+    struct KernelArgs args = {0};
+    constexpr uint32_t numBlocks = 1;
+    hello_world<<<numBlocks, nullptr, stream>>>(&args, sizeof(KernelArgs));
+    aclrtSynchronizeStream(stream);
+
+    aclrtDestroyStream(stream);
+    aclrtResetDevice(deviceId);
+    aclFinalize();
+    return 0;
+}
diff --git a/tests/projects/ascendc/mixed_binary/test.lua b/tests/projects/ascendc/mixed_binary/test.lua
new file mode 100644
index 000000000..0a36e62e4
--- /dev/null
+++ b/tests/projects/ascendc/mixed_binary/test.lua
@@ -0,0 +1,9 @@
+function main(t)
+    if not is_host("linux") then
+        return t:skip("wrong host platform")
+    end
+    if not os.getenv("ASCEND_HOME_PATH") then
+        return t:skip("ASCEND_HOME_PATH not set")
+    end
+    t:build()
+end
diff --git a/tests/projects/ascendc/mixed_binary/xmake.lua b/tests/projects/ascendc/mixed_binary/xmake.lua
new file mode 100644
index 000000000..046da8404
--- /dev/null
+++ b/tests/projects/ascendc/mixed_binary/xmake.lua
@@ -0,0 +1,6 @@
+add_rules("mode.debug", "mode.release")
+
+target("ascendc_mixed")
+    set_kind("binary")
+    add_files("src/main.asc", "src/helper.aicpu")
+    add_ascnpuarchs("dav-2201")
diff --git a/xmake/core/tool/toolchain.lua b/xmake/core/tool/toolchain.lua
index f6704cf28..3bca7e37a 100644
--- a/xmake/core/tool/toolchain.lua
+++ b/xmake/core/tool/toolchain.lua
@@ -514,6 +514,10 @@ function _instance:_description(toolkind)
             cu         = "the cuda compiler",
             culd       = "the cuda linker",
             cuccbin    = "the cuda host c++ compiler",
+            asc        = "the ascend c compiler",
+            aicpu      = "the ascend c ai-cpu compiler",
+            ascld      = "the ascend c linker",
+            ascsh      = "the ascend c shared library linker",
             nc         = "the nim compiler",
             ncld       = "the nim linker",
             ncsh       = "the nim shared library linker",
diff --git a/xmake/languages/ascendc/check_main.lua b/xmake/languages/ascendc/check_main.lua
new file mode 100644
index 000000000..0588ba48e
--- /dev/null
+++ b/xmake/languages/ascendc/check_main.lua
@@ -0,0 +1,38 @@
+--!A cross-platform build utility based on Lua
+--
+-- Licensed under the Apache License, Version 2.0 (the "License");
+-- you may not use this file except in compliance with the License.
+-- You may obtain a copy of the License at
+--
+--     http://www.apache.org/licenses/LICENSE-2.0
+--
+-- Unless required by applicable law or agreed to in writing, software
+-- distributed under the License is distributed on an "AS IS" BASIS,
+-- WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+-- See the License for the specific language governing permissions and
+-- limitations under the License.
+--
+-- Copyright (C) 2015-present, Xmake Open Source Community.
+--
+-- @author      wuzhenqing
+-- @file        check_main.lua
+--
+
+-- check it
+function main(sourcefile)
+
+    -- load source code
+    local sourcecode = io.readfile(sourcefile)
+
+    -- remove comment first
+    sourcecode = sourcecode:gsub("/%*.-%*/", "")
+    sourcecode = sourcecode:gsub("//[^\n]*", "")
+
+    -- find int main(int argc, char** argv) {}
+    if sourcecode:find("%s+main%s*%(.-%)") then
+        return true
+    end
+
+    -- no main function
+    return false
+end
diff --git a/xmake/languages/ascendc/load.lua b/xmake/languages/ascendc/load.lua
new file mode 100644
index 000000000..bdcca0eba
--- /dev/null
+++ b/xmake/languages/ascendc/load.lua
@@ -0,0 +1,97 @@
+--!A cross-platform build utility based on Lua
+--
+-- Licensed under the Apache License, Version 2.0 (the "License");
+-- you may not use this file except in compliance with the License.
+-- You may obtain a copy of the License at
+--
+--     http://www.apache.org/licenses/LICENSE-2.0
+--
+-- Unless required by applicable law or agreed to in writing, software
+-- distributed under the License is distributed on an "AS IS" BASIS,
+-- WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+-- See the License for the specific language governing permissions and
+-- limitations under the License.
+--
+-- Copyright (C) 2015-present, Xmake Open Source Community.
+--
+-- @author      wuzhenqing
+-- @file        load.lua
+--
+
+-- get apis
+function _get_apis()
+    local apis = {}
+    apis.values = {
+        -- target.add_xxx
+        "target.add_links"
+    ,   "target.add_syslinks"
+    ,   "target.add_ascnpuarchs"
+    ,   "target.add_ascflags"
+    ,   "target.add_aicpuflags"
+    ,   "target.add_ldflags"
+    ,   "target.add_arflags"
+    ,   "target.add_shflags"
+    ,   "target.add_defines"
+    ,   "target.add_undefines"
+        -- @note do not translate path, it's usually an absolute path or contains
+        -- $ORIGIN/@loader_path
+    ,   "target.add_rpathdirs"
+        -- option.add_xxx
+    ,   "option.add_links"
+    ,   "option.add_syslinks"
+    ,   "option.add_ascnpuarchs"
+    ,   "option.add_ascflags"
+    ,   "option.add_aicpuflags"
+    ,   "option.add_ldflags"
+    ,   "option.add_arflags"
+    ,   "option.add_shflags"
+    ,   "option.add_defines"
+    ,   "option.add_undefines"
+    ,   "option.add_rpathdirs"
+        -- package.add_xxx
+    ,   "package.add_links"
+    ,   "package.add_syslinks"
+    ,   "package.add_ascflags"
+    ,   "package.add_aicpuflags"
+    ,   "package.add_ldflags"
+    ,   "package.add_arflags"
+    ,   "package.add_shflags"
+    ,   "package.add_defines"
+    ,   "package.add_undefines"
+    ,   "package.add_rpathdirs"
+        -- toolchain.add_xxx
+    ,   "toolchain.add_links"
+    ,   "toolchain.add_syslinks"
+    ,   "toolchain.add_ascnpuarchs"
+    ,   "toolchain.add_ascflags"
+    ,   "toolchain.add_aicpuflags"
+    ,   "toolchain.add_ldflags"
+    ,   "toolchain.add_arflags"
+    ,   "toolchain.add_shflags"
+    ,   "toolchain.add_defines"
+    ,   "toolchain.add_undefines"
+    ,   "toolchain.add_rpathdirs"
+    }
+    apis.groups = {
+        -- target.add_xxx
+        "target.add_linkorders"
+    ,   "target.add_linkgroups"
+        -- package.add_xxx
+    ,   "package.add_linkorders"
+    ,   "package.add_linkgroups"
+    }
+    apis.paths = {
+        -- target.add_xxx
+        "target.add_linkdirs"
+    ,   "target.add_includedirs"
+    ,   "target.add_sysincludedirs"
+        -- option.add_xxx
+    ,   "option.add_linkdirs"
+    ,   "option.add_sysincludedirs"
+    }
+    return apis
+end
+
+function main()
+    return {apis = _get_apis()}
+end
diff --git a/xmake/languages/ascendc/xmake.lua b/xmake/languages/ascendc/xmake.lua
new file mode 100644
index 000000000..27774c16a
--- /dev/null
+++ b/xmake/languages/ascendc/xmake.lua
@@ -0,0 +1,112 @@
+--!A cross-platform build utility based on Lua
+--
+-- Licensed under the Apache License, Version 2.0 (the "License");
+-- you may not use this file except in compliance with the License.
+-- You may obtain a copy of the License at
+--
+--     http://www.apache.org/licenses/LICENSE-2.0
+--
+-- Unless required by applicable law or agreed to in writing, software
+-- distributed under the License is distributed on an "AS IS" BASIS,
+-- WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+-- See the License for the specific language governing permissions and
+-- limitations under the License.
+--
+-- Copyright (C) 2015-present, Xmake Open Source Community.
+--
+-- @author      wuzhenqing
+-- @file        xmake.lua
+--
+
+language("ascendc")
+    add_rules("ascendc")
+    set_sourcekinds {asc = ".asc", aicpu = ".aicpu"}
+    set_sourceflags {asc = "ascflags", aicpu = "aicpuflags"}
+    set_targetkinds {binary = "ascld", static = "ar", shared = "ascsh"}
+    set_targetflags {binary = "ldflags", static = "arflags", shared = "shflags"}
+    set_langkinds {ascendc = "asc"}
+    set_mixingkinds("asc", "aicpu", "cc", "cxx", "as")
+
+    on_load("load")
+    on_check_main("check_main")
+
+    set_nameflags {
+        object = {
+            "config.includedirs"
+        ,   "target.runtimes"
+        ,   "target.symbols"
+        ,   "target.warnings"
+        ,   "target.optimize:check"
+        ,   "target.vectorexts:check"
+        ,   "target.languages"
+        ,   "target.includedirs"
+        ,   "target.defines"
+        ,   "target.undefines"
+        ,   "toolchain.includedirs"
+        ,   "toolchain.defines"
+        ,   "toolchain.undefines"
+        ,   "target.sysincludedirs"
+        ,   "toolchain.sysincludedirs"
+        }
+    ,   binary = {
+            "target.runtimes"
+        ,   "config.linkdirs"
+        ,   "target.linkdirs"
+        ,   "target.rpathdirs"
+        ,   "target.strip"
+        ,   "target.symbols"
+        ,   "target.optimize:check"
+        ,   "toolchain.linkdirs"
+        ,   "toolchain.rpathdirs"
+        ,   "config.links"
+            -- move it before target.links to keep package/dependency order correct
+        ,   "target.linkgroups"
+        ,   "target.links"
+        ,   "toolchain.links"
+        ,   "config.syslinks"
+        ,   "target.syslinks"
+        ,   "toolchain.syslinks"
+        }
+    ,   shared = {
+            "target.runtimes"
+        ,   "config.linkdirs"
+        ,   "target.linkdirs"
+        ,   "target.rpathdirs"
+        ,   "target.strip"
+        ,   "target.symbols"
+        ,   "target.optimize:check"
+        ,   "toolchain.linkdirs"
+        ,   "toolchain.rpathdirs"
+        ,   "config.links"
+        ,   "target.links"
+        ,   "target.linkgroups"
+        ,   "toolchain.links"
+        ,   "config.syslinks"
+        ,   "target.syslinks"
+        ,   "toolchain.syslinks"
+        }
+    ,   static = {
+            "target.strip"
+        ,   "target.symbols"
+        }
+    }
+
+    set_menu {
+                config =
+                {
+                    {category = "Cross Compilation Configuration/Compiler Flags Configuration"                       }
+                ,   {nil, "ascflags",      "kv", nil,          "The Ascend C Kernel Compiler Flags"                 }
+                ,   {nil, "aicpuflags",    "kv", nil,          "The Ascend C AI-CPU Compiler Flags"                 }
+
+                ,   {category = "Cross Compilation Configuration/Linker Flags Configuration"                         }
+                ,   {nil, "ldflags",       "kv", nil,          "The Binary Linker Flags"                            }
+                ,   {nil, "arflags",       "kv", nil,          "The Static Library Linker Flags"                    }
+                ,   {nil, "shflags",       "kv", nil,          "The Shared Library Linker Flags"                    }
+
+                ,   {category = "Cross Compilation Configuration/Builtin Flags Configuration"                        }
+                ,   {nil, "links",         "kv", nil,          "The Link Libraries"                                 }
+                ,   {nil, "syslinks",      "kv", nil,          "The System Link Libraries"                          }
+                ,   {nil, "linkdirs",      "kv", nil,          "The Link Search Directories"                        }
+                ,   {nil, "includedirs",   "kv", nil,          "The Include Search Directories"                     }
+                }
+            }
diff --git a/xmake/modules/core/tools/bisheng.lua b/xmake/modules/core/tools/bisheng.lua
new file mode 100644
index 000000000..5d3726e66
--- /dev/null
+++ b/xmake/modules/core/tools/bisheng.lua
@@ -0,0 +1,71 @@
+--!A cross-platform build utility based on Lua
+--
+-- Licensed under the Apache License, Version 2.0 (the "License");
+-- you may not use this file except in compliance with the License.
+-- You may obtain a copy of the License at
+--
+--     http://www.apache.org/licenses/LICENSE-2.0
+--
+-- Unless required by applicable law or agreed to in writing, software
+-- distributed under the License is distributed on an "AS IS" BASIS,
+-- WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+-- See the License for the specific language governing permissions and
+-- limitations under the License.
+--
+-- Copyright (C) 2015-present, Xmake Open Source Community.
+--
+-- @author      wuzhenqing
+-- @file        bisheng.lua
+--
+
+inherit("clang")
+
+function init(self)
+    _super.init(self)
+    self:add("shared.ascflags", "-fPIC")
+    self:add("shared.aicpuflags", "-fPIC")
+    self:add("ascflags", "-Qunused-arguments")
+    self:add("aicpuflags", "-Qunused-arguments")
+    self:set("ascshflags", "-shared")
+end
+
+-- make the language flag
+function nf_language(self, stdname)
+    if _g.cxxmaps == nil then
+        _g.cxxmaps =
+        {
+            cxx03       = "-std=c++03"
+        ,   cxx11       = "-std=c++11"
+        ,   cxx14       = "-std=c++14"
+        ,   cxx17       = "-std=c++17"
+        ,   cxx20       = "-std=c++20"
+        ,   cxxlatest   = {"-std=c++20", "-std=c++17", "-std=c++14", "-std=c++11", "-std=c++03"}
+        }
+        local cxxmaps2 = {}
+        for k, v in pairs(_g.cxxmaps) do
+            cxxmaps2[k:gsub("xx", "++")] = v
+        end
+        table.join2(_g.cxxmaps, cxxmaps2)
+    end
+    local result = _g.cxxmaps[stdname]
+    if type(result) == "table" then
+        local flagkind = self:kind() == "aicpu" and "aicpuflags" or "ascflags"
+        for _, flag in ipairs(result) do
+            if self:has_flags(flag, flagkind) then
+                result = flag
+                _g.cxxmaps[stdname] = result
+                break
+            end
+        end
+    end
+    return result
+end
+
+-- add source kind override flags (e.g. for .cce files compiled as asc)
+function add_sourceflags(self, sourcefile, fileconfig, target, targetkind)
+    local sourcekind = fileconfig.sourcekind
+    if sourcekind then
+        local maps = {asc = "-x asc", aicpu = "-x aicpu"}
+        return maps[sourcekind]
+    end
+end
diff --git a/xmake/modules/core/tools/bisheng/has_flags.lua b/xmake/modules/core/tools/bisheng/has_flags.lua
new file mode 100644
index 000000000..626c58c51
--- /dev/null
+++ b/xmake/modules/core/tools/bisheng/has_flags.lua
@@ -0,0 +1,166 @@
+--!A cross-platform build utility based on Lua
+--
+-- Licensed under the Apache License, Version 2.0 (the "License");
+-- you may not use this file except in compliance with the License.
+-- You may obtain a copy of the License at
+--
+--     http://www.apache.org/licenses/LICENSE-2.0
+--
+-- Unless required by applicable law or agreed to in writing, software
+-- distributed under the License is distributed on an "AS IS" BASIS,
+-- WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+-- See the License for the specific language governing permissions and
+-- limitations under the License.
+--
+-- Copyright (C) 2015-present, Xmake Open Source Community.
+--
+-- @author      wuzhenqing
+-- @file        has_flags.lua
+--
+
+-- Note: This module is derived from gcc/has_flags.lua with modifications
+-- for Ascend C specific source kinds (.asc, .aicpu) and --npu-arch handling.
+-- Inherit cannot be used here because xmake's tool module inheritance does
+-- not support overriding individual local functions.
+
+-- imports
+import("core.cache.detectcache")
+import("core.language.language")
+
+-- is linker?
+function _islinker(flags, opt)
+
+    -- the flags is "-Wl,<arg>" or "-Xlinker <arg>"?
+    local flags_str = table.concat(flags, " ")
+    if flags_str:startswith("-Wl,") or
+       flags_str:startswith("-Xlinker ") then
+        return true
+    end
+
+    -- the tool kind is ld or sh?
+    local toolkind = opt.toolkind or ""
+    return toolkind == "ld" or
+           toolkind == "sh" or
+           toolkind == "ascld" or
+           toolkind == "ascsh" or
+           toolkind:endswith("-ld") or
+           toolkind:endswith("-sh") or
+           toolkind:endswith("ld") or
+           toolkind:endswith("sh")
+end
+
+-- try running
+function _try_running(program, argv, opt)
+    local errors = nil
+    return try {
+        function ()
+            os.runv(program, argv, opt)
+            return true
+        end,
+        catch {
+            function (errs)
+                errors = (errs or ""):trim()
+            end
+        }
+    }, errors
+end
+
+-- get source kind
+function _get_sourcekind(opt, islinker)
+    if not islinker then
+        local toolkind = opt.toolkind or "asc"
+        if language.sourcekinds()[toolkind] then
+            return toolkind
+        end
+    end
+    return "asc"
+end
+
+function _map_sourcekind(sourcekind)
+    local map = {cxx = "c++"}
+    return map[sourcekind] or sourcekind
+end
+
+-- get extension
+function _get_extension(sourcekind)
+    return table.wrap(language.sourcekinds()[sourcekind])[1] or ".asc"
+end
+
+-- has npu arch?
+function _has_npu_arch(flags)
+    for _, flag in ipairs(flags) do
+        if flag:startswith("--npu-arch=") then
+            return true
+        end
+    end
+end
+
+-- attempt to check it from the argument list
+function _check_from_arglist(flags, opt, islinker)
+    local flag = flags[1]
+    if not flag then
+        return false
+    end
+
+    -- check for the builtin flag=value
+    if flag:startswith("--npu-arch=") then
+        return true
+    end
+
+    -- check from the `--help` menu, only for compile flags
+    if islinker or #flags > 1 then
+        return
+    end
+
+    local key = "core.tools.bisheng.has_flags"
+    local flagskey = opt.program .. "_" .. (opt.programver or "")
+    local allflags = detectcache:get2(key, flagskey)
+    if not allflags then
+        allflags = {}
+        local arglist = try {function () return os.iorunv(opt.program, {"--help"}, {envs = opt.envs}) end}
+        if arglist then
+            for arg in arglist:gmatch("%s+(%-[%-%a%d]+)%s+") do
+                allflags[arg] = true
+            end
+        end
+        detectcache:set2(key, flagskey, allflags)
+    end
+    return allflags[flag]
+end
+
+-- try running to check flags
+function _check_try_running(flags, opt, islinker)
+    local snippet = opt.snippet or
+        "int main(int argc, char** argv)\n{return 0;}\n"
+    local sourcekind = _get_sourcekind(opt, islinker)
+    local sourcefile = os.tmpfile("bisheng_has_flags") ..
+        _get_extension(sourcekind)
+    if not os.isfile(sourcefile) then
+        io.writefile(sourcefile, snippet)
+    end
+
+    local args = table.join("-o", os.tmpfile(), sourcefile)
+    if not islinker then
+        table.insert(args, 1, "-c")
+    end
+    if sourcekind == "asc" and not _has_npu_arch(flags) then
+        sourcekind = "c++"
+    end
+    sourcekind = _map_sourcekind(sourcekind)
+    table.insert(args, 1, sourcekind)
+    table.insert(args, 1, "-x")
+    return _try_running(opt.program, table.join(flags, args), opt)
+end
+
+function _has_flags(flags, opt)
+    opt = opt or {}
+    local islinker = _islinker(flags, opt)
+    if not opt.tryrun and _check_from_arglist(flags, opt, islinker) then
+        return true
+    end
+    return _check_try_running(flags, opt, islinker)
+end
+
+function main(...)
+    return _has_flags(...)
+end
diff --git a/xmake/modules/detect/tools/find_bisheng.lua b/xmake/modules/detect/tools/find_bisheng.lua
new file mode 100644
index 000000000..f5832b09e
--- /dev/null
+++ b/xmake/modules/detect/tools/find_bisheng.lua
@@ -0,0 +1,47 @@
+--!A cross-platform build utility based on Lua
+--
+-- Licensed under the Apache License, Version 2.0 (the "License");
+-- you may not use this file except in compliance with the License.
+-- You may obtain a copy of the License at
+--
+--     http://www.apache.org/licenses/LICENSE-2.0
+--
+-- Unless required by applicable law or agreed to in writing, software
+-- distributed under the License is distributed on an "AS IS" BASIS,
+-- WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+-- See the License for the specific language governing permissions and
+-- limitations under the License.
+--
+-- Copyright (C) 2015-present, Xmake Open Source Community.
+--
+-- @author      wuzhenqing
+-- @file        find_bisheng.lua
+--
+
+-- imports
+import("lib.detect.find_program")
+import("lib.detect.find_programver")
+
+-- find bisheng
+--
+-- @param opt   the argument options, e.g. {version = true}
+--
+-- @return      program, version
+--
+-- @code
+--
+-- local bisheng = find_bisheng()
+-- local bisheng, version = find_bisheng({version = true})
+--
+-- @endcode
+--
+function main(opt)
+    opt = opt or {}
+    opt.norunfile = true
+    local program = find_program(opt.program or "bisheng", opt)
+    local version = nil
+    if program and opt and opt.version then
+        version = find_programver(program, opt)
+    end
+    return program, version
+end
diff --git a/xmake/platforms/linux/xmake.lua b/xmake/platforms/linux/xmake.lua
index b411dc4d3..c4d09880e 100644
--- a/xmake/platforms/linux/xmake.lua
+++ b/xmake/platforms/linux/xmake.lua
@@ -36,7 +36,8 @@ platform("linux")
     --
     -- TODO Perhaps we should handle it better, or remove the cross toolchain.
     set_toolchains("envs", "gcc", "clang",
-        "cross", "yasm", "nasm", "fasm", "cuda", "go", "rust", "swift", "gfortran", "zig", "fpc", "nim", "dotnet")
+        "cross", "yasm", "nasm", "fasm", "cuda", "ascendc", "go", "rust", "swift",
+        "gfortran", "zig", "fpc", "nim", "dotnet")
 
     set_menu {
                 config =
@@ -63,5 +64,3 @@ platform("linux")
                 ,   {nil, "vcpkg",          "kv", "auto",       "The Vcpkg Directory"               }
                 }
             }
-
-
diff --git a/xmake/rules/ascendc/xmake.lua b/xmake/rules/ascendc/xmake.lua
new file mode 100644
index 000000000..312230b07
--- /dev/null
+++ b/xmake/rules/ascendc/xmake.lua
@@ -0,0 +1,53 @@
+--!A cross-platform build utility based on Lua
+--
+-- Licensed under the Apache License, Version 2.0 (the "License");
+-- you may not use this file except in compliance with the License.
+-- You may obtain a copy of the License at
+--
+--     http://www.apache.org/licenses/LICENSE-2.0
+--
+-- Unless required by applicable law or agreed to in writing, software
+-- distributed under the License is distributed on an "AS IS" BASIS,
+-- WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+-- See the License for the specific language governing permissions and
+-- limitations under the License.
+--
+-- Copyright (C) 2015-present, Xmake Open Source Community.
+--
+-- @author      wuzhenqing
+-- @file        xmake.lua
+--
+
+-- define rule: ascendc.env
+rule("ascendc.env")
+    after_load(function (target)
+        -- set default C++ standard (user can override via set_languages)
+        target:add("languages", "c++17")
+    end)
+
+-- define rule: ascendc.npuarchs
+rule("ascendc.npuarchs")
+    on_config(function (target)
+        for _, arch in ipairs(table.unique(table.wrap(target:get("ascnpuarchs")))) do
+            target:add("ascflags", "--npu-arch=" .. arch)
+            target:add("shflags", "--npu-arch=" .. arch)
+            target:add("ldflags", "--npu-arch=" .. arch)
+        end
+    end)
+
+-- define rule: ascendc.build.asc
+rule("ascendc.build.asc")
+    set_sourcekinds("asc")
+    add_deps("ascendc.npuarchs")
+    on_build_files("private.action.build.object", {jobgraph = true, batch = true})
+
+-- define rule: ascendc.build.aicpu
+rule("ascendc.build.aicpu")
+    set_sourcekinds("aicpu")
+    add_deps("ascendc.npuarchs")
+    on_build_files("private.action.build.object", {jobgraph = true, batch = true})
+
+-- define rule: ascendc
+rule("ascendc")
+    add_deps("ascendc.env", "ascendc.build.asc", "ascendc.build.aicpu")
+    add_deps("utils.inherit.links")
diff --git a/xmake/toolchains/ascendc/check.lua b/xmake/toolchains/ascendc/check.lua
new file mode 100644
index 000000000..60df001db
--- /dev/null
+++ b/xmake/toolchains/ascendc/check.lua
@@ -0,0 +1,90 @@
+--!A cross-platform build utility based on Lua
+--
+-- Licensed under the Apache License, Version 2.0 (the "License");
+-- you may not use this file except in compliance with the License.
+-- You may obtain a copy of the License at
+--
+--     http://www.apache.org/licenses/LICENSE-2.0
+--
+-- Unless required by applicable law or agreed to in writing, software
+-- distributed under the License is distributed on an "AS IS" BASIS,
+-- WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+-- See the License for the specific language governing permissions and
+-- limitations under the License.
+--
+-- Copyright (C) 2015-present, Xmake Open Source Community.
+--
+-- @author      wuzhenqing
+-- @file        check.lua
+--
+
+-- imports
+import("lib.detect.find_tool")
+
+-- map host architecture to CANN host tool directory
+function _host_archdir(arch)
+    local host_archdirs = {
+        x86_64 = "x86_64-linux"
+    ,   x64 = "x86_64-linux"
+    ,   arm64 = "aarch64-linux"
+    ,   aarch64 = "aarch64-linux"
+    }
+    return host_archdirs[arch]
+end
+
+-- check the ascendc toolchain
+function main(toolchain)
+    if not toolchain:is_plat("linux") then
+        return false
+    end
+
+    -- resolve sdkdir: --sdk= > ASCEND_HOME_PATH > ASCEND_TOOLKIT_HOME
+    local sdkroot = toolchain:sdkdir()
+    if not sdkroot then
+        sdkroot = os.getenv("ASCEND_HOME_PATH") or os.getenv("ASCEND_TOOLKIT_HOME")
+    end
+    if not sdkroot or not os.isdir(sdkroot) then
+        return false
+    end
+    sdkroot = path.absolute(sdkroot)
+
+    -- map host arch to CANN host directory
+    local host_arch = os.arch()
+    local host_archdir = _host_archdir(host_arch)
+    if not host_archdir then
+        return false
+    end
+
+    local hostroot = path.join(sdkroot, host_archdir)
+    if not os.isdir(hostroot) then
+        return false
+    end
+
+    -- check required executables
+    local bindir = path.join(hostroot, "bin")
+    local bisheng_bin = path.join(bindir, "bisheng")
+    local llvm_ar = path.join(bindir, "llvm-ar")
+    if not os.isexec(bisheng_bin) or not os.isexec(llvm_ar) then
+        return false
+    end
+
+    -- ensure bisheng can load its own shared libraries during version check
+    local host_libdir = path.join(hostroot, "lib64")
+    local ld_library_path = os.getenv("LD_LIBRARY_PATH") or ""
+    local envs = {
+        LD_LIBRARY_PATH = ld_library_path ~= "" and
+            (host_libdir .. path.envsep() .. ld_library_path) or host_libdir
+    }
+
+    -- use find_tool (unified interface) instead of find_bisheng directly
+    local result = find_tool("bisheng", {program = bisheng_bin, version = true, envs = envs})
+    if not result or not result.program then
+        return false
+    end
+
+    toolchain:config_set("sdkdir", sdkroot)
+    toolchain:config_set("bindir", bindir)
+    toolchain:config_set("hostroot", hostroot)
+    cprint("checking for Huawei Ascend C Toolchain (host: %s) ... ${color.success}${text.success}", host_arch)
+    return true
+end
diff --git a/xmake/toolchains/ascendc/load.lua b/xmake/toolchains/ascendc/load.lua
new file mode 100644
index 000000000..696541c9d
--- /dev/null
+++ b/xmake/toolchains/ascendc/load.lua
@@ -0,0 +1,52 @@
+--!A cross-platform build utility based on Lua
+--
+-- Licensed under the Apache License, Version 2.0 (the "License");
+-- you may not use this file except in compliance with the License.
+-- You may obtain a copy of the License at
+--
+--     http://www.apache.org/licenses/LICENSE-2.0
+--
+-- Unless required by applicable law or agreed to in writing, software
+-- distributed under the License is distributed on an "AS IS" BASIS,
+-- WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+-- See the License for the specific language governing permissions and
+-- limitations under the License.
+--
+-- Copyright (C) 2015-present, Xmake Open Source Community.
+--
+-- @author      wuzhenqing
+-- @file        load.lua
+--
+
+function main(toolchain)
+    local sdkdir = toolchain:config("sdkdir")
+    local hostroot = toolchain:config("hostroot")
+    if not sdkdir or not hostroot then
+        raise("ascendc toolchain not checked")
+    end
+
+    -- add run environments
+    toolchain:add("runenvs", "LD_LIBRARY_PATH", path.join(hostroot, "lib64"))
+
+    -- add toolchain-level search paths (picked up by language nameflags)
+    toolchain:add("includedirs", path.join(hostroot, "include"))
+    toolchain:add("includedirs", path.join(hostroot, "asc", "include"))
+    toolchain:add("includedirs", path.join(hostroot, "asc", "include", "aicpu_api"))
+    toolchain:add("includedirs", path.join(sdkdir, "include"))
+    toolchain:add("linkdirs", path.join(hostroot, "lib64"))
+    toolchain:add("linkdirs", path.join(hostroot, "devlib"))
+    toolchain:add("linkdirs", path.join(sdkdir, "lib64"))
+    toolchain:add("rpathdirs", path.join(hostroot, "lib64"))
+    toolchain:add("rpathdirs", path.join(hostroot, "devlib"))
+    toolchain:add("rpathdirs", path.join(sdkdir, "lib64"))
+    toolchain:add("syslinks", "ascendcl", "acl_rt", "dl")
+
+    -- aicpu-specific compiler flags (SDK-path-dependent)
+    toolchain:add("aicpuflags", "-fvisibility-inlines-hidden")
+    toolchain:add("aicpuflags", "-D_GLIBCXX_USE_CXX11_ABI=0")
+    toolchain:add("aicpuflags", "-D_GNU_SOURCE")
+    toolchain:add("aicpuflags", "--cce-aicpu-L" .. path.join(hostroot, "lib64", "device", "lib64"))
+    toolchain:add("aicpuflags", "--cce-aicpu-laicpu_api")
+    toolchain:add("aicpuflags", "--cce-aicpu-toolkit-path=" .. path.join(sdkdir, "toolkit", "toolchain", "hcc", "bin"))
+    toolchain:add("aicpuflags", "--cce-aicpu-sysroot=" .. path.join(sdkdir, "toolkit", "toolchain", "hcc", "sysroot"))
+end
diff --git a/xmake/toolchains/ascendc/xmake.lua b/xmake/toolchains/ascendc/xmake.lua
new file mode 100644
index 000000000..6071a6840
--- /dev/null
+++ b/xmake/toolchains/ascendc/xmake.lua
@@ -0,0 +1,46 @@
+--!A cross-platform build utility based on Lua
+--
+-- Licensed under the Apache License, Version 2.0 (the "License");
+-- you may not use this file except in compliance with the License.
+-- You may obtain a copy of the License at
+--
+--     http://www.apache.org/licenses/LICENSE-2.0
+--
+-- Unless required by applicable law or agreed to in writing, software
+-- distributed under the License is distributed on an "AS IS" BASIS,
+-- WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+-- See the License for the specific language governing permissions and
+-- limitations under the License.
+--
+-- Copyright (C) 2015-present, Xmake Open Source Community.
+--
+-- @author      wuzhenqing
+-- @file        xmake.lua
+--
+
+-- Huawei Ascend C / bisheng toolchain
+--
+-- @code
+-- target("example")
+--     set_toolchains("ascendc")
+--     add_ascnpuarchs("dav-2201")
+-- @endcode
+--
+toolchain("ascendc")
+
+    -- set homepage
+    set_homepage("https://www.hiascend.com/")
+    set_description("Huawei Ascend C (bisheng compiler driver)")
+
+    -- set toolset
+    set_toolset("asc",  "bisheng")
+    set_toolset("aicpu", "bisheng")
+    set_toolset("ascld", "bisheng")
+    set_toolset("ascsh", "bisheng")
+    set_toolset("ar",   "llvm-ar", "ar")
+
+    -- check toolchain
+    on_check("check")
+
+    -- on load
+    on_load("load")