MCC User Manual

This document describes how to compile MUSA code with mcc and gives some details about mcc's implementations.

This document assumes a basic familiarity with the CUDA programming language. Information about CUDA programming can be found in the CUDA programming guide.

Overview

MUSA is a CUDA-like programming language for concurrency programming on MTGPU devices.

Prerequisites

mcc currently supports building MUSA which is compatible with CUDA 11.4. Before you build MUSA code, you'll need to have installed the MUSA Toolkit. mcc will attempt to find MUSA installation under the default path "/usr/local/musa". Users also can specify a different location of MUSA Toolkit location via the option "--musa-path".

MUSA compilation is only supported on Linux currently. The compiler on Windows is coming soon.

Purpose of MCC

A MUSA application consists of a mixture of conventional C++ host code and GPU device functions, which makes the compilation trajectory involve several steps of compilation and linkage for each MUSA source file. mcc, the MUSA programming language compiler driver, helps hide the complexity of these phases of building a source code.

Compilation phases

Mostly, the compilation of MUSA in mcc is like conventional C++ by clang++:

1. Preprocess the source code to deal with macros and include directives by macro expansion and header including.
2. Compile to llvm ir with a series of organized transformation and optimization passes.
3. Code generation: lower to target specific code and generate object files.
4. Link libraries to generate executable binary, including static/shared library and standalone executable file.

The difference is that MUSA code is mixed with the host-side part and device-side part. Basically, mcc will compile the source twice to generate the device code object and host code object, then try to combine them together. So the above phases will repeat twice at least for both the device-side code and host-side part.

In detail, the binary value of the device code object will be assigned to a global constant variable when compiling host-side code, before code generation. Phases of host-side code are almost the same as conventional C++'s.

When the kernel is to be launched to GPU devices, MUSA runtime and driver will help extract the device code object, which is outside of the scope of this document.

Compilation trajectory

Compilation driven by mcc follows the trajectory as below:

Supported input file suffixes

Each phase of compilation driven by mcc accepts specific input files. mcc will identify them by file suffixes and compile options.

Input file suffix	Description
.mu	A MUSA source file, containing host code and device functions.
.c	A C source file.
.cc, .cxx, .cpp	A C++ source file.
.ll	A LLVM IR file.
.o, .obj	An object file.
.a	A library file.
.so	A shared library file.

And other file suffixes are supported by clang/clang++.

Compiler command-line options

mcc is the compiler driver for the MUSA language inspired by llvm-project 14.0.0. Most of the compiler options work fine just as Clang and LLVM toolchain. Additionally, it provides specific options for the MUSA language and MTGPU target.

Options for help information

--help/-help

Displays help for available options.

--help-hidden

Displays help for hidden options.

- v

Displays verbose output for sub-commands.

--version

Prints version information.

Options for input file and specifications

-o [file]

Defines the name of the output file.

- x [language]

Specifies the language for the input files(e.g. -x musa), and subsequent input files will recognize as having type.

- D [macro]=[value]

Defines [macro] to [value] (or 1 if [value] omitted).

Options for path specifications

--musa-path= value

Specifies MUSA Toolkit installation path.

--musa-path-ignore-env

Ignores environment variables to detect MUSA Toolkit installation.

--musa-inc-path= value

Specifies the path of the header file in MUSA Toolkit, the default path is musa-path/include.

--musa-libdevice-path= value

Specifies the path of libdevice.bc in MUSA Toolkit, the default path is musa-path/mtgpu/bitcode/.

--musa-rt-lib-path= value

Specifies MUSA runtime library path.

Options for including and linking

- I dir

Adds a directory to the end of the list of include search paths.

- include file

Includes the file before parsing.

- l lib-name

Links with a library.

- L dir

Adds a directory to the library search path.

- nogpuinc

Does not add include paths for CUDA/HIP/MUSA and does not include the default CUDA/HIP/MUSA wrapper headers.

- nogpulib

Does not link the device library for CUDA/HIP/MUSA device compilation.

Options for compilation phase control

- c

Only runs preprocess, compile, and assemble steps.

- emit-llvm

Uses the LLVM representation for assembler and object files.

--musa-compile-host-device

Compiles MUSA code for both host and device (default). Has no effect on non-MUSA compilations.

--musa-device-only

Compiles MUSA code for device only.

--musa-host-only

Compiles MUSA code for host only. Has no effect on non-MUSA compilations.

--offload-arch= value

MUSA offloading target ID in the form of a device architecture followed by target ID features string delimited by a colon. Each target ID feature is a pre-defined string followed by a plus or minus sign (e.g. gfx908:xnack+:sramecc-). This command may be specified more than once to enable multi-arch compilation, e.g.mcc -o axpy axpy.mu -lmusart --offload-arch=mp_21 --offload-arch=mp_22, this command includes two pairs of options that specify arch mp_21 and mp_22 to enable compiler to build both binaries of mp_21 and mp_22 into a single fat-binary. Compiler will build mp_21 and mp_22 by default without any offload arch specifying option added to the command.

MUSA offloading target ID: mp_10, mp_21, mp_22.

- mtgpu

mcc recognizes MUSA source code by suffix .mu, needs to be hinted to control the compilation compile the source to MTGPU target by this option. While mcc accepts source files with suffix .cu, in which case mcc needs option -mtgpu to recognize the source file with suffix .cu as MUSA language source file.

- cuda_wrapper

mcc can work in CUDA wrapper mode, supporting developers to use CUDA APIs instead of MUSA's.

More compiler options can be referred to from Clang 14.0.0 document.

Options for frontend

- Wno-shared-var-init

MCC supports to use "-Wno-shared-var-init" to ignore compilation errors caused by initializing shared memory, users need to ensure that the initialization operation is safe.

Options for backend

Each backend option starts with "-mllvm" to indicate that the following option is for backend.

- mtgpu-maxregcnt= value

Sets the max temp register count for the MTGPU target.

- mtgpu-if-convert

This option will enable an optimization: "if" block will be converted to predicate execution for small size of the block.

- mtgpu-internalize-symbols= value

This option is true by default and will try to internalize all the symbols including global variables and functions.

- mtgpu-enable-const-calc

This option is disabled by default and will enable constant calculation program generation optimization.

- mtgpu-load-store-opt

This option is disabled by default and will merge shared memory load store instructions.

- enable-ldma-index

This option is disabled by default and will enable ldma instruction to use a non-zero imm-offset index for the instruction selection stage.

- mtgpu-prera-internal-opt

This option is disabled by default and will enable optimization by using internal registers to reduce the register pressure of temp register before register allocation.This option CANNOT BE ENABLED WITH - mtgpu-postra-internal-opt AT THE SAME TIME.

- mtgpu-postra-internal-opt

This option is disabled by default and will enable optimization by using internal registers to reduce the register pressure of temp register after register allocation. This option CANNOT BE ENABLED WITH - mtgpu-prera-internal-opt AT THE SAME TIME.

- internal-opt-avoid-use-nearest

This option is disabled by default and will enlarge the interval of optimization when doing the internal register optimization. Usually used with above 2 options.

- mtgpu-load-cluster-mutation

This option is disabled by default and will enable "load cluster mutation" on the instruction order after register allocation.

- mtgpu-memory-sched-mutation

This option is disabled by default and will enable "memory specific schedule mutation" on the instruction order after register allocation.

- mtgpu-tiny-offset-hint

This option is disabled by default and will hint compiler that code will not use large offset for memory access which will exceed 2^32 to achieve more optimization opportunities.

Options for special control

- mathx-disable-fp64

No hardware support on SUDI and all FP64 operations are implemented by libdevice, which will expand the code size and cause a long time to compile. This option will disable any FP64 operations and speed up the compilation for SUDI(mp_10).

Example

A MUSA source code file: axpy.mu, use the following command line to compile to an executable:

mcc axpy.mu -lmusart - L/usr/local/musa/lib -O2 - o axpy

In which:

• "axpy.mu": is the input source file.
• "-lmusart": links the MUSA runtime library, libmusart.so.
• "-L/usr/local/musa/lib": hints compiler to find libraries under path "/usr/local/musa/lib".
• "-O2": indicates the optimization level. Currently mcc only supports O2 and above opt. level. And this is the default opt. level for mcc when compiling device side code.
• "-o axpy": specifies the output filename, "axpy", and the output location under where the command is executed.

If the source file does not end up with .mu, a pair of options are required: -x musa, the command as follows:

mcc - x musa axpy.cu -lmusart - L/usr/local/musa/lib -O2 - o axpy

And if code in the source file uses API from CUDA, like cudaMemcpy, cudaMemsetetc, please add an extra option -cuda_wrapper and link libcuda2musa.so instead of libmusart.so, the command as follows:

mcc axpy.cu -mtgpu - cuda_wrapper - lcuda2musa - L/usr/local/musa/lib -O2 - o axpy

Extended attribute

Kernel function attribute

mtgpu_num_usreg

MCC supports the __attribute__((mtgpu_num_usreg(num_usreg))) attributes for the MTGPU target. These attributes may be attached to a kernel function definition and are an optimization hint.

If these attributes are specified, then the MTGPU target backend will attempt to limit the number of TEMP regs and/or ATTR regs used to the specified value(s). The number of used TEMP regs and/or ATTR regs may further be rounded up to satisfy the allocation requirements or constraints of the sub-target. Passing 0 as num_usreg implies the default behavior (no limits).

Usage:

__attribute__((mtgpu_num_usreg(256))) void KernelName(...) {...}

mtgpu_unroll_threshold

MCC supports the __attribute__((mtgpu_unroll_threshold(unroll_threshold))) attributes for the MTGPU target. These attributes may be attached to a kernel function definition and are an optimization hint.

Usage:

__attribute__((mtgpu_unroll_threshold(400))) void KernelName(...) {...}

mtgpu_tiny_offset

MCC supports the __attribute__((mtgpu_tiny_offset)) attribute for the MTGPU target. This attribute may be attached to a kernel function definition to hint the compiler that the offsets of memory access operations in this kernel will not exceed INT_MAX to create more optimization opportunities.

Usage:

__global__ __attribute__((mtgpu_tiny_offset)) void kernel_name(...) {...}