GetUserWorkspace
函数原型
__aicore__ inline GM_ADDR GetUserWorkspace(GM_ADDR workspace)
参数说明
参数名称 |
输入/输出 |
描述 |
---|---|---|
workspace |
输入 |
传入workspace的指针,包括系统workspace和用户使用的workspace。 |
支持的型号
Atlas 训练系列产品
Atlas推理系列产品AI Core
Atlas A2训练系列产品/Atlas 800I A2推理产品
注意事项
无
返回值
用户使用workspace指针。
调用示例
在Host侧的Tiling函数中,先通过GetWorkspaceSizes接口获取workspace大小的存放位置,再设置worksapce的块数和workspace的size。
// 用户自定义的tiling函数 static ge::graphStatus TilingFunc(gert::TilingContext* context) { AddApiTiling tiling; ... size_t usrSize = 256; // 设置用户需要使用的workspace大小。 // 如需要使用系统workspace需要调用GetLibApiWorkSpaceSize获取系统workspace的大小。 auto ascendcPlatform = platform_ascendc:: PlatformAscendC(context->GetPlatformInfo()); uint32_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize(); size_t *currentWorkspace = context->GetWorkspaceSizes(1); // 通过框架获取workspace的指针,GetWorkspaceSizes入参为所需workspace的块数。当前限制使用一块。 currentWorkspace[0] = usrSize + sysWorkspaceSize; // 设置总的workspace的数值大小,总的workspace空间由框架来申请并管理。 ... }
在device侧调用本接口获取用户的workspace指针:
// 用户写的Kernel函数,核函数必须包括GM_ADDR workspace入参,位置需要放在tiling之前 extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling) { ... GM_ADDR usrWorkspace = AscendC::GetUserWorkspace(workspace); // 获取用户workspace指针。 ... }
父主题: workspace