昇腾社区首页
中文
注册

快速上手

本节介绍异构编程环境配置与编译器使用的简单用例,方便用户快速验证环境信息,熟悉毕昇编译器的使用。

安装和环境配置

毕昇编译器跟随CANN软件包一起发布。安装完CANN包后,毕昇编译器所在目录为:${INSTALL_DIR}/compiler/ccec_compiler。

${INSTALL_DIR}请替换为CANN软件安装后文件存储路径。若安装的Ascend-cann-toolkit软件包,以root安装举例,则安装后文件存储路径为:/usr/local/Ascend/ascend-toolkit/latest。

编程开始前,需要配置毕昇编译器二进制程序相关环境变量,有配置CANN环境变量和设置PATH环境变量两种方式:

  • 方式一:配置CANN环境变量。
    CANN软件包提供进程级环境变量设置脚本,供用户在进程中引用,以自动完成环境变量设置。执行命令参考如下,以下示例为root或非root用户默认安装路径,请以实际安装路径为准。
    # 以root用户安装toolkit包后配置环境变量
    source /usr/local/Ascend/ascend-toolkit/set_env.sh 
    # 以非root用户安装toolkit包后配置环境变量
    source ${HOME}/Ascend/ascend-toolkit/set_env.sh 
  • 方式二:设置PATH环境变量。
    # 获取CANN包中的毕昇编译器安装目录,举例如下:
    $ export PATH=${INSTALL_DIR}/compiler/ccec_compiler/bin/:$PATH

简单的异构程序编译示例

本示例简单演示一个异构程序,启动4个block(核)的kernel函数,每个block写一份自己的数据,host侧使用ACL Runtime接口进行运行时管理。
// 文件名QuickStartDemo.cce
#include "acl/acl.h"
#include <stdio.h>
#include <stdlib.h>

#ifdef ASCENDC_CPU_DEBUG
#define __aicore__
#else
#define __aicore__ [aicore]
#endif

#define BLOCKS 4
#define CACHELINE_SZ 64

// Define a kernel
__global__ __aicore__ void foo(__gm__ uint8_t *Out, int Stride) {
    Out[block_idx * Stride] = block_idx;
}

int main(int argc, char *argv[]) {
    aclInit(nullptr);
    aclrtSetDevice(0);
    aclrtStream stream;
    aclrtCreateStream(&stream);

    uint8_t ExpectedValue[] = {0, 1, 2, 3};
    uint8_t *OutputValue = nullptr;
    aclrtMalloc((void **)&OutputValue, BLOCKS, ACL_MEM_MALLOC_HUGE_FIRST);

    uint8_t InitValue[BLOCKS] = {0};
    aclrtMemcpyAsync((void *)OutputValue, sizeof(InitValue), InitValue,
                     sizeof(InitValue), ACL_MEMCPY_HOST_TO_DEVICE, stream);
    aclrtSynchronizeStream(stream);

    // Invoke a kernel
    foo<<<BLOCKS, nullptr, stream>>>(OutputValue, CACHELINE_SZ);

    uint8_t *OutHost = nullptr;
    aclrtMallocHost((void **)&OutHost, BLOCKS * CACHELINE_SZ);
    aclrtMemcpyAsync(OutHost, BLOCKS * CACHELINE_SZ, OutputValue,
                     BLOCKS * CACHELINE_SZ, ACL_MEMCPY_DEVICE_TO_HOST, stream);
    aclrtSynchronizeStream(stream);

    for (int I = 0; I < sizeof(ExpectedValue) / sizeof(uint8_t); I++) {
        printf("i%d\t Expect: 0x%04x\t\t\t\tResult: 0x%04x\n", I, ExpectedValue[I],
                OutHost[I * CACHELINE_SZ]);
    }

    aclrtFreeHost(OutHost);
    aclrtFree(OutputValue);

    aclrtDestroyStream(stream);
    aclrtResetDevice(0);
    aclFinalize();
    return 0;
}

编译命令如下,编译选项的具体介绍请参考编译选项

# CANN软件包中的runtime路径
export RT_INC=${INSTALL_DIR}/runtime/include
export RT_LIB=${INSTALL_DIR}/runtime/lib64

# 功能:Host & Device代码混合编译,生成可执行文件,需链接libascendcl.so 和 libruntime.so
# 编译选项--cce-soc-version和--cce-soc-core-type指的是编译AscendXXXYY上的Vector核程序
$bisheng -O2 --cce-soc-version=AscendXXXYY --cce-soc-core-type=VecCore  -I$RT_INC -L$RT_LIB -lascendcl -lruntime QuickStartDemo.cce  -o QuickStartDemo

运行结果如下:

1
2
3
4
5
$ ./QuickStartDemo
i0       Expect: 0x0000                         Result: 0x0000
i1       Expect: 0x0001                         Result: 0x0001
i2       Expect: 0x0002                         Result: 0x0002
i3       Expect: 0x0003                         Result: 0x0003