昇腾社区首页
中文
注册
开发者
下载
支持
积分兑换 NEW

约束说明

  • 在同一个编译单元，若存在多个核函数，暂不支持自动推导Kernel类型，需要开发者手动设置Kernel类型。
    • 特别地，针对如下型号，无论是否是同一个编译单元多个核函数的场景，均不支持在开发者未设置Kernel类型时进行自动推导。建议开发者手动设置Kernel类型。
      • Atlas 推理系列产品
    • 针对Atlas 推理系列产品 ，暂不支持设置Kernel类型为KERNEL_TYPE_MIX_VECTOR_CORE。
  • KERNEL_TASK_TYPE_DEFAULT接口需在核函数或者inline的aicore函数中进行调用，禁止在非inline的aicore函数中进行调用。
  • 纯Scalar算子无法实现自动推导

    需手动标记Kernel函数类型，推荐设置为纯Vector类型，添加__vector__ attribute进行标记：

    1
2
3
4
    __global__ __vector__ __aicore__ void func0(__gm__ uint8* Addr) {
    Addr[1] = Addr[0];
    AscendC::printf("Hello world");
}
  • Kernel函数不支持特化

    使用特化核函数调用，目前依然会使用主模板核函数。

     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
    #include "kernel_operator.h"
#include "acl/acl.h"

// 测试__mix__(x, y)中x, y由模板参数控制是否可行
template<int32_t cube, int32_t vec>
__mix__(cube, vec) __global__ void hello_world()
{
    if ASCEND_IS_AIC {
        AscendC::printf("Hello World AIC with cube, vec !!!\n");
    } else {
        AscendC::printf("Hello World AIV with cube, vec !!!\n");
    }
}

// 特化写法，不支持
template<>
__mix__(1, 0) __global__ void hello_world<1, 0>()
{
    if ASCEND_IS_AIC {
        AscendC::printf("1:0 Hello World AIC with vec  !!!\n");
    } else {
        AscendC::printf("1:0 Hello World AIV with vec !!!\n");
    }
}


int32_t main(int argc, char const *argv[])
{
    uint8_t *aHost;
    uint8_t *aDevice;

    aclInit(nullptr);
    int32_t deviceId = 0;
    aclrtSetDevice(deviceId);
    aclrtStream stream = nullptr;
    aclrtCreateStream(&stream);
    aclrtMallocHost((void **)(&aHost), sizeof(uint32_t));
    aclrtMalloc((void **)&aDevice, sizeof(uint32_t), ACL_MEM_MALLOC_HUGE_FIRST);
    *aHost = 12;
    aclrtMemcpy(aDevice, sizeof(uint32_t), aHost, sizeof(uint32_t), ACL_MEMCPY_HOST_TO_DEVICE);
    constexpr uint32_t numBlocks = 8;

    hello_world<1, 2><<<numBlocks, nullptr, stream>>>();
    aclrtSynchronizeStream(stream);
    hello_world<1, 0><<<numBlocks, nullptr, stream>>>();
    aclrtSynchronizeStream(stream);

    aclrtFree(aDevice);
    aclrtFreeHost(aHost);
    aclrtDestroyStream(stream);
    aclrtResetDevice(deviceId);
    aclFinalize();
    return 0;
}
  • bfloat16_t等数据类型在Host侧不支持，使用这些数据类型时，Host和Device不能写在同一个实现文件里。Host侧不支持的数据类型如下：

    Atlas A2 训练系列产品/Atlas A2 推理系列产品：bfloat16_t。

    Atlas A3 训练系列产品/Atlas A3 推理系列产品：bfloat16_t。

  • 不支持可变参数模板和可变参数函数
    1
2
3
    // 不支持以下写法
template <typename... Args>  // Args 是一个“类型参数包”
void func(Args... args);     // args 是一个“函数参数包”
  • 不支持#line 预处理
    1
2
3
    // 不支持以下写法
#line number // 更改行号
#line number "filename" // 更改文件名
父主题： AI Core算子编译