From fa807b737b4b6b3004cccafc0dfdfdd66fcdce2d Mon Sep 17 00:00:00 2001
From: wuzhenqing <wuzhenqing@stu.hit.edu.cn>
Date: Mon, 4 May 2026 11:15:50 +0000
Subject: 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
---
 .../projects/ascendc/mixed_binary/src/helper.aicpu |  6 ++++++
 tests/projects/ascendc/mixed_binary/src/main.asc   | 25 ++++++++++++++++++++++
 2 files changed, 31 insertions(+)
 create mode 100644 tests/projects/ascendc/mixed_binary/src/helper.aicpu
 create mode 100644 tests/projects/ascendc/mixed_binary/src/main.asc

(limited to 'tests/projects/ascendc/mixed_binary/src')

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;
+}
-- 
cgit v1.3.1