---
title: 函数执行空间限定符
description: "函数执行空间限定符（Function Execution Space Qualifier）指示函数是在host侧执行还是在device侧执行，以及它是否可从host侧或device侧调用。"
url: https://www.hiascend.com/document/detail/zh/canncommercial/latest/compiler/BishengCompiler/atlas_bisheng_10_0006.html
sourcePath: /source/zh/canncommercial/900/compiler/BishengCompiler/atlas_bisheng_10_0006.html
indexId: 6f0de346a8cb8ae16fb0a7ae6162155627c000654c6792c0ce0f9aed95cb8ae381
---
# 函数执行空间限定符

函数执行空间限定符（Function Execution Space Qualifier）指示函数是在host侧执行还是在device侧执行，以及它是否可从host侧或device侧调用。

- __global__
__global__执行空间限定符声明一个kernel函数。kernel函数有如下性质：在device上执行；只能被host侧函数调用；__global__只是表示这是device侧函数的入口，并不表示具体的设备类型，具体的设备类型由[aicore]标记。具有如下使用约束：

  - 一个[aicore] __global__函数必须返回void类型，并且不能是class的成员函数。
  - 主机侧调用[aicore] __global__函数必须使用<<<>>>异构调用语法。
  - [aicore] __global__的调用是异步的，意味着函数返回，并不表示kernel函数在device侧已经执行完成，如果需要同步，需要使用Runtime同步接口显式同步，如    aclrtSynchronizeStream(https://www.hiascend.comdocument/detail/zh/canncommercial/900/API/runtimeapi/aclcppdevg_03_0076.html)
接口。
- [aicore]
aicore执行空间限定符声明一个函数，它具有如下属性：

  - 在device侧执行
  - 只能被__global__函数，或者其他aicore函数调用
  - 该限定符与Ascend C中限定符 __aicore__ 含义一致。适用于开发自定义算子，但未使用Ascend C头文件场景。

```
// Only callable from device functions with same kind
// of execution space
[aicore] void bar() {}
// Define a kernel function execute on AICore device
__global__ [aicore] void foo() {
bar(); // OK.
}
```