Module Management

Module Management#

HIP Runtime API Reference: Module Management

Collaboration diagram for Module Management:

Modules
	Cooperative groups kernel launch of Module management.

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	hipModuleLoad (hipModule_t module, const char fname)
	Loads code object from file into a module the currrent context.

hipError_t	hipModuleUnload (hipModule_t module)
	Frees the module.

hipError_t	hipModuleGetFunction (hipFunction_t function, hipModule_t module, const char kname)
	Function with kname will be extracted if present in module.

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	hipGetFuncBySymbol (hipFunction_t functionPtr, const void symbolPtr)
	Gets pointer to device entry function that matches entry function symbolPtr.

hipError_t	hipModuleGetTexRef (textureReference *texRef, hipModule_t hmod, const char name)
	returns the handle of the texture reference with the name from the module.

hipError_t	hipModuleLoadData (hipModule_t module, const void image)
	builds module from code object which resides in host memory. Image is pointer to that location.

hipError_t	hipModuleLoadDataEx (hipModule_t module, const void image, unsigned int numOptions, hipJitOption options, void *optionValues)
	builds module from code object which resides in host memory. Image is pointer to that location. Options are not used. hipModuleLoadData is called.

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.

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 module management functions of HIP runtime API.

Function Documentation

◆ 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

◆ hipGetFuncBySymbol()

hipError_t hipGetFuncBySymbol	(	hipFunction_t *	functionPtr,
		const void *	symbolPtr
	)

Gets pointer to device entry function that matches entry function symbolPtr.

Parameters

[out]	functionPtr	Device entry function
[in]	symbolPtr	Pointer to device entry function to search for

Returns: hipSuccess, hipErrorInvalidDeviceFunction

◆ 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.

◆ hipModuleGetFunction()

hipError_t hipModuleGetFunction	(	hipFunction_t *	function,
		hipModule_t	module,
		const char *	kname
	)

Function with kname will be extracted if present in module.

Parameters

[in]	module	Module to get function from
[in]	kname	Pointer to the name of function
[out]	function	Pointer to function handle

Returns: hipSuccess, hipErrorInvalidValue, hipErrorInvalidContext, hipErrorNotInitialized, hipErrorNotFound,

◆ hipModuleGetTexRef()

hipError_t hipModuleGetTexRef	(	textureReference **	texRef,
		hipModule_t	hmod,
		const char *	name
	)

returns the handle of the texture reference with the name from the module.

Parameters

[in]	hmod	Module
[in]	name	Pointer of name of texture reference
[out]	texRef	Pointer of texture reference

Returns: hipSuccess, hipErrorNotInitialized, hipErrorNotFound, hipErrorInvalidValue

◆ 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

◆ hipModuleLoad()

hipError_t hipModuleLoad	(	hipModule_t *	module,
		const char *	fname
	)

Loads code object from file into a module the currrent context.

Parameters

[in]	fname	Filename of code object to load
[out]	module	Module

Warning: File/memory resources allocated in this function are released only in hipModuleUnload.

Returns: hipSuccess, hipErrorInvalidValue, hipErrorInvalidContext, hipErrorFileNotFound, hipErrorOutOfMemory, hipErrorSharedObjectInitFailed, hipErrorNotInitialized

◆ hipModuleLoadData()

hipError_t hipModuleLoadData	(	hipModule_t *	module,
		const void *	image
	)

builds module from code object which resides in host memory. Image is pointer to that location.

Parameters

[in]	image	The pointer to the location of data
[out]	module	Retuned module

Returns: hipSuccess, hipErrorNotInitialized, hipErrorOutOfMemory, hipErrorNotInitialized

◆ hipModuleLoadDataEx()

hipError_t hipModuleLoadDataEx	(	hipModule_t *	module,
		const void *	image,
		unsigned int	numOptions,
		hipJitOption *	options,
		void **	optionValues
	)

builds module from code object which resides in host memory. Image is pointer to that location. Options are not used. hipModuleLoadData is called.

Parameters

[in]	image	The pointer to the location of data
[out]	module	Retuned module
[in]	numOptions	Number of options
[in]	options	Options for JIT
[in]	optionValues	Option values for JIT

Returns: hipSuccess, hipErrorNotInitialized, hipErrorOutOfMemory, hipErrorNotInitialized

◆ hipModuleUnload()

hipError_t hipModuleUnload ( hipModule_t module )

Frees the module.

Parameters

[in] module Module to free

Returns: hipSuccess, hipErrorInvalidResourceHandle

The module is freed, and the code objects associated with it are destroyed.

Module Management

Module Management#

Modules

Functions

Detailed Description

Function Documentation

◆ hipExtLaunchKernel()

◆ hipExtLaunchKernelGGL()

◆ hipExtLaunchMultiKernelMultiDevice() [1/2]

◆ hipExtLaunchMultiKernelMultiDevice() [2/2]

◆ hipExtModuleLaunchKernel()

◆ hipFuncGetAttribute()

◆ hipFuncGetAttributes()

◆ hipGetFuncBySymbol()

◆ hipHccModuleLaunchKernel()

◆ hipModuleGetFunction()

◆ hipModuleGetTexRef()

◆ hipModuleLaunchKernel()

◆ hipModuleLoad()

◆ hipModuleLoadData()

◆ hipModuleLoadDataEx()

◆ hipModuleUnload()