文档
注册
评分
提单
论坛
小AI

Compare

函数功能

逐元素比较两个tensor大小,如果比较后的结果为真,则输出结果的对应比特位为1,否则为0。

支持多种比较模式:

  • LT:小于(lower-than)
  • GT:大于(greater-than)
  • GE:大于或等于(greater-equal)
  • EQ:等于(equal)
  • NE:不等于(not-equal)
  • LE:小于或等于(lower-equal)

函数原型

  • 整个tensor参与计算
    1
    2
    3
    4
    5
    6
    dstLocal = src0Local < src1Local;
    dstLocal = src0Local > src1Local;
    dstLocal = src0Local <= src1Local;
    dstLocal = src0Local >= src1Local;
    dstLocal = src0Local == src1Local;
    dstLocal = src0Local != src1Local;
    

    Atlas 200/500 A2推理产品暂不支持整个tensor参与计算的运算符重载。

  • tensor前n个数据计算
    1
    2
    template <typename T, typename U>
    __aicore__ inline void Compare(const LocalTensor<U>& dstLocal, const LocalTensor<T>& src0Local, const LocalTensor<T>& src1Local, CMPMODE cmpMode, uint32_t calCount)
    
  • tensor高维切分计算
    • mask逐bit模式
      1
      2
      template <typename T, typename U, bool isSetMask = true>
      __aicore__ inline void Compare(const LocalTensor<U>& dstLocal, const LocalTensor<T>& src0Local, const LocalTensor<T>& src1Local, CMPMODE cmpMode, const uint64_t mask[], uint8_t repeatTimes, const BinaryRepeatParams& repeatParams)
      
    • mask连续模式
      1
      2
      template <typename T, typename U, bool isSetMask = true>
      __aicore__ inline void Compare(const LocalTensor<U>& dstLocal, const LocalTensor<T>& src0Local, const LocalTensor<T>& src1Local, CMPMODE cmpMode, const uint64_t mask, uint8_t repeatTimes, const BinaryRepeatParams& repeatParams)
      

参数说明

表1 模板参数说明

参数名

描述

T

操作数数据类型。

isSetMask

保留参数,保持默认值即可。

表2 接口参数说明

参数名称

输入/输出

含义

dstLocal

输出

目的操作数。

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

LocalTensor的起始地址需要32字节对齐。

dstLocal用于存储比较结果,将dstLocal中uint8_t类型的数据按照bit位展开,由左至右依次表征对应位置的src0和src1的比较结果,如果比较后的结果为真,则对应比特位为1,否则为0。

Atlas 训练系列产品,支持的数据类型为:int8_t/uint8_t

Atlas推理系列产品AI Core,支持的数据类型为:int8_t/uint8_t

Atlas A2训练系列产品/Atlas 800I A2推理产品,支持的数据类型为:int8_t/uint8_t

Atlas 200/500 A2推理产品,支持的数据类型为:int8_t/uint8_t

src0Local、src1Local

输入

源操作数。

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

LocalTensor的起始地址需要32字节对齐。

Atlas 训练系列产品,支持的数据类型为:half/float

Atlas推理系列产品AI Core,支持的数据类型为:half/float

Atlas A2训练系列产品/Atlas 800I A2推理产品,支持的数据类型为:half/float(所有CMPMODE都支持), int32_t(只支持CMPMODE::EQ)

Atlas 200/500 A2推理产品,支持的数据类型为:half/float

cmpMode

输入

CMPMODE类型,表示比较模式,包括EQ,NE,GE,LE,GT,LT。

  • LT: src0小于(lower-than)src1
  • GT: src0大于(greater-than)src1
  • GE:src0大于或等于(greater-equal)src1
  • EQ:src0等于(equal)src1
  • NE:src0不等于(not-equal)src1
  • LE:src0小于或等于(lower-equal)src1

mask

输入

mask用于控制每次迭代内参与计算的元素。

Atlas 训练系列产品,保留参数,设置无效。

Atlas推理系列产品AI Core,保留参数,设置无效。

Atlas A2训练系列产品/Atlas 800I A2推理产品,保留参数,设置无效。

Atlas 200/500 A2推理产品,设置有效。

  • 连续模式:表示前面连续的多少个元素参与计算。取值范围和操作数的数据类型有关,数据类型不同,每次迭代内能够处理的元素个数最大值不同。当操作数为8位时,mask∈[1, 256];当操作数为16位时,mask∈[1, 128];当操作数为32位时,mask∈[1, 64]。
  • 逐bit模式:可以按位控制哪些元素参与计算,bit位的值为1表示参与计算,0表示不参与。参数类型为长度为2或者4的uint64_t类型数组。

    例如,mask=[8, 0],8=0b1000,表示仅第4个元素参与计算。

    参数取值范围和操作数的数据类型有关,数据类型不同,每次迭代内能够处理的元素个数最大值不同。当操作数为8位时,mask[0]、mask[1]、mask[2]、mask[3]∈[0, 264-1]并且不同时为0;当操作数为16位时,mask[0]、mask[1]∈[0, 264-1]并且不同时为0;当操作数为32位时,mask[1]为0,mask[0]∈(0, 264-1]。

repeatTimes

输入

重复迭代次数。矢量计算单元,每次读取连续的256 Bytes数据进行计算,为完成对输入数据的处理,必须通过多次迭代(repeat)才能完成所有数据的读取与计算。repeatTimes表示迭代的次数。

关于该参数的具体描述请参考通用参数说明

repeatParams

输入

控制操作数地址步长的参数。BinaryRepeatParams类型,包含操作数相邻迭代间相同datablock的地址步长,操作数同一迭代内不同datablock的地址步长等参数。

相邻迭代间的地址步长参数说明请参考repeatStride(相邻迭代间相同datablock的地址步长);同一迭代内datablock的地址步长参数说明请参考dataBlockStride(同一迭代内不同datablock的地址步长)

calCount

输入

输入数据元素个数。设置calCount时,需要保证calCount个元素所占空间256字节对齐。

返回值

支持的型号

Atlas 训练系列产品

Atlas推理系列产品AI Core

Atlas A2训练系列产品/Atlas 800I A2推理产品

Atlas 200/500 A2推理产品

注意事项

  • dstLocal支持多种输入数据类型,内部会将其转换为uint8_t类型处理。
  • dstLocal按照小端顺序排序成二进制结果,对应src中相应位置的数据比较结果。
  • 使用整个tensor参与计算的运算符重载功能,src0Local和src1Local需满足256Byte对齐;使用tensor前n个数据参与计算的接口,设置calCount时,需要保证calCount个元素所占空间256字节对齐。

调用示例

本样例中,源操作数src0和src1各存储了256个float类型的数据。样例实现的功能为,逐元素对src0和src1中的数据进行比较,如果src0中的元素小于src1中的元素,dst结果中对应的比特位置1;反之,则置0。dst结果使用uint8_t类型数据存储。

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

  • 整个tensor参与计算
    1
    dstLocal = src0Local < src1Local;
    
  • tensor前n个数据计算
    1
    AscendC::Compare(dstLocal, src0Local, src1Local, AscendC::CMPMODE::LT, srcDataSize);
    

样例模板

#include "kernel_operator.h"
class KernelCmp {
public:
    __aicore__ inline KernelCmp() {}
    __aicore__ inline void Init(__gm__ uint8_t* src0Gm, __gm__ uint8_t* src1Gm, __gm__ uint8_t* dstGm)
    {
        src0Global.SetGlobalBuffer((__gm__ float*)src0Gm);
        src1Global.SetGlobalBuffer((__gm__ float*)src1Gm);
        dstGlobal.SetGlobalBuffer((__gm__ uint8_t*)dstGm);
        pipe.InitBuffer(inQueueSrc0, 1, srcDataSize * sizeof(float));
        pipe.InitBuffer(inQueueSrc1, 1, srcDataSize * sizeof(float));
        pipe.InitBuffer(outQueueDst, 1, dstDataSize * sizeof(uint8_t));
    }
    __aicore__ inline void Process()
    {
        CopyIn();
        Compute();
        CopyOut();
    }
private:
    __aicore__ inline void CopyIn()
    {
        AscendC::LocalTensor<float> src0Local = inQueueSrc0.AllocTensor<float>();
        AscendC::LocalTensor<float> src1Local = inQueueSrc1.AllocTensor<float>();
        AscendC::DataCopy(src0Local, src0Global, srcDataSize);
        AscendC::DataCopy(src1Local, src1Global, srcDataSize);
        inQueueSrc0.EnQue(src0Local);
        inQueueSrc1.EnQue(src1Local);
    }
    __aicore__ inline void Compute()
    {
        AscendC::LocalTensor<float> src0Local = inQueueSrc0.DeQue<float>();
        AscendC::LocalTensor<float> src1Local = inQueueSrc1.DeQue<float>();
        AscendC::LocalTensor<uint8_t> dstLocal = outQueueDst.AllocTensor<uint8_t>();
 
        // 可根据实际使用接口Compare进行替换
        AscendC::Compare(dstLocal, src0Local, src1Local, AscendC::CMPMODE::LT, srcDataSize);
 
        outQueueDst.EnQue<uint8_t>(dstLocal);
        inQueueSrc0.FreeTensor(src0Local);
        inQueueSrc1.FreeTensor(src1Local);
    }
    __aicore__ inline void CopyOut()
    {
        AscendC::LocalTensor<uint8_t> dstLocal = outQueueDst.DeQue<uint8_t>();
        AscendC::DataCopy(dstGlobal, dstLocal, dstDataSize);
        outQueueDst.FreeTensor(dstLocal);
    }
private:
    AscendC::TPipe pipe;
    AscendC::TQue<AscendC::QuePosition::VECIN, 1> inQueueSrc0, inQueueSrc1;
    AscendC::TQue<AscendC::QuePosition::VECOUT, 1> outQueueDst;
    AscendC::GlobalTensor<float> src0Global, src1Global;
    AscendC::GlobalTensor<uint8_t> dstGlobal;
    uint32_t srcDataSize = 256;
    uint32_t dstDataSize = srcDataSize / AscendC::AscendCUtils::GetBitSize(sizeof(uint8_t));
};
extern "C" __global__ __aicore__ void main_cpu_cmp_sel_demo(__gm__ uint8_t* src0Gm, __gm__ uint8_t* src1Gm, __gm__ uint8_t* dstGm)
{
    KernelCmp op;
    op.Init(src0Gm, src1Gm, dstGm);
    op.Process();
}
搜索结果
找到“0”个结果

当前产品无相关内容

未找到相关内容,请尝试其他搜索词