Duplicate

Applicability

Product	Supported/Unsupported
Atlas A3 training products / Atlas A3 inference products	√
Atlas A2 training products / Atlas A2 inference products	√
Atlas 200I/500 A2 inference products	√
Atlas inference product 's AI Core	√
Atlas inference product 's Vector Core	x
Atlas training products	√

Function Usage

Copies a variable or immediate value for multiple times and fills the vector.

Prototype

Computation of the first n data elements of a tensor

Scalar as the source operand

          
               template <typename T>
__aicore__ inline void Duplicate(const LocalTensor<T>& dst, const T& scalarValue, const int32_t& count)

High-dimensional tensor sharding computation

Bitwise mask mode

          
               template <typename T, bool isSetMask = true>
__aicore__ inline void Duplicate(const LocalTensor<T>& dst, const T& scalarValue, uint64_t mask[], const uint8_t repeatTime, const uint16_t dstBlockStride, const uint8_t dstRepeatStride)

Contiguous mask mode

          
               template <typename T, bool isSetMask = true>
__aicore__ inline void Duplicate(const LocalTensor<T>& dst, const T& scalarValue, uint64_t mask, const uint8_t repeatTime, const uint16_t dstBlockStride, const uint8_t dstRepeatStride)

Parameters

**Table 1** Parameters in the template
Parameter	Description
T	Operand data type. Atlas A3 training products / Atlas A3 inference products : The supported data types are int16_t, uint16_t, half, bfloat16_t, int32_t, uint32_t, and float. Atlas A2 training products / Atlas A2 inference products : The supported data types are int16_t, uint16_t, half, bfloat16_t, int32_t, uint32_t, and float. Atlas 200I/500 A2 inference products : The supported data types are int16_t, uint16_t, half, int32_t, uint32_t, and float. Atlas inference product 's AI Core: The supported data types are int16_t, uint16_t, half, int32_t, uint32_t, and float. Atlas training products : The supported data types are int16_t, uint16_t, half, int32_t, uint32_t, and float.
isSetMask	Indicates whether to set mask inside the API. true: sets mask inside the API. false: sets mask outside the API. Developers need to use the SetVectorMask API to set the mask value. In this mode, the mask value in the input parameter of this API must be set to the placeholder MASK_PLACEHOLDER.

**Table 2** Parameters
Parameter	Input/Output	Meaning
dst	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.
scalarValue	Input	Source operand to be copied. The data type must be the same as that of the elements in dst.
count	Input	Number of elements involved in the computation.
mask/mask[]	Input	The mask parameter is used to control the elements involved in computation in each iteration. Bitwise mode: controls the elements that participate in computation by bit. If a bit is set to 1, the corresponding element participates in the computation. If a bit is set to 0, the corresponding element is masked in the computation. The mask is in array form. The array length and the value range of the array elements are related to the data type of the operand. When the operand is 16-bit, the array length is 2. In this case, mask[0] and mask[1] must be in the range of [0, 2⁶⁴ – 1] and cannot be 0 at the same time. When the operand is 32-bit, the array length is 1. In this case, mask[0] must be in the range of (0, 2⁶⁴ – 1]. When the operand is 64-bit, the array length is 1. In this case, mask[0] must be in the range of (0, 2³² – 1]. For example, if mask = [0, 8] and 8 = 0b1000, only the fourth element participates in computation. Contiguous mode: indicates the number of contiguous elements that participate in computation. The value range is related to the operand data type. The maximum number of elements that can be processed in each repeat varies according to the data type. When the operand is 16-bit, mask ∈ [1, 128]. When the operand is 32-bit, mask ∈ [1, 64]. When the operand is 64-bit, mask ∈ [1, 32].
repeatTime	Input	The Vector Unit reads 8 data blocks (32 bytes each and 256 bytes in total) of contiguous data each time, and has to go through several repeats before all data can be read and computed. repeatTime indicates the number of repeats.
dstBlockStride	Input	Address stride of the vector destination operand between different data blocks in a single repeat
dstRepeatStride	Input	Address stride of the vector destination operand for the same data block between adjacent repeats

Restrictions

For details about the operand address alignment requirements, see General Address Alignment Restrictions.

Returns

None

Examples

This example shows only part of the code involved in the computation process (Compute). To run the code, copy the code block and paste it to the corresponding position in the Compute function in Template Samples.

Example of high-dimensional tensor sharding computation (contiguous mask mode)

        
             uint64_t mask = 128;
half scalar = 18.0;
// repeatTime = 2, 128 elements one repeat, 256 elements total
// dstBlkStride = 1, no gap between blocks in one repeat
// dstRepStride = 8, no gap between repeats
AscendC::Duplicate(dstLocal, scalar, mask, 2, 1, 8 );

Example of high-dimensional tensor sharding computation (bitwise mask mode)

        
             uint64_t mask[2] = { UINT64_MAX, UINT64_MAX };
half scalar = 18.0;
// repeatTime = 2, 128 elements one repeat, 256 elements total
// dstBlkStride = 1, no gap between blocks in one repeat
// dstRepStride = 8, no gap between repeats
AscendC::Duplicate(dstLocal, scalar, mask, 2, 1, 8 );

Example of computing the first n pieces of data in a tensor, where the source operand is a scalar

        
             half inputVal(18.0);
AscendC::Duplicate<half>(dstLocal, inputVal, srcDataSize);

Result example:

scalar: 18.0
dstLocal: [18.0 18.0 18.0 ... 18.0 18.0]

More Samples

You can refer to the following examples to learn how to use the high-dimensional tensor sharding computation APIs of the Duplicate instruction to perform more flexible operations and implement more advanced functions. This example shows only part of the code in the computation process. To run the code, copy and paste the code snippet to the corresponding position of the Compute function in Template Samples.

Use the contiguous mask mode of a high-dimensional tensor sharding computation API to implement discontinuous data calculation.

        
             uint64_t mask = 64;  // Only the first 64 bits are calculated in each repeat.
half scalar = 18.0;
// repeatTime = 2, 128 elements one repeat, 256 elements total
// dstBlkStride = 1, dstRepStride = 8 
AscendC::Duplicate(dstLocal, scalar, mask, 2, 1, 8 );

Result example:

[18.0 18.0 18.0 ... 18.0  undefined ... undefined 
 18.0 18.0 18.0 ... 18.0 undefined ... undefined ](The length of each segment of the computed result or undefined data is 64.)

Use the bitwise mask mode of a high-dimensional tensor sharding computation API to implement discontinuous data calculation.

        
             uint64_t mask[2] = { UINT64_MAX, 0 }; // mask[0] is set to max, mask[1] is set to empty, and only the first 64 bits are calculated each time.
half scalar = 18.0;
// repeatTime = 2, 128 elements one repeat, 512 elements total
// dstBlkStride = 1, no gap between blocks in one repeat
// dstRepStride = 8, no gap between repeats
AscendC::Duplicate(dstLocal, scalar, mask, 2, 1, 8);

Result example:

Input (src0Local): [1.0 2.0 3.0... 256.0]
Input (src1Local): half scalar = 18.0;
Output (dstLocal):
[18.0 18.0 18.0 ... 18.0 undefined ... undefined
 18.0 18.0 18.0 ... 18.0 undefined ... undefined](The length of each segment of the computed result or undefined data is 64.)

Set the repeatStride parameter of a high-dimensional tensor sharding computation API to implement discontinuous data computation.

        
             uint64_t mask = 128;
half scalar = 18.0;
// repeatTime = 1, 128 elements one repeat, 256 elements total
// dstBlkStride = 2, 1 block gap between blocks in one repeat
// dstRepStride = 0, repeatTime = 1
AscendC::Duplicate(dstLocal, scalar, mask, 1, 2, 0);

Result example:

Input (src0Local): [1.0 2.0 3.0... 256.0]
Input (src1Local): half scalar = 18.0;
Output (dstLocal):
[18.0 18.0 18.0 ... 18.0 undefined ... undefined
 18.0 18.0 18.0 ... 18.0 undefined ... undefined
 18.0 18.0 18.0 ... 18.0 undefined ... undefined
 18.0 18.0 18.0 ... 18.0 undefined ... undefined
 18.0 18.0 18.0 ... 18.0 undefined ... undefined
 18.0 18.0 18.0 ... 18.0 undefined ... undefined
 18.0 18.0 18.0 ... 18.0 undefined ... undefined
 18.0 18.0 18.0 ... 18.0 undefined ... undefined](The length of each segment of the computed result is 16.)

Set the repeatStride parameter of a high-dimensional tensor sharding computation API to implement discontinuous data computation.

        
             uint64_t mask = 64;
half scalar = 18.0;
// repeatTime = 2, 128 elements one repeat, 256 elements total
// dstBlkStride = 1, no gap between blocks in one repeat
// dstRepStride = 12, 4 blocks gap between repeats
AscendC::Duplicate(dstLocal, scalar, mask, 2, 1, 12);

Result example:

Input (src0Local): [1.0 2.0 3.0... 256.0]
Input (src1Local): half scalar = 18.0;
Output (dstLocal):
[18.0 18.0 18.0 ... 18.0 undefined ... undefined 18.0 18.0 18.0 ... 18.0](The length of each segment of the computed result is 64, and that of undefined data is 128.)

Template Samples

      
       
         
         
           #include "kernel_operator.h"
class KernelDuplicate {
public:
    __aicore__ inline KernelDuplicate() {}
    __aicore__ inline void Init(__gm__ uint8_t* src, __gm__ uint8_t* dstGm)
    {
        srcGlobal.SetGlobalBuffer((__gm__ half*)src);
        dstGlobal.SetGlobalBuffer((__gm__ half*)dstGm);
        pipe.InitBuffer(inQueueSrc, 1, srcDataSize * sizeof(half));
        pipe.InitBuffer(outQueueDst, 1, dstDataSize * sizeof(half));
    }
    __aicore__ inline void Process()
    {
        CopyIn();
        Compute();
        CopyOut();
    }
private:
    __aicore__ inline void CopyIn()
    {
        AscendC::LocalTensor<half> srcLocal = inQueueSrc.AllocTensor<half>();
        AscendC::DataCopy(srcLocal, srcGlobal, srcDataSize);
        inQueueSrc.EnQue(srcLocal);
    }
    __aicore__ inline void Compute()
    {
        AscendC::LocalTensor<half> srcLocal = inQueueSrc.DeQue<half>();
        AscendC::LocalTensor<half> dstLocal = outQueueDst.AllocTensor<half>();
        half inputVal(18.0);
        AscendC::Duplicate<half>(dstLocal, inputVal, srcDataSize);
        outQueueDst.EnQue<half>(dstLocal);
        inQueueSrc.FreeTensor(srcLocal);
    }
    __aicore__ inline void CopyOut()
    {
        AscendC::LocalTensor<half> dstLocal = outQueueDst.DeQue<half>();
        AscendC::DataCopy(dstGlobal, dstLocal, dstDataSize);
        outQueueDst.FreeTensor(dstLocal);
    }
private:
    AscendC::TPipe pipe;
    AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueueSrc;
    AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueueDst;
    AscendC::GlobalTensor<half> srcGlobal, dstGlobal;
    int srcDataSize = 256;
    int dstDataSize = 256;
};
extern "C" __global__ __aicore__ void duplicate_kernel(__gm__ uint8_t* src, __gm__ uint8_t* dstGm)
{
    KernelDuplicate op;
    op.Init(src, dstGm);
    op.Process();
}

          

        

      
     

Parent topic: Data Padding