3.5 内核（函数）

在.cu文件中，用_global_标识了内核。在CUDA运行时中，它们可以使用由三对尖括号（<<>>）构成的单行代码启动。在第7章给出了内核怎样被启动和怎样在GPU上执行的细节。

模块中的GPU可执行代码以内核形式呈现，或者使用CUDA运行时的语言集成特性启动（<<>>符号）或者使用驱动程序API中的cuLaunchKernel()函数启动。在本书写作时，CUDA还没有任何CUDA模块中可执行代码的动态驻留管理。当模块加载时，内核整个地加载进设备内存。

一旦模块被加载，内核可使用cuModuleGetFunction()函数查询，内核的属性使用cuFuncGetAttribute()函数查询，内核通过cuLaunchKernel()函数启动，cuLaunchKernel()包含了一整套已被废弃的API入口点：像cuFuncSetBlockShape()指定下次内核启动使用的线程块大小；函数cuParamSetv()指定下次内核启动传入的参数；cuLaunch()、cuLaunchGrid()和cuLaunchGridAsync()使用此前设置的状态启动内核。这些API很低效，因为这花费了太多的调用来启动一次内核，并且很多参数，像线程块大小，最好在每次请求内核启动时统一指定。

函数cuFuncGetAttribute()可以被用来查询指定的属性，例如：

·每线程块最大线程数
·静态分配共享内存的数量
用户分配的常量内存大小
每个函数使用的本地内存数量

• 函数中每个线程使用的寄存器数量

• 被编译函数的虚拟（PTX）与二进制架构版本

当使用驱动程序API，最好使用extern C来禁止C++中的默认名称改编行为。否则，你必须对cuModuleGetFunction()使用改编的名称。

1. CUDA运行时

在CUDA运行时创建的可执行文件被加载的同时，它会在主机内存创建一个全局数据结构体，来描述当CUDA设备创建后应分配的CUDA资源。一旦CUDA设备被初始化，这些全局数据结构体就被用来一次性创建CUDA资源。因为这些全局数据在CUDA运行时中属于进程级共享，所以使用CUDA运行时不能递增的加载或卸载CUDA模块。

因为CUDA运行时用C++语言集成，在API函数（像CUDAFuncGetAttributes()和CUDAMemcpy-ToSymbol())内，内核和符号必须用名字指定（即，不是字符串字面值）。

2. 缓存设置

在费米架构中，流处理器簇有一级缓存，它可以被分割为16KB大小的共享内存/48KB一级缓存或48KB的共享内存/16KB的一级缓存 [1]。最初，CUDA允许以每个内核为基础配置缓存，在CUDA运行时中使用CUDAFuncSetCacheConfig()或驱动程序API中使用cuFuncSetCacheConfig()设置。后来，这一状态变得全局化：使用cuCtxSetCacheConfig()/cudaDeviceSetCacheConfig()函数设定默认缓存配置。

[1] SM 3.X添加了均分缓存的能力（32KB/32KB）。