Log

功能说明

按元素以e、2、10为底做对数运算，计算公式如下，其中PAR表示矢量计算单元一个迭代能够处理的元素个数 ：

定义原型

参数说明

表1 2级接口参数说明

参数名

输入/输出

描述

dstLocal

输出

目的操作数。

类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。

Atlas A2训练系列产品，支持的数据类型为：half/float

srcLocal

输入

源操作数。

类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。

源操作数的数据类型需要与目的操作数保持一致。

Atlas A2训练系列产品，支持的数据类型为：half/float

tmpTensor

输入

中间操作数，Log2指令执行期间，需要一块临时空间用于存储中间结果。

类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。

若输入为half类型，临时空间的大小请参考定义原型。

若输入为float类型，不需要临时空间。

Atlas A2训练系列产品，支持的数据类型为：half/float

isReuseSource

输入

是否允许修改源操作数。该参数预留，传入默认值false即可。

isBasicBlock

输入

是否使用基本块进行计算。

True: 若输入源数据数量大于基本块大小(1024)，则会对输入数据进行分块，以块为单位进行处理。

False: 根据输入源数据量，分配相应的额外临时空间进行处理。

返回值

无

支持的型号

Atlas A2训练系列产品

约束说明

调用示例

本样例中只展示Compute流程中的部分代码。如果您需要运行样例代码，请将该代码段拷贝并替换样例模板中Compute函数的部分代码即可。

样例模板

#include "kernel_operator.h"
namespace AscendC {
template <typename srcType>
class KernelLog {
public:
    __aicore__ inline KernelLog()
    {}
    __aicore__ inline void Init(GM_ADDR src_gm, GM_ADDR dst_gm, uint32_t srcSize)
    {
        src_global.SetGlobalBuffer(reinterpret_cast<__gm__ srcType *>(src_gm), srcSize);
        dst_global.SetGlobalBuffer(reinterpret_cast<__gm__ srcType *>(dst_gm), srcSize);
        pipe.InitBuffer(inQueueX, 1, srcSize * sizeof(srcType));
        pipe.InitBuffer(outQueue, 1, srcSize * sizeof(srcType));
        bufferSize = srcSize;
    }
    __aicore__ inline void Process()
    {
        CopyIn();
        Compute();
        CopyOut();
    }
private:
    __aicore__ inline void CopyIn()
    {
        LocalTensor<srcType> srcLocal = inQueueX.AllocTensor<srcType>();
        DataCopy(srcLocal, src_global, bufferSize);
        inQueueX.EnQue(srcLocal);
    }
    __aicore__ inline void Compute()
    {
        LocalTensor<srcType> dstLocal = outQueue.AllocTensor<srcType>();
        LocalTensor<srcType> srcLocal = inQueueX.DeQue<srcType>();
        Log(dstLocal, srcLocal);
        //Log10(dstLocal, srcLocal);
        //Log2<srcType, false, false>(dstLocal, srcLocal);
        outQueue.EnQue<srcType>(dstLocal);
        inQueueX.FreeTensor(srcLocal);
    }
    __aicore__ inline void CopyOut()
    {
        LocalTensor<srcType> dstLocal = outQueue.DeQue<srcType>();
        DataCopy(dst_global, dstLocal, bufferSize);
        outQueue.FreeTensor(dstLocal);
    }
private:
    GlobalTensor<srcType> src_global;
    GlobalTensor<srcType> dst_global;
    TPipe pipe;
    TQue<QuePosition::VECIN, 1> inQueueX;
    TQue<QuePosition::VECOUT, 1> outQueue;
    uint32_t bufferSize = 0;
};
template <typename dataType>
__aicore__ void kernel_log_operator(GM_ADDR src_gm, GM_ADDR dst_gm, uint32_t srcSize)
{
    KernelLog<dataType> op;
    op.Init(src_gm, dst_gm, srcSize);
    op.Process();
}
} // namespace AscendC
                                                                    
extern "C" __global__ __aicore__ void kernel_log_operator(GM_ADDR src_gm, GM_ADDR dst_gm, uint32_t srcSize) 
{                                                                                                                 
    AscendC::kernel_log_operator<half>(src_gm, dst_gm, srcSize);  //传入类型和大小                                                       
}
父主题： 数学库