【2023·CANN训练营第一季】TIK C++算子执行的不同模式

390 阅读2分钟

TIK C++算子执行的不同模式

TIK C++算子可用**CPU模式NPU模式**执行

CPU模式: 算子功能调试用,可以模拟在NPU上的计算行为,不需要依赖昇腾设备

NPU模式: 算子功能/性能调试用,可以使用NPU的强大算力进行运算加速

使用方式

使用**内置宏 __CCE_KT_TEST__ **标识被宏包括的代码在特定的模式下编译

#ifdef __CCE_KT_TEST__ 表示在CPU模式下会编译该段代码

#ifndef __CCE_KT_TEST__ 表示在NPU模式下会编译该段代码

不同模式区别

CPU模式NPU模式
依赖昇腾设备不依赖依赖
使用的编译器g++编译器自研ccec编译器
使用场景算子功能调试算子功能/性能调试
推荐开发顺序
模式限制只限于精度调试调试较为复杂

编写思路

1. 运行CPU模式包含的头文件

#include “tikicpulib.h"
#define_aicore_

2. 运行NPU模式包含的头文件

#include"acl/acl.h"
#define_aicore [aicore]

3. 核函数的定义

extern "C" __global__ __aicore__ void Hellowor1d(__gm__ uint&_t*.foo){}

4. 设备侧执行逻辑

void Hellowor1d(__gm__ uint&_t*.foo){}

5. 主机侧执行逻辑

负责数据在主机侧内存的申请,主机到设备的拷贝,核函数执行同步和回收资源的工作

int32_t main(int32_t argc, char* argv[]){
  size_t fooSize = 256;
  uint32_t blockDim = 8;
#ifdef __CCE_KT_TEST__

内置宏 __CCE_KT_TEST__

是区分运行CPU模式或NPU模式逻辑的标志

  • 主机侧执行CPU模式逻辑

    使用封装的执行宏ICPU_RUN_KF

    uint8_t *foo = (uint8_t *)tik2::GmAl1oc(fooSize);
    ICPU_RUN_KF(Hellolorld, blockDim, foo);
    tik2::GmFree((void*)foo);
    
  • 主机侧执行NPU模式逻辑

    使用内核调用符<<<…>>>

    #else
      aclInit (nullptr);
      acirtStream stream = nullptr;
      aclrtCreatestream(&stream);
    
      uint8_t *fooDevice;
      aclrtMalloc((void**)&fooDevice, fooSize, ACL_MEM_MALLOC_HUGE_FIRST);
      Hellolorld<<<blockDim, nullptr, stream>>>(fooDevice);
      aclrtSynchronizestream(stream);
    ​
      aclrtFree(fooDevice);
      aclrtDestroystream(stream);
      aclFinalize();
    ```
    
#endif
  return 0;  
}

编译程序

  • 设置环境变量

    install_path=/usr/local/Ascend/ascend-toolkit/latest
    source ${install_path}/bin/setenv.bash
    # 或 
    source ${install_path}/../set_env.sh
    
  • CPU模式的流程命令

    (根据不同的芯片版本而异)

    # 编译命令
    g++ hello_world.cpp                                           \
      -D__CCE_KT_TEST__=1 -D__CCE_AICORE__=100 -D__DAV_C100__     \
      -I${install_path}/tools/tikicpulib/lib/include              \
      -L${install_path}/toolkit/tools/simulator/Ascend910A/lib    \
      -L${install_path}/tools/tikicpulib/lib -ltikicpulib_stubreg \
      -L${install_path}/tools/tikicpulib/lib/Ascend910 -lstdc++   \
      -O2 -std=c++17 -o hello_world_cpu
    
    # 执行命令
    export LD_LIBRARY_PATH=${install_path}/tools/tikicpulib/lib:$LD_LIBRARY_PATH
    ./hello_world_cpu
    

image.png

CPU模式

  • NPU模式的流程命令

    (根据不同的芯片版本而异)

    # 编译命令
      ccec -x cce hello_world.cpp             \
         --cce-aicore-arch=dav-c100         \
           -I${install_path}/acllib/include   \
         -L${install_path}/atc/lib64        \
           -lruntime -lascendcl -lstdc++      \
         -O2 -std=c++17 -o hello_world_npu
           
    # 执行命令
      ./hello_world_npu
    

image.png

NPU模式