Module Management

Module Management#

HIP Runtime API Reference: Module Management

Collaboration diagram for Module Management:

Functions
HIP_PUBLIC_API 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=nullptr, hipEvent_t stopEvent=nullptr, uint32_t flags=0)
	Launches kernel with parameters and shared memory on stream with arguments passed to kernel params or extra arguments.

HIP_PUBLIC_API 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=nullptr, hipEvent_t stopEvent=nullptr) __attribute__((deprecated("use hipExtModuleLaunchKernel instead")))
	This HIP API is deprecated, please use hipExtModuleLaunchKernel() instead.

hipError_t	hipModuleLoad (hipModule_t module, const char fname)
	Loads code object from file into a hipModule_t.

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

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.

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.

Detailed Description

This section describes the module management functions of HIP runtime API.

Function Documentation

◆ hipExtLaunchMultiKernelMultiDevice()

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]	hipLaunchParams	List of launch parameters, one per device.
[in]	numDevices	Size of the launchParamsList array.
[in]	flags	Flags to control launch behavior.

Returns: hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue

◆ hipExtModuleLaunchKernel()

HIP_PUBLIC_API 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 = `nullptr`,
		hipEvent_t	stopEvent = `nullptr`,
		uint32_t	flags = `0`
	)

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]	gridDimX	X grid dimension specified in work-items.
[in]	gridDimY	Y grid dimension specified in work-items.
[in]	gridDimZ	Z grid dimension specified in work-items.
[in]	blockDimX	X block dimension 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. 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]	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.
[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, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue.

Warning: kernellParams argument is not yet implemented in HIP, use extra instead. Please refer to hip_porting_driver_api.md for sample usage. 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.

◆ hipFuncGetAttribute()

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

Find out a specific attribute for a given function.

Parameters

[out]	value
[in]	attrib
[in]	hfunc

Returns: hipSuccess, hipErrorInvalidValue, hipErrorInvalidDeviceFunction

◆ hipFuncGetAttributes()

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

Find out attributes for a given function.

Parameters

[out]	attr
[in]	func

Returns: hipSuccess, hipErrorInvalidValue, hipErrorInvalidDeviceFunction

◆ hipHccModuleLaunchKernel()

HIP_PUBLIC_API 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 = `nullptr`,
		hipEvent_t	stopEvent = `nullptr`
	)

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

◆ hipLaunchCooperativeKernel()

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

Parameters

[in]	f	Kernel to launch.
[in]	gridDim	Grid dimensions specified as multiple of blockDim.
[in]	blockDim	Block dimensions specified in work-items
[in]	kernelParams	A list of kernel arguments
[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.

Please note, HIP does not support kernel launch with total work items defined in dimension with size gridDim x blockDim >= 2^32.

Returns: hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue, hipErrorCooperativeLaunchTooLarge

◆ hipLaunchCooperativeKernelMultiDevice()

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

[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, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue, hipErrorCooperativeLaunchTooLarge

◆ 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
[in]	kname
[out]	function

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
[in]	name
[out]	texRef

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

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, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue

Warning: kernellParams argument is not yet implemented in HIP. Please use extra instead. Please refer to hip_porting_driver_api.md for sample usage.

◆ hipModuleLoad()

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

Loads code object from file into a hipModule_t.

Parameters

[in]	fname
[out]	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
[out]	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
[out]	module
[in]	number	of options
[in]	options	for JIT
[in]	option	values for JIT

Returns: hipSuccess, hipErrorNotInitialized, hipErrorOutOfMemory, hipErrorNotInitialized

◆ hipModuleUnload()

hipError_t hipModuleUnload ( hipModule_t module )

Frees the module.

Parameters

[in] module

Returns: hipSuccess, hipInvalidValue module is freed and the code objects associated with it are destroyed

Module Management

Module Management#

Functions

Detailed Description

Function Documentation

◆ hipExtLaunchMultiKernelMultiDevice()

◆ hipExtModuleLaunchKernel()

◆ hipFuncGetAttribute()

◆ hipFuncGetAttributes()

◆ hipHccModuleLaunchKernel()

◆ hipLaunchCooperativeKernel()

◆ hipLaunchCooperativeKernelMultiDevice()

◆ hipModuleGetFunction()

◆ hipModuleGetTexRef()

◆ hipModuleLaunchKernel()

◆ hipModuleLoad()

◆ hipModuleLoadData()

◆ hipModuleLoadDataEx()

◆ hipModuleUnload()