SetPadValue(ISASI)

Function Usage

Sets the value filled by DataCopyPad. The following paths are supported:

GM->VECIN/GM->VECOUT

Prototype

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

Parameters

**Table 1** Parameters in the template
Parameter	Input/Output	Description
T	Input	Data type of the padding value, which is the same as the type of the data transferred by DataCopyPad.
pos	Input	Used to specify the destination address to which DataCopyPad transfers data from GM. The default value is TPosition::MAX, which is equivalent to the destination address VECIN/VECOUT. The following values are supported: TPosition::VECIN/TPosition::VECOUT/TPosition::MAX

**Table 2** Parameters
Parameter	Input/Output	Description
paddingValue	Input	Value filled by DataCopyPad. Its data type is the same as the type of the data transferred by DataCopyPad.

Returns

None

Availability

Atlas A2 training products/Atlas A2 inference products

Atlas A3 training products/Atlas A3 inference products

Constraints

None

Example

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

Input (srcGm, shape = [32, 31]): [[1, 1, 1, ..., 1], [1, 1, 1, ..., 1], ..., [1, 1, 1, ..., 1]]
Output (dstGm, shape = [32, 32]): [[1, 1, 1, ..., 1, 37], [1, 1, 1, ..., 1, 37], ..., [1, 1, 1, ..., 1, 37]]

Parent topic: Data Movement