Cooperative groups

Cooperative kernel launches

The following host-side functions are used for cooperative kernel launches.

hipError_t hipModuleLaunchCooperativeKernel(hipFunction_t f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, hipStream_t stream, void **kernelParams)

launches kernel f with launch parameters and shared memory on stream with arguments passed to kernelParams, where thread blocks can cooperate and synchronize as they execute

Please note, HIP does not support kernel launch with total work items defined in dimension with size \( gridDim \cdot blockDim \geq 2^{32} \).

Parameters:

• f – [in] Kernel to launch.

• gridDimX – [in] X grid dimension specified as multiple of blockDimX.

• gridDimY – [in] Y grid dimension specified as multiple of blockDimY.

• gridDimZ – [in] Z grid dimension specified as multiple of blockDimZ.

• blockDimX – [in] X block dimension specified in work-items.

• blockDimY – [in] Y block dimension specified in work-items.

• blockDimZ – [in] Z block dimension specified in work-items.

• sharedMemBytes – [in] Amount of dynamic shared memory to allocate for this kernel. The HIP-Clang compiler provides support for extern shared declarations.

• stream – [in] Stream where the kernel should be dispatched. May be 0, in which case the default stream is used with associated synchronization rules.

• kernelParams – [in] A list of kernel arguments.

Returns:

hipSuccess, hipErrorDeinitialized, hipErrorNotInitialized, hipErrorInvalidContext, hipErrorInvalidHandle, hipErrorInvalidImage, hipErrorInvalidValue, hipErrorInvalidConfiguration, hipErrorLaunchFailure, hipErrorLaunchOutOfResources, hipErrorLaunchTimeOut, hipErrorCooperativeLaunchTooLarge, hipErrorSharedObjectInitFailed

hipError_t hipModuleLaunchCooperativeKernelMultiDevice(hipFunctionLaunchParams *launchParamsList, unsigned int numDevices, unsigned int flags)

Launches kernels on multiple devices where thread blocks can cooperate and synchronize as they execute.

Parameters:

• launchParamsList – [in] List of launch parameters, one per device.

• numDevices – [in] Size of the launchParamsList array.

• flags – [in] Flags to control launch behavior.

Returns:

hipSuccess, hipErrorDeinitialized, hipErrorNotInitialized, hipErrorInvalidContext, hipErrorInvalidHandle, hipErrorInvalidImage, hipErrorInvalidValue, hipErrorInvalidConfiguration, hipErrorInvalidResourceHandle, hipErrorLaunchFailure, hipErrorLaunchOutOfResources, hipErrorLaunchTimeOut, hipErrorCooperativeLaunchTooLarge, hipErrorSharedObjectInitFailed

hipError_t hipLaunchCooperativeKernel(const void *f, dim3 gridDim, dim3 blockDimX, void **kernelParams, unsigned int sharedMemBytes, hipStream_t stream)

Launches kernel f with launch parameters and shared memory on stream with arguments passed to kernelparams or extra, where thread blocks can cooperate and synchronize as they execute.

Please note, HIP does not support kernel launch with total work items defined in dimension with size \( gridDim \cdot blockDim \geq 2^{32} \).

Parameters:

• f – [in] - Kernel to launch.

• gridDim – [in] - Grid dimensions specified as multiple of blockDim.

• blockDimX – [in] - Block dimensions specified in work-items

• kernelParams – [in] - Pointer of arguments passed to the kernel. If the kernel has multiple parameters, ‘kernelParams’ should be array of pointers, each points the corresponding argument.

• sharedMemBytes – [in] - Amount of dynamic shared memory to allocate for this kernel. The HIP-Clang compiler provides support for extern shared declarations.

• stream – [in] - Stream where the kernel should be dispatched. May be 0, in which case th default stream is used with associated synchronization rules.

Returns:

hipSuccess, hipErrorNotInitialized, hipErrorInvalidValue, hipErrorCooperativeLaunchTooLarge

hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams *launchParamsList, int numDevices, unsigned int flags)

Launches kernels on multiple devices where thread blocks can cooperate and synchronize as they execute.

Parameters:

• launchParamsList – [in] List of launch parameters, one per device.

• numDevices – [in] Size of the launchParamsList array.

• flags – [in] Flags to control launch behavior.

Returns:

hipSuccess, hipErrorNotInitialized, hipErrorInvalidValue, hipErrorCooperativeLaunchTooLarge

template<class T>
inline hipError_t hipLaunchCooperativeKernel(T f, dim3 gridDim, dim3 blockDim, void **kernelParams, unsigned int sharedMemBytes, hipStream_t stream)

Launches a device function.

Template Parameters:

T – The type of the kernel function.

Parameters:

• f – [in] Kernel function to launch.

• gridDim – [in] Grid dimensions specified as multiple of blockDim.

• blockDim – [in] Block dimensions specified in work-items.

• kernelParams – [in] A list of kernel arguments.

• sharedMemBytes – [in] Amount of dynamic shared memory to allocate for this kernel. The HIP-Clang compiler provides support for extern shared declarations.

• stream – [in] Stream which on the kernel launched.

Returns:

hipSuccess, hipErrorLaunchFailure, hipErrorInvalidValue, hipErrorInvalidResourceHandle

template<class T>
inline hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams *launchParamsList, unsigned int numDevices, unsigned int flags = 0)

Launches kernel function on multiple devices, where thread blocks can cooperate and synchronize on execution.

Parameters:

• launchParamsList – [in] List of kernel launch parameters, one per device.

• numDevices – [in] Size of launchParamsList array.

• flags – [in] Flag to handle launch behavior.

Returns:

hipSuccess, hipErrorLaunchFailure, hipErrorInvalidValue, hipErrorInvalidResourceHandle

Cooperative groups classes

The following cooperative groups classes can be used on the device side.

Warning

doxygenclass: Cannot find class “cooperative_groups::thread_group” in doxygen xml output for project “HIP 7.13.0 Documentation” from directory: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hip/checkouts/develop/projects/hip/docs/doxygen/xml

Warning

doxygenclass: Cannot find class “cooperative_groups::thread_block” in doxygen xml output for project “HIP 7.13.0 Documentation” from directory: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hip/checkouts/develop/projects/hip/docs/doxygen/xml

Warning

doxygenclass: Cannot find class “cooperative_groups::grid_group” in doxygen xml output for project “HIP 7.13.0 Documentation” from directory: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hip/checkouts/develop/projects/hip/docs/doxygen/xml

Warning

doxygenclass: Cannot find class “cooperative_groups::multi_grid_group” in doxygen xml output for project “HIP 7.13.0 Documentation” from directory: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hip/checkouts/develop/projects/hip/docs/doxygen/xml

Warning

doxygenclass: Cannot find class “cooperative_groups::thread_block_tile” in doxygen xml output for project “HIP 7.13.0 Documentation” from directory: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hip/checkouts/develop/projects/hip/docs/doxygen/xml

Warning

doxygenclass: Cannot find class “cooperative_groups::coalesced_group” in doxygen xml output for project “HIP 7.13.0 Documentation” from directory: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hip/checkouts/develop/projects/hip/docs/doxygen/xml

Cooperative groups construct functions

The following functions are used to construct different group-type instances on the device side.

Warning

doxygengroup: Cannot find group “CooperativeGConstruct” in doxygen xml output for project “HIP 7.13.0 Documentation” from directory: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hip/checkouts/develop/projects/hip/docs/doxygen/xml

Cooperative groups exposed API functions

The following functions are the exposed API for different group-type instances on the device side.

Warning

doxygengroup: Cannot find group “CooperativeGAPI” in doxygen xml output for project “HIP 7.13.0 Documentation” from directory: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hip/checkouts/develop/projects/hip/docs/doxygen/xml