Trunc【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit产品支持情况产品是否支持Ascend 950PR/Ascend 950DT√Atlas A3 训练系列产品 / Atlas A3 推理系列产品√Atlas A2 训练系列产品 / Atlas A2 推理系列产品√功能说明按元素做浮点数截断即向零取整操作。计算公式如下举例如下Trunc(3.9) 3Trunc(-3.9) -3函数原型通过sharedTmpBuffer入参传入临时空间源操作数Tensor全部/部分参与计算template typename T, bool isReuseSource false __aicore__ inline void Trunc(const LocalTensorT dstTensor, const LocalTensorT srcTensor, const LocalTensoruint8_t sharedTmpBuffer, const uint32_t calCount)源操作数Tensor全部参与计算template typename T, bool isReuseSource false __aicore__ inline void Trunc(const LocalTensorT dstTensor, const LocalTensorT srcTensor, const LocalTensoruint8_t sharedTmpBuffer)接口框架申请临时空间源操作数Tensor全部/部分参与计算template typename T, bool isReuseSource false __aicore__ inline void Trunc(const LocalTensorT dstTensor, const LocalTensorT srcTensor, const uint32_t calCount)源操作数Tensor全部参与计算template typename T, bool isReuseSource false __aicore__ inline void Trunc(const LocalTensorT dstTensor, const LocalTensorT srcTensor)由于该接口的内部实现中涉及精度转换。需要额外的临时空间来存储计算过程中的中间变量。临时空间支持接口框架申请和开发者通过sharedTmpBuffer入参传入两种方式。接口框架申请临时空间开发者无需申请但是需要预留临时空间的大小。通过sharedTmpBuffer入参传入使用该tensor作为临时空间进行处理接口框架不再申请。该方式开发者可以自行管理sharedTmpBuffer内存空间并在接口调用完成后复用该部分内存内存不会反复申请释放灵活性较高内存利用率也较高。接口框架申请的方式开发者需要预留临时空间通过sharedTmpBuffer传入的情况开发者需要为tensor申请空间。临时空间大小BufferSize的获取方式如下通过GetTruncMaxMinTmpSize中提供的GetTruncMaxMinTmpSize接口获取需要预留空间的大小。参数说明表 1模板参数说明参数名描述T操作数的数据类型。Ascend 950PR/Ascend 950DT支持的数据类型为half、float。Atlas A3 训练系列产品 / Atlas A3 推理系列产品支持的数据类型为half、float。Atlas A2 训练系列产品 / Atlas A2 推理系列产品支持的数据类型为half、float。isReuseSource是否允许修改源操作数。该参数预留传入默认值false即可。表 2接口参数说明参数名输入/输出描述dstTensor输出目的操作数。类型为LocalTensor支持的TPosition为VECIN/VECCALC/VECOUT。srcTensor输入源操作数。类型为LocalTensor支持的TPosition为VECIN/VECCALC/VECOUT。sharedTmpBuffer输入临时缓存。类型为LocalTensor支持的TPosition为VECIN/VECCALC/VECOUT。用于Trunc内部复杂计算时存储中间变量由开发者提供。临时空间大小BufferSize的获取方式请参考GetTruncMaxMinTmpSize。calCount输入参与计算的元素个数。返回值说明无约束说明支持源操作数与目的操作数地址重叠。不支持sharedTmpBuffer与源操作数和目的操作数地址重叠。操作数地址对齐要求请参见通用地址对齐约束。调用示例#include kernel_operator.h template typename srcType class KernelTrunc { public: __aicore__ inline KernelTrunc(){} __aicore__ inline void Init(GM_ADDR srcGm, GM_ADDR dstGm, uint32_t srcSize) { src_global.SetGlobalBuffer(reinterpret_cast__gm__ srcType *(srcGm), srcSize); dst_global.SetGlobalBuffer(reinterpret_cast__gm__ srcType *(dstGm), 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() { AscendC::LocalTensorsrcType srcLocal inQueueX.AllocTensorsrcType(); AscendC::DataCopy(srcLocal, src_global, bufferSize); inQueueX.EnQue(srcLocal); } __aicore__ inline void Compute() { AscendC::LocalTensorsrcType dstLocal outQueue.AllocTensorsrcType(); AscendC::LocalTensorsrcType srcLocal inQueueX.DeQuesrcType(); AscendC::TruncsrcType, false(dstLocal, srcLocal); outQueue.EnQuesrcType(dstLocal); inQueueX.FreeTensor(srcLocal); } __aicore__ inline void CopyOut() { AscendC::LocalTensorsrcType dstLocal outQueue.DeQuesrcType(); AscendC::DataCopy(dst_global, dstLocal, bufferSize); outQueue.FreeTensor(dstLocal); } private: AscendC::GlobalTensorsrcType src_global; AscendC::GlobalTensorsrcType dst_global; AscendC::TPipe pipe; AscendC::TQueAscendC::TPosition::VECIN, 1 inQueueX; AscendC::TQueAscendC::TPosition::VECOUT, 1 outQueue; uint32_t bufferSize 0; }; template typename dataType __aicore__ void kernel_trunc_operator(GM_ADDR srcGm, GM_ADDR dstGm, uint32_t srcSize) { KernelTruncdataType op; op.Init(srcGm, dstGm, srcSize); op.Process(); } extern C __global__ __aicore__ void trunc_operator(GM_ADDR srcGm, GM_ADDR dstGm, uint32_t srcSize) { kernel_trunc_operatorhalf(srcGm, dstGm, srcSize); }结果示例如下输入数据(srcLocal): [ 0.5317103 -6.37912032 5.53408647 ... 11.11059642 -11.67860335] 输出数据(dstLocal): [ 0.0 -6.0 5.0 ... 11.0 -11.0 ]【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考