Add Huawei Ascend C toolchain support

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

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());
+}

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;
+}

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

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")

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;
+}

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;
+}

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

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")

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",

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

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