CUDA Driver API :: CUDA Toolkit Documentation

6.22. 执行控制

本节介绍底层CUDA驱动程序应用程序编程接口的执行控制功能。

Functions

CUresult cuFuncGetAttribute ( int* pi, CUfunction_attribute attrib, CUfunction hfunc ): Returns information about a function.
CUresult cuFuncGetModule ( CUmodule* hmod, CUfunction hfunc ): Returns a module handle.
CUresult cuFuncGetName ( const char** name, CUfunction hfunc ): Returns the function name for a CUfunction handle.
CUresult cuFuncGetParamInfo ( CUfunction func, size_t paramIndex, size_t* paramOffset, size_t* paramSize ): Returns the offset and size of a kernel parameter in the device-side parameter layout.
CUresult cuFuncIsLoaded ( CUfunctionLoadingState* state, CUfunction function ): Returns if the function is loaded.
CUresult cuFuncLoad ( CUfunction function ): Loads a function.
CUresult cuFuncSetAttribute ( CUfunction hfunc, CUfunction_attribute attrib, int value ): Sets information about a function.
CUresult cuFuncSetCacheConfig ( CUfunction hfunc, CUfunc_cache config ): Sets the preferred cache configuration for a device function.
CUresult cuLaunchCooperativeKernel ( CUfunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, CUstream hStream, void** kernelParams ): Launches a CUDA function CUfunction or a CUDA kernel CUkernel where thread blocks can cooperate and synchronize as they execute.
CUresult cuLaunchCooperativeKernelMultiDevice ( CUDA_LAUNCH_PARAMS* launchParamsList, unsigned int numDevices, unsigned int flags ): Launches CUDA functions on multiple devices where thread blocks can cooperate and synchronize as they execute.
CUresult cuLaunchHostFunc ( CUstream hStream, CUhostFn fn, void* userData ): Enqueues a host function call in a stream.
CUresult cuLaunchKernel ( CUfunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, CUstream hStream, void** kernelParams, void** extra ): Launches a CUDA function CUfunction or a CUDA kernel CUkernel.
CUresult cuLaunchKernelEx ( const CUlaunchConfig* config, CUfunction f, void** kernelParams, void** extra ): Launches a CUDA function CUfunction or a CUDA kernel CUkernel with launch-time configuration.

Functions

CUresult cuFuncGetAttribute ( int* pi, CUfunction_attribute attrib, CUfunction hfunc )

返回关于函数的信息。

参数

pi: - Returned attribute value
attrib: - Attribute requested
hfunc: - Function to query attribute of

CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE, CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_FUNCTION_NOT_LOADED

描述

在*pi中返回由hfunc指定的内核上属性attrib的整数值。支持的属性包括：

CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK: 每个块的最大线程数，超过此数值将导致函数启动失败。该数值取决于函数本身以及当前加载该函数的设备。
CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES: 该函数每个块所需的静态分配共享内存大小（以字节为单位）。这不包括用户在运行时请求的动态分配共享内存。
CU_FUNC_ATTRIBUTE_CONST_SIZE_BYTES: 该函数所需的用户分配常量内存大小（以字节为单位）。
CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES: 该函数每个线程使用的本地内存大小（以字节为单位）。
CU_FUNC_ATTRIBUTE_NUM_REGS: 该函数每个线程使用的寄存器数量。
CU_FUNC_ATTRIBUTE_PTX_VERSION: 该函数编译所针对的PTX虚拟架构版本。该值的计算方式为主版本号*10 + 次版本号，例如PTX 1.3版本函数将返回值13。请注意：对于CUDA 3.0之前编译的cubin文件，该属性可能返回未定义值0。
CU_FUNC_ATTRIBUTE_BINARY_VERSION: 函数编译所针对的二进制架构版本。该值为主版本号*10 + 次版本号，因此二进制版本1.3的函数将返回值13。注意：对于未正确编码二进制架构版本的旧版cubin文件，此属性将返回值10。
CU_FUNC_CACHE_MODE_CA: 该属性用于指示函数是否已使用用户指定的编译选项"-Xptxas --dlcm=ca"进行编译。
CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES: 动态分配的共享内存的最大大小（以字节为单位）。
CU_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT: 首选共享内存与L1缓存的分配比例，以共享内存总量的百分比表示。
CU_FUNC_ATTRIBUTE_CLUSTER_SIZE_MUST_BE_SET: 如果设置了此属性，则内核启动时必须指定有效的集群大小。
CU_FUNC_ATTRIBUTE_REQUIRED_CLUSTER_WIDTH: 以块为单位的所需集群宽度。
CU_FUNC_ATTRIBUTE_REQUIRED_CLUSTER_HEIGHT: 以块为单位的必需簇高度。
CU_FUNC_ATTRIBUTE_REQUIRED_CLUSTER_DEPTH: 以块为单位的所需集群深度。
CU_FUNC_ATTRIBUTE_NON_PORTABLE_CLUSTER_SIZE_ALLOWED: 表示该函数是否允许以非便携式集群大小启动。1表示允许，0表示禁止。非便携式集群大小可能仅在程序测试过的特定SKU上有效，若在不同硬件平台上运行可能导致启动失败。CUDA API提供了cudaOccupancyMaxActiveClusters来帮助检查当前设备是否支持所需集群大小。便携式集群大小保证在所有高于目标计算能力的硬件上都可用，sm_90架构的便携式集群大小为每集群8个块。未来计算能力版本可能会增加该数值。特定硬件单元可能支持更高的非便携式集群大小。
CU_FUNC_ATTRIBUTE_CLUSTER_SCHEDULING_POLICY_PREFERENCE: 函数的块调度策略。值类型为CUclusterSchedulingPolicy。

除少数例外情况外，函数属性也可以通过cuModuleEnumerateFunctions返回的未加载函数句柄进行查询。如果属性需要完全加载的函数但函数尚未加载，则会返回CUDA_ERROR_FUNCTION_NOT_LOADED。可以通过cuFuncIsloaded查询函数的加载状态。在查询以下需要函数已加载的属性之前，可以调用cuFuncLoad显式加载函数：

Note:

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

另请参阅：

cuCtxGetCacheConfig, cuCtxSetCacheConfig, cuFuncSetCacheConfig, cuLaunchKernel, cudaFuncGetAttributes, cudaFuncSetAttribute, cuFuncIsLoaded, cuFuncLoad, cuKernelGetAttribute

CUresult cuFuncGetModule ( CUmodule* hmod, CUfunction hfunc )

返回一个模块句柄。

参数

hmod: - Returned module handle
hfunc: - Function to retrieve module for

CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_NOT_FOUND

描述

在*hmod中返回函数hfunc所在模块的句柄。该模块的生命周期与其加载的上下文生命周期一致，或直到模块被显式卸载为止。

CUDA运行时管理其自身加载到主上下文中的模块。如果此API返回的句柄指向由CUDA运行时加载的模块，则对该模块调用cuModuleUnload()将导致未定义行为。

Note:

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

CUresult cuFuncGetName ( const char** name, CUfunction hfunc )

返回CUfunction句柄对应的函数名称。

参数

name: - The returned name of the function
hfunc: - The function handle to retrieve the name for

CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE,

描述

返回与函数句柄hfunc关联的函数名称到**name中。函数名称以空终止字符串形式返回。返回的名称仅在函数句柄有效时有效。如果模块被卸载或重新加载，必须再次调用API以获取更新后的名称。如果函数未声明为具有C链接，此API可能返回一个混淆的名称。如果**name或hfunc为NULL，则返回CUDA_ERROR_INVALID_VALUE。

Note:

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

CUresult cuFuncGetParamInfo ( CUfunction func, size_t paramIndex, size_t* paramOffset, size_t* paramSize )

返回设备端参数布局中内核参数的偏移量和大小。

参数

func: - The function to query
paramIndex: - The parameter index to query
paramOffset: - Returns the offset into the device-side parameter layout at which the parameter resides
paramSize: - Optionally returns the size of the parameter in the device-side parameter layout

CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE,

描述

查询func参数列表中位于paramIndex处的内核参数，并分别在paramOffset和paramSize中返回该参数在设备端参数布局中的偏移量和大小。此信息可用于通过cudaGraphKernelNodeSetParam()和cudaGraphKernelNodeUpdatesApply()从设备更新内核节点参数。paramIndex必须小于func接受的参数数量。如果仅需要参数偏移量，可以将paramSize设置为NULL。

Note:

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

另请参阅：

cuKernelGetParamInfo

CUresult cuFuncIsLoaded ( CUfunctionLoadingState* state, CUfunction function )

返回函数是否已加载。

参数

state: - returned loading state
function: - the function to check

CUDA_SUCCESS, CUDA_ERROR_INVALID_HANDLE, CUDA_ERROR_INVALID_VALUE

描述

返回function在state中的加载状态。

另请参阅：

cuFuncLoad, cuModuleEnumerateFunctions

CUresult cuFuncLoad ( CUfunction function )

加载一个函数。

参数

function: - the function to load

CUDA_SUCCESS, CUDA_ERROR_INVALID_HANDLE, CUDA_ERROR_INVALID_VALUE

描述

完成对function的函数加载。对一个已完全加载的函数调用此API不会有任何效果。

另请参阅：

cuModuleEnumerateFunctions, cuFuncIsLoaded

CUresult cuFuncSetAttribute ( CUfunction hfunc, CUfunction_attribute attrib, int value )

设置关于函数的信息。

参数

hfunc: - Function to query attribute of
attrib: - Attribute requested
value: - The value to set

CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE, CUDA_ERROR_INVALID_VALUE

描述

此调用将指定内核hfunc上的属性attrib设置为由val指定的整数值。如果属性新值成功设置，此函数将返回CUDA_SUCCESS。如果设置失败，此调用将返回错误。并非所有属性都可以设置值。尝试为只读属性设置值将导致错误(CUDA_ERROR_INVALID_VALUE)

cuFuncSetAttribute 调用支持的属性有：

CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES: 动态分配的共享内存的最大字节数。该值应包含请求的动态分配共享内存的最大大小。此值与函数属性CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES的总和不得超过设备属性CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK_OPTIN。可请求的动态共享内存的最大大小可能因GPU架构而异。
CU_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT: 在L1缓存和共享内存使用相同硬件资源的设备上，此属性设置共享内存的预留偏好比例（占总共享内存的百分比）。参见CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_MULTIPROCESSOR。这仅是一个提示，驱动程序可根据函数执行需要选择不同的比例。
CU_FUNC_ATTRIBUTE_REQUIRED_CLUSTER_WIDTH: 所需的块级集群宽度。宽度、高度和深度值必须全为0或全为正数。集群维度的有效性会在启动时进行检查。如果在编译时设置了该值，则不能在运行时设置。在运行时设置该值将返回CUDA_ERROR_NOT_PERMITTED。
CU_FUNC_ATTRIBUTE_REQUIRED_CLUSTER_HEIGHT: 以块为单位的必需集群高度。宽度、高度和深度值必须全部为0或全部为正数。集群维度的有效性将在启动时检查。如果在编译时设置该值，则无法在运行时设置。在运行时设置将返回CUDA_ERROR_NOT_PERMITTED。
CU_FUNC_ATTRIBUTE_REQUIRED_CLUSTER_DEPTH: 所需的块集群深度。宽度、高度和深度值必须全部为0或全部为正数。集群维度的有效性将在启动时检查。如果在编译时设置该值，则不能在运行时设置。在运行时设置将返回CUDA_ERROR_NOT_PERMITTED。
CU_FUNC_ATTRIBUTE_NON_PORTABLE_CLUSTER_SIZE_ALLOWED: 表示该函数是否可以使用非便携式集群大小启动。1表示允许，0表示禁止。
CU_FUNC_ATTRIBUTE_CLUSTER_SCHEDULING_POLICY_PREFERENCE: 函数的块调度策略。值类型为CUclusterSchedulingPolicy。

Note:

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

另请参阅：

cuCtxGetCacheConfig, cuCtxSetCacheConfig, cuFuncSetCacheConfig, cuLaunchKernel, cudaFuncGetAttributes, cudaFuncSetAttribute, cuKernelSetAttribute

CUresult cuFuncSetCacheConfig ( CUfunction hfunc, CUfunc_cache config )

为设备函数设置首选缓存配置。

参数

hfunc: - Kernel to configure cache for
config: - Requested cache configuration

CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_INVALID_CONTEXT

描述

在使用相同硬件资源的L1缓存和共享内存的设备上，这通过config为设备函数hfunc设置首选的缓存配置。这只是一个偏好设置。驱动程序会尽可能使用请求的配置，但如果需要执行hfunc，也可以自由选择不同的配置。通过cuCtxSetCacheConfig()设置的任何上下文范围内的偏好将被此每个函数的设置覆盖，除非每个函数的设置是CU_FUNC_CACHE_PREFER_NONE。在这种情况下，将使用当前上下文范围内的设置。

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

使用与最近偏好设置不同的偏好启动内核可能会插入一个设备端同步点。

支持的缓存配置包括：

CU_FUNC_CACHE_PREFER_NONE: 对共享内存或L1缓存无偏好（默认设置）
CU_FUNC_CACHE_PREFER_SHARED: 优先分配更大的共享内存和更小的L1缓存
CU_FUNC_CACHE_PREFER_L1: 优先使用更大的L1缓存和更小的共享内存
CU_FUNC_CACHE_PREFER_EQUAL: 倾向于使用大小相等的L1缓存和共享内存

Note:

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

另请参阅：

cuCtxGetCacheConfig, cuCtxSetCacheConfig, cuFuncGetAttribute, cuLaunchKernel, cudaFuncSetCacheConfig, cuKernelSetCacheConfig

CUresult cuLaunchCooperativeKernel ( CUfunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, CUstream hStream, void** kernelParams )

启动一个CUDA函数CUfunction或CUDA内核CUkernel，在执行过程中线程块可以协作和同步。

参数

f: - Function CUfunction or Kernel CUkernel to launch
gridDimX: - Width of grid in blocks
gridDimY: - Height of grid in blocks
gridDimZ: - Depth of grid in blocks
blockDimX: - X dimension of each thread block
blockDimY: - Y dimension of each thread block
blockDimZ: - Z dimension of each thread block
sharedMemBytes: - Dynamic shared-memory size per thread block in bytes
hStream: - Stream identifier
kernelParams: - Array of pointers to kernel parameters

CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE, CUDA_ERROR_INVALID_IMAGE, CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_LAUNCH_FAILED, CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES, CUDA_ERROR_LAUNCH_TIMEOUT, CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING, CUDA_ERROR_COOPERATIVE_LAUNCH_TOO_LARGE, CUDA_ERROR_SHARED_OBJECT_INIT_FAILED, CUDA_ERROR_NOT_FOUND

描述

调用函数CUfunction或内核CUkernelf在一个gridDimX x gridDimY x gridDimZ的块网格上执行。每个块包含blockDimX x blockDimY x blockDimZ个线程。

sharedMemBytes 设置每个线程块可用的动态共享内存大小。

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

启动的块总数不能超过cuOccupancyMaxActiveBlocksPerMultiprocessor（或cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags）返回的每个多处理器最大块数乘以设备属性CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT指定的多处理器数量。

该内核无法利用CUDA动态并行功能。

内核参数必须通过kernelParams指定。如果函数f有N个参数，那么kernelParams需要是一个包含N个指针的数组。kernelParams[0]到kernelParams[N-1]中的每个指针都必须指向一个内存区域，实际的内核参数将从该区域复制。内核参数的数量及其偏移量和大小不需要特别指定，因为这些信息可以直接从内核映像中获取。

调用cuLaunchCooperativeKernel()会设置持久性函数状态，该状态与通过cuLaunchKernel API设置的函数状态相同

当通过cuLaunchCooperativeKernel()启动内核f时，之前与f关联的块形状、共享大小和参数信息将被覆盖。

请注意，要使用cuLaunchCooperativeKernel()，内核函数f必须满足以下条件之一：使用3.2或更高版本的工具链编译以包含内核参数信息，或者不包含任何内核参数。如果这两个条件都不满足，cuLaunchCooperativeKernel()将返回CUDA_ERROR_INVALID_IMAGE。

请注意，该API也可用于启动无上下文的核函数CUkernel，具体方法是通过cuLibraryGetKernel()查询句柄，然后将其强制转换为CUfunction传递给API。在这种情况下，启动核函数的上下文将取自指定的流hStream，若流为NULL则使用当前上下文。

Note:

此函数使用标准的默认流语义。
请注意，此函数也可能返回之前异步启动的错误代码。

另请参阅：

cuCtxGetCacheConfig, cuCtxSetCacheConfig, cuFuncSetCacheConfig, cuFuncGetAttribute, cuLaunchCooperativeKernelMultiDevice, cudaLaunchCooperativeKernel, cuLibraryGetKernel, cuKernelSetCacheConfig, cuKernelGetAttribute, cuKernelSetAttribute

CUresult cuLaunchCooperativeKernelMultiDevice ( CUDA_LAUNCH_PARAMS* launchParamsList, unsigned int numDevices, unsigned int flags )

在多个设备上启动CUDA函数，使线程块在执行过程中能够协作和同步。

参数

launchParamsList: - List of launch parameters, one per device
numDevices: - Size of the launchParamsList array
flags: - Flags to control launch behavior

已弃用

该函数自CUDA 11.3起已弃用。

描述

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

无法在同一设备上启动两个内核。此多设备启动所针对的所有设备必须完全相同。所有设备的设备属性CU_DEVICE_ATTRIBUTE_COOPERATIVE_MULTI_DEVICE_LAUNCH必须具有非零值。

所有启动的内核在编译代码方面必须完全相同。请注意，每个设备上启动的内核所属模块中的任何__device__、__constant__或__managed__变量，都会在每个设备上独立实例化。应用程序需负责确保这些变量被正确初始化和使用。

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

用于启动这些内核的流必须通过cuStreamCreate或cuStreamCreateWithPriority创建。不能使用NULL流、CU_STREAM_LEGACY或CU_STREAM_PER_THREAD。

每个内核启动的块总数不能超过cuOccupancyMaxActiveBlocksPerMultiprocessor（或cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags）返回的每个多处理器最大活动块数乘以设备属性CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT指定的多处理器数量。由于每个设备启动的块总数必须在所有设备上保持一致，因此每个设备可启动的最大块数将受限于具有最少多处理器数量的设备。

内核无法利用CUDA动态并行功能。

CUDA_LAUNCH_PARAMS 结构体定义如下：

‎        typedef struct CUDA_LAUNCH_PARAMS_st
              {
                  CUfunction function;
                  unsigned int gridDimX;
                  unsigned int gridDimY;
                  unsigned int gridDimZ;
                  unsigned int blockDimX;
                  unsigned int blockDimY;
                  unsigned int blockDimZ;
                  unsigned int sharedMemBytes;
                  CUstream hStream;
                  void **kernelParams;
              } CUDA_LAUNCH_PARAMS;

where:

CUDA_LAUNCH_PARAMS::function 指定要启动的内核函数。所有函数在编译代码层面必须完全相同。注意您也可以通过cuLibraryGetKernel()查询句柄并转换为CUfunction来指定无上下文的内核CUkernel。这种情况下，内核启动的上下文将从指定的流CUDA_LAUNCH_PARAMS::hStream中获取。
CUDA_LAUNCH_PARAMS::gridDimX 表示网格的宽度（以块为单位）。所有启动的内核必须保持该值一致。
CUDA_LAUNCH_PARAMS::gridDimY 表示网格的高度（以块为单位）。所有启动的内核必须保持该值一致。
CUDA_LAUNCH_PARAMS::gridDimZ 表示网格在块中的深度。所有启动的内核必须保持一致。
CUDA_LAUNCH_PARAMS::blockDimX 是每个线程块的X维度。所有启动的内核必须保持一致。
CUDA_LAUNCH_PARAMS::blockDimX 是每个线程块的Y维度。所有启动的内核必须保持一致。
CUDA_LAUNCH_PARAMS::blockDimZ 是每个线程块的Z维度。所有启动的内核必须保持一致。
CUDA_LAUNCH_PARAMS::sharedMemBytes 表示每个线程块的动态共享内存大小（以字节为单位）。所有启动的内核必须保持该值一致。
CUDA_LAUNCH_PARAMS::hStream 是用于执行启动操作的流句柄。该流不能是NULL流或CU_STREAM_LEGACY或CU_STREAM_PER_THREAD。与此流关联的CUDA上下文必须与CUDA_LAUNCH_PARAMS::function关联的上下文匹配。
CUDA_LAUNCH_PARAMS::kernelParams 是指向内核参数的指针数组。如果 CUDA_LAUNCH_PARAMS::function 有N个参数，那么 CUDA_LAUNCH_PARAMS::kernelParams 需要是一个包含N个指针的数组。CUDA_LAUNCH_PARAMS::kernelParams[0] 到 CUDA_LAUNCH_PARAMS::kernelParams[N-1] 中的每个指针都必须指向一个内存区域，实际的内核参数将从该区域复制。内核参数的数量及其偏移量和大小不需要特别指定，因为这些信息会直接从内核映像中获取。

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

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

调用cuLaunchCooperativeKernelMultiDevice()会设置持久性函数状态，该状态与通过cuLaunchKernel API为launchParamsList中的每个元素单独调用时设置的函数状态相同。

当通过cuLaunchCooperativeKernelMultiDevice()启动内核时，launchParamsList中每个CUDA_LAUNCH_PARAMS::function关联的原有块形状、共享大小和参数信息将被覆盖。

请注意，要使用cuLaunchCooperativeKernelMultiDevice()，内核必须使用3.2或更高版本的工具链编译，以便包含内核参数信息，或者没有内核参数。如果这两个条件都不满足，那么cuLaunchCooperativeKernelMultiDevice()将返回CUDA_ERROR_INVALID_IMAGE。

Note:

此函数使用标准的默认流语义。
请注意，此函数也可能返回之前异步启动的错误代码。

另请参阅：

cuCtxGetCacheConfig, cuCtxSetCacheConfig, cuFuncSetCacheConfig, cuFuncGetAttribute, cuLaunchCooperativeKernel, cudaLaunchCooperativeKernelMultiDevice

CUresult cuLaunchHostFunc ( CUstream hStream, CUhostFn fn, void* userData )

在流中排队一个主机函数调用。

参数

hStream: - Stream to enqueue function call in
fn: - The function to call once preceding stream operations are complete
userData: - User-specified data to be passed to the function

CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE, CUDA_ERROR_NOT_SUPPORTED

描述

将一个主机函数加入流中排队执行。该函数将在当前已排队的工作完成后被调用，并会阻塞之后添加的工作。

主机函数不得进行任何CUDA API调用。尝试使用CUDA API可能会导致CUDA_ERROR_NOT_PERMITTED错误，但这不是必然的。主机函数不得执行任何可能依赖于未完成CUDA工作的同步操作，除非这些工作被明确要求在之前运行。没有明确执行顺序要求的主机函数（例如在独立流中）将以未定义的顺序执行，并可能被串行化。

就统一内存管理而言，执行过程提供以下保证：

在函数执行期间，该流被视为空闲状态。因此，例如，该函数可以始终使用附加到其入队流的内存。
函数开始执行的效果等同于在同一个流中同步记录在函数之前的事件。因此，它会同步那些在函数执行前已被"连接"的流。
在任何流中添加设备工作不会使该流变为活动状态，直到所有先前的主机函数和流回调执行完毕。因此，例如，即使工作已添加到另一个流中，如果该工作通过事件被排序在函数调用之后，函数仍可能使用全局附加内存。
函数的完成不会导致流变为活动状态，除非如上所述。如果没有后续的设备工作，流将保持空闲状态，并且在连续的宿主函数或流回调之间没有设备工作时也将保持空闲。因此，例如，可以通过在流结束时从宿主函数发出信号来完成流同步。

请注意，与cuStreamAddCallback不同，如果CUDA上下文中发生错误，该函数将不会被调用。

Note:

此函数使用标准的默认流语义。
请注意，此函数也可能返回之前异步启动的错误代码。

另请参阅：

cuStreamCreate, cuStreamQuery, cuStreamSynchronize, cuStreamWaitEvent, cuStreamDestroy, cuMemAllocManaged, cuStreamAttachMemAsync, cuStreamAddCallback

CUresult cuLaunchKernel ( CUfunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, CUstream hStream, void** kernelParams, void** extra )

启动一个CUDA函数CUfunction或CUDA内核CUkernel。

参数

f: - Function CUfunction or Kernel CUkernel to launch
gridDimX: - Width of grid in blocks
gridDimY: - Height of grid in blocks
gridDimZ: - Depth of grid in blocks
blockDimX: - X dimension of each thread block
blockDimY: - Y dimension of each thread block
blockDimZ: - Z dimension of each thread block
sharedMemBytes: - Dynamic shared-memory size per thread block in bytes
hStream: - Stream identifier
kernelParams: - Array of pointers to kernel parameters
extra: - Extra options

描述

调用函数CUfunction或内核CUkernelf在一个gridDimX x gridDimY x gridDimZ的块网格上执行。每个块包含blockDimX x blockDimY x blockDimZ个线程。

sharedMemBytes 设置每个线程块可用的动态共享内存大小。

可以以下两种方式之一指定内核参数到f：

1) 内核参数可以通过kernelParams指定。如果函数f有N个参数，那么kernelParams需要是一个包含N个指针的数组。kernelParams[0]到kernelParams[N-1]中的每个指针都必须指向一个内存区域，实际的内核参数将从该区域复制。内核参数的数量及其偏移量和大小不需要特别指定，因为这些信息可以直接从内核映像中获取。

2) 内核参数也可以由应用程序打包到单个缓冲区中，通过extra参数传入。这种方式需要应用程序自行了解每个内核参数在缓冲区中的大小和对齐/填充情况。以下是通过extra参数实现这种方式的示例：

‎    size_t argBufferSize;
          char argBuffer[256];
      
          // populate argBuffer and argBufferSize
      
          void *config[] = {
              CU_LAUNCH_PARAM_BUFFER_POINTER, argBuffer,
              CU_LAUNCH_PARAM_BUFFER_SIZE,    &argBufferSize,
              CU_LAUNCH_PARAM_END
          };
          status = cuLaunchKernel(f, gx, gy, gz, bx, by, bz, sh, s, NULL, config);

extra参数的存在是为了让cuLaunchKernel能够接收一些不太常用的额外参数。extra指定了一个包含额外设置名称及其对应值的列表。每个额外设置名称后面必须紧跟其对应的值。该列表必须以NULL或CU_LAUNCH_PARAM_END作为结束标志。

CU_LAUNCH_PARAM_END，用于表示extra数组的结束；
CU_LAUNCH_PARAM_BUFFER_POINTER，指定extra中的下一个值将是指向包含内核f启动所需所有参数的缓冲区的指针；
CU_LAUNCH_PARAM_BUFFER_SIZE，用于指定extra中的下一个值将是一个指向size_t的指针，该size_t包含由CU_LAUNCH_PARAM_BUFFER_POINTER指定的缓冲区大小；

如果同时使用kernelParams和extra指定内核参数（即kernelParams和extra均不为NULL），将返回错误CUDA_ERROR_INVALID_VALUE。

调用cuLaunchKernel()会使通过以下已弃用API设置的持久函数状态失效：cuFuncSetBlockShape()、cuFuncSetSharedSize()、cuParamSetSize()、cuParamSeti()、cuParamSetf()、cuParamSetv()。

请注意，要使用cuLaunchKernel()，内核f必须使用工具链版本3.2或更高版本编译，以便包含内核参数信息，或者没有内核参数。如果这两个条件都不满足，那么cuLaunchKernel()将返回CUDA_ERROR_INVALID_IMAGE。

Note:

此函数使用标准的默认流语义。
请注意，此函数也可能返回之前异步启动的错误代码。

另请参阅：

cuCtxGetCacheConfig, cuCtxSetCacheConfig, cuFuncSetCacheConfig, cuFuncGetAttribute, cudaLaunchKernel, cuLibraryGetKernel, cuKernelSetCacheConfig, cuKernelGetAttribute, cuKernelSetAttribute

CUresult cuLaunchKernelEx ( const CUlaunchConfig* config, CUfunction f, void** kernelParams, void** extra )

启动一个CUDA函数CUfunction或CUDA内核CUkernel，并带有启动时配置。

参数

config: - Config to launch
f: - Function CUfunction or Kernel CUkernel to launch
kernelParams: - Array of pointers to kernel parameters
extra: - Extra options

描述

调用函数CUfunction或内核CUkernelf，并使用指定的启动配置config。

CUlaunchConfig 结构体定义为：

‎       typedef struct CUlaunchConfig_st {
           unsigned int gridDimX;
           unsigned int gridDimY;
           unsigned int gridDimZ;
           unsigned int blockDimX;
           unsigned int blockDimY;
           unsigned int blockDimZ;
           unsigned int sharedMemBytes;
           CUstream hStream;
           CUlaunchAttribute *attrs;
           unsigned int numAttrs;
       } CUlaunchConfig;

其中：

CUlaunchConfig::gridDimX 表示网格的宽度（以块为单位）。
CUlaunchConfig::gridDimY 表示网格的高度（以块为单位）。
CUlaunchConfig::gridDimZ 表示网格在块中的深度。
CUlaunchConfig::blockDimX 是每个线程块的X维度。
CUlaunchConfig::blockDimX 是每个线程块的Y维度。
CUlaunchConfig::blockDimZ 是每个线程块的Z维度。
CUlaunchConfig::sharedMemBytes 是每个线程块的动态共享内存大小，单位为字节。
CUlaunchConfig::hStream 是执行启动操作的流句柄。与此流关联的CUDA上下文必须与函数f关联的上下文匹配。
CUlaunchConfig::attrs 是一个由 CUlaunchConfig::numAttrs 个连续的 CUlaunchAttribute 元素组成的数组。如果 CUlaunchConfig::numAttrs 为零，则不考虑此指针的值。但在这种情况下，建议将指针设置为 NULL。
CUlaunchConfig::numAttrs 表示填充CUlaunchConfig::attrs数组前CUlaunchConfig::numAttrs个位置的属性数量。

启动时配置通过向CUlaunchConfig::attrs添加条目来指定。每个条目包含一个属性ID和对应的属性值。

CUlaunchAttribute 结构体的定义如下：

‎       typedef struct CUlaunchAttribute_st {
           CUlaunchAttributeID id;
           CUlaunchAttributeValue value;
       } CUlaunchAttribute;

where:

CUlaunchAttribute::id 是一个唯一标识该属性的枚举值。
CUlaunchAttribute::value 是一个包含属性值的联合体。

使用config参数的示例：

‎       CUlaunchAttribute coopAttr = {.id = CU_LAUNCH_ATTRIBUTE_COOPERATIVE,
                                     .value = 1};
       CUlaunchConfig config = {... // set block and grid dimensions
                              .attrs = &coopAttr,
                              .numAttrs = 1};
      
       cuLaunchKernelEx(&config, kernel, NULL, NULL);

CUlaunchAttributeID 枚举的定义如下：

‎       typedef enum CUlaunchAttributeID_enum {
           CU_LAUNCH_ATTRIBUTE_IGNORE = 0,
           CU_LAUNCH_ATTRIBUTE_ACCESS_POLICY_WINDOW   = 1,
           CU_LAUNCH_ATTRIBUTE_COOPERATIVE            = 2,
           CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY = 3,
           CU_LAUNCH_ATTRIBUTE_CLUSTER_DIMENSION                    = 4,
           CU_LAUNCH_ATTRIBUTE_CLUSTER_SCHEDULING_POLICY_PREFERENCE = 5,
           CU_LAUNCH_ATTRIBUTE_PROGRAMMATIC_STREAM_SERIALIZATION    = 6,
           CU_LAUNCH_ATTRIBUTE_PROGRAMMATIC_EVENT                   = 7,
           CU_LAUNCH_ATTRIBUTE_PRIORITY               = 8,
           CU_LAUNCH_ATTRIBUTE_MEM_SYNC_DOMAIN_MAP    = 9,
           CU_LAUNCH_ATTRIBUTE_MEM_SYNC_DOMAIN        = 10,
           CU_LAUNCH_ATTRIBUTE_PREFERRED_CLUSTER_DIMENSION = 11,
           CU_LAUNCH_ATTRIBUTE_LAUNCH_COMPLETION_EVENT = 12,
           CU_LAUNCH_ATTRIBUTE_DEVICE_UPDATABLE_KERNEL_NODE = 13,
       } CUlaunchAttributeID;

以及对应的 CUlaunchAttributeValue 联合体如下：

‎       typedef union CUlaunchAttributeValue_union {
           CUaccessPolicyWindow accessPolicyWindow;
           int cooperative;
           CUsynchronizationPolicy syncPolicy;
           struct {
               unsigned int x;
               unsigned int y;
               unsigned int z;
           } clusterDim;
           CUclusterSchedulingPolicy clusterSchedulingPolicyPreference;
           int programmaticStreamSerializationAllowed;
           struct {
               CUevent event;
               int flags;
               int triggerAtBlockStart;
           } programmaticEvent;
           int priority;
           CUlaunchMemSyncDomainMap memSyncDomainMap;
           CUlaunchMemSyncDomain memSyncDomain;
           struct {
               unsigned int x;
               unsigned int y;
               unsigned int z;
           } preferredClusterDim;
           struct {
               CUevent event;
               int flags;
           } launchCompletionEvent;
           struct {
               int deviceUpdatable;
               CUgraphDeviceNode devNode;
           } deviceUpdatableKernelNode;
       } CUlaunchAttributeValue;

将CU_LAUNCH_ATTRIBUTE_COOPERATIVE设置为非零值会使内核启动成为协作式启动，其使用方式和语义与cuLaunchCooperativeKernel完全相同。

将CU_LAUNCH_ATTRIBUTE_PROGRAMMATIC_STREAM_SERIALIZATION设置为非零值会使内核使用编程方式解决其流依赖关系——如果前一个内核请求重叠执行，则使CUDA运行时有机会允许该网格的执行与流中的前一个内核重叠。

CU_LAUNCH_ATTRIBUTE_PROGRAMMATIC_EVENT 记录一个与内核启动相关联的事件。通过此启动属性记录的事件确保仅在关联内核中的所有块触发该事件后才会触发。一个块可以通过PTX launchdep.release或CUDA内置函数cudaTriggerProgrammaticLaunchCompletion()来触发事件。如果triggerAtBlockStart设置为非0值，还可以在每个块执行开始时插入触发器。请注意，依赖项（包括调用cuEventSynchronize()的CPU线程）不保证在事件释放时立即观察到释放。例如，cuEventSynchronize()可能仅在关联内核完成很久之后才观察到事件触发。此记录类型主要用于在设备任务之间建立程序化依赖关系。提供的事件不能是进程间或互操作事件。该事件必须禁用计时（即创建时设置了CU_EVENT_DISABLE_TIMING标志）。

CU_LAUNCH_ATTRIBUTE_LAUNCH_COMPLETION_EVENT 记录一个与内核启动相关联的事件。理论上，该事件会在内核的所有块开始执行后被触发。目前这是一个尽力而为的机制。如果内核B在内核A的启动完成事件上有依赖关系，B可能会等待A完全执行完毕。或者，B的块也可能在A的所有块开始之前就开始执行，例如：

如果B可以申请A无法获得的执行资源，例如它们运行在不同的GPU上。
如果B的优先级高于A。

如果此类顺序反转可能导致死锁，请谨慎操作。提供的事件不能是进程间或互操作事件。该事件必须禁用计时（即必须设置CU_EVENT_DISABLE_TIMING标志来创建）。

在捕获的启动中将CU_LAUNCH_ATTRIBUTE_DEVICE_UPDATABLE_KERNEL_NODE设置为1会使生成的内核节点可被设备更新。此属性专用于图形，将其传递到非捕获流中的启动会导致错误。不允许传递除0或1以外的值。

成功时，将通过CUlaunchAttributeValue::deviceUpdatableKernelNode::devNode返回一个句柄，该句柄可传递给各种设备端更新函数，以便在另一个内核中更新节点的内核参数。有关可进行的设备更新类型及其相关限制的更多信息，请参阅cudaGraphKernelNodeUpdatesApply。

与常规内核节点相比，支持设备更新的内核节点存在额外限制。首先，无法通过cuGraphDestroyNode从图中移除设备可更新节点。此外，一旦启用此功能，节点将无法退出该模式，任何尝试将属性设置为0的操作都将导致错误。包含一个或多个设备可更新节点的图也不允许多重实例化。

CU_LAUNCH_ATTRIBUTE_PREFERRED_CLUSTER_DIMENSION 允许内核启动指定首选替代集群维度。线程块可以根据此属性指定的维度（分组为"首选替代集群"）或CU_LAUNCH_ATTRIBUTE_CLUSTER_DIMENSION属性指定的维度（分组为"常规集群"）进行分组。"首选替代集群"的集群维度必须是常规集群维度的正整数倍。设备将尽最大努力优先将线程块分组为首选集群而非常规集群。当设备认为必要时（主要是当设备暂时缺乏物理资源来启动更大的首选集群时），设备可能会转而启动常规集群，以尽可能充分利用物理设备资源。

每种类型的集群都会有自己的枚举/坐标设置，就好像网格仅由该类型的集群组成。例如，如果首选替代集群的尺寸是常规集群的两倍，那么可能会同时存在一个索引为(1,0,0)的常规集群和一个索引为(1,0,0)的首选集群。在这个例子中，首选替代集群(1,0,0)会取代常规集群(2,0,0)和(3,0,0)并将它们的区块组合起来。

此属性仅在指定了常规集群维度时生效。首选替代集群维度必须是常规集群维度的正整数倍，并且必须能够整除网格。如果内核的`__launch_bounds__`中设置了`maxBlocksPerCluster`，则其值不得超过该限制。否则，其值必须小于驱动程序支持的最大值。此外，允许将此属性设置为物理上无法适配任何特定设备的值。

其他属性的效果与通过持久化API设置时的效果一致。

参见cuStreamSetAttribute获取

参见cuFuncSetAttribute获取

可以像使用cuLaunchKernel一样的方式指定内核参数到f。

请注意，该API也可用于启动无上下文的核函数CUkernel，具体方法是通过cuLibraryGetKernel()查询句柄，然后将其强制转换为CUfunction传递给API。在这种情况下，启动核函数的上下文将取自指定的流CUlaunchConfig::hStream，若流为NULL则使用当前上下文。

Note:

此函数使用标准的默认流语义。
请注意，此函数也可能返回之前异步启动的错误代码。

另请参阅：

cuCtxGetCacheConfig, cuCtxSetCacheConfig, cuFuncSetCacheConfig, cuFuncGetAttribute, cudaLaunchKernel, cudaLaunchKernelEx, cuLibraryGetKernel, cuKernelSetCacheConfig, cuKernelGetAttribute, cuKernelSetAttribute