Execution Control

Execution Control#

HIP Runtime API Reference: Execution Control

Collaboration diagram for Execution Control:

Functions
hipError_t	hipExtModuleLaunchKernel (hipFunction_t f, uint32_t globalWorkSizeX, uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ, uint32_t localWorkSizeX, uint32_t localWorkSizeY, uint32_t localWorkSizeZ, size_t sharedMemBytes, hipStream_t hStream, void kernelParams, void extra, hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags)
	Launches kernel with parameters and shared memory on stream with arguments passed to kernel params or extra arguments.

hipError_t	hipHccModuleLaunchKernel (hipFunction_t f, uint32_t globalWorkSizeX, uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ, uint32_t localWorkSizeX, uint32_t localWorkSizeY, uint32_t localWorkSizeZ, size_t sharedMemBytes, hipStream_t hStream, void kernelParams, void extra, hipEvent_t startEvent, hipEvent_t stopEvent)
	This HIP API is deprecated, please use hipExtModuleLaunchKernel() instead.

hipError_t	hipExtLaunchKernel (const void function_address, dim3 numBlocks, dim3 dimBlocks, void *args, size_t sharedMemBytes, hipStream_t stream, hipEvent_t startEvent, hipEvent_t stopEvent, int flags)
	Launches kernel from the pointer address, with arguments and shared memory on stream.

template<typename... Args, typename F = void (*)(Args...)>
void	hipExtLaunchKernelGGL (F kernel, const dim3 &numBlocks, const dim3 &dimBlocks, std::uint32_t sharedMemBytes, hipStream_t stream, hipEvent_t startEvent, hipEvent_t stopEvent, std::uint32_t flags, Args... args)
	Launches kernel with dimention parameters and shared memory on stream with templated kernel and arguments.

hipError_t	hipFuncSetAttribute (const void *func, hipFuncAttribute attr, int value)
	Set attribute for a specific function.

hipError_t	hipFuncSetCacheConfig (const void *func, hipFuncCache_t config)
	Set Cache configuration for a specific function.

hipError_t	hipFuncSetSharedMemConfig (const void *func, hipSharedMemConfig config)
	Set shared memory configuation for a specific function.

hipError_t	hipFuncGetAttributes (struct hipFuncAttributes attr, const void func)
	Find out attributes for a given function.

hipError_t	hipFuncGetAttribute (int *value, hipFunction_attribute attrib, hipFunction_t hfunc)
	Find out a specific attribute for a given function.

hipError_t	hipModuleLaunchKernel (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, void extra)
	launches kernel f with launch parameters and shared memory on stream with arguments passed to kernelparams or extra

hipError_t	hipExtLaunchMultiKernelMultiDevice (hipLaunchParams *launchParamsList, int numDevices, unsigned int flags)
	Launches kernels on multiple devices and guarantees all specified kernels are dispatched on respective streams before enqueuing any other work on the specified streams from any other threads.

hipError_t	hipLaunchKernelExC (const hipLaunchConfig_t config, const void fPtr, void **args)
	Launches a HIP kernel using a generic function pointer and the specified configuration.

hipError_t	hipDrvLaunchKernelEx (const HIP_LAUNCH_CONFIG config, hipFunction_t f, void params, void *extra)
	Launches a HIP kernel using the driver API with the specified configuration.

template<class T >
hipError_t	hipExtLaunchMultiKernelMultiDevice (hipLaunchParams *launchParamsList, unsigned int numDevices, unsigned int flags=0)
	Launches kernels on multiple devices and guarantees all specified kernels are dispatched on respective streams before enqueuing any other work on the specified streams from any other threads.

Detailed Description

This section describes the execution control functions of HIP runtime API.

Function Documentation

◆ hipDrvLaunchKernelEx()

hipError_t hipDrvLaunchKernelEx	(	const HIP_LAUNCH_CONFIG *	config,
		hipFunction_t	f,
		void **	params,
		void **	extra
	)

Launches a HIP kernel using the driver API with the specified configuration.

This function dispatches the device kernel represented by a HIP function object. It passes both the kernel parameters and any extra configuration arguments to the kernel launch.

Parameters

[in]	config	Pointer to the kernel launch configuration structure.
[in]	f	HIP function object representing the device kernel to be launched.
[in]	params	Array of pointers to the kernel parameters.
[in]	extra	Array of pointers for additional launch parameters or extra configuration data.

Returns: hipSuccess if the kernel is launched successfully, otherwise an appropriate error code.

◆ hipExtLaunchKernel()

hipError_t hipExtLaunchKernel	(	const void *	function_address,
		dim3	numBlocks,
		dim3	dimBlocks,
		void **	args,
		size_t	sharedMemBytes,
		hipStream_t	stream,
		hipEvent_t	startEvent,
		hipEvent_t	stopEvent,
		int	flags
	)

Launches kernel from the pointer address, with arguments and shared memory on stream.

Parameters

[in]	function_address	pointer to the Kernel to launch.
[in]	numBlocks	number of blocks.
[in]	dimBlocks	dimension of a block.
[in]	args	pointer to kernel arguments.
[in]	sharedMemBytes	Amount of dynamic shared memory to allocate for this kernel. HIP-Clang compiler provides support for extern shared declarations.
[in]	stream	Stream where the kernel should be dispatched. May be 0, in which case the default stream is used with associated synchronization rules.
[in]	startEvent	If non-null, specified event will be updated to track the start time of the kernel launch. The event must be created before calling this API.
[in]	stopEvent	If non-null, specified event will be updated to track the stop time of the kernel launch. The event must be created before calling this API.
[in]	flags	The value of hipExtAnyOrderLaunch, signifies if kernel can be launched in any order.

Returns: hipSuccess, hipInvalidDeviceId, hipErrorNotInitialized, hipErrorInvalidValue.

◆ hipExtLaunchKernelGGL()

template<typename... Args, typename F = void (*)(Args...)>

void hipExtLaunchKernelGGL	(	F	kernel,
		const dim3 &	numBlocks,
		const dim3 &	dimBlocks,
		std::uint32_t	sharedMemBytes,
		hipStream_t	stream,
		hipEvent_t	startEvent,
		hipEvent_t	stopEvent,
		std::uint32_t	flags,
		Args...	args
	)

inline

Launches kernel with dimention parameters and shared memory on stream with templated kernel and arguments.

Parameters

[in]	kernel	Kernel to launch.
[in]	numBlocks	const number of blocks.
[in]	dimBlocks	const dimension of a block.
[in]	sharedMemBytes	Amount of dynamic shared memory to allocate for this kernel. HIP-Clang compiler provides support for extern shared declarations.
[in]	stream	Stream where the kernel should be dispatched. May be 0, in which case the default stream is used with associated synchronization rules.
[in]	startEvent	If non-null, specified event will be updated to track the start time of the kernel launch. The event must be created before calling this API.
[in]	stopEvent	If non-null, specified event will be updated to track the stop time of the kernel launch. The event must be created before calling this API.
[in]	flags	The value of hipExtAnyOrderLaunch, signifies if kernel can be launched in any order.
[in]	args	templated kernel arguments.

◆ hipExtLaunchMultiKernelMultiDevice() [1/2]

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

Launches kernels on multiple devices and guarantees all specified kernels are dispatched on respective streams before enqueuing any other work on the specified streams from any other threads.

Parameters

[in]	launchParamsList	List of launch parameters, one per device.
[in]	numDevices	Size of the launchParamsList array.
[in]	flags	Flags to control launch behavior.

Returns: hipSuccess, hipErrorNotInitialized, hipErrorInvalidValue

◆ hipExtLaunchMultiKernelMultiDevice() [2/2]

template<class T >

hipError_t hipExtLaunchMultiKernelMultiDevice	(	hipLaunchParams *	launchParamsList,
		unsigned int	numDevices,
		unsigned int	flags = `0`
	)

inline

Launches kernels on multiple devices and guarantees all specified kernels are dispatched on respective streams before enqueuing any other work on the specified streams from any other threads.

Parameters

[in]	launchParamsList	List of launch parameters, one per device.
[in]	numDevices	Size of the launchParamsList array.
[in]	flags	Flags to control launch behavior.

Returns: hipSuccess, hipErrorInvalidValue

◆ hipExtModuleLaunchKernel()

hipError_t hipExtModuleLaunchKernel	(	hipFunction_t	f,
		uint32_t	globalWorkSizeX,
		uint32_t	globalWorkSizeY,
		uint32_t	globalWorkSizeZ,
		uint32_t	localWorkSizeX,
		uint32_t	localWorkSizeY,
		uint32_t	localWorkSizeZ,
		size_t	sharedMemBytes,
		hipStream_t	hStream,
		void **	kernelParams,
		void **	extra,
		hipEvent_t	startEvent,
		hipEvent_t	stopEvent,
		uint32_t	flags
	)

Launches kernel with parameters and shared memory on stream with arguments passed to kernel params or extra arguments.

Parameters

[in]	f	Kernel to launch.
[in]	globalWorkSizeX	X grid dimension specified in work-items.
[in]	globalWorkSizeY	Y grid dimension specified in work-items.
[in]	globalWorkSizeZ	Z grid dimension specified in work-items.
[in]	localWorkSizeX	X block dimension specified in work-items.
[in]	localWorkSizeY	Y block dimension specified in work-items.
[in]	localWorkSizeZ	Z block dimension specified in work-items.
[in]	sharedMemBytes	Amount of dynamic shared memory to allocate for this kernel. HIP-Clang compiler provides support for extern shared declarations.
[in]	hStream	Stream where the kernel should be dispatched. May be 0, in which case the default stream is used with associated synchronization rules.
[in]	kernelParams	pointer to kernel parameters.
[in]	extra	Pointer to kernel arguments. These are passed directly to the kernel and must be in the memory layout and alignment expected by the kernel. All passed arguments must be naturally aligned according to their type. The memory address of each argument should be a multiple of its size in bytes. Please refer to hip_porting_driver_api.md for sample usage.
[in]	startEvent	If non-null, specified event will be updated to track the start time of the kernel launch. The event must be created before calling this API.
[in]	stopEvent	If non-null, specified event will be updated to track the stop time of the kernel launch. The event must be created before calling this API.
[in]	flags	The value of hipExtAnyOrderLaunch, signifies if kernel can be launched in any order.

Returns: hipSuccess, hipInvalidDeviceId, hipErrorNotInitialized, hipErrorInvalidValue.

HIP/ROCm actually updates the start event when the associated kernel completes. Currently, timing between startEvent and stopEvent does not include the time it takes to perform a system scope release/cache flush - only the time it takes to issues writes to cache.

Note: For this HIP API, the flag 'hipExtAnyOrderLaunch' is not supported on AMD GFX9xx boards.

◆ hipFuncGetAttribute()

hipError_t hipFuncGetAttribute	(	int *	value,
		hipFunction_attribute	attrib,
		hipFunction_t	hfunc
	)

Find out a specific attribute for a given function.

Parameters

[out]	value	Pointer to the value
[in]	attrib	Attributes of the given funtion
[in]	hfunc	Function to get attributes from

Returns: hipSuccess, hipErrorInvalidValue, hipErrorInvalidDeviceFunction

◆ hipFuncGetAttributes()

hipError_t hipFuncGetAttributes	(	struct hipFuncAttributes *	attr,
		const void *	func
	)

Find out attributes for a given function.

Parameters

[out]	attr	Attributes of funtion
[in]	func	Pointer to the function handle

Returns: hipSuccess, hipErrorInvalidValue, hipErrorInvalidDeviceFunction

◆ hipFuncSetAttribute()

hipError_t hipFuncSetAttribute	(	const void *	func,
		hipFuncAttribute	attr,
		int	value
	)

Set attribute for a specific function.

Parameters

[in]	func	Pointer of the function
[in]	attr	Attribute to set
[in]	value	Value to set

Returns: hipSuccess, hipErrorInvalidDeviceFunction, hipErrorInvalidValue

Note: AMD devices and some Nvidia GPUS do not support shared cache banking, and the hint is ignored on those architectures.

◆ hipFuncSetCacheConfig()

hipError_t hipFuncSetCacheConfig	(	const void *	func,
		hipFuncCache_t	config
	)

Set Cache configuration for a specific function.

Parameters

[in]	func	Pointer of the function.
[in]	config	Configuration to set.

Returns: hipSuccess, hipErrorNotInitialized Note: AMD devices and some Nvidia GPUS do not support reconfigurable cache. This hint is ignored on those architectures.

◆ hipFuncSetSharedMemConfig()

hipError_t hipFuncSetSharedMemConfig	(	const void *	func,
		hipSharedMemConfig	config
	)

Set shared memory configuation for a specific function.

Parameters

[in]	func	Pointer of the function
[in]	config	Configuration

Returns: hipSuccess, hipErrorInvalidDeviceFunction, hipErrorInvalidValue

Note: AMD devices and some Nvidia GPUS do not support shared cache banking, and the hint is ignored on those architectures.

◆ hipHccModuleLaunchKernel()

hipError_t hipHccModuleLaunchKernel	(	hipFunction_t	f,
		uint32_t	globalWorkSizeX,
		uint32_t	globalWorkSizeY,
		uint32_t	globalWorkSizeZ,
		uint32_t	localWorkSizeX,
		uint32_t	localWorkSizeY,
		uint32_t	localWorkSizeZ,
		size_t	sharedMemBytes,
		hipStream_t	hStream,
		void **	kernelParams,
		void **	extra,
		hipEvent_t	startEvent,
		hipEvent_t	stopEvent
	)

This HIP API is deprecated, please use hipExtModuleLaunchKernel() instead.

◆ hipLaunchKernelExC()

hipError_t hipLaunchKernelExC	(	const hipLaunchConfig_t *	config,
		const void *	fPtr,
		void **	args
	)

Launches a HIP kernel using a generic function pointer and the specified configuration.

This function is equivalent to hipLaunchKernelEx but accepts the kernel as a generic function pointer.

Parameters

[in]	config	Pointer to the kernel launch configuration structure.
[in]	fPtr	Pointer to the device kernel function.
[in]	args	Array of pointers to the kernel arguments.

Returns: hipSuccess if the kernel is launched successfully, otherwise an appropriate error code.

◆ hipModuleLaunchKernel()

hipError_t hipModuleLaunchKernel	(	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,
		void **	extra
	)

launches kernel f with launch parameters and shared memory on stream with arguments passed to kernelparams or extra

Parameters

[in]	f	Kernel to launch.
[in]	gridDimX	X grid dimension specified as multiple of blockDimX.
[in]	gridDimY	Y grid dimension specified as multiple of blockDimY.
[in]	gridDimZ	Z grid dimension specified as multiple of blockDimZ.
[in]	blockDimX	X block dimensions specified in work-items
[in]	blockDimY	Y grid dimension specified in work-items
[in]	blockDimZ	Z grid dimension specified in work-items
[in]	sharedMemBytes	Amount of dynamic shared memory to allocate for this kernel. The HIP-Clang compiler provides support for extern shared declarations.
[in]	stream	Stream where the kernel should be dispatched. May be 0, in which case th default stream is used with associated synchronization rules.
[in]	kernelParams	Kernel parameters to launch
[in]	extra	Pointer to kernel arguments. These are passed directly to the kernel and must be in the memory layout and alignment expected by the kernel. All passed arguments must be naturally aligned according to their type. The memory address of each argument should be a multiple of its size in bytes. Please refer to hip_porting_driver_api.md for sample usage.

Please note, HIP does not support kernel launch with total work items defined in dimension with size gridDim x blockDim >= 2^32. So gridDim.x * blockDim.x, gridDim.y * blockDim.y and gridDim.z * blockDim.z are always less than 2^32.

Returns: hipSuccess, hipErrorNotInitialized, hipErrorInvalidValue

Execution Control

Execution Control#

Functions

Detailed Description

Function Documentation

◆ hipDrvLaunchKernelEx()

◆ hipExtLaunchKernel()

◆ hipExtLaunchKernelGGL()

◆ hipExtLaunchMultiKernelMultiDevice() [1/2]

◆ hipExtLaunchMultiKernelMultiDevice() [2/2]

◆ hipExtModuleLaunchKernel()

◆ hipFuncGetAttribute()

◆ hipFuncGetAttributes()

◆ hipFuncSetAttribute()

◆ hipFuncSetCacheConfig()

◆ hipFuncSetSharedMemConfig()

◆ hipHccModuleLaunchKernel()

◆ hipLaunchKernelExC()

◆ hipModuleLaunchKernel()