Restrictions

If there are multiple kernel functions in the same compilation unit, the kernel type cannot be automatically inferred. You need to manually set the kernel type.
- Particularly, for the following models, automatic inference is not supported when the kernel type is not set, regardless of whether there are multiple kernel functions in the same compilation unit. You are advised to manually set the kernel type.
  - Atlas 350 Accelerator Card
  - Atlas inference product
- For the Atlas inference product, the kernel type cannot be set to KERNEL_TYPE_MIX_VECTOR_CORE.
The KERNEL_TASK_TYPE_DEFAULT API needs to be called in the kernel function.
Automatic inference cannot be implemented for scalar-only operators.
You need to manually mark the kernel function type. It is recommended that the kernel function type be set to the vector-only type and __vector__ attribute be added for marking.
1 2 3 4
__global__ __vector__ __aicore__ void func0(__gm__ uint8* Addr) { Addr[1] = Addr[0]; AscendC::printf("Hello world"); }

The kernel function does not support specialization.

When a specialized kernel function is used for calling, the primary template kernel function is still used.

#include "kernel_operator.h"
#include "acl/acl.h"

// Test whether x and y in __mix__(x, y) can be controlled by template parameters.
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");
    }
}

// Specialization syntax, not supported
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;
}

Data types such as bfloat16_t are not supported on the host. When these data types are used, the host and device cannot be written in the same implementation file. Data types that are not supported on the host are as follows:
Atlas 350 Accelerator Card: bfloat16_t, hifloat8_t, fp8_e5m2_t, fp8_e4m3fn_t, fp8_e8m0_t, fp4x2_e2m1_t, fp4x2_e1m2_t, int4x2_t

Atlas A2 training product/Atlas A2 inference product: bfloat16_t

Atlas A3 training product/Atlas A3 inference product: bfloat16_t

Variable parameter templates and functions are not supported.

// The following syntax is not supported:
template <typename... Args>  // Args is a type parameter package.
void func(Args... args);     // args is a function parameter package.

The #line preprocessing is not supported.

// The following syntax is not supported:
#line number // Change the line number.
#line number "filename" // Change the file name.

Parent topic: AI Core SIMD Programming