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

9 files changed, 132 insertions, 0 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")