一、什么是TIK C++？

????????TIK C++是一種使用C/C++作為前端語言的算子開發(fā)工具，通過四層接口抽象、并行編程范式、孿生調(diào)試等技術(shù)，極大提高算子開發(fā)效率，助力AI開發(fā)者低成本完成算子開發(fā)和模型調(diào)優(yōu)部署。

使用TIK C++開發(fā)自定義算子的優(yōu)勢：

C/C++原語編程

編程模型屏蔽硬件差異，編程范式提高開發(fā)效率

多層級API封裝，從簡單到靈活，兼顧易用與高效

孿生調(diào)試，CPU側(cè)模擬NPU側(cè)的行為，可先在CPU側(cè)調(diào)試

二、核函數(shù)的概念和使用

????????核函數(shù)（Kernel Function）是TIK C++算子設(shè)備側(cè)的入口。TIK C++允許用戶使用核函數(shù)這種C/C++函數(shù)的語法擴展來管理設(shè)備側(cè)的運行代碼，用戶在核函數(shù)中實現(xiàn)算子邏輯的編寫，例如自定義算子類及其成員函數(shù)以實現(xiàn)該算子的所有功能。核函數(shù)是主機側(cè)和設(shè)備側(cè)連接的橋梁。

????????核函數(shù)是直接在設(shè)備側(cè)執(zhí)行的代碼。在核函數(shù)中，需要為在一個核上執(zhí)行的代碼規(guī)定要進行的數(shù)據(jù)訪問和計算操作，當(dāng)核函數(shù)被調(diào)用時，多個核將并行執(zhí)行同一個計算任務(wù)

2.1 如何編寫核函數(shù)？

1、使用函數(shù)類型限定符

????????除了需要按照C/C++函數(shù)聲明的方式定義核函數(shù)之外，還要為核函數(shù)加上額外的函數(shù)類型限定符，包含__global__和__aicore__ 使用__global__函數(shù)類型限定符來標(biāo)識它是一個核函數(shù)，可以被<<<...>>>調(diào)用；使用__aicore__函數(shù)類型限定符來標(biāo)識該核函數(shù)在設(shè)備側(cè)AI Core上執(zhí)行：

其他規(guī)則：

__global__ __aicore__voidkernel_name(argument list)

2、使用變量類型限定符

?????????指針入?yún)⒆兞拷y(tǒng)一的類型定義為__gm__ uint8_t* 用戶可統(tǒng)一使用uint8_t類型的指針，并在使用時轉(zhuǎn)化為實際的指針類型；亦可直接傳入實際的指針類型。

必須具有void返回類型

使用extern "C"

僅支持入?yún)橹羔橆愋突駽/C++內(nèi)置數(shù)據(jù)類型(Primitive Data Types)，如：half* s0、float* s1、int32_t c

2.2?如何調(diào)用核函數(shù)？

常見的C/C++函數(shù)調(diào)用方式是如下的形式：

function_name(argument list)

核函數(shù)使用內(nèi)核調(diào)用符<<<...>>>這種語法形式，來規(guī)定核函數(shù)的執(zhí)行配置：

kernel_name<<<blockDim, l2ctrl, stream>>>(argument list);

使用內(nèi)核調(diào)用符<<<…>>>調(diào)用核函數(shù)：

HelloWorld<<<8,nullptr, stream>>>(fooDevice);

blockDim設(shè)置為8，表示在8個核上調(diào)用了HelloWorld核函數(shù)，每個核都會獨立且并行地執(zhí)行該核函數(shù) Stream可以通過aclrtCreateStream來創(chuàng)建，它的作用是在當(dāng)前進程或線程中顯式創(chuàng)建一個aclrtStream argument list設(shè)置為fooDevice這1個入?yún)?/p>

其中：

blockDim，規(guī)定了核函數(shù)將會在幾個核上執(zhí)行，每個執(zhí)行該核函數(shù)的核會被分配一個邏輯ID，表現(xiàn)為內(nèi)置變量block_idx，編號從0開始，可為不同的邏輯核定義不同的行為，可以在算子實現(xiàn)中使用。

l2ctrl，保留參數(shù)，暫時設(shè)置為固定值nullptr。

stream，類型為aclrtStream，stream是一個任務(wù)隊列，應(yīng)用程序通過stream來管理任務(wù)的并行。

核函數(shù)的調(diào)用是異步的，核函數(shù)的調(diào)用結(jié)束后，控制權(quán)立刻返回給主機側(cè) 強制主機側(cè)程序等待所有核函數(shù)執(zhí)行完畢的API（阻塞應(yīng)用程序運行，直到指定Stream中的所有任務(wù)都完成，同步接口）為aclrtSynchronizeStream

三、TIK C++的helloworld樣例

3.1、算子執(zhí)行的不同模式

CPU模式：算子功能調(diào)試用，可以模擬在NPU上的計算行為，不需要依賴昇騰設(shè)備

NPU模式：算子功能/性能調(diào)試，可以使用NPU的強大算力進行運算加速

使用內(nèi)置宏__CCE_KT_TEST__標(biāo)識被宏包括的代碼在特定的模式下編譯

#ifdef __CCE_KT_TEST__ 表示在CPU模式下會編譯該段代碼

#ifndef __CCE_KT_TEST__ 表示在NPU模式下會編譯該段代碼

?上述步驟：

3.2、樣例展示及各模塊功能

# 設(shè)置環(huán)境變量

install_path=/usr/local/Ascend/ascend-toolkit/latest

source ${install_path}/bin/setenv.bash

或 source ${install_path}/../set_env.sh

(根據(jù)不同的芯片版本而異)

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

CPU模式的流程命令（# 執(zhí)行命令）

exportLD_LIBRARY_PATH=${install_path}/tools/tikicpulib/lib:$LD_LIBRARY_PATH./hello_world_cpu

CPU模式的流程命令執(zhí)行結(jié)果：

?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

NPU模式的流程命令（# 執(zhí)行命令）

./hello_world_npu

NPU模式的流程命令執(zhí)行結(jié)果