Module management

Module management#

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.

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.

Parameters:

  • f – [in] Kernel to launch.

  • globalWorkSizeX – [in] X grid dimension specified in work-items.

  • globalWorkSizeY – [in] Y grid dimension specified in work-items.

  • globalWorkSizeZ – [in] Z grid dimension specified in work-items.

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

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

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

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

  • hStream – [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] pointer to kernel parameters.

  • extra – [in] 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.

  • startEvent – [in] 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.

  • stopEvent – [in] 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.

  • flags – [in] The value of hipExtAnyOrderLaunch, signifies if kernel can be launched in any order.

Returns:

hipSuccess, hipInvalidDeviceId, hipErrorNotInitialized, hipErrorInvalidValue.

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.

Parameters:

  • function_address – [in] pointer to the Kernel to launch.

  • numBlocks – [in] number of blocks.

  • dimBlocks – [in] dimension of a block.

  • args – [in] pointer to kernel arguments.

  • sharedMemBytes – [in] Amount of dynamic shared memory to allocate for this kernel. 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.

  • startEvent – [in] 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.

  • stopEvent – [in] 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.

  • flags – [in] The value of hipExtAnyOrderLaunch, signifies if kernel can be launched in any order.

Returns:

hipSuccess, hipInvalidDeviceId, hipErrorNotInitialized, hipErrorInvalidValue.

template<typename ...Args, typename F = void (*)(Args...)>
inline 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.

Parameters:

  • kernel – [in] Kernel to launch.

  • numBlocks – [in] const number of blocks.

  • dimBlocks – [in] const dimension of a block.

  • sharedMemBytes – [in] Amount of dynamic shared memory to allocate for this kernel. 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.

  • startEvent – [in] 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.

  • stopEvent – [in] 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.

  • flags – [in] The value of hipExtAnyOrderLaunch, signifies if kernel can be launched in any order.

  • args – [in] templated kernel arguments.

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

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

Warning

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

Parameters:

  • fname – [in] Filename of code object to load

  • module – [out] Module

Returns:

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

hipError_t hipModuleUnload(hipModule_t module)#

Frees the module.

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

Parameters:

module – [in] Module to free

Returns:

hipSuccess, hipErrorInvalidResourceHandle

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

Function with kname will be extracted if present in module.

Parameters:

  • module – [in] Module to get function from

  • kname – [in] Pointer to the name of function

  • function – [out] Pointer to function handle

Returns:

hipSuccess, hipErrorInvalidValue, hipErrorInvalidContext, hipErrorNotInitialized, hipErrorNotFound,

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

Find out attributes for a given function.

Parameters:

  • attr – [out] Attributes of funtion

  • func – [in] Pointer to the function handle

Returns:

hipSuccess, hipErrorInvalidValue, hipErrorInvalidDeviceFunction

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

Find out a specific attribute for a given function.

Parameters:

  • value – [out] Pointer to the value

  • attrib – [in] Attributes of the given funtion

  • hfunc – [in] Function to get attributes from

Returns:

hipSuccess, hipErrorInvalidValue, hipErrorInvalidDeviceFunction

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

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

Parameters:

  • functionPtr – [out] Device entry function

  • symbolPtr – [in] Pointer to device entry function to search for

Returns:

hipSuccess, hipErrorInvalidDeviceFunction

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:

  • hmod – [in] Module

  • name – [in] Pointer of name of texture reference

  • texRef – [out] Pointer of texture reference

Returns:

hipSuccess, hipErrorNotInitialized, hipErrorNotFound, hipErrorInvalidValue

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:

  • image – [in] The pointer to the location of data

  • module – [out] Retuned module

Returns:

hipSuccess, hipErrorNotInitialized, hipErrorOutOfMemory, hipErrorNotInitialized

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:

  • image – [in] The pointer to the location of data

  • module – [out] Retuned module

  • numOptions – [in] Number of options

  • options – [in] Options for JIT

  • optionValues – [in] Option values for JIT

Returns:

hipSuccess, hipErrorNotInitialized, hipErrorOutOfMemory, hipErrorNotInitialized

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

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.

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 dimensions specified in work-items

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

  • blockDimZ – [in] Z grid 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 th default stream is used with associated synchronization rules.

  • kernelParams – [in] Kernel parameters to launch

  • extra – [in] 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.

Returns:

hipSuccess, hipErrorNotInitialized, hipErrorInvalidValue

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:

  • 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

template<class T>
inline 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.

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

Contents