目录

CreateVecIndex

Function Description

Creates the vector index with firstValue as the start value.

Function Prototype

Computation of the first n data elements of a tensor:

template <typename T> 
__aicore__ inline void CreateVecIndex(LocalTensor<T> dstLocal, const T &firstValue, uint32_t calCount)

Parameters

Table 1 Parameters

Parameter Input/Output Description
dstLocal Output

Destination operand.

The type is LocalTensor, and the supported TPosition is VECIN, VECCALC, or VECOUT.

The start address of the LocalTensor must be 32-byte aligned.

For Kirin 9020, Kirin 9030, and Kirin X90 series processors, the supported data types include int8, int16, int32, float16, and float.
firstValue Input The first value of the index. The data type must be the same as that of the element in dstLocal.
calCount Input Number of elements of the input data.

Supported Models

Kirin 9020 series processors

Kirin 9030 series processors

Kirin X90 series processors

Precautions

Return Value

None

Calling Example

This example shows only part of the code used in the computation process (Compute). To run the sample code, use it to replace the code snippet related to the Compute function in Sample Template.

Example of computing the first n data elements of a tensor:

AscendC::CreateVecIndex(dstLocal, (T)0, 128);

An example of the result is as follows:

Input (firstValue): 0
Output data (dstLocal): [0 1 2... 127]

Sample Template

#include "kernel_operator.h"
template <typename T> 
class CreateVecIndexTest {
public:
    __aicore__ inline CreateVecIndexTest() {}
    __aicore__ inline void Init(GM_ADDR dstGm, uint64_t mask, uint8_t repeatTimes,
        uint16_t dstBlkStride, uint8_t dstRepStride)
    {
        m_mask = mask;
        m_repeatTimes = repeatTimes;
        m_dstBlkStride = dstBlkStride;
        m_dstRepStride = dstRepStride;
        m_elementCount = m_dstBlkStride * m_dstRepStride * 32 * m_repeatTimes / sizeof(T);
        m_dstGlobal.SetGlobalBuffer((__gm__ T*)dstGm);
        m_pipe.InitBuffer(m_queOut, 1, m_dstBlkStride * m_dstRepStride * 32 * m_repeatTimes);
        m_pipe.InitBuffer(m_queTmp, 1, 1024);
    }
    __aicore__ inline void Process()
    {
        // CopyIn is not required.
        Compute();
        CopyOut();
    }
private:

    __aicore__ inline void Compute()
    {
        AscendC::LocalTensor<T> dstLocal = m_queOut.AllocTensor<T>();
        AscendC::LocalTensor<uint8_t> tmpLocal = m_queTmp.AllocTensor<uint8_t>();
        AscendC::Duplicate(dstLocal, (T)0, m_elementCount);
        AscendC::PipeBarrier<PIPE_ALL>();
        AscendC::CreateVecIndex(dstLocal, (T)0, m_repeatTimes * 256 / sizeof(T));
        m_queOut.EnQue(dstLocal);
        m_queTmp.FreeTensor(tmpLocal);
    }
    __aicore__ inline void CopyOut()
    {
        AscendC::LocalTensor<T> dstLocal = m_queOut.DeQue<T>();
 
        AscendC::DataCopy(m_dstGlobal, dstLocal, m_elementCount);
        m_queOut.FreeTensor(dstLocal);
    }
private:
    AscendC::TPipe m_pipe;
    uint32_t m_elementCount;
    uint32_t m_mask;
    uint32_t m_repeatTimes;
    uint32_t m_dstBlkStride;
    uint32_t m_dstRepStride;
    AscendC::GlobalTensor<T> m_dstGlobal;
    AscendC::TQue<AscendC::QuePosition::VECOUT, 1> m_queOut;
    AscendC::TQue<AscendC::QuePosition::VECIN, 1> m_queTmp;
}; // class CreateVecIndexTest
template <typename T> 
__global__ __aicore__ void testCreateVecIndex(GM_ADDR dstGm, uint64_t mask, uint8_t repeatTimes,
        uint16_t dstBlkStride, uint8_t dstRepStride)
{
    CreateVecIndexTest<T> op;
    op.Init(dstGm, mask, repeatTimes, dstBlkStride, dstRepStride);
    op.Process();
}