Executable

Classes
struct	hsa_code_object_reader_s
	Code object reader handle. More...
struct	hsa_executable_s
	Struct containing an opaque handle to an executable, which contains ISA for finalized kernels and indirect functions together with the allocated global or readonly segment variables they reference. More...
struct	hsa_loaded_code_object_s
	Loaded code object handle. More...
struct	hsa_executable_symbol_s
	Executable symbol handle. More...

Typedefs
typedef struct hsa_code_object_reader_s	hsa_code_object_reader_t
	Code object reader handle.
typedef struct hsa_executable_s	hsa_executable_t
	Struct containing an opaque handle to an executable, which contains ISA for finalized kernels and indirect functions together with the allocated global or readonly segment variables they reference.
typedef struct hsa_loaded_code_object_s	hsa_loaded_code_object_t
	Loaded code object handle.
typedef struct hsa_executable_symbol_s	hsa_executable_symbol_t
	Executable symbol handle.

Enumerations
enum	hsa_executable_state_t { HSA_EXECUTABLE_STATE_UNFROZEN = 0 , HSA_EXECUTABLE_STATE_FROZEN = 1 }
	Executable state. More...
enum	hsa_executable_info_t { HSA_EXECUTABLE_INFO_PROFILE = 1 , HSA_EXECUTABLE_INFO_STATE = 2 , HSA_EXECUTABLE_INFO_DEFAULT_FLOAT_ROUNDING_MODE = 3 }
	Executable attributes. More...
enum	hsa_symbol_kind_t { HSA_SYMBOL_KIND_VARIABLE = 0 , HSA_SYMBOL_KIND_KERNEL = 1 , HSA_SYMBOL_KIND_INDIRECT_FUNCTION = 2 }
	Symbol type. More...
enum	hsa_symbol_linkage_t { HSA_SYMBOL_LINKAGE_MODULE = 0 , HSA_SYMBOL_LINKAGE_PROGRAM = 1 }
	Linkage type of a symbol. More...
enum	hsa_variable_allocation_t { HSA_VARIABLE_ALLOCATION_AGENT = 0 , HSA_VARIABLE_ALLOCATION_PROGRAM = 1 }
	Allocation type of a variable. More...
enum	hsa_variable_segment_t { HSA_VARIABLE_SEGMENT_GLOBAL = 0 , HSA_VARIABLE_SEGMENT_READONLY = 1 }
	Memory segment associated with a variable. More...
enum	hsa_executable_symbol_info_t {   HSA_EXECUTABLE_SYMBOL_INFO_TYPE = 0 , HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH = 1 , HSA_EXECUTABLE_SYMBOL_INFO_NAME = 2 , HSA_EXECUTABLE_SYMBOL_INFO_MODULE_NAME_LENGTH = 3 ,   HSA_EXECUTABLE_SYMBOL_INFO_MODULE_NAME = 4 , HSA_EXECUTABLE_SYMBOL_INFO_AGENT = 20 , HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS = 21 , HSA_EXECUTABLE_SYMBOL_INFO_LINKAGE = 5 ,   HSA_EXECUTABLE_SYMBOL_INFO_IS_DEFINITION = 17 , HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ALLOCATION = 6 , HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SEGMENT = 7 , HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ALIGNMENT = 8 ,   HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE = 9 , HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_IS_CONST = 10 , HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT = 22 , HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE = 11 ,   HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_ALIGNMENT = 12 , HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE = 13 , HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE = 14 , HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_DYNAMIC_CALLSTACK = 15 ,   HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_CALL_CONVENTION = 18 , HSA_EXECUTABLE_SYMBOL_INFO_INDIRECT_FUNCTION_OBJECT = 23 , HSA_EXECUTABLE_SYMBOL_INFO_INDIRECT_FUNCTION_CALL_CONVENTION = 16 }
	Executable symbol attributes. More...

Functions
hsa_status_t HSA_API	hsa_code_object_reader_create_from_file (hsa_file_t file, hsa_code_object_reader_t *code_object_reader)
	Create a code object reader to operate on a file.
hsa_status_t HSA_API	hsa_code_object_reader_create_from_memory (const void *code_object, size_t size, hsa_code_object_reader_t *code_object_reader)
	Create a code object reader to operate on memory.
hsa_status_t HSA_API	hsa_code_object_reader_destroy (hsa_code_object_reader_t code_object_reader)
	Destroy a code object reader.
hsa_status_t HSA_API HSA_DEPRECATED	hsa_executable_create (hsa_profile_t profile, hsa_executable_state_t executable_state, const char *options, hsa_executable_t *executable)
	Create an empty executable.
hsa_status_t HSA_API	hsa_executable_create_alt (hsa_profile_t profile, hsa_default_float_rounding_mode_t default_float_rounding_mode, const char *options, hsa_executable_t *executable)
	Create an empty executable.
hsa_status_t HSA_API	hsa_executable_destroy (hsa_executable_t executable)
	Destroy an executable.
hsa_status_t HSA_API	hsa_executable_load_program_code_object (hsa_executable_t executable, hsa_code_object_reader_t code_object_reader, const char *options, hsa_loaded_code_object_t *loaded_code_object)
	Load a program code object into an executable.
hsa_status_t HSA_API	hsa_executable_load_agent_code_object (hsa_executable_t executable, hsa_agent_t agent, hsa_code_object_reader_t code_object_reader, const char *options, hsa_loaded_code_object_t *loaded_code_object)
	Load an agent code object into an executable.
hsa_status_t HSA_API	hsa_executable_freeze (hsa_executable_t executable, const char *options)
	Freeze the executable.
hsa_status_t HSA_API	hsa_executable_get_info (hsa_executable_t executable, hsa_executable_info_t attribute, void *value)
	Get the current value of an attribute for a given executable.
hsa_status_t HSA_API	hsa_executable_global_variable_define (hsa_executable_t executable, const char *variable_name, void *address)
	Define an external global variable with program allocation.
hsa_status_t HSA_API	hsa_executable_agent_global_variable_define (hsa_executable_t executable, hsa_agent_t agent, const char *variable_name, void *address)
	Define an external global variable with agent allocation.
hsa_status_t HSA_API	hsa_executable_readonly_variable_define (hsa_executable_t executable, hsa_agent_t agent, const char *variable_name, void *address)
	Define an external readonly variable.
hsa_status_t HSA_API	hsa_executable_validate (hsa_executable_t executable, uint32_t *result)
	Validate an executable.
hsa_status_t HSA_API	hsa_executable_validate_alt (hsa_executable_t executable, const char *options, uint32_t *result)
	Validate an executable.
hsa_status_t HSA_API HSA_DEPRECATED	hsa_executable_get_symbol (hsa_executable_t executable, const char *module_name, const char *symbol_name, hsa_agent_t agent, int32_t call_convention, hsa_executable_symbol_t *symbol)
	Get the symbol handle for a given a symbol name.
hsa_status_t HSA_API	hsa_executable_get_symbol_by_name (hsa_executable_t executable, const char *symbol_name, const hsa_agent_t *agent, hsa_executable_symbol_t *symbol)
	Retrieve the symbol handle corresponding to a given a symbol name.
hsa_status_t HSA_API	hsa_executable_symbol_get_info (hsa_executable_symbol_t executable_symbol, hsa_executable_symbol_info_t attribute, void *value)
	Get the current value of an attribute for a given executable symbol.
hsa_status_t HSA_API HSA_DEPRECATED	hsa_executable_iterate_symbols (hsa_executable_t executable, hsa_status_t(*callback)(hsa_executable_t exec, hsa_executable_symbol_t symbol, void *data), void *data)
	Iterate over the symbols in a executable, and invoke an application-defined callback on every iteration.
hsa_status_t HSA_API	hsa_executable_iterate_agent_symbols (hsa_executable_t executable, hsa_agent_t agent, hsa_status_t(*callback)(hsa_executable_t exec, hsa_agent_t agent, hsa_executable_symbol_t symbol, void *data), void *data)
	Iterate over the kernels, indirect functions, and agent allocation variables in an executable for a given agent, and invoke an application- defined callback on every iteration.
hsa_status_t HSA_API	hsa_executable_iterate_program_symbols (hsa_executable_t executable, hsa_status_t(*callback)(hsa_executable_t exec, hsa_executable_symbol_t symbol, void *data), void *data)
	Iterate over the program allocation variables in an executable, and invoke an application-defined callback on every iteration.

Detailed Description

Typedef Documentation

hsa_code_object_reader_t

typedef struct hsa_code_object_reader_s hsa_code_object_reader_t

Code object reader handle.

A code object reader is used to load a code object from file (when created using hsa_code_object_reader_create_from_file), or from memory (if created using hsa_code_object_reader_create_from_memory).

hsa_executable_symbol_t

typedef struct hsa_executable_symbol_s hsa_executable_symbol_t

Executable symbol handle.

The lifetime of an executable object symbol matches that of the executable associated with it. An operation on a symbol whose associated executable has been destroyed results in undefined behavior.

hsa_executable_t

typedef struct hsa_executable_s hsa_executable_t

Struct containing an opaque handle to an executable, which contains ISA for finalized kernels and indirect functions together with the allocated global or readonly segment variables they reference.

hsa_loaded_code_object_t

typedef struct hsa_loaded_code_object_s hsa_loaded_code_object_t

Loaded code object handle.

Enumeration Type Documentation

hsa_executable_info_t

enum hsa_executable_info_t

Executable attributes.

Enumerator
HSA_EXECUTABLE_INFO_PROFILE	Profile this executable is created for. The type of this attribute is hsa_profile_t.
HSA_EXECUTABLE_INFO_STATE	Executable state. The type of this attribute is hsa_executable_state_t.
HSA_EXECUTABLE_INFO_DEFAULT_FLOAT_ROUNDING_MODE	Default floating-point rounding mode specified when executable was created. The type of this attribute is hsa_default_float_rounding_mode_t.

Definition at line 4401 of file hsa.h.

hsa_executable_state_t

enum hsa_executable_state_t

Executable state.

Enumerator
HSA_EXECUTABLE_STATE_UNFROZEN	Executable state, which allows the user to load code objects and define external variables. Variable addresses, kernel code handles, and indirect function code handles are not available in query operations until the executable is frozen (zero always returned).
HSA_EXECUTABLE_STATE_FROZEN	Executable state, which allows the user to query variable addresses, kernel code handles, and indirect function code handles using query operations. Loading new code objects, as well as defining external variables, is not allowed in this state.

Definition at line 4119 of file hsa.h.

hsa_executable_symbol_info_t

enum hsa_executable_symbol_info_t

Executable symbol attributes.

Enumerator
HSA_EXECUTABLE_SYMBOL_INFO_TYPE	The kind of the symbol. The type of this attribute is hsa_symbol_kind_t.
HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH	The length of the symbol name in bytes, not including the NUL terminator. The type of this attribute is uint32_t.
HSA_EXECUTABLE_SYMBOL_INFO_NAME	The name of the symbol. The type of this attribute is character array with the length equal to the value of HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH attribute.
HSA_EXECUTABLE_SYMBOL_INFO_MODULE_NAME_LENGTH	Deprecated The length of the module name in bytes (not including the NUL terminator) to which this symbol belongs if this symbol has module linkage, otherwise 0 is returned. The type of this attribute is uint32_t.
HSA_EXECUTABLE_SYMBOL_INFO_MODULE_NAME	Deprecated The module name to which this symbol belongs if this symbol has module linkage, otherwise an empty string is returned. The type of this attribute is character array with the length equal to the value of HSA_EXECUTABLE_SYMBOL_INFO_MODULE_NAME_LENGTH attribute.
HSA_EXECUTABLE_SYMBOL_INFO_AGENT	Deprecated Agent associated with this symbol. If the symbol is a variable, the value of this attribute is only defined if HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ALLOCATION is HSA_VARIABLE_ALLOCATION_AGENT. The type of this attribute is hsa_agent_t.
HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS	The address of the variable. The value of this attribute is undefined if the symbol is not a variable. The type of this attribute is uint64_t. If executable's state is HSA_EXECUTABLE_STATE_UNFROZEN, then 0 is returned.
HSA_EXECUTABLE_SYMBOL_INFO_LINKAGE	The linkage kind of the symbol. The type of this attribute is hsa_symbol_linkage_t.
HSA_EXECUTABLE_SYMBOL_INFO_IS_DEFINITION	Indicates whether the symbol corresponds to a definition. The type of this attribute is bool.
HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ALLOCATION	Deprecated The allocation kind of the variable. The value of this attribute is undefined if the symbol is not a variable. The type of this attribute is hsa_variable_allocation_t.
HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SEGMENT	Deprecated The segment kind of the variable. The value of this attribute is undefined if the symbol is not a variable. The type of this attribute is hsa_variable_segment_t.
HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ALIGNMENT	Deprecated Alignment of the symbol in memory. The value of this attribute is undefined if the symbol is not a variable. The type of this attribute is uint32_t. The current alignment of the variable in memory may be greater than the value specified in the source program variable declaration.
HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE	Deprecated Size of the variable. The value of this attribute is undefined if the symbol is not a variable. The type of this attribute is uint32_t. A value of 0 is returned if the variable is an external variable and has an unknown dimension.
HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_IS_CONST	Deprecated Indicates whether the variable is constant. The value of this attribute is undefined if the symbol is not a variable. The type of this attribute is bool.
HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT	Kernel object handle, used in the kernel dispatch packet. The value of this attribute is undefined if the symbol is not a kernel. The type of this attribute is uint64_t. If the state of the executable is HSA_EXECUTABLE_STATE_UNFROZEN, then 0 is returned.
HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE	Size of kernarg segment memory that is required to hold the values of the kernel arguments, in bytes. Must be a multiple of 16. The value of this attribute is undefined if the symbol is not a kernel. The type of this attribute is uint32_t.
HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_ALIGNMENT	Alignment (in bytes) of the buffer used to pass arguments to the kernel, which is the maximum of 16 and the maximum alignment of any of the kernel arguments. The value of this attribute is undefined if the symbol is not a kernel. The type of this attribute is uint32_t.
HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE	Size of static group segment memory required by the kernel (per work-group), in bytes. The value of this attribute is undefined if the symbol is not a kernel. The type of this attribute is uint32_t. The reported amount does not include any dynamically allocated group segment memory that may be requested by the application when a kernel is dispatched.
HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE	Size of static private, spill, and arg segment memory required by this kernel (per work-item), in bytes. The value of this attribute is undefined if the symbol is not a kernel. The type of this attribute is uint32_t. If the value of HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_DYNAMIC_CALLSTACK is true, the kernel may use more private memory than the reported value, and the application must add the dynamic call stack usage to private_segment_size when populating a kernel dispatch packet.
HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_DYNAMIC_CALLSTACK	Dynamic callstack flag. The value of this attribute is undefined if the symbol is not a kernel. The type of this attribute is bool. If this flag is set (the value is true), the kernel uses a dynamically sized call stack. This can happen if recursive calls, calls to indirect functions, or the HSAIL alloca instruction are present in the kernel.
HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_CALL_CONVENTION	Deprecated Call convention of the kernel. The value of this attribute is undefined if the symbol is not a kernel. The type of this attribute is uint32_t.
HSA_EXECUTABLE_SYMBOL_INFO_INDIRECT_FUNCTION_OBJECT	Indirect function object handle. The value of this attribute is undefined if the symbol is not an indirect function, or the associated agent does not support the Full Profile. The type of this attribute depends on the machine model: the type is uint32_t for small machine model, and uint64_t for large model. If the state of the executable is HSA_EXECUTABLE_STATE_UNFROZEN, then 0 is returned.
HSA_EXECUTABLE_SYMBOL_INFO_INDIRECT_FUNCTION_CALL_CONVENTION	Deprecated Call convention of the indirect function. The value of this attribute is undefined if the symbol is not an indirect function, or the associated agent does not support the Full Profile. The type of this attribute is uint32_t.

Definition at line 4807 of file hsa.h.

hsa_symbol_kind_t

enum hsa_symbol_kind_t

Symbol type.

Enumerator
HSA_SYMBOL_KIND_VARIABLE	Variable.
HSA_SYMBOL_KIND_KERNEL	Kernel.
HSA_SYMBOL_KIND_INDIRECT_FUNCTION	Indirect function.

Definition at line 4743 of file hsa.h.

hsa_symbol_linkage_t

enum hsa_symbol_linkage_t

Linkage type of a symbol.

Enumerator
HSA_SYMBOL_LINKAGE_MODULE	Module linkage.
HSA_SYMBOL_LINKAGE_PROGRAM	Program linkage.

Definition at line 4762 of file hsa.h.

hsa_variable_allocation_t

enum hsa_variable_allocation_t

Allocation type of a variable.

Enumerator
HSA_VARIABLE_ALLOCATION_AGENT	Agent allocation.
HSA_VARIABLE_ALLOCATION_PROGRAM	Program allocation.

Definition at line 4777 of file hsa.h.

hsa_variable_segment_t

enum hsa_variable_segment_t

Memory segment associated with a variable.

Enumerator
HSA_VARIABLE_SEGMENT_GLOBAL	Global memory segment.
HSA_VARIABLE_SEGMENT_READONLY	Readonly memory segment.

Definition at line 4792 of file hsa.h.

Function Documentation

hsa_code_object_reader_create_from_file()

hsa_status_t HSA_API hsa_code_object_reader_create_from_file	(	hsa_file_t	file,
		hsa_code_object_reader_t *	code_object_reader )

Create a code object reader to operate on a file.

Parameters

[in]

file

File descriptor. The file must have been opened by application with at least read permissions prior calling this function. The file must contain a vendor-specific code object.

The file is owned and managed by the application; the lifetime of the file descriptor must exceed that of any associated code object reader.

Parameters

[out]

code_object_reader

Memory location to store the newly created code object reader handle. Must not be NULL.

Return values

HSA_STATUS_SUCCESS	The function has been executed successfully.
HSA_STATUS_ERROR_NOT_INITIALIZED	The HSA runtime has not been initialized.
HSA_STATUS_ERROR_INVALID_FILE	file is invalid.
HSA_STATUS_ERROR_OUT_OF_RESOURCES	The HSA runtime failed to allocate the required resources.
HSA_STATUS_ERROR_INVALID_ARGUMENT	code_object_reader is NULL.

hsa_code_object_reader_create_from_memory()

hsa_status_t HSA_API hsa_code_object_reader_create_from_memory	(	const void *	code_object,
		size_t	size,
		hsa_code_object_reader_t *	code_object_reader )

Create a code object reader to operate on memory.

Parameters

[in]	code_object	Memory buffer that contains a vendor-specific code object. The buffer is owned and managed by the application; the lifetime of the buffer must exceed that of any associated code object reader.
[in]	size	Size of the buffer pointed to by code_object. Must not be 0.
[out]	code_object_reader	Memory location to store newly created code object reader handle. Must not be NULL.

Return values

HSA_STATUS_SUCCESS	The function has been executed successfully.
HSA_STATUS_ERROR_NOT_INITIALIZED	The HSA runtime has not been initialized.
HSA_STATUS_ERROR_OUT_OF_RESOURCES	The HSA runtime failed to allocate the required resources.
HSA_STATUS_ERROR_INVALID_ARGUMENT	code_object is NULL, size is zero, or code_object_reader is NULL.

hsa_code_object_reader_destroy()

hsa_status_t HSA_API hsa_code_object_reader_destroy

(

hsa_code_object_reader_t

code_object_reader

)

Destroy a code object reader.

The code object reader handle becomes invalid after completion of this function. Any file or memory used to create the code object read is not closed, removed, or deallocated by this function.

Parameters

[in]

code_object_reader

Code object reader to destroy.

Return values

HSA_STATUS_SUCCESS	The function has been executed successfully.
HSA_STATUS_ERROR_NOT_INITIALIZED	The HSA runtime has not been initialized.
HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER	code_object_reader is invalid.

hsa_executable_agent_global_variable_define()

hsa_status_t HSA_API hsa_executable_agent_global_variable_define	(	hsa_executable_t	executable,
		hsa_agent_t	agent,
		const char *	variable_name,
		void *	address )

Define an external global variable with agent allocation.

This function allows the application to provide the definition of a variable in the global segment memory with agent allocation. The variable must be defined before loading a code object into an executable. In addition, code objects loaded must not define the variable.

Parameters

[in]	executable	Executable. Must not be in frozen state.
[in]	agent	Agent for which the variable is being defined.
[in]	variable_name	Name of the variable. The Programmer's Reference Manual describes the standard name mangling scheme.
[in]	address	Address where the variable is defined. This address must have been previously allocated using hsa_memory_allocate in a global region that is only visible to agent. The application cannot deallocate the buffer pointed by address before executable is destroyed.

Return values

HSA_STATUS_SUCCESS	The function has been executed successfully.
HSA_STATUS_ERROR_NOT_INITIALIZED	The HSA runtime has not been initialized.
HSA_STATUS_ERROR_OUT_OF_RESOURCES	The HSA runtime failed to allocate the required resources.
HSA_STATUS_ERROR_INVALID_EXECUTABLE	The executable is invalid.
HSA_STATUS_ERROR_INVALID_AGENT	agent is invalid.
HSA_STATUS_ERROR_VARIABLE_ALREADY_DEFINED	The variable is already defined.
HSA_STATUS_ERROR_INVALID_SYMBOL_NAME	There is no variable with the variable_name.
HSA_STATUS_ERROR_FROZEN_EXECUTABLE	executable is frozen.
HSA_STATUS_ERROR_INVALID_ARGUMENT	variable_name is NULL.

hsa_executable_create()

hsa_status_t HSA_API HSA_DEPRECATED hsa_executable_create	(	hsa_profile_t	profile,
		hsa_executable_state_t	executable_state,
		const char *	options,
		hsa_executable_t *	executable )

Create an empty executable.

Deprecated

Use hsa_executable_create_alt instead, which allows the application to specify the default floating-point rounding mode of the executable and assumes an unfrozen initial state.

Parameters

[in]	profile	Profile used in the executable.
[in]	executable_state	Executable state. If the state is HSA_EXECUTABLE_STATE_FROZEN, the resulting executable is useless because no code objects can be loaded, and no variables can be defined.
[in]	options	Standard and vendor-specific options. Unknown options are ignored. A standard option begins with the "-hsa_" prefix. Options beginning with the "-hsa_ext_<extension_name>_" prefix are reserved for extensions. A vendor-specific option begins with the "-<vendor_name>_" prefix. Must be a NUL-terminated string. May be NULL.
[out]	executable	Memory location where the HSA runtime stores the newly created executable handle.

Return values

HSA_STATUS_SUCCESS	The function has been executed successfully.
HSA_STATUS_ERROR_NOT_INITIALIZED	The HSA runtime has not been initialized.
HSA_STATUS_ERROR_OUT_OF_RESOURCES	The HSA runtime failed to allocate the required resources.
HSA_STATUS_ERROR_INVALID_ARGUMENT	profile is invalid, or executable is NULL.

hsa_executable_create_alt()

hsa_status_t HSA_API hsa_executable_create_alt	(	hsa_profile_t	profile,
		hsa_default_float_rounding_mode_t	default_float_rounding_mode,
		const char *	options,
		hsa_executable_t *	executable )

Create an empty executable.

Parameters

[in]	profile	Profile used in the executable.
[in]	default_float_rounding_mode	Default floating-point rounding mode used in the executable. Allowed rounding modes are near and zero (default is not allowed).
[in]	options	Standard and vendor-specific options. Unknown options are ignored. A standard option begins with the "-hsa_" prefix. Options beginning with the "-hsa_ext_<extension_name>_" prefix are reserved for extensions. A vendor-specific option begins with the "-<vendor_name>_" prefix. Must be a NUL-terminated string. May be NULL.
[out]	executable	Memory location where the HSA runtime stores newly created executable handle. The initial state of the executable is HSA_EXECUTABLE_STATE_UNFROZEN.

Return values

HSA_STATUS_SUCCESS	The function has been executed successfully.
HSA_STATUS_ERROR_NOT_INITIALIZED	The HSA runtime has not been initialized.
HSA_STATUS_ERROR_OUT_OF_RESOURCES	The HSA runtime failed to allocate the required resources.
HSA_STATUS_ERROR_INVALID_ARGUMENT	profile is invalid, or executable is NULL.

hsa_executable_destroy()

hsa_status_t HSA_API hsa_executable_destroy

(

hsa_executable_t

executable

)

Destroy an executable.

An executable handle becomes invalid after the executable has been destroyed. Code object handles that were loaded into this executable are still valid after the executable has been destroyed, and can be used as intended. Resources allocated outside and associated with this executable (such as external global or readonly variables) can be released after the executable has been destroyed.

Executable should not be destroyed while kernels are in flight.

Parameters

[in]

executable

Executable.

Return values

HSA_STATUS_SUCCESS	The function has been executed successfully.
HSA_STATUS_ERROR_NOT_INITIALIZED	The HSA runtime has not been initialized.
HSA_STATUS_ERROR_INVALID_EXECUTABLE	The executable is invalid.

hsa_executable_freeze()

hsa_status_t HSA_API hsa_executable_freeze	(	hsa_executable_t	executable,
		const char *	options )

Freeze the executable.

No modifications to executable can be made after freezing: no code objects can be loaded to the executable, and no external variables can be defined. Freezing the executable does not prevent querying the executable's attributes. The application must define all the external variables in an executable before freezing it.

Parameters

[in]	executable	Executable.
[in]	options	Standard and vendor-specific options. Unknown options are ignored. A standard option begins with the "-hsa_" prefix. Options beginning with the "-hsa_ext_<extension_name>_" prefix are reserved for extensions. A vendor-specific option begins with the "-<vendor_name>_" prefix. Must be a NUL-terminated string. May be NULL.

Return values

HSA_STATUS_SUCCESS	The function has been executed successfully.
HSA_STATUS_ERROR_NOT_INITIALIZED	The HSA runtime has not been initialized.
HSA_STATUS_ERROR_INVALID_EXECUTABLE	The executable is invalid.
HSA_STATUS_ERROR_VARIABLE_UNDEFINED	One or more variables are undefined in the executable.
HSA_STATUS_ERROR_FROZEN_EXECUTABLE	executable is already frozen.

hsa_executable_get_info()

hsa_status_t HSA_API hsa_executable_get_info	(	hsa_executable_t	executable,
		hsa_executable_info_t	attribute,
		void *	value )

Get the current value of an attribute for a given executable.

Parameters

[in]	executable	Executable.
[in]	attribute	Attribute to query.
[out]	value	Pointer to an application-allocated buffer where to store the value of the attribute. If the buffer passed by the application is not large enough to hold the value of attribute, the behavior is undefined.

Return values

HSA_STATUS_SUCCESS	The function has been executed successfully.
HSA_STATUS_ERROR_NOT_INITIALIZED	The HSA runtime has not been initialized.
HSA_STATUS_ERROR_INVALID_EXECUTABLE	The executable is invalid.
HSA_STATUS_ERROR_INVALID_ARGUMENT	attribute is an invalid executable attribute, or value is NULL.

hsa_executable_get_symbol()

hsa_status_t HSA_API HSA_DEPRECATED hsa_executable_get_symbol	(	hsa_executable_t	executable,
		const char *	module_name,
		const char *	symbol_name,
		hsa_agent_t	agent,
		int32_t	call_convention,
		hsa_executable_symbol_t *	symbol )

Get the symbol handle for a given a symbol name.

Deprecated

Use hsa_executable_get_symbol_by_name instead.

Parameters

[in]	executable	Executable.
[in]	module_name	Module name. Must be NULL if the symbol has program linkage.
[in]	symbol_name	Symbol name.
[in]	agent	Agent associated with the symbol. If the symbol is independent of any agent (for example, a variable with program allocation), this argument is ignored.
[in]	call_convention	Call convention associated with the symbol. If the symbol does not correspond to an indirect function, this argument is ignored.
[out]	symbol	Memory location where the HSA runtime stores the symbol handle.

Return values

HSA_STATUS_SUCCESS	The function has been executed successfully.
HSA_STATUS_ERROR_NOT_INITIALIZED	The HSA runtime has not been initialized.
HSA_STATUS_ERROR_INVALID_EXECUTABLE	The executable is invalid.
HSA_STATUS_ERROR_INVALID_SYMBOL_NAME	There is no symbol with a name that matches symbol_name.
HSA_STATUS_ERROR_INVALID_ARGUMENT	symbol_name is NULL, or symbol is NULL.

hsa_executable_get_symbol_by_name()

hsa_status_t HSA_API hsa_executable_get_symbol_by_name	(	hsa_executable_t	executable,
		const char *	symbol_name,
		const hsa_agent_t *	agent,
		hsa_executable_symbol_t *	symbol )

Retrieve the symbol handle corresponding to a given a symbol name.

Parameters

[in]	executable	Executable.
[in]	symbol_name	Symbol name. Must be a NUL-terminated character array. The Programmer's Reference Manual describes the standard name mangling scheme.
[in]	agent	Pointer to the agent for which the symbol with the given name is defined. If the symbol corresponding to the given name has program allocation, agent must be NULL.
[out]	symbol	Memory location where the HSA runtime stores the symbol handle. Must not be NULL.

Return values

HSA_STATUS_SUCCESS	The function has been executed successfully.
HSA_STATUS_ERROR_NOT_INITIALIZED	The HSA runtime has not been initialized.
HSA_STATUS_ERROR_INVALID_EXECUTABLE	The executable is invalid.
HSA_STATUS_ERROR_INVALID_SYMBOL_NAME	There is no symbol with a name that matches symbol_name.
HSA_STATUS_ERROR_INVALID_ARGUMENT	symbol_name is NULL, or symbol is NULL.

hsa_executable_global_variable_define()

hsa_status_t HSA_API hsa_executable_global_variable_define	(	hsa_executable_t	executable,
		const char *	variable_name,
		void *	address )

Define an external global variable with program allocation.

This function allows the application to provide the definition of a variable in the global segment memory with program allocation. The variable must be defined before loading a code object into an executable. In addition, code objects loaded must not define the variable.

Parameters

[in]	executable	Executable. Must not be in frozen state.
[in]	variable_name	Name of the variable. The Programmer's Reference Manual describes the standard name mangling scheme.
[in]	address	Address where the variable is defined. This address must be in global memory and can be read and written by any agent in the system. The application cannot deallocate the buffer pointed by address before executable is destroyed.

Return values

HSA_STATUS_SUCCESS	The function has been executed successfully.
HSA_STATUS_ERROR_NOT_INITIALIZED	The HSA runtime has not been initialized.
HSA_STATUS_ERROR_OUT_OF_RESOURCES	The HSA runtime failed to allocate the required resources.
HSA_STATUS_ERROR_INVALID_EXECUTABLE	The executable is invalid.
HSA_STATUS_ERROR_VARIABLE_ALREADY_DEFINED	The variable is already defined.
HSA_STATUS_ERROR_INVALID_SYMBOL_NAME	There is no variable with the variable_name.
HSA_STATUS_ERROR_FROZEN_EXECUTABLE	executable is frozen.
HSA_STATUS_ERROR_INVALID_ARGUMENT	variable_name is NULL.

hsa_executable_iterate_agent_symbols()

hsa_status_t HSA_API hsa_executable_iterate_agent_symbols	(	hsa_executable_t	executable,
		hsa_agent_t	agent,
		hsa_status_t(*	callback )(hsa_executable_t exec, hsa_agent_t agent, hsa_executable_symbol_t symbol, void *data),
		void *	data )

Iterate over the kernels, indirect functions, and agent allocation variables in an executable for a given agent, and invoke an application- defined callback on every iteration.

Parameters

[in]	executable	Executable.
[in]	agent	Agent.
[in]	callback	Callback to be invoked once per executable symbol. The HSA runtime passes three arguments to the callback: the executable, a symbol, and the application data. If callback returns a status other than HSA_STATUS_SUCCESS for a particular iteration, the traversal stops and hsa_executable_iterate_symbols returns that status value.
[in]	data	Application data that is passed to callback on every iteration. May be NULL.

Return values

HSA_STATUS_SUCCESS	The function has been executed successfully.
HSA_STATUS_ERROR_NOT_INITIALIZED	The HSA runtime has not been initialized.
HSA_STATUS_ERROR_INVALID_EXECUTABLE	The executable is invalid.
HSA_STATUS_ERROR_INVALID_ARGUMENT	callback is NULL.

hsa_executable_iterate_program_symbols()

hsa_status_t HSA_API hsa_executable_iterate_program_symbols	(	hsa_executable_t	executable,
		hsa_status_t(*	callback )(hsa_executable_t exec, hsa_executable_symbol_t symbol, void *data),
		void *	data )

Iterate over the program allocation variables in an executable, and invoke an application-defined callback on every iteration.

Parameters

[in]	executable	Executable.
[in]	callback	Callback to be invoked once per executable symbol. The HSA runtime passes three arguments to the callback: the executable, a symbol, and the application data. If callback returns a status other than HSA_STATUS_SUCCESS for a particular iteration, the traversal stops and hsa_executable_iterate_symbols returns that status value.
[in]	data	Application data that is passed to callback on every iteration. May be NULL.

Return values

HSA_STATUS_SUCCESS	The function has been executed successfully.
HSA_STATUS_ERROR_NOT_INITIALIZED	The HSA runtime has not been initialized.
HSA_STATUS_ERROR_INVALID_EXECUTABLE	The executable is invalid.
HSA_STATUS_ERROR_INVALID_ARGUMENT	callback is NULL.

hsa_executable_iterate_symbols()

hsa_status_t HSA_API HSA_DEPRECATED hsa_executable_iterate_symbols	(	hsa_executable_t	executable,
		hsa_status_t(*	callback )(hsa_executable_t exec, hsa_executable_symbol_t symbol, void *data),
		void *	data )

Iterate over the symbols in a executable, and invoke an application-defined callback on every iteration.

Deprecated

Parameters

[in]	executable	Executable.
[in]	callback	Callback to be invoked once per executable symbol. The HSA runtime passes three arguments to the callback: the executable, a symbol, and the application data. If callback returns a status other than HSA_STATUS_SUCCESS for a particular iteration, the traversal stops and hsa_executable_iterate_symbols returns that status value.
[in]	data	Application data that is passed to callback on every iteration. May be NULL.

Return values

HSA_STATUS_SUCCESS	The function has been executed successfully.
HSA_STATUS_ERROR_NOT_INITIALIZED	The HSA runtime has not been initialized.
HSA_STATUS_ERROR_INVALID_EXECUTABLE	The executable is invalid.
HSA_STATUS_ERROR_INVALID_ARGUMENT	callback is NULL.

hsa_executable_load_agent_code_object()

hsa_status_t HSA_API hsa_executable_load_agent_code_object	(	hsa_executable_t	executable,
		hsa_agent_t	agent,
		hsa_code_object_reader_t	code_object_reader,
		const char *	options,
		hsa_loaded_code_object_t *	loaded_code_object )

Load an agent code object into an executable.

The agent code object contains all defined agent allocation variables, functions, indirect functions, and kernels in a given program for a given instruction set architecture.

Any module linkage declaration must have been defined either by a define variable or by loading a code object that has a symbol with module linkage definition.

The default floating-point rounding mode of the code object associated with code_object_reader must match that of the executable (HSA_EXECUTABLE_INFO_DEFAULT_FLOAT_ROUNDING_MODE), or be default (in which case the value of HSA_EXECUTABLE_INFO_DEFAULT_FLOAT_ROUNDING_MODE is used). If the agent code object uses extensions, the implementation and the agent must support them for this operation to return successfully.

Parameters

[in]	executable	Executable.
[in]	agent	Agent to load code object for. A code object can be loaded into an executable at most once for a given agent. The instruction set architecture of the code object must be supported by the agent.
[in]	code_object_reader	A code object reader that holds the code object to load. If a code object reader is destroyed before all the associated executables are destroyed, the behavior is undefined.
[in]	options	Standard and vendor-specific options. Unknown options are ignored. A standard option begins with the "-hsa_" prefix. Options beginning with the "-hsa_ext_<extension_name>_" prefix are reserved for extensions. A vendor-specific option begins with the "-<vendor_name>_" prefix. Must be a NUL-terminated string. May be NULL.
[out]	loaded_code_object	Pointer to a memory location where the HSA runtime stores the loaded code object handle. May be NULL.

Return values

HSA_STATUS_SUCCESS	The function has been executed successfully.
HSA_STATUS_ERROR_NOT_INITIALIZED	The HSA runtime has not been initialized.
HSA_STATUS_ERROR_OUT_OF_RESOURCES	The HSA runtime failed to allocate the required resources.
HSA_STATUS_ERROR_INVALID_EXECUTABLE	The executable is invalid.
HSA_STATUS_ERROR_FROZEN_EXECUTABLE	The executable is frozen.
HSA_STATUS_ERROR_INVALID_AGENT	The agent is invalid.
HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER	code_object_reader is invalid.
HSA_STATUS_ERROR_INCOMPATIBLE_ARGUMENTS	The code object read by code_object_reader is not compatible with the agent (for example, the agent does not support the instruction set architecture of the code object), the executable (for example, there is a default floating-point mode mismatch between the two), or the implementation.

hsa_executable_load_program_code_object()

hsa_status_t HSA_API hsa_executable_load_program_code_object	(	hsa_executable_t	executable,
		hsa_code_object_reader_t	code_object_reader,
		const char *	options,
		hsa_loaded_code_object_t *	loaded_code_object )

Load a program code object into an executable.

A program code object contains information about resources that are accessible by all kernel agents that run the executable, and can be loaded at most once into an executable.

If the program code object uses extensions, the implementation must support them for this operation to return successfully.

Parameters

[in]	executable	Executable.
[in]	code_object_reader	A code object reader that holds the program code object to load. If a code object reader is destroyed before all the associated executables are destroyed, the behavior is undefined.
[in]	options	Standard and vendor-specific options. Unknown options are ignored. A standard option begins with the "-hsa_" prefix. Options beginning with the "-hsa_ext_<extension_name>_" prefix are reserved for extensions. A vendor-specific option begins with the "-<vendor_name>_" prefix. Must be a NUL-terminated string. May be NULL.
[out]	loaded_code_object	Pointer to a memory location where the HSA runtime stores the loaded code object handle. May be NULL.

Return values

HSA_STATUS_SUCCESS	The function has been executed successfully.
HSA_STATUS_ERROR_NOT_INITIALIZED	The HSA runtime has not been initialized.
HSA_STATUS_ERROR_OUT_OF_RESOURCES	The HSA runtime failed to allocate the required resources.
HSA_STATUS_ERROR_INVALID_EXECUTABLE	The executable is invalid.
HSA_STATUS_ERROR_FROZEN_EXECUTABLE	The executable is frozen.
HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER	code_object_reader is invalid.
HSA_STATUS_ERROR_INCOMPATIBLE_ARGUMENTS	The program code object is not compatible with the executable or the implementation (for example, the code object uses an extension that is not supported by the implementation).

hsa_executable_readonly_variable_define()

hsa_status_t HSA_API hsa_executable_readonly_variable_define	(	hsa_executable_t	executable,
		hsa_agent_t	agent,
		const char *	variable_name,
		void *	address )

Define an external readonly variable.

This function allows the application to provide the definition of a variable in the readonly segment memory. The variable must be defined before loading a code object into an executable. In addition, code objects loaded must not define the variable.

Parameters

[in]	executable	Executable. Must not be in frozen state.
[in]	agent	Agent for which the variable is being defined.
[in]	variable_name	Name of the variable. The Programmer's Reference Manual describes the standard name mangling scheme.
[in]	address	Address where the variable is defined. This address must have been previously allocated using hsa_memory_allocate in a readonly region associated with agent. The application cannot deallocate the buffer pointed by address before executable is destroyed.
[in]	address	Address where the variable is defined. The buffer pointed by address is owned by the application, and cannot be deallocated before executable is destroyed.

Return values

HSA_STATUS_SUCCESS	The function has been executed successfully.
HSA_STATUS_ERROR_NOT_INITIALIZED	The HSA runtime has not been initialized.
HSA_STATUS_ERROR_OUT_OF_RESOURCES	The HSA runtime failed to allocate the required resources.
HSA_STATUS_ERROR_INVALID_EXECUTABLE	Executable is invalid.
HSA_STATUS_ERROR_INVALID_AGENT	agent is invalid.
HSA_STATUS_ERROR_VARIABLE_ALREADY_DEFINED	The variable is already defined.
HSA_STATUS_ERROR_INVALID_SYMBOL_NAME	There is no variable with the variable_name.
HSA_STATUS_ERROR_FROZEN_EXECUTABLE	executable is frozen.
HSA_STATUS_ERROR_INVALID_ARGUMENT	variable_name is NULL.

hsa_executable_symbol_get_info()

hsa_status_t HSA_API hsa_executable_symbol_get_info	(	hsa_executable_symbol_t	executable_symbol,
		hsa_executable_symbol_info_t	attribute,
		void *	value )

Get the current value of an attribute for a given executable symbol.

Parameters

[in]	executable_symbol	Executable symbol.
[in]	attribute	Attribute to query.
[out]	value	Pointer to an application-allocated buffer where to store the value of the attribute. If the buffer passed by the application is not large enough to hold the value of attribute, the behavior is undefined.

Return values

HSA_STATUS_SUCCESS	The function has been executed successfully.
HSA_STATUS_ERROR_NOT_INITIALIZED	The HSA runtime has not been initialized.
HSA_STATUS_ERROR_INVALID_EXECUTABLE_SYMBOL	The executable symbol is invalid.
HSA_STATUS_ERROR_INVALID_ARGUMENT	attribute is an invalid executable symbol attribute, or value is NULL.

hsa_executable_validate()

hsa_status_t HSA_API hsa_executable_validate	(	hsa_executable_t	executable,
		uint32_t *	result )

Validate an executable.

Checks that all code objects have matching machine model, profile, and default floating-point rounding mode. Checks that all declarations have definitions. Checks declaration-definition compatibility (see the HSA Programming Reference Manual for compatibility rules). Invoking this function is equivalent to invoking hsa_executable_validate_alt with no options.

Parameters

[in]	executable	Executable. Must be in frozen state.
[out]	result	Memory location where the HSA runtime stores the validation result. If the executable passes validation, the result is 0.

Return values

HSA_STATUS_SUCCESS	The function has been executed successfully.
HSA_STATUS_ERROR_NOT_INITIALIZED	The HSA runtime has not been initialized.
HSA_STATUS_ERROR_INVALID_EXECUTABLE	executable is invalid.
HSA_STATUS_ERROR_INVALID_ARGUMENT	result is NULL.

hsa_executable_validate_alt()

hsa_status_t HSA_API hsa_executable_validate_alt	(	hsa_executable_t	executable,
		const char *	options,
		uint32_t *	result )

Validate an executable.

Checks that all code objects have matching machine model, profile, and default floating-point rounding mode. Checks that all declarations have definitions. Checks declaration-definition compatibility (see the HSA Programming Reference Manual for compatibility rules).

Parameters

[in]	executable	Executable. Must be in frozen state.
[in]	options	Standard and vendor-specific options. Unknown options are ignored. A standard option begins with the "-hsa_" prefix. Options beginning with the "-hsa_ext_<extension_name>_" prefix are reserved for extensions. A vendor-specific option begins with the "-<vendor_name>_" prefix. Must be a NUL-terminated string. May be NULL.
[out]	result	Memory location where the HSA runtime stores the validation result. If the executable passes validation, the result is 0.

Return values

HSA_STATUS_SUCCESS	The function has been executed successfully.
HSA_STATUS_ERROR_NOT_INITIALIZED	The HSA runtime has not been initialized.
HSA_STATUS_ERROR_INVALID_EXECUTABLE	executable is invalid.
HSA_STATUS_ERROR_INVALID_ARGUMENT	result is NULL.