workspace

功能介绍

workspace是设备侧Global Memory上的一块内存。您需要使用Global Memory的内存时，可以考虑使用workspace。具体使用场景如下：

需要使用Unified Buffer和L1 Buffer上空间且空间不够用时，可以将数据暂存至workspace上。
调用SyncAll等API接口时，可以使用workspace作为入参。
其他需要使用Global Memory上内存空间的场景。

使用方法

在Host侧的Tiling函数中，先通过GetWorkspaceSizes接口获取workspace大小的存放位置，再设置workspace的块数和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侧调用GetUserWorkspace获取用户的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指针。
    ...
}

使用约束

针对Atlas A2训练系列产品/Atlas 800I A2推理产品，必须设置系统workspace。
使用workspace，必须设置系统workspace大小。
Ascend C当前限制仅能使用一块workspace。

父主题： 内存管理