CUDA 运行时 API :: CUDA 工具包文档

6.10. 占用率

本节介绍 CUDA 运行时应用程序编程接口的占用率计算函数。

除了占用率计算器函数（cudaOccupancyMaxActiveBlocksPerMultiprocessor 和 cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags）之外，还有 C++ 独有的基于占用率的启动配置函数，这些函数在 C++ API 例程模块中进行了文档化。

请参阅 cudaOccupancyMaxPotentialBlockSize ( C++ API), cudaOccupancyMaxPotentialBlockSize ( C++ API), cudaOccupancyMaxPotentialBlockSizeVariableSMem ( C++ API), cudaOccupancyMaxPotentialBlockSizeVariableSMem ( C++ API) cudaOccupancyAvailableDynamicSMemPerBlock (C++ API),

函数

__host__ cudaError_t cudaOccupancyAvailableDynamicSMemPerBlock ( size_t* dynamicSmemSize, const void* func, int numBlocks, int blockSize ): 返回在 SM 上启动 numBlocks 个块时每个块可用的动态共享内存。
__host__ __device__ cudaError_t cudaOccupancyMaxActiveBlocksPerMultiprocessor ( int* numBlocks, const void* func, int blockSize, size_t dynamicSMemSize ): 返回设备函数的占用率。
__host__ cudaError_t cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags ( int* numBlocks, const void* func, int blockSize, size_t dynamicSMemSize, unsigned int flags ): 返回带有指定标志的设备函数的占用率。
__host__ cudaError_t cudaOccupancyMaxActiveClusters ( int* numClusters, const void* func, const cudaLaunchConfig_t* launchConfig ): 给定内核函数 (func) 和启动配置 (config)，返回目标设备上可以共存的最大簇数量，结果存于 *numClusters 中。
__host__ cudaError_t cudaOccupancyMaxPotentialClusterSize ( int* clusterSize, const void* func, const cudaLaunchConfig_t* launchConfig ): 给定内核函数 (func) 和启动配置 (config)，返回最大簇大小，结果存于 *clusterSize 中。

函数

__host__ cudaError_t cudaOccupancyAvailableDynamicSMemPerBlock ( size_t* dynamicSmemSize, const void* func, int numBlocks, int blockSize )

返回在启动时每个块可用的动态共享内存numBlocksSM 上的块。

参数

dynamicSmemSize: - 返回的最大动态共享内存
func: - 用于计算占用率的内核函数
numBlocks: - 适合 SM 的块数
blockSize: - 块的大小

返回值

cudaSuccess, cudaErrorInvalidDevice, cudaErrorInvalidDeviceFunction, cudaErrorInvalidValue, cudaErrorUnknown,

描述

在*dynamicSmemSize中返回允许的最大动态共享内存大小numBlocks每个 SM 的块。

注意

请注意，此函数也可能返回来自先前异步启动的错误代码。
请注意，如果此调用尝试初始化内部 CUDA RT 状态，此函数也可能返回 cudaErrorInitializationError、cudaErrorInsufficientDriver 或 cudaErrorNoDevice。
请注意，如 cudaStreamAddCallback 所指定，不得从回调中调用任何 CUDA 函数。cudaErrorNotPermitted 可能会（但不能保证）在这种情况下作为诊断返回。
该 API 也可以与内核 cudaKernel_t 一起使用，方法是使用 cudaLibraryGetKernel() 或 cudaGetKernel 查询句柄，然后将其强制转换为 void* 传递给 API。符号entryFuncAddr传递给 cudaGetKernel 的符号应是使用相同 CUDA 运行时实例注册的符号。
传递属于不同运行时实例的符号将导致未定义的行为。唯一可以可靠地传递给不同运行时实例的类型是 cudaKernel_t

另请参阅

cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags, cudaOccupancyMaxPotentialBlockSize ( C++ API), cudaOccupancyMaxPotentialBlockSizeWithFlags ( C++ API), cudaOccupancyMaxPotentialBlockSizeVariableSMem ( C++ API), cudaOccupancyMaxPotentialBlockSizeVariableSMemWithFlags ( C++ API), cudaOccupancyAvailableDynamicSMemPerBlock

__host__ __device__ cudaError_t cudaOccupancyMaxActiveBlocksPerMultiprocessor ( int* numBlocks, const void* func, int blockSize, size_t dynamicSMemSize )

返回设备函数的占用率。

参数

numBlocks: - 返回的占用率
func: - 用于计算占用率的内核函数
blockSize: - 内核计划启动的块大小
dynamicSMemSize: - 计划使用的每个块的动态共享内存大小，以字节为单位

返回值

cudaSuccess, cudaErrorInvalidDevice, cudaErrorInvalidDeviceFunction, cudaErrorInvalidValue, cudaErrorUnknown,

描述

在*numBlocks设备函数的每个流式多处理器 (SM) 的最大活动块数。

注意

请注意，此函数也可能返回来自先前异步启动的错误代码。
请注意，如果此调用尝试初始化内部 CUDA RT 状态，此函数也可能返回 cudaErrorInitializationError、cudaErrorInsufficientDriver 或 cudaErrorNoDevice。
请注意，如 cudaStreamAddCallback 所指定，不得从回调中调用任何 CUDA 函数。cudaErrorNotPermitted 可能会（但不能保证）在这种情况下作为诊断返回。
该 API 也可以与内核 cudaKernel_t 一起使用，方法是使用 cudaLibraryGetKernel() 或 cudaGetKernel 查询句柄，然后将其强制转换为 void* 传递给 API。符号entryFuncAddr传递给 cudaGetKernel 的符号应是使用相同 CUDA 运行时实例注册的符号。
传递属于不同运行时实例的符号将导致未定义的行为。唯一可以可靠地传递给不同运行时实例的类型是 cudaKernel_t

另请参阅

__host__ cudaError_t cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags ( int* numBlocks, const void* func, int blockSize, size_t dynamicSMemSize, unsigned int flags )

返回带有指定标志的设备函数的占用率。

参数

numBlocks: - 返回的占用率
func: - 用于计算占用率的内核函数
blockSize: - 内核计划启动的块大小
dynamicSMemSize: - 计划使用的每个块的动态共享内存大小，以字节为单位
flags: - 占用率计算器的请求行为

返回值

cudaSuccess, cudaErrorInvalidDevice, cudaErrorInvalidDeviceFunction, cudaErrorInvalidValue, cudaErrorUnknown,

描述

在*numBlocks设备函数的每个流式多处理器 (SM) 的最大活动块数。

此flags参数控制特殊情况的处理方式。有效标志包括

cudaOccupancyDefault：保持与 cudaOccupancyMaxActiveBlocksPerMultiprocessor 相同的默认行为

cudaOccupancyDisableCachingOverride：此标志禁止在全局缓存影响占用率的平台上的默认行为。在这些平台上，如果启用了缓存，但每个块的 SM 资源使用会导致零占用率，则占用率计算器将计算占用率，就好像禁用了缓存一样。设置此标志使占用率计算器在这种情况下返回 0。有关此功能的更多信息，请参阅 Maxwell 调优指南的“统一 L1/纹理缓存”部分。

注意

请注意，此函数也可能返回来自先前异步启动的错误代码。
请注意，如果此调用尝试初始化内部 CUDA RT 状态，此函数也可能返回 cudaErrorInitializationError、cudaErrorInsufficientDriver 或 cudaErrorNoDevice。
请注意，如 cudaStreamAddCallback 所指定，不得从回调中调用任何 CUDA 函数。cudaErrorNotPermitted 可能会（但不能保证）在这种情况下作为诊断返回。
该 API 也可以与内核 cudaKernel_t 一起使用，方法是使用 cudaLibraryGetKernel() 或 cudaGetKernel 查询句柄，然后将其强制转换为 void* 传递给 API。符号entryFuncAddr传递给 cudaGetKernel 的符号应是使用相同 CUDA 运行时实例注册的符号。
传递属于不同运行时实例的符号将导致未定义的行为。唯一可以可靠地传递给不同运行时实例的类型是 cudaKernel_t

另请参阅

cudaOccupancyMaxActiveBlocksPerMultiprocessor, cudaOccupancyMaxPotentialBlockSize ( C++ API), cudaOccupancyMaxPotentialBlockSizeWithFlags ( C++ API), cudaOccupancyMaxPotentialBlockSizeVariableSMem ( C++ API), cudaOccupancyMaxPotentialBlockSizeVariableSMemWithFlags ( C++ API), cudaOccupancyAvailableDynamicSMemPerBlock (C++ API), cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags

__host__ cudaError_t cudaOccupancyMaxActiveClusters ( int* numClusters, const void* func, const cudaLaunchConfig_t* launchConfig )

给定内核函数（funcfunc）和启动配置（config），返回目标设备上可以共存的最大簇数量，结果存于.

参数

*numClusters: 参数
func: numClusters
- 返回的目标设备上可以共存的最大簇数量

返回值

func

描述

- 用于计算最大簇数量的内核函数

launchConfig

注意

请注意，此函数也可能返回来自先前异步启动的错误代码。
请注意，如果此调用尝试初始化内部 CUDA RT 状态，此函数也可能返回 cudaErrorInitializationError、cudaErrorInsufficientDriver 或 cudaErrorNoDevice。
请注意，如 cudaStreamAddCallback 所指定，不得从回调中调用任何 CUDA 函数。cudaErrorNotPermitted 可能会（但不能保证）在这种情况下作为诊断返回。
该 API 也可以与内核 cudaKernel_t 一起使用，方法是使用 cudaLibraryGetKernel() 或 cudaGetKernel 查询句柄，然后将其强制转换为 void* 传递给 API。符号entryFuncAddr传递给 cudaGetKernel 的符号应是使用相同 CUDA 运行时实例注册的符号。
传递属于不同运行时实例的符号将导致未定义的行为。唯一可以可靠地传递给不同运行时实例的类型是 cudaKernel_t

另请参阅

- 启动配置

返回值

给定内核函数（funcfunc）和启动配置（cudaSuccess, cudaErrorInvalidDeviceFunction, cudaErrorInvalidValue, cudaErrorInvalidClusterSize, cudaErrorUnknown,如果函数已设置所需的簇大小（请参阅 cudaFuncGetAttributes），则 config 中的簇大小必须未指定或与所需大小匹配。如果没有所需大小，则必须在 config 中指定簇大小，否则函数将返回错误。.

参数

请注意，内核函数的各种属性可能会影响占用率计算。运行时环境可能会影响硬件调度簇的方式，因此无法保证计算出的占用率可以实现。: 另请参阅 cudaFuncGetAttributes cudaOccupancyMaxActiveClusters (C++ API), cuOccupancyMaxActiveClusters
func: __host__ cudaError_t cudaOccupancyMaxPotentialClusterSize ( int* clusterSize, const void* func, const cudaLaunchConfig_t* launchConfig )
- 返回的目标设备上可以共存的最大簇数量

返回值

给定内核函数（

描述

func）和启动配置（），返回最大簇大小，结果存于如果函数已设置所需的簇大小（请参阅 cudaFuncGetAttributes），则 config 中的簇大小必须未指定或与所需大小匹配。如果没有所需大小，则必须在 config 中指定簇大小，否则函数将返回错误。*clusterSize

参数

clusterSize

注意

请注意，此函数也可能返回来自先前异步启动的错误代码。
请注意，如果此调用尝试初始化内部 CUDA RT 状态，此函数也可能返回 cudaErrorInitializationError、cudaErrorInsufficientDriver 或 cudaErrorNoDevice。
请注意，如 cudaStreamAddCallback 所指定，不得从回调中调用任何 CUDA 函数。cudaErrorNotPermitted 可能会（但不能保证）在这种情况下作为诊断返回。
该 API 也可以与内核 cudaKernel_t 一起使用，方法是使用 cudaLibraryGetKernel() 或 cudaGetKernel 查询句柄，然后将其强制转换为 void* 传递给 API。符号entryFuncAddr传递给 cudaGetKernel 的符号应是使用相同 CUDA 运行时实例注册的符号。
传递属于不同运行时实例的符号将导致未定义的行为。唯一可以可靠地传递给不同运行时实例的类型是 cudaKernel_t

另请参阅

- 返回的给定内核函数和启动配置可以启动的最大簇大小