6.22. 执行控制

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

函数

CUresult cuFuncGetAttribute ( int* pi, CUfunction_attribute attrib, CUfunction hfunc )
返回有关函数的信息。
CUresult cuFuncGetModule ( CUmodule* hmod, CUfunction hfunc )
返回模块句柄。
CUresult cuFuncGetName ( const char** name, CUfunction hfunc )
返回 CUfunction 句柄的函数名称。
CUresult cuFuncGetParamInfo ( CUfunction func, size_t paramIndex, size_t* paramOffset, size_t* paramSize )
返回设备端参数布局中内核参数的偏移量和大小。
CUresult cuFuncIsLoaded ( CUfunctionLoadingState* state, CUfunction function )
返回函数是否已加载。
CUresult cuFuncLoad ( CUfunction function )
加载函数。
CUresult cuFuncSetAttribute ( CUfunction hfunc, CUfunction_attribute attrib, int  value )
设置有关函数的信息。
CUresult cuFuncSetCacheConfig ( CUfunction hfunc, CUfunc_cache config )
为设备函数设置首选缓存配置。
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，其中线程块可以在执行时进行协作和同步。
CUresult cuLaunchCooperativeKernelMultiDevice ( CUDA_LAUNCH_PARAMS* launchParamsList, unsigned int  numDevices, unsigned int  flags )
在多个设备上启动 CUDA 函数，其中线程块可以在执行时进行协作和同步。
CUresult cuLaunchHostFunc ( CUstream hStream, CUhostFn fn, void* userData )
在流中加入主机函数调用队列。
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。
CUresult cuLaunchKernelEx ( const CUlaunchConfig* config, CUfunction f, void** kernelParams, void** extra )
使用启动时配置启动 CUDA 函数 CUfunction 或 CUDA 内核 CUkernel。

函数

CUresult cuFuncGetAttribute ( int* pi, CUfunction_attribute attrib, CUfunction hfunc )
返回有关函数的信息。
参数
pi
- 返回的属性值
attrib
- 请求的属性
hfunc
- 要查询属性的函数
返回值

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中返回属性的整数值attrib在给定的内核上hfunc。 支持的属性包括

除少数例外情况外，也可以在从 cuModuleEnumerateFunctions 返回的未加载函数句柄上查询函数属性。 如果属性需要完全加载的函数，但该函数未加载，则返回 CUDA_ERROR_FUNCTION_NOT_LOADED。 可以使用 cuFuncIsloaded 查询函数的加载状态。 可以调用 cuFuncLoad 以在查询以下需要加载函数的属性之前显式加载函数

注意

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

另请参阅

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

CUresult cuFuncGetModule ( CUmodule* hmod, CUfunction hfunc )
返回模块句柄。
参数
hmod
- 返回的模块句柄
hfunc
- 要检索模块的函数
返回值

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() 将导致未定义的行为。

注意

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

CUresult cuFuncGetName ( const char** name, CUfunction hfunc )
返回 CUfunction 句柄的函数名称。
参数
name
- 返回的函数名称
hfunc
hfunc
返回值

- 要检索名称的函数句柄

描述

在CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE,**namehfunc与函数句柄关联的函数名称CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE,。 函数名称作为空终止字符串返回。 返回的名称仅在函数句柄有效时有效。 如果模块被卸载或重新加载，则必须再次调用 API 才能获取更新的名称。 如果函数未声明为具有 C 链接，则此 API 可能会返回损坏的名称。 如果hfunc或

注意

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

为 NULL，则返回 CUDA_ERROR_INVALID_VALUE。
CUresult cuFuncGetParamInfo ( CUfunction func, size_t paramIndex, size_t* paramOffset, size_t* paramSize )
参数
返回设备端参数布局中内核参数的偏移量和大小。
func
- 要查询的函数
paramIndex
- 要查询的参数索引
paramOffset
- 返回参数所在的设备端参数布局中的偏移量
paramSize
返回值

- 要检索名称的函数句柄

描述

- 可选地返回设备端参数布局中参数的大小- 要查询的函数查询内核参数，位于进入func 的- 要查询的参数索引参数列表，并在- 返回参数所在的设备端参数布局中的偏移量和- 要查询的函数中分别返回参数在设备端参数布局中的偏移量和大小。 此信息可用于通过 cudaGraphKernelNodeSetParam() 和 cudaGraphKernelNodeUpdatesApply() 从设备更新内核节点参数。返回设备端参数布局中内核参数的偏移量和大小。必须小于- 返回参数所在的设备端参数布局中的偏移量采用的参数数量。

注意

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

另请参阅

如果只需要参数偏移量，则可以将

设置为 NULL。
另请参阅
参数
CUresult cuFuncIsLoaded ( CUfunctionLoadingState* state, CUfunction function )
返回函数是否已加载。
state
- 返回的加载状态
返回值

function

描述

在 CUresult cuFuncIsLoaded ( CUfunctionLoadingState* state, CUfunction function ) - 要检查的函数state.

另请参阅

CUDA_SUCCESS, CUDA_ERROR_INVALID_HANDLE, CUDA_ERROR_INVALID_VALUE

返回
cuFuncLoad, cuModuleEnumerateFunctions
参数
state
CUresult cuFuncLoad ( CUfunction function )
返回值

function

描述

加载函数。statefunction

另请参阅

- 要加载的函数

完成
的函数加载。 对完全加载的函数调用此 API 无效。
参数
hfunc
- 要查询属性的函数
attrib
- 请求的属性
cuModuleEnumerateFunctions, cuFuncIsLoaded
CUresult cuFuncSetAttribute ( CUfunction hfunc, CUfunction_attribute attrib, int  value )
返回值

设置有关函数的信息。

描述

valueattrib在给定的内核上hfunc- 要设置的值CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE, CUDA_ERROR_INVALID_VALUE此调用将指定属性

的值设置为

注意

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

另请参阅

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

CUresult cuFuncSetCacheConfig ( CUfunction hfunc, CUfunc_cache config )
为设备函数设置首选缓存配置。
参数
hfunc
- 要为其配置缓存的内核
config
- 请求的缓存配置
返回值

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 缓存和共享内存大小固定的设备上，此设置不起作用。

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

支持的缓存配置包括

注意

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

另请参阅

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
- 要启动的函数 CUfunction 或内核 CUkernel
gridDimX
- 块网格的宽度
gridDimY
- 块网格的高度
gridDimZ
- 块网格的深度
blockDimX
- 每个线程块的 X 维度
blockDimY
- 每个线程块的 Y 维度
blockDimZ
- 每个线程块的 Z 维度
sharedMemBytes
- 每个线程块的动态共享内存大小（以字节为单位）
hStream
- 流标识符
kernelParams
- 指向内核参数的指针数组
返回值

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

描述

在fon agridDimXxgridDimYxgridDimZ个块的网格上调用函数 CUfunction 或内核 CUkernel。blockDimXxblockDimYxblockDimZ每个块包含线程。

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

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

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

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

内核参数必须通过kernelParams指定。f如果kernelParams有 N 个参数，则kernelParams需要是指向 N 个指针的数组。kernelParams[N-1] 中的每一个都必须指向一个内存区域，从中复制实际的内核参数。内核参数的数量及其偏移量和大小不需要指定，因为这些信息直接从内核的镜像中检索。

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

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

请注意，要使用 cuLaunchCooperativeKernel()，内核f必须已使用工具链版本 3.2 或更高版本编译，以便它将包含内核参数信息，或者没有内核参数。如果未满足这些条件中的任何一个，则 cuLaunchCooperativeKernel() 将返回 CUDA_ERROR_INVALID_IMAGE。

请注意，API 也可以用于启动无上下文内核 CUkernel，方法是使用 cuLibraryGetKernel() 查询句柄，然后通过强制转换为 CUfunction 将其传递给 API。 在这里，要在其上启动内核的上下文将从指定的流hStream或在 NULL 流的情况下从当前上下文中获取。

注意

  • 此函数使用标准的 默认流 语义。

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

另请参阅

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

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

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 11.3 开始弃用。

描述

调用

数组中指定的内核，其中数组的每个元素都指定执行单个内核启动所需的所有参数。 这些内核可以在执行时进行协作和同步。 数组的大小由launchParamsListarray where each element of the array specifies all the parameters required to perform a single kernel launch. These kernels can cooperate and synchronize as they execute. The size of the array is specified bynumDevices.

指定。 不能在同一设备上启动两个内核。 此多设备启动所针对的所有设备必须相同。 所有设备都必须具有设备属性 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;
，其中

默认情况下，内核在所有指定的流中的所有先前工作完成之前，不会在任何 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() 启动内核时，先前与 CUDA_LAUNCH_PARAMS::function 关联的块形状、共享大小和参数信息将被覆盖。launchParamsList关联的块形状、共享大小和参数信息将被覆盖。

请注意，要使用 cuLaunchCooperativeKernelMultiDevice()，内核必须已使用工具链版本 3.2 或更高版本编译，以便它将包含内核参数信息，或者没有内核参数。 如果未满足这些条件中的任何一个，则 cuLaunchCooperativeKernelMultiDevice() 将返回 CUDA_ERROR_INVALID_IMAGE。

注意

  • 此函数使用标准的 默认流 语义。

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

另请参阅

cuCtxGetCacheConfig, cuCtxSetCacheConfig, cuFuncSetCacheConfig, cuFuncGetAttribute, cuLaunchCooperativeKernel, cudaLaunchCooperativeKernelMultiDevice

CUresult cuLaunchHostFunc ( CUstream hStream, CUhostFn fn, void* userData )
在流中排队主机函数调用。
参数
hStream
- 要在其中排队函数调用的流
fn
- 一旦先前的流操作完成，要调用的函数
userData
- 要传递给函数的、用户指定的数据
返回值

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 上下文中发生错误，将不会调用该函数。

注意

  • 此函数使用标准的 默认流 语义。

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

另请参阅

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
- 要启动的函数 CUfunction 或内核 CUkernel
gridDimX
- 块网格的宽度
gridDimY
- 块网格的高度
gridDimZ
- 块网格的深度
blockDimX
- 每个线程块的 X 维度
blockDimY
- 每个线程块的 Y 维度
blockDimZ
- 每个线程块的 Z 维度
sharedMemBytes
- 每个线程块的动态共享内存大小（以字节为单位）
hStream
- 流标识符
kernelParams
- 指向内核参数的指针数组
extra
- 额外的选项
返回值

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_SHARED_OBJECT_INIT_FAILED, CUDA_ERROR_NOT_FOUND

描述

在fon agridDimXxgridDimYxgridDimZ个块的网格上调用函数 CUfunction 或内核 CUkernel。blockDimXxblockDimYxblockDimZ每个块包含线程。

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

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

1) 内核参数可以通过kernelParams指定。f如果kernelParams有 N 个参数，则kernelParams需要是指向 N 个指针的数组。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 结尾。

如果内核参数同时使用以下两者指定，则将返回错误 CUDA_ERROR_INVALID_VALUEkernelParams参数列表，并在extra(即两者kernelParams参数列表，并在extra均为非 NULL）。

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

请注意，要使用 cuLaunchKernel()，内核f必须使用工具链版本 3.2 或更高版本编译，以便它包含内核参数信息，或者不包含内核参数。如果未满足任一条件，则 cuLaunchKernel() 将返回 CUDA_ERROR_INVALID_IMAGE。

请注意，API 也可以用于启动无上下文内核 CUkernel，方法是使用 cuLibraryGetKernel() 查询句柄，然后通过强制转换为 CUfunction 将其传递给 API。 在这里，要在其上启动内核的上下文将从指定的流hStream或在 NULL 流的情况下从当前上下文中获取。

注意

  • 此函数使用标准的 默认流 语义。

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

另请参阅

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
- 要启动的配置
f
- 要启动的函数 CUfunction 或内核 CUkernel
kernelParams
- 指向内核参数的指针数组
extra
- 额外的选项
返回值

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

描述

在f使用指定的启动时配置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::attrs 来指定。每个条目都是一个属性 ID 和一个相应的属性值。

The CUlaunchAttribute 结构定义如下

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

使用示例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`，则它也不能超过 `maxBlocksPerCluster`。否则，它必须小于驱动程序可以支持的最大值。否则，允许将此属性设置为物理上无法容纳在任何特定设备上的值。

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

有关详细信息，请参见 cuStreamSetAttribute

有关详细信息，请参见 cuFuncSetAttribute

内核参数到f可以使用与使用 cuLaunchKernel 相同的方式来指定它们。

请注意，该 API 还可以用于启动无上下文内核 CUkernel，方法是使用 cuLibraryGetKernel() 查询句柄，然后通过强制转换为 CUfunction 将其传递给 API。在这里，要启动内核的上下文将从指定的流 CUlaunchConfig::hStream 或 NULL 流情况下的当前上下文中获取。

注意

  • 此函数使用标准的 默认流 语义。

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

另请参阅

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