---
title: 如何使用workspace
description: "workspace是设备侧Global Memory上的一块内存。workspace内存分为两部分：系统workspace和用户workspace。"
url: https://www.hiascend.com/document/detail/zh/canncommercial/latest/programug/Ascendcopdevg/atlas_ascendc_10_0092.html
sourcePath: /source/zh/canncommercial/900/programug/Ascendcopdevg/atlas_ascendc_10_0092.html
indexId: cb2b1bcc22c6951dcb136ab2e6156ea6c920e162f775bb1f762440362a2a5d7980
---
# 如何使用workspace

workspace是设备侧Global Memory上的一块内存。workspace内存分为两部分：系统workspace和用户workspace。

- 系统workspace：Ascend C API需要预留的workspace内存
  API在计算过程需要一些workspace内存作为缓存，因此算子需要为API预留workspace内存，预留内存大小通过GetLibApiWorkSpaceSize(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_1036.html)接口获取。

- 用户workspace：算子实现使用到的workspace内存
算子内部需要通过额外的device内存进行数据交换或者缓存的时候才需要分配，根据实际情况自行分配。使用场景如下：

  - 需要使用Unified Buffer和L1 Buffer上的空间且空间不够用时，可以将数据暂存至workspace上。
  - 调用    SyncAll(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_0204.html)
等API接口时，需要workspace作为入参。
  - 其他需要使用Global Memory上内存空间的场景。

不同开发方式下，具体的使用方法如下：

- 工程化算子开发方式
  在tiling函数中先通过GetWorkspaceSizes(https://www.hiascend.comdocument/detail/zh/canncommercial/900/maintenref/basicdataapi/atlasopapi_07_00244.html)接口获取workspace大小的存放位置，再设置workspace的大小，框架侧会为其申请对应大小的设备侧Global Memory，在对应的算子kernel侧实现时可以使用这块workspace内存。在使用Matmul Kernel侧接口(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_0613.html)等需要系统workspace的高阶API时，设置的workspace空间大小为系统workspace和用户workspace之和。 1 2 3 4 5 6 7 8 9 10 11 12 13 // 用户自定义的tiling函数 static ge::graphStatus TilingFunc(gert::TilingContext* context) { AddApiTiling tiling; ... size_t usrSize = 256; // 设置用户需要使用的workspace大小为256字节。 // 如需要使用系统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侧kernel入口处的workspace为用户的workspace指针： 1 2 3 4 5 6 // 用户写的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) { ... }

- Kernel直调算子开发场景
  需要使用workspace空间时，建议开启编译选项HAVE_WORKSPACE。host侧开发者仍需要自行申请workspace的空间，并传入。在使用Matmul Kernel侧接口(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_0613.html)等需要系统workspace的高阶API时，设置的workspace空间大小为系统workspace和用户workspace之和。系统workspace大小可以通过PlatformAscendCManager(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_1039.html)的GetLibApiWorkSpaceSize接口获取。开启HAVE_WORKSPACE后，开发者在kernel侧入参处获取的workspace为偏移了系统workspace后的用户workspace。