如何申请workspace作为临时内存

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

• 系统workspace：AscendC API需要预留的workspace内存

  API在计算过程需要一些workspace内存作为缓存，因此算子需要为API预留workspace内存，预留内存大小通过GetLibApiWorkSpaceSize接口获取。

• 开发者workspace：算子实现使用到的workspace内存

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

  • 需要使用Unified Buffer和L1 Buffer上空间且空间不够用时，可以将数据暂存至workspace上。
  • 其他需要使用Global Memory上内存空间的场景。

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

• 工程化算子开发方式

  在tiling函数中先通过GetWorkspaceSizes接口获取workspace大小的存放位置，再设置workspace的大小，框架侧会为其在申请对应大小的设备侧Global Memory，在对应的算子kernel侧实现时可以使用这块workspace内存。在使用Matmul等需要系统workspace的高阶API时，设置的workspace空间大小为系统workspace和开发者workspace之和。

  // 开发者自定义的tiling函数
  static ge::graphStatus TilingFunc(gert::TilingContext* context)
  {
      AddApiTiling tiling;
      // ...
      size_t usrSize = 256; // 设置开发者需要使用的workspace大小。
      // 如需要使用系统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指针：

  // 开发者写的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)
   {
       // ...
        
   }