A workspace is a memory block in global memory on the device. The workspace consists of the system workspace and the developer 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.
Developer 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. The application scenarios are as follows:
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 developer workspace.
// 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; // Set 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 developer workspace pointer.
// Kernel function written by the developer. 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)
{
// ...
}