summaryrefslogtreecommitdiff
path: root/tests/projects/ascendc/mixed_binary/src
diff options
context:
space:
mode:
authorwuzhenqing <[email protected]>2026-05-04 11:15:50 +0000
committerwuzhenqing <[email protected]>2026-05-04 11:18:09 +0000
commitfa807b737b4b6b3004cccafc0dfdfdd66fcdce2d (patch)
tree558dc6bc9e5749180fa97a569aab8bb6c2e54fec /tests/projects/ascendc/mixed_binary/src
parent8bbbb2860fdd1910db466fe318ee8d465728ca32 (diff)
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
Diffstat (limited to 'tests/projects/ascendc/mixed_binary/src')
-rw-r--r--tests/projects/ascendc/mixed_binary/src/helper.aicpu6
-rw-r--r--tests/projects/ascendc/mixed_binary/src/main.asc25
2 files changed, 31 insertions, 0 deletions
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;
+}