6.8. 执行控制

本节介绍 CUDA 运行时应用程序编程接口的执行控制功能。

某些函数具有重载的 C++ API 模板版本，在 C++ API 例程 模块中单独记录。

函数

__host__ ​ __device__ ​cudaError_t cudaFuncGetAttributes ( cudaFuncAttributes* attr, const void* func )
查找给定函数的属性。
__host__ ​cudaError_t cudaFuncGetName ( const char** name, const void* func )
返回设备入口函数指针的函数名称。
__host__ ​cudaError_t cudaFuncGetParamInfo ( const void* func, size_t paramIndex, size_t* paramOffset, size_t* paramSize )
返回内核参数在设备端参数布局中的偏移量和大小。
__host__ ​cudaError_t cudaFuncSetAttribute ( const void* func, cudaFuncAttribute attr, int  value )
设置给定函数的属性。
__host__ ​cudaError_t cudaFuncSetCacheConfig ( const void* func, cudaFuncCache cacheConfig )
为设备函数设置首选缓存配置。
__device__ ​ void* cudaGetParameterBuffer ( size_t alignment, size_t size )
获取参数缓冲区。
__device__ ​ void cudaGridDependencySynchronize ( void )
程序化网格依赖同步。
__host__ ​cudaError_t cudaLaunchCooperativeKernel ( const void* func, dim3 gridDim, dim3 blockDim, void** args, size_t sharedMem, cudaStream_t stream )
启动设备函数，其中线程块可以在执行时进行协作和同步。
__host__ ​cudaError_t cudaLaunchCooperativeKernelMultiDevice ( cudaLaunchParams* launchParamsList, unsigned int  numDevices, unsigned int  flags = 0 )
在多个设备上启动设备函数，其中线程块可以在执行时进行协作和同步。
__device__ ​cudaError_t cudaLaunchDevice ( void* func, void* parameterBuffer, dim3 gridDimension, dim3 blockDimension, unsigned int  sharedMemSize, cudaStream_t stream )
启动指定的内核。
__host__ ​cudaError_t cudaLaunchHostFunc ( cudaStream_t stream, cudaHostFn_t fn, void* userData )
在流中排队主机函数调用。
__host__ ​cudaError_t cudaLaunchKernel ( const void* func, dim3 gridDim, dim3 blockDim, void** args, size_t sharedMem, cudaStream_t stream )
启动设备函数。
__host__ ​cudaError_t cudaLaunchKernelExC ( const cudaLaunchConfig_t* config, const void* func, void** args )
使用启动时配置启动 CUDA 函数。
__host__ ​cudaError_t cudaSetDoubleForDevice ( double* d )
转换双精度参数以在设备上执行。
__host__ ​cudaError_t cudaSetDoubleForHost ( double* d )
转换在设备上执行后的双精度参数。
__device__ ​ void cudaTriggerProgrammaticLaunchCompletion ( void )
程序化依赖触发器。

函数

__host__ ​ __device__ ​cudaError_t cudaFuncGetAttributes ( cudaFuncAttributes* attr, const void* func )
查找给定函数的属性。
参数
attr
- 返回指向函数属性的指针
func
- 设备函数符号
返回值

cudaSuccess, cudaErrorInvalidDeviceFunction

描述

此函数获取通过以下方式指定的函数的属性func. func是设备函数符号，必须声明为__global__函数。获取的属性放置在attr中。如果指定的函数不存在，则假定它是一个 cudaKernel_t 并按原样使用。对于模板函数，请按以下方式传递函数符号：func_name<template_arg_0,...,template_arg_N>

请注意，某些函数属性（例如 maxThreadsPerBlock）可能会因当前使用的设备而异。

注意

  • 请注意，此函数也可能返回先前异步启动的错误代码。

  • 使用字符串命名函数作为func参数在 CUDA 4.1 中已弃用，并在 CUDA 5.0 中移除。

  • 请注意，如果此调用尝试初始化内部 CUDA RT 状态，此函数也可能返回 cudaErrorInitializationError、cudaErrorInsufficientDriver 或 cudaErrorNoDevice。

  • 请注意，根据 cudaStreamAddCallback 的规定，不得从回调中调用任何 CUDA 函数。cudaErrorNotPermitted 可能会（但不能保证）作为诊断在这种情况下返回。

  • API 也可以与内核 cudaKernel_t 一起使用，方法是使用 cudaLibraryGetKernel() 或 cudaGetKernel 查询句柄，然后通过强制转换为 void* 将其传递给 API。符号entryFuncAddr传递给 cudaGetKernel 的符号应该是使用相同 CUDA 运行时实例注册的符号。

  • 传递属于不同运行时实例的符号将导致未定义的行为。唯一可以可靠地传递到不同运行时实例的类型是 cudaKernel_t

另请参阅

cudaFuncSetCacheConfig (C API), cudaFuncGetAttributes (C++ API), cudaLaunchKernel (C API), cuFuncGetAttribute

__host__ ​cudaError_t cudaFuncGetName ( const char** name, const void* func )
返回设备入口函数指针的函数名称。
参数
name
- 返回的函数名称
func
- 要检索名称的函数指针
返回值

cudaSuccess, cudaErrorInvalidValue, cudaErrorInvalidDeviceFunction

描述

在**name中返回与符号func关联的函数名称。函数名称以空字符结尾的字符串形式返回。如果函数未声明为具有 C 链接，则此 API 可能会返回损坏的名称。如果**name为 NULL，则返回 cudaErrorInvalidValue。如果func不是设备入口函数，则假定它是一个 cudaKernel_t 并按原样使用。

注意

cudaFuncGetName (C++ API)

__host__ ​cudaError_t cudaFuncGetParamInfo ( const void* func, size_t paramIndex, size_t* paramOffset, size_t* paramSize )
返回内核参数在设备端参数布局中的偏移量和大小。
参数
func
- 要查询的函数
paramIndex
- 要查询的参数索引
paramOffset
- 参数所在的设备端参数布局中的偏移量
paramSize
- 参数在设备端参数布局中的大小（字节）
返回值

CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE,

描述

查询paramIndex中func 的参数列表中的参数，并通过paramOffset和paramSize. paramOffset返回参数信息，返回参数在设备端参数布局中的偏移量。paramSize返回参数的大小（以字节为单位）。此信息可用于通过 cudaGraphKernelNodeSetParam() 和 cudaGraphKernelNodeUpdatesApply() 从设备更新内核节点参数。paramIndex必须小于func所接受的参数数量。

注意

  • 请注意，此函数也可能返回先前异步启动的错误代码。

  • API 也可以与内核 cudaKernel_t 一起使用，方法是使用 cudaLibraryGetKernel() 或 cudaGetKernel 查询句柄，然后通过强制转换为 void* 将其传递给 API。符号entryFuncAddr传递给 cudaGetKernel 的符号应该是使用相同 CUDA 运行时实例注册的符号。

  • 传递属于不同运行时实例的符号将导致未定义的行为。唯一可以可靠地传递到不同运行时实例的类型是 cudaKernel_t

__host__ ​cudaError_t cudaFuncSetAttribute ( const void* func, cudaFuncAttribute attr, int  value )
设置给定函数的属性。
参数
func
- 要获取属性的函数
attr
- 要设置的属性
value
- 要设置的值
返回值

cudaSuccess, cudaErrorInvalidDeviceFunction, cudaErrorInvalidValue

描述

此函数设置通过以下方式指定的函数的属性func。参数func必须是指向在设备上执行的函数的指针。由func指定的参数必须声明为__global__函数。由attr定义的枚举设置为由value定义的值。如果指定的函数不存在，则假定它是一个 cudaKernel_t 并按原样使用。如果指定的属性无法写入，或者值不正确，则返回 cudaErrorInvalidValue。

的有效值attr是

注意

cudaLaunchKernel (C++ API), cudaFuncSetCacheConfig (C++ API), cudaFuncGetAttributes (C API),

__host__ ​cudaError_t cudaFuncSetCacheConfig ( const void* func, cudaFuncCache cacheConfig )
为设备函数设置首选缓存配置。
参数
func
- 设备函数符号
cacheConfig
- 请求的缓存配置
返回值

cudaSuccess, cudaErrorInvalidDeviceFunction

描述

在 L1 缓存和共享内存使用相同硬件资源的设备上，这通过cacheConfig为通过func指定的函数设置首选缓存配置。这只是一个首选项。运行时将在可能的情况下使用请求的配置，但它可以自由选择不同的配置（如果执行func.

func是设备函数符号，必须声明为__global__函数需要）。如果指定的函数不存在，则返回 cudaErrorInvalidDeviceFunction。对于模板函数，请按以下方式传递函数符号：func_name<template_arg_0,...,template_arg_N>

在 L1 缓存和共享内存大小固定的设备上，此设置不起作用。

使用与最近的首选项设置不同的首选项启动内核可能会插入设备端同步点。

支持的缓存配置包括

注意

另请参阅

cudaFuncSetCacheConfig (C++ API), cudaFuncGetAttributes (C API), cudaLaunchKernel (C API), cuFuncSetCacheConfig

__device__ ​ void* cudaGetParameterBuffer ( size_t alignment, size_t size )
获取参数缓冲区。
参数
alignment
- 指定参数缓冲区的对齐要求
size
- 指定大小要求（以字节为单位）
返回值

返回指向已分配 parameterBuffer 的指针

描述

获取一个参数缓冲区，该缓冲区可以填充内核启动的参数。传递给 cudaLaunchDevice 的参数必须通过此函数分配。

这是一个底层 API，只能从并行线程执行 (PTX) 访问。CUDA 用户代码应使用 <<< >>> 来启动内核。

注意

请注意，此函数也可能返回先前异步启动的错误代码。

另请参阅

cudaLaunchDevice

__device__ ​ void cudaGridDependencySynchronize ( void ) [inline]
程序化网格依赖同步。
描述

此设备函数将阻塞线程，直到所有直接网格依赖项完成。此 API 旨在与程序化/启动事件/依赖项结合使用。有关更多信息，请参阅 cudaLaunchAttributeID::cudaLaunchAttributeProgrammaticStreamSerialization 和 cudaLaunchAttributeID::cudaLaunchAttributeProgrammaticEvent。

__host__ ​cudaError_t cudaLaunchCooperativeKernel ( const void* func, dim3 gridDim, dim3 blockDim, void** args, size_t sharedMem, cudaStream_t stream )
启动一个设备函数，其中线程块可以在执行时进行协作和同步。
参数
func
- 设备函数符号
gridDim
- 网格维度
blockDim
- 块维度
args
- 参数
sharedMem
- 共享内存
stream
- 流标识符
返回值

cudaSuccess, cudaErrorInvalidDeviceFunction, cudaErrorInvalidConfiguration, cudaErrorLaunchFailure, cudaErrorLaunchTimeout, cudaErrorLaunchOutOfResources, cudaErrorCooperativeLaunchTooLarge, cudaErrorSharedObjectInitFailed

描述

此函数调用内核func在gridDim (gridDim.xgridDim.ygridDim.z) 网格的块中。每个块包含blockDim (blockDim.xblockDim.yblockDim.z) 线程。

调用此内核的设备必须具有设备属性 cudaDevAttrCooperativeLaunch 的非零值。

启动的块总数不能超过每个多处理器最大块数（由 cudaOccupancyMaxActiveBlocksPerMultiprocessor 返回）或 cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags 返回，再乘以设备属性 cudaDevAttrMultiProcessorCount 指定的多处理器数量。

内核不能使用 CUDA 动态并行性。

如果内核有 N 个参数，则args应该指向 N 个指针的数组。每个指针，从args[0]到args[N - 1], 指向将从中复制实际参数的内存区域。

对于模板函数，请按以下方式传递函数符号：func_name<template_arg_0,...,template_arg_N>

sharedMem设置每个线程块可用的动态共享内存量。

stream指定调用关联到的流。

注意

另请参阅

cudaLaunchCooperativeKernel (C++ API), cudaLaunchCooperativeKernelMultiDevice, cuLaunchCooperativeKernel

__host__ ​cudaError_t cudaLaunchCooperativeKernelMultiDevice ( cudaLaunchParams* launchParamsList, unsigned int  numDevices, unsigned int  flags = 0 )
在多个设备上启动设备函数，其中线程块可以在执行时进行协作和同步。
参数
launchParamsList
- 启动参数列表，每个设备一个
numDevices
- 大小为launchParamsList数组
flags
- 用于控制启动行为的标志
返回值

cudaSuccess, cudaErrorInvalidDeviceFunction, cudaErrorInvalidConfiguration, cudaErrorLaunchFailure, cudaErrorLaunchTimeout, cudaErrorLaunchOutOfResources, cudaErrorCooperativeLaunchTooLarge, cudaErrorSharedObjectInitFailed

已弃用

此函数从 CUDA 11.3 版本开始已弃用。

描述

调用在launchParamsList数组中指定的内核，其中数组的每个元素都指定执行单个内核启动所需的所有参数。这些内核可以在执行时进行协作和同步。数组的大小由numDevices.

不能在同一设备上启动两个内核。此多设备启动所针对的所有设备必须是相同的。所有设备都必须具有设备属性 cudaDevAttrCooperativeMultiDeviceLaunch 的非零值。

必须在所有设备上启动相同的内核。请注意，任何 __device__ 或 __constant__ 变量都在每个设备上独立实例化。应用程序有责任确保这些变量被正确初始化和使用。

网格的大小（以块为单位指定）、块本身的大小以及每个线程块使用的共享内存量也必须在所有启动的内核之间匹配。

用于启动这些内核的流必须通过 cudaStreamCreate 或 cudaStreamCreateWithPriority 或 cudaStreamCreateWithPriority 创建。不能使用 NULL 流或 cudaStreamLegacy 或 cudaStreamPerThread。

每个内核启动的块总数不能超过每个多处理器最大块数（由 cudaOccupancyMaxActiveBlocksPerMultiprocessor 返回）或 cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags 返回，再乘以设备属性 cudaDevAttrMultiProcessorCount 指定的多处理器数量。由于每个设备启动的块总数必须在所有设备之间匹配，因此每个设备可以启动的最大块数将受到多处理器数量最少的设备的限制。

内核不能使用 CUDA 动态并行性。

cudaLaunchParams 结构定义如下

‎        struct cudaLaunchParams
              {
                  void *func;
                  dim3 gridDim;
                  dim3 blockDim;
                  void **args;
                  size_t sharedMem;
                  cudaStream_t 
                  stream;
              };
其中

默认情况下，内核在所有指定流中的所有先前工作完成之前，不会在任何 GPU 上开始执行。可以通过指定标志 cudaCooperativeLaunchMultiDeviceNoPreSync 来覆盖此行为。指定此标志后，每个内核将仅等待与该 GPU 对应的流中的先前工作完成，然后才开始执行。

同样，默认情况下，在所有 GPU 上的内核完成之前，在任何指定流中推送的任何后续工作都不会开始执行。可以通过指定标志 cudaCooperativeLaunchMultiDeviceNoPostSync 来覆盖此行为。指定此标志后，在任何指定流中推送的任何后续工作将仅等待在与该流对应的 GPU 上启动的内核完成，然后才开始执行。

注意

另请参阅

cudaLaunchCooperativeKernel (C++ API), cudaLaunchCooperativeKernel, cuLaunchCooperativeKernelMultiDevice

__device__ ​cudaError_t cudaLaunchDevice ( void* func, void* parameterBuffer, dim3 gridDimension, dim3 blockDimension, unsigned int  sharedMemSize, cudaStream_t stream )
启动指定的内核。
参数
func
- 指向要启动的内核的指针
parameterBuffer
- 保存启动内核的参数。parameterBuffer 可以为 NULL。（可选）
gridDimension
- 指定网格维度
blockDimension
- 指定块维度
sharedMemSize
- 指定共享内存大小
stream
- 指定要使用的流
返回值

cudaSuccess, cudaErrorInvalidDevice, cudaErrorLaunchMaxDepthExceeded, cudaErrorInvalidConfiguration, cudaErrorStartupFailure, cudaErrorLaunchPendingCountExceeded, cudaErrorLaunchOutOfResources

描述

使用指定的参数缓冲区启动指定的内核。可以通过调用 cudaGetParameterBuffer() 获取参数缓冲区。

这是一个底层 API，只能从并行线程执行 (PTX) 访问。CUDA 用户代码应使用 <<< >>> 来启动内核。

注意

请注意，此函数也可能返回先前异步启动的错误代码。

有关启动配置和参数布局的详细描述，请参阅 CUDA 编程指南中的“执行配置”和“参数缓冲区布局”部分。

另请参阅

cudaGetParameterBuffer

__host__ ​cudaError_t cudaLaunchHostFunc ( cudaStream_t stream, cudaHostFn_t fn, void* userData )
在流中入队一个主机函数调用。
参数
stream
fn
- 在先前的流操作完成后要调用的函数
userData
- 要传递给函数的、用户指定的数据
返回值

cudaSuccess, cudaErrorInvalidResourceHandle, cudaErrorInvalidValue, cudaErrorNotSupported

描述

在流中入队一个主机函数以运行。该函数将在当前入队的工作之后被调用，并将阻止在其之后添加的工作。

主机函数不得进行任何 CUDA API 调用。尝试使用 CUDA API 可能会导致 cudaErrorNotPermitted 错误，但这并非强制要求。主机函数不得执行任何可能依赖于尚未强制要求提前运行的未完成 CUDA 工作的同步。没有强制顺序的主机函数（例如在独立流中）以未定义的顺序执行，并且可能会被序列化。

为了统一内存的目的，执行提供了一些保证：

  • 在函数执行期间，该流被视为空闲。因此，例如，该函数始终可以使用附加到其入队流的内存。

  • 函数执行的开始与同步在函数紧邻之前记录在同一流中的事件具有相同的效果。因此，它同步了在函数之前“加入”的流。

  • 向任何流添加设备工作不会使该流变为活动状态，直到所有先前的宿主函数和流回调都已执行完毕。因此，例如，即使已将工作添加到另一个流中，函数也可能使用全局附加内存，如果该工作已通过事件在函数调用之后排序。

  • 函数的完成不会导致流变为活动状态，除非如上所述。如果函数之后没有设备工作，则该流将保持空闲，并且在连续的宿主函数或流回调之间（没有设备工作的情况下）也将保持空闲。因此，例如，可以通过从流末尾的宿主函数发出信号来完成流同步。

请注意，与 cuStreamAddCallback 相比，如果在 CUDA 上下文中发生错误，则不会调用该函数。

注意

另请参阅

cudaStreamCreate, cudaStreamQuery, cudaStreamSynchronize, cudaStreamWaitEvent, cudaStreamDestroy, cudaMallocManaged, cudaStreamAttachMemAsync, cudaStreamAddCallback, cuLaunchHostFunc

__host__ ​cudaError_t cudaLaunchKernel ( const void* func, dim3 gridDim, dim3 blockDim, void** args, size_t sharedMem, cudaStream_t stream )
启动一个设备函数。
参数
func
- 设备函数符号
gridDim
- 网格维度
blockDim
- 块维度
args
- 参数
sharedMem
- 共享内存
stream
- 流标识符
返回值

cudaSuccess, cudaErrorInvalidDeviceFunction, cudaErrorInvalidConfiguration, cudaErrorLaunchFailure, cudaErrorLaunchTimeout, cudaErrorLaunchOutOfResources, cudaErrorSharedObjectInitFailed, cudaErrorInvalidPtx, cudaErrorUnsupportedPtxVersion, cudaErrorNoKernelImageForDevice, cudaErrorJitCompilerNotFound, cudaErrorJitCompilationDisabled

描述

此函数调用内核func在gridDim (gridDim.xgridDim.ygridDim.z) 网格的块中。每个块包含blockDim (blockDim.xblockDim.yblockDim.z) 线程。

如果内核有 N 个参数，则args应该指向 N 个指针的数组。每个指针，从args[0]到args[N - 1], 指向将从中复制实际参数的内存区域。

对于模板函数，请按以下方式传递函数符号：func_name<template_arg_0,...,template_arg_N>

sharedMem设置每个线程块可用的动态共享内存量。

stream指定调用关联到的流。

注意

另请参阅

cudaLaunchKernel (C++ API), cuLaunchKernel

__host__ ​cudaError_t cudaLaunchKernelExC ( const cudaLaunchConfig_t* config, const void* func, void** args )
使用启动时配置启动 CUDA 函数。
参数
config
- 启动配置
func
- 要启动的内核
args
- 指向内核参数的指针数组
返回值

cudaSuccess, cudaErrorInvalidDeviceFunction, cudaErrorInvalidConfiguration, cudaErrorLaunchFailure, cudaErrorLaunchTimeout, cudaErrorLaunchOutOfResources, cudaErrorSharedObjectInitFailed, cudaErrorInvalidPtx, cudaErrorUnsupportedPtxVersion, cudaErrorNoKernelImageForDevice, cudaErrorJitCompilerNotFound, cudaErrorJitCompilationDisabled

描述

请注意，功能等效的可变参数模板 cudaLaunchKernelEx 适用于 C++11 及更高版本。

调用内核func在config->gridDim (config->gridDim.xconfig->gridDim.yconfig->gridDim.z) 网格的块中。每个块包含config->blockDim (config->blockDim.xconfig->blockDim.yconfig->blockDim.z) 线程。

config->dynamicSmemBytes设置每个线程块可用的动态共享内存量。

config->stream指定调用关联到的流。

除了网格和块维度、动态共享内存大小和流之外的配置，还可以通过以下两个字段提供：config

config->attrs是一个数组config->numAttrs连续的 cudaLaunchAttribute 元素。如果以下情况，则不考虑此指针的值：config->numAttrs为零。但是，在这种情况下，建议将指针设置为 NULL。config->numAttrs是填充第一个位置的属性数量config->numAttrs的位置config->attrs数组。

如果内核有 N 个参数，则args应该指向 N 个指针的数组。每个指针，从args[0]到args[N - 1], 指向将从中复制实际参数的内存区域。

注意：此函数如此命名是为了避免无意中调用模板版本，cudaLaunchKernelEx，对于接受单个 void** 或 void* 参数的内核。

注意

另请参阅

cudaLaunchKernelEx(const cudaLaunchConfig_t *config, void (*kernel)(ExpTypes...), ActTypes &&... args) "cudaLaunchKernelEx (C++ API)", cuLaunchKernelEx

__host__ ​cudaError_t cudaSetDoubleForDevice ( double* d )
转换双精度参数以便在设备上执行。
参数
d
- 要转换的双精度浮点数
返回值

cudaSuccess

已弃用

此函数从 CUDA 7.5 版本开始已弃用

描述

转换双精度值d如果设备不支持双精度算术，则转换为内部浮点表示形式。如果设备本身支持双精度，则此函数不执行任何操作。

注意

另请参阅

cudaFuncSetCacheConfig ( C API), cudaFuncGetAttributes ( C API), cudaSetDoubleForHost

__host__ ​cudaError_t cudaSetDoubleForHost ( double* d )
在设备上执行后转换双精度参数。
参数
d
- 要转换的双精度浮点数
返回值

cudaSuccess

已弃用

此函数从 CUDA 7.5 版本开始已弃用

描述

转换双精度值d如果设备不支持双精度算术，则从潜在的内部浮点表示形式转换。如果设备本身支持双精度，则此函数不执行任何操作。

注意

另请参阅

cudaFuncSetCacheConfig ( C API), cudaFuncGetAttributes ( C API), cudaSetDoubleForDevice

__device__ ​ void cudaTriggerProgrammaticLaunchCompletion ( void ) [inline]
程序化依赖触发器。
描述

此设备函数确保程序化启动完成边/事件得到满足。有关更多信息，请参阅 cudaLaunchAttributeID::cudaLaunchAttributeProgrammaticStreamSerialization 和 cudaLaunchAttributeID::cudaLaunchAttributeProgrammaticEvent。只有当网格中的每个 CTA 都已退出或至少调用过此函数一次时，事件/边沿触发才会发生，否则触发会在所有 warp 完成执行后但在网格完成之前自动发生。触发仅启用辅助内核的调度。它本身不提供内存可见性保证。用户可以通过插入正确范围的内存栅栏来强制执行内存可见性。