Duplicate

Applicability

Product	Supported
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

Copies a variable or an immediate value multiple times and fills the result into a vector.

Prototype

Computation of the first n data elements of a tensor

When the source operand is a scalar:

          
               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** Template parameters
Parameter	Description
T	Operand data type. For 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. For 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. For Atlas 200I/500 A2 inference products , the supported data types are int16_t, uint16_t, half, int32_t, uint32_t, and float. For the Atlas inference product 's AI Core, the supported data types are int16_t, uint16_t, half, int32_t, uint32_t, and float. For 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. Its data type must match that of the elements in dst.
count	Input	Number of elements involved in the computation.
mask/mask[]	Input	mask is used to control the elements that participate in computation in each iteration. Bitwise mode: controls which elements are involved in computation bit by bit. A bit value of 1 means the corresponding element participates in computation, while 0 means it does not. The mask value is an array. The array length and the value range of the array elements are related to the operand data type. When the operand is 16-bit, the array length is 2, mask[0] and mask[1] ∈ [0, 2⁶⁴ -1] and cannot be 0 at the same time. When the operand is 32-bit, the array length is 1 and mask[0] ∈ (0, 2⁶⁴ – 1]. When the operand is 64-bit, the array length is 1 and mask[0] ∈ (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 compute unit reads 8 consecutive data blocks (32 bytes per block, 256 bytes in total) for computation each time. To process the input data, the data needs to be read and computed over multiple repeats. 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

Example

This example only shows code snippets for the computation process. To run the code, copy and paste these code snippets into the corresponding position of the Compute function in Template Sample.

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 data elements of a tensor, with the source operand being 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 only shows code snippets for the computation process. To run the code, copy and paste these code snippets into the corresponding position of the Compute function in Template Sample.

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 DataBlock Stride parameter of a high-dimensional tensor sharding computation API to implement non-contiguous 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 Repeat stride parameter of a high-dimensional tensor sharding computation API to implement non-contiguous 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 Sample

      
       
         
         
           #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