fix ascendc toolchain

wuzhenqing · wuzhenqing · commit 5dae47e33b09 · 2026-05-01T10:57:41.000Z
diff --git a/tests/modules/language/ascendc/test.lua b/tests/modules/language/ascendc/test.lua
@@ -1,6 +1,6 @@
 import("core.language.language")
 
-local function _write_asc_file(name, source)
+function _write_asc_file(name, source)
     local sourcefile = os.tmpfile(name) .. ".asc"
     io.writefile(sourcefile, source)
     return sourcefile
@@ -16,8 +16,7 @@ function test_check_main(t)
     local mainfile = _write_asc_file("ascendc_check_main", [[
 #include "acl/acl.h"
 
-int32_t main(int argc, char const *argv[])
-{
+int32_t main(int argc, char const *argv[]) {
     return 0;
 }
 ]])
@@ -28,8 +27,7 @@ int32_t main(int argc, char const *argv[])
 /*
 int main() { return 1; }
 */
-__global__ __vector__ void kernel()
-{
+__global__ __vector__ void kernel() {
 }
 ]])
     t:are_equal(check_main(commentfile), false)
diff --git a/tests/projects/ascendc/libs/src/lib.asc b/tests/projects/ascendc/libs/src/lib.asc
@@ -1,6 +1,5 @@
 #include "kernel_operator.h"
 
-__global__ __vector__ void hello_world()
-{
+__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
@@ -3,8 +3,7 @@
 
 extern __global__ __vector__ void hello_world();
 
-int32_t main(int argc, char const *argv[])
-{
+int32_t main(int argc, char const *argv[]) {
     aclInit(nullptr);
     int32_t deviceId = 0;
     aclrtSetDevice(deviceId);
diff --git a/tests/projects/ascendc/libs/xmake.lua b/tests/projects/ascendc/libs/xmake.lua
@@ -2,19 +2,16 @@ add_rules("mode.debug", "mode.release")
 
 target("ascendc_static")
     set_kind("static")
-    set_toolchains("ascendc")
     add_files("src/lib.asc")
     add_ascnpuarchs("dav-2201")
 
 target("ascendc_shared")
     set_kind("shared")
-    set_toolchains("ascendc")
     add_files("src/lib.asc")
     add_ascnpuarchs("dav-2201")
 
 target("ascendc_libs_bin")
     add_deps("ascendc_static", "ascendc_shared")
     set_kind("binary")
-    set_toolchains("ascendc")
     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
@@ -1,7 +1,6 @@
 #include "aicpu_api.h"
 
-__global__ __aicpu__ uint32_t hello_world(void *args)
-{
+__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
@@ -6,8 +6,7 @@ struct KernelArgs {
 
 extern __global__ __aicpu__ uint32_t hello_world(void *args);
 
-int32_t main(int argc, char const *argv[])
-{
+int32_t main(int argc, char const *argv[]) {
     aclInit(nullptr);
     int32_t deviceId = 0;
     aclrtSetDevice(deviceId);
diff --git a/xmake/languages/ascendc/check_main.lua b/xmake/languages/ascendc/check_main.lua
@@ -26,7 +26,7 @@ function main(sourcefile)
 
     -- remove comment first
     sourcecode = sourcecode:gsub("/%*.-%*/", "")
-    sourcecode = sourcecode:gsub("//.-\n", "\n")
+    sourcecode = sourcecode:gsub("//[^\n]*", "")
 
     -- find int main(int argc, char** argv) {}
     if sourcecode:find("%s+main%s*%(.-%)") then
diff --git a/xmake/modules/core/tools/bisheng.lua b/xmake/modules/core/tools/bisheng.lua
@@ -27,6 +27,11 @@ function init(self)
     self:add("ascflags", "-Qunused-arguments")
     self:add("aicpuflags", "-Qunused-arguments")
     self:set("ascshflags", "-shared")
+    -- default C++ standard and source kind for bisheng compiler driver
+    self:add("ascflags", "-std=c++17", {force = true})
+    self:add("aicpuflags", "-std=c++17", {force = true})
+    self:add("ascflags", "-x", "asc", {force = true})
+    self:add("aicpuflags", "-x", "aicpu", {force = true})
 end
 
 -- make the language flag
diff --git a/xmake/modules/core/tools/bisheng/has_flags.lua b/xmake/modules/core/tools/bisheng/has_flags.lua
@@ -133,7 +133,7 @@ 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:" .. snippet) ..
+    local sourcefile = os.tmpfile("bisheng_has_flags") ..
         _get_extension(sourcekind)
     if not os.isfile(sourcefile) then
         io.writefile(sourcefile, snippet)
diff --git a/xmake/platforms/linux/xmake.lua b/xmake/platforms/linux/xmake.lua
@@ -45,8 +45,6 @@ platform("linux")
                     {category = "Cuda SDK Configuration"                                            }
                 ,   {nil, "cuda",           "kv", "auto",       "The Cuda SDK Directory"            }
                 ,   {nil, "cuda_sdkver",    "kv", "auto",       "The Cuda SDK Version"              }
-                ,   {category = "Ascend CANN SDK Configuration"                                      }
-                ,   {nil, "ascendc",        "kv", "auto",       "The Ascend CANN SDK Directory"     }
                 ,   {category = "Qt SDK Configuration"                                              }
                 ,   {nil, "qt",             "kv", "auto",       "The Qt SDK Directory"              }
                 ,   {nil, "qt_host",        "kv", "auto",       "The Qt Host SDK Directory"         }
@@ -59,8 +57,6 @@ platform("linux")
                 {
                     {category = "Cuda SDK Configuration"                                            }
                 ,   {nil, "cuda",           "kv", "auto",       "The Cuda SDK Directory"            }
-                ,   {category = "Ascend CANN SDK Configuration"                                      }
-                ,   {nil, "ascendc",        "kv", "auto",       "The Ascend CANN SDK Directory"     }
                 ,   {category = "Qt SDK Configuration"                                              }
                 ,   {nil, "qt",             "kv", "auto",       "The Qt SDK Directory"              }
                 ,   {nil, "qt_host",        "kv", "auto",       "The Qt Host SDK Directory"         }
diff --git a/xmake/rules/ascendc/xmake.lua b/xmake/rules/ascendc/xmake.lua
@@ -18,46 +18,65 @@
 -- @file        xmake.lua
 --
 
--- define rule: ascendc.npuarchs
-rule("ascendc.npuarchs")
+-- define rule: ascendc.env
+rule("ascendc.env")
     on_config(function (target)
-        local function nf_ascnpuarch(arch)
-            if type(arch) ~= "string" then
-                return
-            end
-            arch = arch:trim()
-            if arch == "" then
-                return
-            end
-            if arch:startswith("--npu-arch=") then
-                return arch
-            end
-            return "--npu-arch=" .. arch
+        local toolchain_inst = target:toolchain("ascendc")
+        if not toolchain_inst then
+            return
         end
 
-        local function add_flag_once(name, flag)
-            if not table.contains(table.wrap(target:get(name)), flag) then
-                target:add(name, flag, {force = true})
-            end
+        local sdkdir = toolchain_inst:config("sdkdir")
+        if not sdkdir then
+            return
         end
 
-        local npuarchs = table.wrap(target:get("ascnpuarchs"))
-        for _, opt in ipairs(target:orderopts()) do
-            table.join2(npuarchs, opt:get("ascnpuarchs"))
-        end
-        for _, toolchain_inst in ipairs(target:toolchains()) do
-            if toolchain_inst:name() == "ascendc" then
-                table.join2(npuarchs, toolchain_inst:get("ascnpuarchs"))
-            end
+        -- resolve host architecture directory
+        local host_arch = os.arch()
+        local host_archdirs = {
+            x86_64 = "x86_64-linux"
+        ,   x64 = "x86_64-linux"
+        ,   arm64 = "aarch64-linux"
+        ,   aarch64 = "aarch64-linux"
+        }
+        local host_archdir = host_archdirs[host_arch]
+        if not host_archdir then
+            return
         end
+        local hostroot = path.join(sdkdir, host_archdir)
 
-        for _, arch in ipairs(table.unique(npuarchs)) do
-            local flag = nf_ascnpuarch(arch)
-            if flag then
-                add_flag_once("ascflags", flag)
-                add_flag_once("shflags", flag)
-                add_flag_once("ldflags", flag)
-            end
+        -- aicpu-specific flags (SDK-path-dependent)
+        target:add("aicpuflags", "-O2")
+        target:add("aicpuflags", "-fvisibility-inlines-hidden", {force = true})
+        target:add("aicpuflags", "-D_GLIBCXX_USE_CXX11_ABI=0", {force = true})
+        target:add("aicpuflags", "-D_GNU_SOURCE", {force = true})
+        target:add("aicpuflags", "--cce-aicpu-L" .. path.join(hostroot, "lib64", "device", "lib64"), {force = true})
+        target:add("aicpuflags", "--cce-aicpu-laicpu_api", {force = true})
+        local toolkit_bindir = path.join(sdkdir, "toolkit", "toolchain", "hcc", "bin")
+        local toolkit_sysroot = path.join(sdkdir, "toolkit", "toolchain", "hcc", "sysroot")
+        target:add("aicpuflags", "--cce-aicpu-toolkit-path=" .. toolkit_bindir, {force = true})
+        target:add("aicpuflags", "--cce-aicpu-sysroot=" .. toolkit_sysroot, {force = true})
+
+        -- common include/lib directories
+        target:add("includedirs", path.join(hostroot, "include"))
+        target:add("includedirs", path.join(hostroot, "asc", "include", "aicpu_api"))
+        target:add("includedirs", path.join(sdkdir, "include"))
+        target:add("linkdirs", path.join(hostroot, "lib64"))
+        target:add("linkdirs", path.join(hostroot, "devlib"))
+        target:add("linkdirs", path.join(sdkdir, "lib64"))
+        target:add("rpathdirs", path.join(hostroot, "lib64"))
+        target:add("rpathdirs", path.join(hostroot, "devlib"))
+        target:add("rpathdirs", path.join(sdkdir, "lib64"))
+        target:add("syslinks", "ascendcl", "acl_rt", "dl")
+    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, {force = true})
+            target:add("shflags", "--npu-arch=" .. arch, {force = true})
+            target:add("ldflags", "--npu-arch=" .. arch, {force = true})
         end
     end)
 
@@ -77,7 +96,7 @@ rule("ascendc.build.aicpu")
 rule("ascendc")
 
     -- add build rules
-    add_deps("ascendc.build.asc", "ascendc.build.aicpu")
+    add_deps("ascendc.env", "ascendc.build.asc", "ascendc.build.aicpu")
 
     -- inherit links and linkdirs of all dependent targets by default
     add_deps("utils.inherit.links")
diff --git a/xmake/toolchains/ascendc/check.lua b/xmake/toolchains/ascendc/check.lua
@@ -19,26 +19,38 @@
 --
 
 -- imports
-import("detect.tools.find_bisheng")
-import("toolchains.ascendc.common", {rootdir = os.programdir()})
+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
 
-    local sdkroot = common.sdkdir(toolchain)
+    -- resolve sdkdir: --sdk= > ASCEND_HOME_PATH > ASCEND_TOOLKIT_HOME
+    local sdkroot = toolchain:sdkdir()
     if not sdkroot then
-        return false
+        sdkroot = os.getenv("ASCEND_HOME_PATH") or os.getenv("ASCEND_TOOLKIT_HOME")
     end
-    sdkroot = path.absolute(sdkroot)
-    if not os.isdir(sdkroot) then
+    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 = common.host_archdir(host_arch)
+    local host_archdir = _host_archdir(host_arch)
     if not host_archdir then
         return false
     end
@@ -48,13 +60,11 @@ function main(toolchain)
         return false
     end
 
+    -- check required executables
     local bindir = path.join(hostroot, "bin")
-    local bisheng = path.join(bindir, "bisheng")
+    local bisheng_bin = path.join(bindir, "bisheng")
     local llvm_ar = path.join(bindir, "llvm-ar")
-    if not os.isexec(bisheng) then
-        return false
-    end
-    if not os.isexec(llvm_ar) then
+    if not os.isexec(bisheng_bin) or not os.isexec(llvm_ar) then
         return false
     end
 
@@ -66,11 +76,9 @@ function main(toolchain)
             (host_libdir .. path.envsep() .. ld_library_path) or host_libdir
     }
 
-    local program, version = find_bisheng({program = bisheng, version = true, envs = envs})
-    if not program then
-        return false
-    end
-    if not version then
+    -- 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
 
diff --git a/xmake/toolchains/ascendc/common.lua b/xmake/toolchains/ascendc/common.lua
diff --git a/xmake/toolchains/ascendc/load.lua b/xmake/toolchains/ascendc/load.lua
diff --git a/xmake/toolchains/ascendc/xmake.lua b/xmake/toolchains/ascendc/xmake.lua

Original file line number	Diff line number	Diff line change
`@@ -1,6 +1,5 @@`
`1`	`1`	`#include "kernel_operator.h"`
`2`	`2`
`3`		`-__global__ __vector__ void hello_world()`
`4`		`-{`
	`3`	`+__global__ __vector__ void hello_world() {`
`5`	`4`	`AscendC::printf("[Block (%lu/%lu)]: Hello World!!!\n", AscendC::GetBlockIdx(), AscendC::GetBlockNum());`
`6`	`5`	`}`
Original file line number	Diff line number	Diff line change
`@@ -1,7 +1,6 @@`
`1`	`1`	`#include "aicpu_api.h"`
`2`	`2`
`3`		`-__global__ __aicpu__ uint32_t hello_world(void *args)`
`4`		`-{`
	`3`	`+__global__ __aicpu__ uint32_t hello_world(void *args) {`
`5`	`4`	`AscendC::printf("Hello World!!!\n");`
`6`	`5`	`return 0;`
`7`	`6`	`}`