基于Ascend C开发的矢量计算相关的算子可以运行在Vector Core上,该接口用于使能Vector Core。基于Vector Core架构进行算子编程的详细介绍请参考基于VectorCore编程。
KERNEL_TASK_TYPE_DEFAULT(value)
参数 |
输入/输出 |
说明 |
---|---|---|
value |
输入 |
取值为KERNEL_TYPE_MIX_VECTOR_CORE,设置kernel类型为Vector Core。 算子执行时会同时启动AI Core和Vector Core, 此时AI Core会当成Vector Core使用,比如用户在host侧设置block_dim为10,则会启动总数为10的AI Core和Vector Core。 通过SetBlockDim设置核数时, 需要大于AI Core的核数,否则不会启动VectorCore。 |
Atlas 推理系列产品
无
extern "C" __global__ __aicore__ void add_custom(__gm__ uint8_t *x, __gm__ uint8_t *y, __gm__ uint8_t *z, __gm__ uint8_t *workspace, __gm__ uint8_t *tiling)
{
GET_TILING_DATA(tilingData, tiling);
if (workspace == nullptr) {
return;
}
SetSysWorkspace(workspace);
GM_ADDR usr = GetUserWorkspace(workspace);
KernelAdd op;
op.Init(x, y, z, tilingData.blockDim, tilingData.totalLength, tilingData.tileNum);
KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_VECTOR_CORE); // 使能VectorCore
if (TILING_KEY_IS(1)) {
op.Process1();
} else if (TILING_KEY_IS(2)) {
op.Process2();
}
// ...
}
// 配套的host侧tiling函数示例: ge::graphStatus TilingFunc(gert::TilingContext* context) { // 使能VectorCore,将block_dim置为AI Core中vector核数 + Vector Core中的vector核数 auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo); auto totalCoreNum = ascendcPlatform.GetCoreNumAiv(); if (ascendcPlatform.GetSocVersion() == platform_ascendc::SocVersion::ASCEND310P) { totalCoreNum = totalCoreNum + ascendcPlatform.GetCoreNumVector(); } context->SetBlockDim(totalCoreNum); }