How to 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 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 global memory of the corresponding size on the device to the workspace. 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 Kernel APIs), 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 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.

  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 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 size of the system workspace can be obtained through the GetLibApiWorkSpaceSize API of PlatformAscendCManager. After HAVE_WORKSPACE is enabled, the workspace obtained from the input parameter on the kernel is the user workspace after the system workspace is offset.