news 2026/5/11 7:39:42

CANN/asc-devkit异或运算API文档

作者头像

张小明

前端开发工程师

1.2k 24
文章封面图
CANN/asc-devkit异或运算API文档

Xor

【免费下载链接】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 推理系列产品

√

功能说明

按元素执行Xor运算，Xor（异或）的概念和运算规则如下：

  • 概念：参加运算的两个数据，按二进制位进行“异或”运算。
  • 运算规则：0^0=0；0^1=1；1^0=1；1^1=0；即：参加运算的两个对象，如果两个相应位为“异”（值不同），则该位结果为1，否则为 0【同0异1】。

计算公式如下：

例如：3^5=6，即0000 0011^0000 0101 = 0000 0110

函数原型

  • 通过sharedTmpBuffer入参传入临时空间

    • 源操作数Tensor全部/部分参与计算

      template <typename T, bool isReuseSource = false> __aicore__ inline void Xor(const LocalTensor<T>& dstTensor, const LocalTensor<T>& src0Tensor, const LocalTensor<T>& src1Tensor, const LocalTensor<uint8_t>& sharedTmpBuffer, const uint32_t calCount)

    • 源操作数Tensor全部参与计算

      template <typename T, bool isReuseSource = false> __aicore__ inline void Xor(const LocalTensor<T>& dstTensor, const LocalTensor<T>& src0Tensor, const LocalTensor<T>& src1Tensor, const LocalTensor<uint8_t>& sharedTmpBuffer)

  • 接口框架申请临时空间

    • 源操作数Tensor全部/部分参与计算

      template <typename T, bool isReuseSource = false> __aicore__ inline void Xor(const LocalTensor<T>& dstTensor, const LocalTensor<T>& src0Tensor, const LocalTensor<T>& src1Tensor, const uint32_t calCount)

    • 源操作数Tensor全部参与计算

      template <typename T, bool isReuseSource = false> __aicore__ inline void Xor(const LocalTensor<T>& dstTensor, const LocalTensor<T>& src0Tensor, const LocalTensor<T>& src1Tensor)

由于该接口的内部实现中涉及复杂的数学计算，需要额外的临时空间来存储计算过程中的中间变量。临时空间支持开发者通过sharedTmpBuffer入参传入和接口框架申请两种方式。

  • 通过sharedTmpBuffer入参传入，使用该tensor作为临时空间进行处理，接口框架不再申请。该方式开发者可以自行管理sharedTmpBuffer内存空间，并在接口调用完成后，复用该部分内存，内存不会反复申请释放，灵活性较高，内存利用率也较高。
  • 接口框架申请临时空间，开发者无需申请，但是需要预留临时空间的大小。

通过sharedTmpBuffer传入的情况，开发者需要为tensor申请空间；接口框架申请的方式，开发者需要预留临时空间。临时空间大小BufferSize的获取方式如下：通过GetXorMaxMinTmpSize中提供的接口获取需要预留空间范围的大小。

参数说明

表 1模板参数说明

参数名

描述

T

操作数的数据类型。

Ascend 950PR/Ascend 950DT，支持的数据类型为：int16_t、uint16_t。

Atlas A3 训练系列产品 / Atlas A3 推理系列产品，支持的数据类型为：int16_t、uint16_t。

Atlas A2 训练系列产品 / Atlas A2 推理系列产品，支持的数据类型为：int16_t、uint16_t。

isReuseSource

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

表 2接口参数说明

参数名

输入/输出

描述

dstTensor

输出

目的操作数。

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

src0Tensor

输入

源操作数0。

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

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

src1Tensor

输入

源操作数1。

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

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

sharedTmpBuffer

输入

临时缓存。

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

用于Xor内部复杂计算时存储中间变量，由开发者提供。

临时空间大小BufferSize的获取方式请参考GetXorMaxMinTmpSize。

calCount

输入

参与计算的元素个数。

返回值说明

无

约束说明

  • 不支持源操作数与目的操作数地址重叠。
  • 当前仅支持ND格式的输入，不支持其他格式。
  • calCount需要保证小于或等于src0Tensor和src1Tensor和dstTensor存储的元素范围。
  • 对于不带calCount参数的接口，需要保证src0Tensor和src1Tensor的shape大小相等。
  • 不支持sharedTmpBuffer与源操作数和目的操作数地址重叠。
  • 操作数地址对齐要求请参见通用地址对齐约束。

调用示例

调用样例kernel侧xor_custom.cpp

#include "kernel_operator.h" constexpr int32_t BUFFER_NUM = 1; class KernelXor { public: __aicore__ inline KernelXor() {} __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength, uint32_t totalLength2, uint32_t tilenum, uint32_t tmpSize, uint32_t mcount) { this->totalLength = totalLength; this->blockLength = totalLength / AscendC::GetBlockNum(); this->blockLength2 = totalLength2 / AscendC::GetBlockNum(); this->tilenum = tilenum; this->tmpSize = tmpSize; this->mcount = mcount; this->tileLength = this->blockLength / tilenum / BUFFER_NUM; this->tileLength2 = this->blockLength2 / tilenum / BUFFER_NUM; xGm.SetGlobalBuffer((__gm__ int16_t *)x + this->blockLength * AscendC::GetBlockIdx(), this->blockLength); yGm.SetGlobalBuffer((__gm__ int16_t *)y + this->blockLength2 * AscendC::GetBlockIdx(), this->blockLength2); zGm.SetGlobalBuffer((__gm__ int16_t *)z + this->blockLength * AscendC::GetBlockIdx(), this->blockLength); if (this->tmpSize != 0) { pipe.InitBuffer(tmpQueue, BUFFER_NUM, this->tmpSize); } pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(int16_t)); pipe.InitBuffer(inQueueY, BUFFER_NUM, this->tileLength2 * sizeof(int16_t)); pipe.InitBuffer(outQueueZ, BUFFER_NUM, this->tileLength * sizeof(int16_t)); } __aicore__ inline void Process() { int32_t loopCount = this->tilenum * BUFFER_NUM; for (int32_t i = 0; i < loopCount; i++) { CopyIn(i); Compute(i); CopyOut(i); } } private: __aicore__ inline void CopyIn(int32_t progress) { AscendC::LocalTensor<int16_t> xLocal = inQueueX.AllocTensor<int16_t>(); AscendC::DataCopy(xLocal, xGm[progress * this->tileLength], this->tileLength); inQueueX.EnQue(xLocal); AscendC::LocalTensor<int16_t> yLocal = inQueueY.AllocTensor<int16_t>(); AscendC::DataCopy(yLocal, yGm[progress * this->tileLength2], this->tileLength2); inQueueY.EnQue(yLocal); } __aicore__ inline void Compute(int32_t progress) { AscendC::LocalTensor<int16_t> xLocal = inQueueX.DeQue<int16_t>(); AscendC::LocalTensor<int16_t> yLocal = inQueueY.DeQue<int16_t>(); AscendC::LocalTensor<int16_t> zLocal = outQueueZ.AllocTensor<int16_t>(); if (this->tmpSize != 0) { AscendC::LocalTensor<uint8_t> tmpLocal = tmpQueue.AllocTensor<uint8_t>(); if (this->mcount != this->totalLength) { AscendC::Xor(zLocal, xLocal, yLocal, tmpLocal, this->mcount); } else { AscendC::Xor(zLocal, xLocal, yLocal, tmpLocal); } tmpQueue.FreeTensor(tmpLocal); } else { if (this->mcount != this->totalLength) { AscendC::Xor(zLocal, xLocal, yLocal, this->mcount); } else { AscendC::Xor(zLocal, xLocal, yLocal); } } outQueueZ.EnQue<int16_t>(zLocal); inQueueX.FreeTensor(xLocal); inQueueY.FreeTensor(yLocal); } __aicore__ inline void CopyOut(int32_t progress) { AscendC::LocalTensor<int16_t> zLocal = outQueueZ.DeQue<int16_t>(); AscendC::DataCopy(zGm[progress * this->tileLength], zLocal, this->tileLength); outQueueZ.FreeTensor(zLocal); } private: AscendC::TPipe pipe; AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueX; AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueY; AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> tmpQueue; AscendC::TQue<AscendC::TPosition::VECOUT, BUFFER_NUM> outQueueZ; AscendC::GlobalTensor<int16_t> xGm; AscendC::GlobalTensor<int16_t> yGm; AscendC::GlobalTensor<int16_t> zGm; uint32_t blockLength; uint32_t blockLength2; uint32_t tilenum; uint32_t tileLength; uint32_t tileLength2; uint32_t tmpSize; uint32_t mcount; uint32_t totalLength; }; extern "C" __global__ __aicore__ void xor_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling) { GET_TILING_DATA(tilingData, tiling); KernelXor op; op.Init(x, y, z, tilingData.totalLength, tilingData.totalLength2, tilingData.tilenum, tilingData.tmpSize, tilingData.mcount); if (TILING_KEY_IS(1)) { op.Process(); } }

host侧xor_custom_tiling.h

#include "register/op_def_registry.h" #include "register/tilingdata_base.h" namespace optiling { BEGIN_TILING_DATA_DEF(XorCustomTilingData) TILING_DATA_FIELD_DEF(uint32_t, totalLength); TILING_DATA_FIELD_DEF(uint32_t, totalLength2); TILING_DATA_FIELD_DEF(uint32_t, tmpSize); TILING_DATA_FIELD_DEF(uint32_t, tilenum); TILING_DATA_FIELD_DEF(uint32_t, mcount); END_TILING_DATA_DEF; REGISTER_TILING_DATA_CLASS(XorCustom, XorCustomTilingData) }

host侧xor_custom.cpp

#include "xor_custom_tiling.h" #include "register/op_def_registry.h" #include "tiling/tiling_api.h" namespace optiling { static ge::graphStatus TilingFunc(gert::TilingContext *context) { XorCustomTilingData tiling; const gert::RuntimeAttrs *xorAttrs = context->GetAttrs(); const uint32_t tilenum = *(xorAttrs->GetAttrPointer<uint32_t>(0)); const uint32_t numBlocks = *(xorAttrs->GetAttrPointer<uint32_t>(1)); const uint32_t sizeflag = *(xorAttrs->GetAttrPointer<uint32_t>(2)); const uint32_t countflag = *(xorAttrs->GetAttrPointer<uint32_t>(3)); uint32_t totalLength = context->GetInputTensor(0)->GetShapeSize(); uint32_t totalLength2 = context->GetInputTensor(1)->GetShapeSize(); context->SetBlockDim(numBlocks); tiling.set_totalLength(totalLength); tiling.set_totalLength2(totalLength2); tiling.set_tilenum(tilenum); if (countflag == 0) { tiling.set_mcount(totalLength2); } else if (countflag == 1) { tiling.set_mcount(totalLength); } std::vector<int64_t> shapeVec = {totalLength}; ge::Shape srcShape(shapeVec); uint32_t typeSize = sizeof(int16_t); uint32_t maxValue = 0; uint32_t minValue = 0; bool isReuseSource = false; AscendC::GetXorMaxMinTmpSize(srcShape, typeSize, isReuseSource, maxValue, minValue); // sizeflag 0：代表取最小的tempBuffer 1：取最大的tempBuffer if (sizeflag == 0) { tiling.set_tmpSize(minValue); } else if (sizeflag == 1) { tiling.set_tmpSize(maxValue); } else if (sizeflag == 2) { tiling.set_tmpSize(0); } tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity()); context->GetRawTilingData()->SetDataSize(tiling.GetDataSize()); context->SetTilingKey(1); size_t *currentWorkspace = context->GetWorkspaceSizes(1); currentWorkspace[0] = 0; return ge::GRAPH_SUCCESS; } } namespace ge { static ge::graphStatus InferShape(gert::InferShapeContext *context) { const gert::Shape *xShape = context->GetInputShape(0); gert::Shape *yShape = context->GetOutputShape(0); *yShape = *xShape; return GRAPH_SUCCESS; } } namespace ops { class XorCustom : public OpDef { public: explicit XorCustom(const char *name) : OpDef(name) { this->Input("x") .ParamType(REQUIRED) .DataType({ge::DT_INT16}) .Format({ge::FORMAT_ND}); this->Input("y") .ParamType(REQUIRED) .DataType({ge::DT_INT16}) .Format({ge::FORMAT_ND}); this->Output("z") .ParamType(REQUIRED) .DataType({ge::DT_INT16}) .Format({ge::FORMAT_ND}); this->SetInferShape(ge::InferShape); this->Attr("tilenum") .AttrType(REQUIRED) .Int(0); this->Attr("numBlocks") .AttrType(REQUIRED) .Int(0); this->Attr("sizeflag") .AttrType(REQUIRED) .Int(0); this->Attr("countflag") .AttrType(REQUIRED) .Int(0); this->AICore() .SetTiling(optiling::TilingFunc); this->AICore().AddConfig("ascendxxx"); // ascendxxx请修改为对应的AI处理器型号。 } }; OP_ADD(XorCustom); } // namespace ops

结果示例如下：

输入输出的数据类型为int16_t，一维向量包含32个数。例如向量中第一个数据进行异或：(-5753) xor 18745 = -24386 输入数据(src0Local): [-5753 28501 20334 -5845 ... -20817 3403 21261 22241] 输入数据(src1Local): [18745 -24448 20873 10759 ... 21940 -26342 9251 31019] 输出数据(dstLocal): [-24386 -12331 7911 -15572 ... -1253 -27567 30510 12234]

【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言，原生支持C和C++标准规范，主要由类库和语言扩展层构成，提供多层级API，满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit

创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考

版权声明: 本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若内容造成侵权/违法违规/事实不符，请联系邮箱：809451989@qq.com进行投诉反馈，一经查实，立即删除！
网站建设 2026/5/11 7:35:33

CANN ops-nn MaxPool3D算子

MaxPool3D 【免费下载链接】ops-nn 本项目是CANN提供的神经网络类计算算子库&#xff0c;实现网络在NPU上加速计算。 项目地址: https://gitcode.com/cann/ops-nn 产品支持情况 产品是否支持Ascend 950PR/Ascend 950DT√Atlas A3 训练系列产品/Atlas A3 推理系列产品√…

作者头像 李华
网站建设 2026/5/11 7:25:53

【信息科学与工程学】【人工智能】【数字孪生】【游戏科学】主要数学模型-第九篇 计算神经科学

认知神经科学的几何、拓扑与计算建模框架 这是一个深度交叉领域的问题，我将从几何表示、拓扑结构、动力学模型和仿真算法四个维度，系统梳理从神经元到全脑的计算神经科学建模方法。 一、神经元与连接的几何表示模型 神经元形态的表示： a) 线表示：将神经元的树突和轴突表示…

作者头像 李华
网站建设 2026/5/11 7:24:42

【信息科学工程学】【社会科学】 第五十五篇 人的利益规则05——行业篇02

编号 行业门类 细分子行业和细分公司类型 行业的上游 行业的下游 产品的上游 产品的下游 利益形态（与各主体） 各类利益形态和利益模型（含周期性） 利益流动模式 人的利益规则 公司内的利益规则 公司间的利益规则 各类利益链和利益流动链条的数学建模框架 行业…

作者头像 李华
网站建设 2026/5/11 7:17:33

DISTINCT 带 WHERE 仍全表扫描？两层优化刀法拆解

DISTINCT 带 WHERE 仍全表扫描&#xff1f;两层优化刀法拆解 引言&#xff1a;一个看似多余的 DISTINCT&#xff0c;藏着性能陷阱 几乎每个写过 SQL 的人都用过 DISTINCT。它的语义很简单——去掉重复行。但"简单"不等于"快"。在一个客户的生产环境中&…

作者头像 李华
网站建设 2026/5/11 7:12:33

uniCloud服务空间从阿里云搬家迁移至支付宝云

阿里云现在新增了小时级低消&#xff0c;大家的开销一下子涨了不少。今天就给大家讲讲&#xff0c;怎么从阿里云迁移到支付宝云。 就拿咱们之前做的商城项目来举例&#xff0c;迁移之前&#xff0c;大家一定要先检查好&#xff1a;项目能不能正常跑起来&#xff0c;还有所有数据…

作者头像 李华