---
title: 核函数
description: "Ascend C支持开发者自定义核函数来扩展C++，核函数在AI处理器上执行时实际是若干线程在并行执行，每个线程有独立的寄存器和栈，共同完成数据计算任务。"
url: https://www.hiascend.com/document/detail/zh/canncommercial/latest/programug/Ascendcopdevg/atlas_ascendc_10_10067.html
sourcePath: /source/zh/canncommercial/900/programug/Ascendcopdevg/atlas_ascendc_10_10067.html
indexId: b65f73532df0bee8fc923aa7facb6ca1d892336e7bb36aa1bb2be8f31daea52c81
---
# 核函数

#### 核函数的定义

Ascend C支持开发者自定义核函数来扩展C++，核函数在AI处理器上执行时实际是若干线程在并行执行，每个线程有独立的寄存器和栈，共同完成数据计算任务。

核函数的定义示例如下：

```
__global__ void kernel_name(argument list)
```

定义核函数时需要遵循以下规则：


- 使用函数类型限定符__global__， 标识它是一个核函数。
- 核函数必须具有void返回类型。
- 函数入参支持的类型如下：

  - 基础数据类型，如int32_t、float等。
  - 基础数据类型的指针类型，如int32_t*、float*等，实际上这些指针指向的是GlobalMemory内存。

#### 核函数的调用

算子程序中的函数可以分为三类：host侧执行函数、核函数（device侧执行）、device侧执行函数（除核函数之外）。下图以Kernel直调算子开发方式为例，描述三者的调用关系：

- host侧执行函数可以调用其它host执行函数，也就是通用C/C++编程中的函数调用；也可以通过<<<...>>>调用核函数。
- 核函数可以调用除核函数之外的其它device侧执行函数。
- device侧执行函数（除核函数之外）使用类型限定符__aicore__标识，可以调用同类的其它device侧执行函数。

图1 算子程序中三种函数间的调用关系

Host侧通过内核调用符<<<...>>>的语法形式调用核函数，如下所示：

```
kernel_name<<<numBlocks, threadsPerBlock, dynUBufSize, stream>>>(args...)
```

执行配置由4个参数决定：- numBlocks：为核函数配置的线程块的个数，即启用的核数, 支持int32_t和dim3类型；
- threadsPerBlock：每个线程块内并发执行的线程数量，支持int32_t和dim3类型；
- dynUBufSize：动态申请内存空间总大小，一般情况设置为0；
- stream：用于host侧和device侧的流同步。