下载
中文
注册
我要评分
文档获取效率
文档正确性
内容完整性
文档易理解
在线提单
论坛求助
昇腾小AI

SetAtomicType

功能说明

通过设置模板参数来设定原子操作不同的数据类型。

函数原型

1
2
template <typename T>
__aicore__ inline void SetAtomicType()

参数说明

表1 模板参数说明

参数名

描述

T

设定不同的数据类型。

Atlas A2训练系列产品/Atlas 800I A2推理产品,支持float/half/int16_t/int32_t/int8_t/bfloat16_t

Atlas推理系列产品AI Core,支持float/half/int16_t

Atlas 200/500 A2推理产品,支持float/half/int16_t/int32_t

返回值

支持的型号

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

Atlas推理系列产品AI Core

Atlas 200/500 A2推理产品

约束说明

需要和SetAtomicAddSetAtomicMaxSetAtomicMin配合使用。

使用完成后,建议清空原子操作的状态(详见SetAtomicNone),以免影响后续相关指令功能。

调用示例

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
// 本演示示例使用DataCopy从VECOUT搬出到外部dstGlobal时进行原子最小,并调用SetAtomicType修改原子最小的数据类型。
#include "kernel_operator.h"

static const int data_size = 256;
template <typename T>
class KernelDataCopyAtomicMin {
public:
    __aicore__ inline KernelDataCopyAtomicMin() {}
    __aicore__ inline void Init(GM_ADDR src0_gm, GM_ADDR src1_gm, GM_ADDR dst_gm)
    {
        src0Global.SetGlobalBuffer((__gm__ T *)src0_gm);
        src1Global.SetGlobalBuffer((__gm__ T *)src1_gm);
        dstGlobal.SetGlobalBuffer((__gm__ T *)dst_gm);

        pipe.InitBuffer(queueSrc, 2, data_size * sizeof(T));
    }
    __aicore__ inline void Process()
    {
        CopyIn();
        Compute();
        CopyOut();
    }
private:
    __aicore__ inline void CopyIn()
    {
        AscendC::LocalTensor<T> src0local = queueSrc.AllocTensor<T>();
        AscendC::LocalTensor<T> src1local = queueSrc.AllocTensor<T>();
        AscendC::DataCopy(src0local, src0Global, data_size);
        queueSrc.EnQue(src0local);
        AscendC::DataCopy(src1local, src1Global, data_size);
        queueSrc.EnQue(src1local);
    }
    __aicore__ inline void Compute()
    {
        AscendC::LocalTensor<T> src0local = queueSrc.DeQue<T>();
        AscendC::LocalTensor<T> src1local = queueSrc.DeQue<T>();
        AscendC::Abs(src0local, src0local, data_size);
        AscendC::Abs(src1local, src1local, data_size);
        queueSrc.EnQue(src0local);
        queueSrc.EnQue(src1local);
    }
    __aicore__ inline void CopyOut()
    {
        AscendC::LocalTensor<T> dst0Local = queueSrc.DeQue<T>();
        AscendC::LocalTensor<T> dst1Local = queueSrc.DeQue<T>();
        AscendC::DataCopy(dstGlobal, dst1Local, data_size);
        AscendC::SetAtomicMin<int8_t>();  // 此处设置的类型可随意,此例中以int8_t为例
        AscendC::SetAtomicType<T>();  // 此处设置真实的数据类型
        AscendC::DataCopy(dstGlobal, dst0Local, data_size);
        queueSrc.FreeTensor(dst0Local);
        AscendC::DataCopy(dstGlobal, dst1Local, data_size);
        queueSrc.FreeTensor(dst1Local);
        AscendC::SetAtomicNone();
    }
private:
    AscendC::TPipe pipe;
    AscendC::TQue<AscendC::QuePosition::VECIN, 2> queueSrc;
    AscendC::GlobalTensor<T> src0Global, src1Global, dstGlobal;
};

extern "C" __global__ __aicore__ void data_copy_atomic_min_kernel(GM_ADDR src0_gm, GM_ADDR src1_gm, GM_ADDR dst_gm)
{
    KernelDataCopyAtomicMin<half> op;
    op.Init(src0_gm, src1_gm, dst_gm);
    op.Process();
}
每个核的输入数据为: 
Src0: [1,1,1,1,1,...,1] // 256个1
Src1: [2,2,2,2,2,...,2] // 256个2
最终输出数据: [1,1,1,1,1,...,1] // 256个1
搜索结果
找到“0”个结果

当前产品无相关内容

未找到相关内容,请尝试其他搜索词