Compare

函数功能

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

支持多种比较模式：

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

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

函数原型

参数说明

返回值

无

支持的型号

Atlas 训练系列产品

Atlas推理系列产品（Ascend 310P处理器）AI Core

Atlas A2训练系列产品

注意事项

• 操作数地址偏移对齐要求请参见通用约束。

• dstLocal支持多种输入数据类型，内部会将其转换为uint8_t类型处理。
• dstLocal按照小端顺序排序成二进制结果，对应src中相应位置的数据比较结果。
• 二级接口和三级接口中，src0Local和src1Local需满足256Byte对齐。

调用示例

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

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

• Compare-3级接口样例

  dstLocal = src0Local < src1Local;

• Compare-2级接口样例

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

• Compare-0级接口样例-mask连续模式

  uint64_t mask = 256/sizeof(float); // 256为每个迭代处理的字节数
  int repeat = 4;
  BinaryRepeatParams repeatParams = { 1, 1, 1, 8, 8, 8 };
  // repeat = 4, 64 elements one repeat, 256 elements total
  // dstBlkStride, src0BlkStride, src1BlkStride = 1, no gap between blocks in one repeat
  // dstRepStride, src0RepStride, src1RepStride = 8, no gap between repeats
  Compare(dstLocal, src0Local, src1Local, CMPMODE::LT, mask, repeat, repeatParams);

• Compare-0级接口样例-mask逐bit模式

  uint64_t mask[2] = { UINT64_MAX, 0};
  int repeat = 4;
  BinaryRepeatParams repeatParams = { 1, 1, 1, 8, 8, 8 };
  // repeat = 4, 64 elements one repeat, 256 elements total
  // srcBlkStride, = 1, no gap between blocks in one repeat
  // dstRepStride, srcRepStride = 8, no gap between repeats
  Compare(dstLocal, src0Local, src1Local, CMPMODE::LT, mask, repeat, repeatParams);

样例模板

#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();
}