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^000^111^011^10即参加运算的两个对象如果两个相应位为“异”值不同则该位结果为1否则为 0【同0异1】。计算公式如下例如3^56即0000 0011^0000 0101 0000 0110函数原型通过sharedTmpBuffer入参传入临时空间源操作数Tensor全部/部分参与计算template typename T, bool isReuseSource false __aicore__ inline void Xor(const LocalTensorT dstTensor, const LocalTensorT src0Tensor, const LocalTensorT src1Tensor, const LocalTensoruint8_t sharedTmpBuffer, const uint32_t calCount)源操作数Tensor全部参与计算template typename T, bool isReuseSource false __aicore__ inline void Xor(const LocalTensorT dstTensor, const LocalTensorT src0Tensor, const LocalTensorT src1Tensor, const LocalTensoruint8_t sharedTmpBuffer)接口框架申请临时空间源操作数Tensor全部/部分参与计算template typename T, bool isReuseSource false __aicore__ inline void Xor(const LocalTensorT dstTensor, const LocalTensorT src0Tensor, const LocalTensorT src1Tensor, const uint32_t calCount)源操作数Tensor全部参与计算template typename T, bool isReuseSource false __aicore__ inline void Xor(const LocalTensorT dstTensor, const LocalTensorT src0Tensor, const LocalTensorT 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::LocalTensorint16_t xLocal inQueueX.AllocTensorint16_t(); AscendC::DataCopy(xLocal, xGm[progress * this-tileLength], this-tileLength); inQueueX.EnQue(xLocal); AscendC::LocalTensorint16_t yLocal inQueueY.AllocTensorint16_t(); AscendC::DataCopy(yLocal, yGm[progress * this-tileLength2], this-tileLength2); inQueueY.EnQue(yLocal); } __aicore__ inline void Compute(int32_t progress) { AscendC::LocalTensorint16_t xLocal inQueueX.DeQueint16_t(); AscendC::LocalTensorint16_t yLocal inQueueY.DeQueint16_t(); AscendC::LocalTensorint16_t zLocal outQueueZ.AllocTensorint16_t(); if (this-tmpSize ! 0) { AscendC::LocalTensoruint8_t tmpLocal tmpQueue.AllocTensoruint8_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.EnQueint16_t(zLocal); inQueueX.FreeTensor(xLocal); inQueueY.FreeTensor(yLocal); } __aicore__ inline void CopyOut(int32_t progress) { AscendC::LocalTensorint16_t zLocal outQueueZ.DeQueint16_t(); AscendC::DataCopy(zGm[progress * this-tileLength], zLocal, this-tileLength); outQueueZ.FreeTensor(zLocal); } private: AscendC::TPipe pipe; AscendC::TQueAscendC::TPosition::VECIN, BUFFER_NUM inQueueX; AscendC::TQueAscendC::TPosition::VECIN, BUFFER_NUM inQueueY; AscendC::TQueAscendC::TPosition::VECIN, BUFFER_NUM tmpQueue; AscendC::TQueAscendC::TPosition::VECOUT, BUFFER_NUM outQueueZ; AscendC::GlobalTensorint16_t xGm; AscendC::GlobalTensorint16_t yGm; AscendC::GlobalTensorint16_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-GetAttrPointeruint32_t(0)); const uint32_t numBlocks *(xorAttrs-GetAttrPointeruint32_t(1)); const uint32_t sizeflag *(xorAttrs-GetAttrPointeruint32_t(2)); const uint32_t countflag *(xorAttrs-GetAttrPointeruint32_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::vectorint64_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），仅供参考