Compare

函数功能

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

支持多种比较模式：

LT：小于（lower-than）
GT：大于（greater-than）

GE：大于或等于（greater-equal）
EQ：等于（equal）
NE：不等于（not-equal）
LE：小于或等于（lower-equal）

函数原型

整个tensor参与计算
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个数据计算
 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模式
   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[2], uint8_t repeatTimes, const BinaryRepeatParams& repeatParams)
- mask连续模式
   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。 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。 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用于控制每次迭代内参与计算的元素。（保留参数，设置无效）
repeatTimes	输入	重复迭代次数。矢量计算单元，每次读取连续的256 Bytes数据进行计算，为完成对输入数据的处理，必须通过多次迭代（repeat）才能完成所有数据的读取与计算。repeatTimes表示迭代的次数。关于该参数的具体描述请参考通用参数说明。
BinaryRepeatParams	输入	控制操作数地址步长的数据结构。结构体内包含操作数相邻迭代间相同block的地址步长，操作数同一迭代内不同block的地址步长等参数。该数据结构的定义请参考BinaryRepeatParams。相邻迭代间的地址步长参数说明请参考Repeat stride（相邻迭代间相同datablock的地址步长）；同一迭代内datablock的地址步长参数说明请参考Block stride（同一迭代内不同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参与计算
```
dstLocal = src0Local < src1Local;
```

tensor前n个数据计算

Compare(dstLocal, src0Local, src1Local, CMPMODE::LT, srcDataSize);

样例模板

#include "kernel_operator.h"
namespace AscendC {
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()
    {
        LocalTensor<float> src0Local = inQueueSrc0.AllocTensor<float>();
        LocalTensor<float> src1Local = inQueueSrc1.AllocTensor<float>();
        DataCopy(src0Local, src0Global, srcDataSize);
        DataCopy(src1Local, src1Global, srcDataSize);
        inQueueSrc0.EnQue(src0Local);
        inQueueSrc1.EnQue(src1Local);
    }
    __aicore__ inline void Compute()
    {
        LocalTensor<float> src0Local = inQueueSrc0.DeQue<float>();
        LocalTensor<float> src1Local = inQueueSrc1.DeQue<float>();
        LocalTensor<uint8_t> dstLocal = outQueueDst.AllocTensor<uint8_t>();

        // 可根据实际使用接口Compare进行替换
       // Compare(dstLocal, src0Local, src1Local, CMPMODE::LT, srcDataSize);

        outQueueDst.EnQue<uint8_t>(dstLocal);
        inQueueSrc0.FreeTensor(src0Local);
        inQueueSrc1.FreeTensor(src1Local);
    }
    __aicore__ inline void CopyOut()
    {
        LocalTensor<uint8_t> dstLocal = outQueueDst.DeQue<uint8_t>();
        DataCopy(dstGlobal, dstLocal, dstDataSize);
        outQueueDst.FreeTensor(dstLocal);
    }
private:
    TPipe pipe;
    TQue<QuePosition::VECIN, 1> inQueueSrc0, inQueueSrc1;
    TQue<QuePosition::VECOUT, 1> outQueueDst;
    GlobalTensor<float> src0Global, src1Global;
    GlobalTensor<uint8_t> dstGlobal;
    uint32_t srcDataSize = 256;
    uint32_t dstDataSize = srcDataSize / AscendCUtils::GetBitSize(sizeof(uint8_t));
};
} // namespace AscendC
extern "C" __global__ __aicore__ void main_cpu_cmp_sel_demo(__gm__ uint8_t* src0Gm, __gm__ uint8_t* src1Gm, __gm__ uint8_t* dstGm)
{
    AscendC::KernelCmp op;
    op.Init(src0Gm, src1Gm, dstGm);
    op.Process();
}

父主题： 比较指令