SetPadValue

函数功能

设置DataCopyPad接口填充的数值。支持的通路如下：

GM->VECIN/GM->VECOUT

函数原型

template <typename T, TPosition pos = TPosition::MAX>
__aicore__ inline void SetPadValue(T paddingValue)

参数说明

表1 模板参数说明
参数名	输入/输出	描述
T	输入	填充值的数据类型，与DataCopyPad接口搬运的数据类型一致。
pos	输入	用于指定DataCopyPad接口搬运过程中从GM搬运数据到哪一个目的地址，默认值为TPosition::MAX，等效于目的地址为VECIN/VECOUT。支持的取值为： TPosition::VECIN/TPosition::VECOUT/TPosition::MAX

表2 参数说明
参数名	输入/输出	描述
paddingValue	输入	DataCopyPad接口填充的数值，数据与DataCopyPad接口搬运的数据类型一致。

返回值

无

支持的型号

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

约束说明

无

调用示例

#include "kernel_operator.h"

template <typename T>
class SetPadValueTest {
public:
    __aicore__ inline SetPadValueTest() {}
    __aicore__ inline void Init(__gm__ uint8_t* dstGm, __gm__ uint8_t* srcGm, uint32_t n1, uint32_t n2)
    {
        m_n1 = n1;
        m_n2 = n2;
        m_n2Align = n2 % 32 == 0 ? n2 : (n2 / 32 + 1) * 32;
        m_srcGlobal.SetGlobalBuffer((__gm__ T*)srcGm);
        m_dstGlobal.SetGlobalBuffer((__gm__ T*)dstGm);

        m_pipe.InitBuffer(m_queInSrc, 1, m_n1 * m_n2Align * sizeof(T));
    }
    __aicore__ inline void Process()
    {
        CopyIn();
        Compute();
        CopyOut();
    }
private:
    __aicore__ inline void CopyIn()
    {
        AscendC::LocalTensor<T> srcLocal = m_queInSrc.AllocTensor<T>();
        AscendC::DataCopyExtParams dataCopyExtParams;
        AscendC::DataCopyPadExtParams<T> padParams;

        dataCopyExtParams.blockCount = m_n1;
        dataCopyExtParams.blockLen = m_n2 * sizeof(T);
        dataCopyExtParams.srcStride = 0;
        dataCopyExtParams.dstStride = 0;

        padParams.isPad = false;
        padParams.leftPadding = 0;
        padParams.rightPadding = 1;

        AscendC::SetPadValue((T)37);
        AscendC::PipeBarrier<PIPE_ALL>();
        AscendC::DataCopyPad(srcLocal, m_srcGlobal, dataCopyExtParams, padParams);
        m_queInSrc.EnQue(srcLocal);
    }
    __aicore__ inline void Compute()
    {
        ;
    }
    __aicore__ inline void CopyOut()
    {
        AscendC::LocalTensor<T> dstLocal = m_queInSrc.DeQue<T>();
        AscendC::DataCopy(m_dstGlobal, dstLocal, m_n1 * m_n2Align);
        m_queInSrc.FreeTensor(dstLocal);
    }
private:
    AscendC::TPipe m_pipe;
    uint32_t m_n1;
    uint32_t m_n2;
    uint32_t m_n2Align;
    AscendC::GlobalTensor<T> m_srcGlobal;
    AscendC::GlobalTensor<T> m_dstGlobal;
    AscendC::TQue<AscendC::QuePosition::VECIN, 1> m_queInSrc;
}; // class SetPadValueTest

template <typename T>
__global__ __aicore__ void testSetPadValue(GM_ADDR dstGm, GM_ADDR srcGm, uint32_t n1, uint32_t n2)
{
    SetPadValueTest<T> op;
    op.Init(dstGm, srcGm, n1, n2);
    op.Process();
}

输入数据（srcGm, shape = [32, 31]）：[[1, 1, 1, ..., 1], [1, 1, 1, ..., 1], ... , [1, 1, 1, ..., 1]]
输出数据（dstGm, shape = [32, 32]）：[[1, 1, 1, ..., 1, 37], [1, 1, 1, ..., 1, 37], ... , [1, 1, 1, ..., 1, 37]]

父主题： 数据搬运