Skip to content

0.2.3: mcpp.rules.ascendc — Ascend C compiled by BiSheng, and no NPU required to do it - #7

Merged
Sunrisepeak merged 1 commit into
mainfrom
feat/rules-ascendc
Sep 7, 2026
Merged

0.2.3: mcpp.rules.ascendc — Ascend C compiled by BiSheng, and no NPU required to do it#7
Sunrisepeak merged 1 commit into
mainfrom
feat/rules-ascendc

Conversation

@Sunrisepeak

Copy link
Copy Markdown
Member

The fifth rule, and the first written against a vendor this project had never
built for. Everything below was measured on CANN 8.5.0, on a machine with
no Ascend hardware.

The compile line

From the toolkit's own CMakeASCInformation.cmake:

bisheng <DEFINES> <INCLUDES> -fPIC <FLAGS> -o <OBJECT> -c -x asc <SOURCE>

and three things that file does not say, each found by running it:

  • ASCEND_HOME_PATH must be in the environment, and points at the arch
    directory, not the toolkit root. Without it: PlugIn Err: can not find ASCEND_HOME_PATH, before the source is looked at.
  • The plugin needs libmmpa.so on the loader path. Without it the plugin
    fails to load and the failure surfaces as unknown type name '__aicore__'
    a message about the source that is not about the source.
  • -DTILING_KEY_VAR=0 and eleven include directories, which is what the
    toolkit's device_intf_pub target carries.

Not --cce-aicore-only

With it, BiSheng emits a Da Vinci object (*unknown arch 0x1029*) the host
linker cannot place. Without it, the same source yields an x86-64 object
carrying both the device binary (registered at load time) and a host-callable
launcher per __global__ function — which is exactly what mcpp's model needs:
a device unit becomes an object and the ordinary link takes it. No registration
JSON, no device-link step.

The launcher is C++-mangled even when the kernel is extern "C", so the
rule tells a project to put an extern "C" wrapper in the same .asc file
rather than depend on two compilers agreeing about mangling.

The link closure, which no project could name

A mixed-mode object needs libascendc_runtime.a; that archive is static, so
its undefined symbols become the program's: libascend_dump, libruntime,
libmmpa, libprofapi, libc_sec. Then -Wl,-rpath-link for the library
directory and the toolkit's driver stubs — GNU ld resolves a shared
library's own DT_NEEDED through -rpath-link, never through -L, and
without it the link fails on memset_s and CheckLogLevel, symbols belonging
to libraries nobody named. Then -Wl,-rpath, because mcpp gives an artifact a
private interpreter and a program that links but cannot start is not progress.

That -rpath-link pair has no declarative spelling — it is what mcpp
2026.9.6.5's link_flag outlet exists for, and why this member's floor is that
release.

Verified end to end

mcpp's examples/09-heterogeneous/cann compiles its kernel with BiSheng, links
the object into an ordinary binary, and resolves every runtime dependency
except libascend_hal.so — the driver, absent on a machine with no NPU and
exactly what should be missing there. Its CPU leg builds and runs
(12 24 36 48).

No tests/ fixture

Same reason the index recipe gives: the toolkit is a 1.1 GB download the CI job
cannot carry, and a test that would never run is worse than none.

…required to do it

The fifth rule, and the first one written against a vendor this project had
never built for. Everything below was measured on CANN 8.5.0 on a machine with
no Ascend hardware.

THE COMPILE LINE comes from the toolkit's own `CMakeASCInformation.cmake`:

    bisheng <DEFINES> <INCLUDES> -fPIC <FLAGS> -o <OBJECT> -c -x asc <SOURCE>

and three things that file does not say, each found by running it:

  * `ASCEND_HOME_PATH` must be in the environment and points at the ARCH
    directory, not the toolkit root. Without it: `PlugIn Err: can not find
    ASCEND_HOME_PATH`, before the source is looked at.
  * The plugin it loads needs `libmmpa.so` on the loader path. Without it the
    plugin fails to load and the failure surfaces as `unknown type name
    '__aicore__'` -- a message about the source that is not about the source.
  * `-DTILING_KEY_VAR=0` and eleven include directories, which is what the
    toolkit's own `device_intf_pub` target carries.

NOT `--cce-aicore-only`, AND THAT IS THE DESIGN DECISION IN THIS RULE. With it,
BiSheng emits a Da Vinci object -- `*unknown arch 0x1029*` -- that a host linker
cannot place. Without it, the same source yields an x86-64 object carrying both
the device binary, registered at load time, and a host-callable launcher per
`__global__` function. The second is what mcpp's model needs: a device unit
becomes an object and the ordinary link takes it. No registration JSON, no
device-link step.

The launcher is C++-MANGLED even when the kernel is declared `extern "C"`, so
the rule's documentation tells a project to put an `extern "C"` wrapper in the
same `.asc` file rather than have its host half depend on two compilers
agreeing about mangling.

THE LINK CLOSURE, which no project could be expected to name. A mixed-mode
object needs `libascendc_runtime.a`; that archive is static, so its undefined
symbols become the program's: `libascend_dump`, `libruntime`, `libmmpa`,
`libprofapi`, `libc_sec`. Then `-Wl,-rpath-link` for both the library directory
and the toolkit's driver stubs, because GNU ld resolves a shared library's own
`DT_NEEDED` entries through `-rpath-link` and never through `-L` -- without it
the link fails on `memset_s` and `CheckLogLevel`, symbols belonging to
libraries nobody named. Then `-Wl,-rpath` for the same directory, because mcpp
gives an artifact a private interpreter and a program that links but cannot
start is not progress.

That is five directives the rule computes, and one of them -- the `-rpath-link`
pair -- has no declarative spelling at all. It is what mcpp 2026.9.6.5's
`link_flag` outlet exists for, and it is why this member's floor is that
release rather than the one whose device-source table carries `.asc`.

VERIFIED END TO END: mcpp's `examples/09-heterogeneous/cann` compiles its
kernel with BiSheng, links the resulting object into an ordinary binary, and
resolves every runtime dependency except `libascend_hal.so` -- the DRIVER,
which is absent on a machine with no NPU and is exactly what should be missing
there. Its CPU leg builds and runs.

NO `tests/` FIXTURE, and the reason is the same one the index recipe gives: the
toolkit is a 1.1 GB download that the CI job cannot carry, and a test that
would never run is worse than none.
@Sunrisepeak
Sunrisepeak merged commit 89e2d1e into main Sep 7, 2026
1 check passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant