How do I use 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
During API compute, some workspace memory is required as the cache. Therefore, the operator needs to reserve workspace memory for the API. The reserved memory 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 called to obtain the storage location of the workspace size, and then the workspace size is set. The framework allocates the global memory of the corresponding size on the device to the workspace. This workspace memory can be used in the implementation of the corresponding operator kernel. When high-level APIs that require the system workspace are used (such as Matmul Kernel APIs), the workspace size is the sum of the system workspace and the user workspace.

        
             // User-defined tiling function
static ge::graphStatus TilingFunc(gert::TilingContext* context)
{
    AddApiTiling tiling;
    ...
    size_t usrSize = 256; // Set the required workspace size to 256 bytes.
    // 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.

        
             // 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 HAVE_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 Kernel APIs), the workspace size is the sum of the system workspace and the user workspace. The GetLibApiWorkSpaceSize API of PlatformAscendCManager can be called to obtain the size of the system workspace. After HAVE_WORKSPACE is enabled, the workspace obtained by the developer from the kernel input parameter is the user workspace with the system workspace offset.

Parent topic: Common Operation