CANN学习资源开源仓的算子调试

在非昇腾设备上，可以利用CPU仿真环境进行算子开发和测试，实际上是通过标准的GCC编译器编译算子Kernel程序，链接CPU调测库，编译生成的可执行文件，完成算子CPU域的运行验证。比如，使用AI Core数量为8核，CPU调试中会创建8个线程模拟并行。

add.cpp 是CPU侧验证工程的核心文件，包含头文件引入与常量定义、算子核函数实现以及CPU侧主程序调用三部分。注意这里忽略tiling函数，然后在Init里面有个printf调试接口打印。

//AscendC算子开发的核心头文件，提供了向量计算、数据搬运、流水线管理等基础API。
#include "kernel_operator.h"
//自定义的数据工具头文件，包含 ReadFile、WriteFile 等辅助函数，用于文件读写。
#include "data_utils.h"
//CPU调试库头文件，仅在CPU仿真环境下使用。
#include "tikicpulib.h"

constexpr int32_t TOTAL_LENGTH = 8 * 2048;                            // total length of data
constexpr int32_t USE_CORE_NUM = 8;                                   // num of core used
constexpr int32_t BLOCK_LENGTH = TOTAL_LENGTH / USE_CORE_NUM;         // length computed of each core
constexpr int32_t TILE_NUM = 8;                                       // split data into 8 tiles for each core
constexpr int32_t BUFFER_NUM = 2;                                     // tensor num for each queue
constexpr int32_t TILE_LENGTH = BLOCK_LENGTH / TILE_NUM / BUFFER_NUM; // separate to 2 parts, due to double buffer

class KernelAdd {
public:
    __aicore__ inline KernelAdd() {}
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z)
    {
        xGm.SetGlobalBuffer((__gm__ half *)x + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH);
        yGm.SetGlobalBuffer((__gm__ half *)y + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH);
        zGm.SetGlobalBuffer((__gm__ half *)z + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH);
        pipe.InitBuffer(inQueueX, BUFFER_NUM, TILE_LENGTH * sizeof(half));
        pipe.InitBuffer(inQueueY, BUFFER_NUM, TILE_LENGTH * sizeof(half));
        pipe.InitBuffer(outQueueZ, BUFFER_NUM, TILE_LENGTH * sizeof(half));
        // 添加打印调试信息
        AscendC::printf("[Block (%lu/%lu)]: BLOCK_LENGTH is %d\n", AscendC::GetBlockIdx(), AscendC::GetBlockNum(), BLOCK_LENGTH);
    }
    __aicore__ inline void Process()
    {
        int32_t loopCount = TILE_NUM * BUFFER_NUM;
        for (int32_t i = 0; i < loopCount; i++) {
            CopyIn(i);
            Compute(i);
            CopyOut(i);
        }
    }

private:
    __aicore__ inline void CopyIn(int32_t progress)
    {
        AscendC::LocalTensor<half> xLocal = inQueueX.AllocTensor<half>();
        AscendC::LocalTensor<half> yLocal = inQueueY.AllocTensor<half>();
        AscendC::DataCopy(xLocal, xGm[progress * TILE_LENGTH], TILE_LENGTH);
        AscendC::DataCopy(yLocal, yGm[progress * TILE_LENGTH], TILE_LENGTH);
        inQueueX.EnQue(xLocal);
        inQueueY.EnQue(yLocal);
    }
    __aicore__ inline void Compute(int32_t progress)
    {
        AscendC::LocalTensor<half> xLocal = inQueueX.DeQue<half>();
        AscendC::LocalTensor<half> yLocal = inQueueY.DeQue<half>();
        AscendC::LocalTensor<half> zLocal = outQueueZ.AllocTensor<half>();
        AscendC::Add(zLocal, xLocal, yLocal, TILE_LENGTH);
        outQueueZ.EnQue<half>(zLocal);
        inQueueX.FreeTensor(xLocal);
        inQueueY.FreeTensor(yLocal);
    }
    __aicore__ inline void CopyOut(int32_t progress)
    {
        AscendC::LocalTensor<half> zLocal = outQueueZ.DeQue<half>();
        AscendC::DataCopy(zGm[progress * TILE_LENGTH], zLocal, TILE_LENGTH);
        outQueueZ.FreeTensor(zLocal);
    }

private:
    AscendC::TPipe pipe;
    AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;
    AscendC::TQue<AscendC::TPosition::VECOUT, BUFFER_NUM> outQueueZ;
    AscendC::GlobalTensor<half> xGm;
    AscendC::GlobalTensor<half> yGm;
    AscendC::GlobalTensor<half> zGm;
};

__global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)
{
    KernelAdd op;
    op.Init(x, y, z);
    op.Process();
}

int32_t main(int32_t argc, char *argv[])
{
    uint32_t numBlocks = 8;
    size_t inputByteSize = 8 * 2048 * sizeof(uint16_t);
    size_t outputByteSize = 8 * 2048 * sizeof(uint16_t);

    //GmAlloc 进行核函数的CPU侧运行验证时，用于创建共享内存：在/tmp目录下创建一个共享文件，并返回该文件的映射指针。
    uint8_t *x = (uint8_t *)AscendC::GmAlloc(inputByteSize);
    uint8_t *y = (uint8_t *)AscendC::GmAlloc(inputByteSize);
    uint8_t *z = (uint8_t *)AscendC::GmAlloc(outputByteSize);

    ReadFile("../input_x.bin", inputByteSize, x, inputByteSize);
    ReadFile("../input_y.bin", inputByteSize, y, inputByteSize);

    // 矢量算子需要设置内核模式为AIV模式
    AscendC::SetKernelMode(KernelMode::AIV_MODE);
    // 调用ICPU_RUN_KF调测宏，完成核函数CPU侧的调用
    ICPU_RUN_KF(add_custom, numBlocks, x, y, z); // use this macro for cpu debug

    // 输出数据写出
    WriteFile("../output.bin", z, outputByteSize);

    // 调用GmFree释放申请的资源
    AscendC::GmFree((void *)x);
    AscendC::GmFree((void *)y);
    AscendC::GmFree((void *)z);
    return 0;
}

data_utils.h的内容略。生成输入数据和标杆结果数据的gen_data.py略。CMakeLists.txt内容：

cmake_minimum_required(VERSION 3.16)
project(cpudebug)

# 查找并加载 tikicpulib（CPU 调试必备），若找不到则终止配置
set(tikicpulib_DIR "/usr/local/Ascend/cann-8.5.0/tools/tikicpulib/lib/cmake")
find_package(tikicpulib REQUIRED)
add_executable(add ${CMAKE_CURRENT_SOURCE_DIR}/add.cpp)
target_compile_options(add PRIVATE
    -g -O2 -std=c++17 -D_GLIBCXX_USE_CXX11_ABI=0 -Wall -Werror
)
# 链接所需库：
target_link_libraries(add PRIVATE
    ascendcl
    tikicpulib::${soc_version}
)

install(TARGETS add
    LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR}
    ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR}
    RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR}
)

编译运行：

cd Sources/cpu_debug && mkdir -p build && cd build && cmake .. -Dsoc_version=Ascend910B4 && make -j 
cd Sources/cpu_debug/build && ./add

然后验证结果verify_result.py略。最后说一下，CPU仿真调试验证我没用过，觉得是多此一举。在开源仓算子项目的基础上要做CPU调试，要改东西，麻烦，我直接上NPU环境去测试验证不好吗？