Documentation

Workspace

A workspace is a memory block on global memory on the device. The workspace consists of the system workspace and the user workspace.

  • System workspace: workspace reserved for Ascend C APIs

    Some workspace needs to be reserved by the operators for API computing. The reserved workspace size can be obtained by calling GetLibApiWorkSpaceSize.

  • User workspace: workspace used for operator implementation

    The user workspace needs to be allocated only when extra device memory is required for data exchange or caching inside the operator. You can allocate it as required. It is supposed to be used in the following scenarios:

    • If the space on the Unified Buffer and L1 Buffer is insufficient, you can temporarily store data in the workspace.
    • When calling APIs such as SyncAll, you can use the workspace as the input parameter.
    • It can be used in other scenarios where the memory space in global memory needs to be used.

The following introduces how to use the workspace in different development modes:

  • Project-based operator development
    In the tiling function, the GetWorkspaceSizes API is used to obtain the location for storing the workspace size, and then set the workspace size. The framework allocates the corresponding size of global memory on the device. The workspace can be used when the corresponding operator is implemented on the kernel. When high-level APIs that require the system workspace are used (such as Matmul), the workspace size is the sum of the system workspace and the user workspace. 
     1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
    		 
    // User-defined tiling function
static ge::graphStatus TilingFunc(gert::TilingContext* context)
{
    AddApiTiling tiling;
    ...
    size_t usrSize = 256; // Set the workspace size as required.
    // To use the system workspace, call GetLibApiWorkSpaceSize to obtain the size of the system workspace.
    auto ascendcPlatform = platform_ascendc:: PlatformAscendC(context->GetPlatformInfo());
    uint32_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
    size_t *currentWorkspace = context->GetWorkspaceSizes(1); // Obtain the workspace pointer through the framework. The input parameter of GetWorkspaceSizes is the number of required workspace blocks. Currently, only one block can be used.
    currentWorkspace[0] = usrSize + sysWorkspaceSize; // Specify the total workspace size. The total workspace space is allocated and managed by the framework.
    ...
}
    The workspace at the kernel entry point on the device is the user workspace pointer. 
    1
2
3
4
5
6
    		 
    // Kernel function written by the user. The kernel function must contain the GM_ADDR workspace input parameter, which must be placed before tiling.
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling)
{
    ...
    
}
  • Kernel launch operator development

    If the workspace is required, you are advised to enable the compilation option -DHAVE_WORKSPACE. You still need to allocate the workspace size on the host and pass it in. When high-level APIs that require the system workspace are used (such as Matmul), the workspace size is the sum of the system workspace and the user workspace. The size of the system workspace can be obtained through the GetLibApiWorkSpaceSize API of PlatformAscendCManager. After -DHAVE_WORKSPACE is enabled, the workspace obtained from the input parameter on the kernel is the user workspace plus the offset system workspace.

Parent topic: Memory Management