约束说明

__global__ __vector__ __aicore__ void func0(__gm__ uint8* Addr) {
    Addr[1] = Addr[0];
    AscendC::printf("Hello world");
} 

#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;
}

// 不支持以下写法
template <typename... Args>  // Args 是一个“类型参数包”
void func(Args... args);     // args 是一个“函数参数包”

// 不支持以下写法
#line number // 更改行号
#line number "filename" // 更改文件名