libomptarget/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h - llvm-project/openmp - Git at Google

 //===--- amdgpu/dynamic_hsa/hsa.h --------------------------------- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // The parts of the hsa api that are presently in use by the amdgpu plugin
 //
 //===----------------------------------------------------------------------===//
 #ifndef HSA_RUNTIME_INC_HSA_H_
 #define HSA_RUNTIME_INC_HSA_H_

 #include <stddef.h>
 #include <stdint.h>

 // Detect and set large model builds.
 #undef HSA_LARGE_MODEL
 #if defined(__LP64__) || defined(_M_X64)
 #define HSA_LARGE_MODEL
 #endif

 #ifdef __cplusplus
 extern "C" {
 #endif

 typedef enum {
   HSA_STATUS_SUCCESS = 0x0,
   HSA_STATUS_INFO_BREAK = 0x1,
   HSA_STATUS_ERROR = 0x1000,
   HSA_STATUS_ERROR_INVALID_CODE_OBJECT = 0x1010,
   HSA_STATUS_ERROR_NOT_INITIALIZED = 0x100B,
 } hsa_status_t;

 hsa_status_t hsa_status_string(hsa_status_t status, const char **status_string);

 typedef struct hsa_dim3_s {
   uint32_t x;
   uint32_t y;
   uint32_t z;
 } hsa_dim3_t;

 hsa_status_t hsa_init();

 hsa_status_t hsa_shut_down();

 typedef struct hsa_agent_s {
   uint64_t handle;
 } hsa_agent_t;

 typedef enum {
   HSA_DEVICE_TYPE_CPU = 0,
   HSA_DEVICE_TYPE_GPU = 1,
   HSA_DEVICE_TYPE_DSP = 2
 } hsa_device_type_t;

 typedef enum {
   HSA_ISA_INFO_NAME_LENGTH = 0,
   HSA_ISA_INFO_NAME = 1
 } hsa_isa_info_t;

 typedef enum {
   HSA_AGENT_INFO_NAME = 0,
   HSA_AGENT_INFO_VENDOR_NAME = 1,
   HSA_AGENT_INFO_FEATURE = 2,
   HSA_AGENT_INFO_PROFILE = 4,
   HSA_AGENT_INFO_WAVEFRONT_SIZE = 6,
   HSA_AGENT_INFO_WORKGROUP_MAX_DIM = 7,
   HSA_AGENT_INFO_WORKGROUP_MAX_SIZE = 8,
   HSA_AGENT_INFO_GRID_MAX_DIM = 9,
   HSA_AGENT_INFO_GRID_MAX_SIZE = 10,
   HSA_AGENT_INFO_FBARRIER_MAX_SIZE = 11,
   HSA_AGENT_INFO_QUEUES_MAX = 12,
   HSA_AGENT_INFO_QUEUE_MIN_SIZE = 13,
   HSA_AGENT_INFO_QUEUE_MAX_SIZE = 14,
   HSA_AGENT_INFO_NODE = 16,
   HSA_AGENT_INFO_DEVICE = 17,
   HSA_AGENT_INFO_CACHE_SIZE = 18,
   HSA_AGENT_INFO_FAST_F16_OPERATION = 24,
 } hsa_agent_info_t;

 typedef enum {
   HSA_SYSTEM_INFO_VERSION_MAJOR = 0,
   HSA_SYSTEM_INFO_VERSION_MINOR = 1,
 } hsa_system_info_t;

 typedef enum {
   HSA_AGENT_FEATURE_KERNEL_DISPATCH = 1,
   HSA_AGENT_FEATURE_AGENT_DISPATCH = 2,
 } hsa_agent_feature_t;

 typedef struct hsa_region_s {
   uint64_t handle;
 } hsa_region_t;

 typedef struct hsa_isa_s {
   uint64_t handle;
 } hsa_isa_t;

 hsa_status_t hsa_system_get_info(hsa_system_info_t attribute, void *value);

 hsa_status_t hsa_agent_get_info(hsa_agent_t agent, hsa_agent_info_t attribute,
                                 void *value);

 hsa_status_t hsa_isa_get_info_alt(hsa_isa_t isa, hsa_isa_info_t attribute,
                                   void *value);

 hsa_status_t hsa_iterate_agents(hsa_status_t (*callback)(hsa_agent_t agent,
                                                          void *data),
                                 void *data);

 hsa_status_t hsa_agent_iterate_isas(hsa_agent_t agent,
                                     hsa_status_t (*callback)(hsa_isa_t isa,
                                                              void *data),
                                     void *data);

 typedef struct hsa_signal_s {
   uint64_t handle;
 } hsa_signal_t;

 #ifdef HSA_LARGE_MODEL
 typedef int64_t hsa_signal_value_t;
 #else
 typedef int32_t hsa_signal_value_t;
 #endif

 hsa_status_t hsa_signal_create(hsa_signal_value_t initial_value,
                                uint32_t num_consumers,
                                const hsa_agent_t *consumers,
                                hsa_signal_t *signal);

 hsa_status_t hsa_amd_signal_create(hsa_signal_value_t initial_value,
                                    uint32_t num_consumers,
                                    const hsa_agent_t *consumers,
                                    uint64_t attributes, hsa_signal_t *signal);

 hsa_status_t hsa_signal_destroy(hsa_signal_t signal);

 void hsa_signal_store_relaxed(hsa_signal_t signal, hsa_signal_value_t value);

 void hsa_signal_store_screlease(hsa_signal_t signal, hsa_signal_value_t value);

 hsa_signal_value_t hsa_signal_load_scacquire(hsa_signal_t signal);

 void hsa_signal_subtract_screlease(hsa_signal_t signal,
                                    hsa_signal_value_t value);

 typedef enum {
   HSA_SIGNAL_CONDITION_EQ = 0,
   HSA_SIGNAL_CONDITION_NE = 1,
 } hsa_signal_condition_t;

 typedef enum {
   HSA_WAIT_STATE_BLOCKED = 0,
   HSA_WAIT_STATE_ACTIVE = 1
 } hsa_wait_state_t;

 hsa_signal_value_t hsa_signal_wait_scacquire(hsa_signal_t signal,
                                              hsa_signal_condition_t condition,
                                              hsa_signal_value_t compare_value,
                                              uint64_t timeout_hint,
                                              hsa_wait_state_t wait_state_hint);

 typedef enum {
   HSA_QUEUE_TYPE_MULTI = 0,
   HSA_QUEUE_TYPE_SINGLE = 1,
 } hsa_queue_type_t;

 typedef enum {
   HSA_QUEUE_FEATURE_KERNEL_DISPATCH = 1,
   HSA_QUEUE_FEATURE_AGENT_DISPATCH = 2
 } hsa_queue_feature_t;

 typedef uint32_t hsa_queue_type32_t;

 typedef struct hsa_queue_s {
   hsa_queue_type32_t type;
   uint32_t features;

 #ifdef HSA_LARGE_MODEL
   void *base_address;
 #elif defined HSA_LITTLE_ENDIAN
   void *base_address;
   uint32_t reserved0;
 #else
   uint32_t reserved0;
   void *base_address;
 #endif
   hsa_signal_t doorbell_signal;
   uint32_t size;
   uint32_t reserved1;
   uint64_t id;
 } hsa_queue_t;

 hsa_status_t hsa_queue_create(hsa_agent_t agent, uint32_t size,
                               hsa_queue_type32_t type,
                               void (*callback)(hsa_status_t status,
                                                hsa_queue_t *source, void *data),
                               void *data, uint32_t private_segment_size,
                               uint32_t group_segment_size, hsa_queue_t **queue);

 hsa_status_t hsa_queue_destroy(hsa_queue_t *queue);

 uint64_t hsa_queue_load_read_index_scacquire(const hsa_queue_t *queue);

 uint64_t hsa_queue_add_write_index_relaxed(const hsa_queue_t *queue,
                                            uint64_t value);

 typedef enum {
   HSA_PACKET_TYPE_KERNEL_DISPATCH = 2,
   HSA_PACKET_TYPE_BARRIER_AND = 3,
 } hsa_packet_type_t;

 typedef enum { HSA_FENCE_SCOPE_SYSTEM = 2 } hsa_fence_scope_t;

 typedef enum {
   HSA_PACKET_HEADER_TYPE = 0,
   HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE = 9,
   HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE = 11
 } hsa_packet_header_t;

 typedef enum {
   HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS = 0
 } hsa_kernel_dispatch_packet_setup_t;

 typedef enum {
   HSA_KERNEL_DISPATCH_PACKET_SETUP_WIDTH_DIMENSIONS = 2
 } hsa_kernel_dispatch_packet_setup_width_t;

 typedef struct hsa_kernel_dispatch_packet_s {
   uint16_t header;
   uint16_t setup;
   uint16_t workgroup_size_x;
   uint16_t workgroup_size_y;
   uint16_t workgroup_size_z;
   uint16_t reserved0;
   uint32_t grid_size_x;
   uint32_t grid_size_y;
   uint32_t grid_size_z;
   uint32_t private_segment_size;
   uint32_t group_segment_size;
   uint64_t kernel_object;
 #ifdef HSA_LARGE_MODEL
   void *kernarg_address;
 #elif defined HSA_LITTLE_ENDIAN
   void *kernarg_address;
   uint32_t reserved1;
 #else
   uint32_t reserved1;
   void *kernarg_address;
 #endif
   uint64_t reserved2;
   hsa_signal_t completion_signal;
 } hsa_kernel_dispatch_packet_t;

 typedef struct hsa_barrier_and_packet_s {
   uint16_t header;
   uint16_t reserved0;
   uint32_t reserved1;
   hsa_signal_t dep_signal[5];
   uint64_t reserved2;
   hsa_signal_t completion_signal;
 } hsa_barrier_and_packet_t;

 typedef enum { HSA_PROFILE_BASE = 0, HSA_PROFILE_FULL = 1 } hsa_profile_t;

 typedef enum {
   HSA_EXECUTABLE_STATE_UNFROZEN = 0,
   HSA_EXECUTABLE_STATE_FROZEN = 1
 } hsa_executable_state_t;

 typedef struct hsa_executable_s {
   uint64_t handle;
 } hsa_executable_t;

 typedef struct hsa_executable_symbol_s {
   uint64_t handle;
 } hsa_executable_symbol_t;

 typedef enum {
   HSA_EXECUTABLE_SYMBOL_INFO_TYPE = 0,
   HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH = 1,
   HSA_EXECUTABLE_SYMBOL_INFO_NAME = 2,
   HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS = 21,
   HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE = 9,
   HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT = 22,
   HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE = 11,
   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_t;

 typedef struct hsa_code_object_s {
   uint64_t handle;
 } hsa_code_object_t;

 typedef enum {
   HSA_SYMBOL_KIND_VARIABLE = 0,
   HSA_SYMBOL_KIND_KERNEL = 1,
   HSA_SYMBOL_KIND_INDIRECT_FUNCTION = 2
 } hsa_symbol_kind_t;

 typedef enum {
   HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT = 0,
   HSA_DEFAULT_FLOAT_ROUNDING_MODE_ZERO = 1,
   HSA_DEFAULT_FLOAT_ROUNDING_MODE_NEAR = 2,
 } hsa_default_float_rounding_mode_t;

 hsa_status_t hsa_memory_copy(void *dst, const void *src, size_t size);

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

 hsa_status_t 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);

 hsa_status_t hsa_executable_destroy(hsa_executable_t executable);

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

 hsa_status_t hsa_executable_validate(hsa_executable_t executable,
                                      uint32_t *result);

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

 hsa_status_t 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);

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

 hsa_status_t hsa_code_object_deserialize(void *serialized_code_object,
                                          size_t serialized_code_object_size,
                                          const char *options,
                                          hsa_code_object_t *code_object);

 hsa_status_t hsa_executable_load_code_object(hsa_executable_t executable,
                                              hsa_agent_t agent,
                                              hsa_code_object_t code_object,
                                              const char *options);

 hsa_status_t hsa_code_object_destroy(hsa_code_object_t code_object);

 typedef bool (*hsa_amd_signal_handler)(hsa_signal_value_t value, void *arg);

 hsa_status_t hsa_amd_signal_async_handler(hsa_signal_t signal,
                                           hsa_signal_condition_t cond,
                                           hsa_signal_value_t value,
                                           hsa_amd_signal_handler handler,
                                           void *arg);

 #ifdef __cplusplus
 }
 #endif

 #endif
	//===--- amdgpu/dynamic_hsa/hsa.h --------------------------------- C++ -*-===//
	//
	// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
	// See https://llvm.org/LICENSE.txt for license information.
	// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
	//
	//===----------------------------------------------------------------------===//
	//
	// The parts of the hsa api that are presently in use by the amdgpu plugin
	//
	//===----------------------------------------------------------------------===//
	#ifndef HSA_RUNTIME_INC_HSA_H_
	#define HSA_RUNTIME_INC_HSA_H_

	#include <stddef.h>
	#include <stdint.h>

	// Detect and set large model builds.
	#undef HSA_LARGE_MODEL
	#if defined(__LP64__) \|\| defined(_M_X64)
	#define HSA_LARGE_MODEL
	#endif

	#ifdef __cplusplus
	extern "C" {
	#endif

	typedef enum {
	HSA_STATUS_SUCCESS = 0x0,
	HSA_STATUS_INFO_BREAK = 0x1,
	HSA_STATUS_ERROR = 0x1000,
	HSA_STATUS_ERROR_INVALID_CODE_OBJECT = 0x1010,
	HSA_STATUS_ERROR_NOT_INITIALIZED = 0x100B,
	} hsa_status_t;

	hsa_status_t hsa_status_string(hsa_status_t status, const char **status_string);

	typedef struct hsa_dim3_s {
	uint32_t x;
	uint32_t y;
	uint32_t z;
	} hsa_dim3_t;

	hsa_status_t hsa_init();

	hsa_status_t hsa_shut_down();

	typedef struct hsa_agent_s {
	uint64_t handle;
	} hsa_agent_t;

	typedef enum {
	HSA_DEVICE_TYPE_CPU = 0,
	HSA_DEVICE_TYPE_GPU = 1,
	HSA_DEVICE_TYPE_DSP = 2
	} hsa_device_type_t;

	typedef enum {
	HSA_ISA_INFO_NAME_LENGTH = 0,
	HSA_ISA_INFO_NAME = 1
	} hsa_isa_info_t;

	typedef enum {
	HSA_AGENT_INFO_NAME = 0,
	HSA_AGENT_INFO_VENDOR_NAME = 1,
	HSA_AGENT_INFO_FEATURE = 2,
	HSA_AGENT_INFO_PROFILE = 4,
	HSA_AGENT_INFO_WAVEFRONT_SIZE = 6,
	HSA_AGENT_INFO_WORKGROUP_MAX_DIM = 7,
	HSA_AGENT_INFO_WORKGROUP_MAX_SIZE = 8,
	HSA_AGENT_INFO_GRID_MAX_DIM = 9,
	HSA_AGENT_INFO_GRID_MAX_SIZE = 10,
	HSA_AGENT_INFO_FBARRIER_MAX_SIZE = 11,
	HSA_AGENT_INFO_QUEUES_MAX = 12,
	HSA_AGENT_INFO_QUEUE_MIN_SIZE = 13,
	HSA_AGENT_INFO_QUEUE_MAX_SIZE = 14,
	HSA_AGENT_INFO_NODE = 16,
	HSA_AGENT_INFO_DEVICE = 17,
	HSA_AGENT_INFO_CACHE_SIZE = 18,
	HSA_AGENT_INFO_FAST_F16_OPERATION = 24,
	} hsa_agent_info_t;

	typedef enum {
	HSA_SYSTEM_INFO_VERSION_MAJOR = 0,
	HSA_SYSTEM_INFO_VERSION_MINOR = 1,
	} hsa_system_info_t;

	typedef enum {
	HSA_AGENT_FEATURE_KERNEL_DISPATCH = 1,
	HSA_AGENT_FEATURE_AGENT_DISPATCH = 2,
	} hsa_agent_feature_t;

	typedef struct hsa_region_s {
	uint64_t handle;
	} hsa_region_t;

	typedef struct hsa_isa_s {
	uint64_t handle;
	} hsa_isa_t;

	hsa_status_t hsa_system_get_info(hsa_system_info_t attribute, void *value);

	hsa_status_t hsa_agent_get_info(hsa_agent_t agent, hsa_agent_info_t attribute,
	void *value);

	hsa_status_t hsa_isa_get_info_alt(hsa_isa_t isa, hsa_isa_info_t attribute,
	void *value);

	hsa_status_t hsa_iterate_agents(hsa_status_t (*callback)(hsa_agent_t agent,
	void *data),
	void *data);

	hsa_status_t hsa_agent_iterate_isas(hsa_agent_t agent,
	hsa_status_t (*callback)(hsa_isa_t isa,
	void *data),
	void *data);

	typedef struct hsa_signal_s {
	uint64_t handle;
	} hsa_signal_t;

	#ifdef HSA_LARGE_MODEL
	typedef int64_t hsa_signal_value_t;
	#else
	typedef int32_t hsa_signal_value_t;
	#endif

	hsa_status_t hsa_signal_create(hsa_signal_value_t initial_value,
	uint32_t num_consumers,
	const hsa_agent_t *consumers,
	hsa_signal_t *signal);

	hsa_status_t hsa_amd_signal_create(hsa_signal_value_t initial_value,
	uint32_t num_consumers,
	const hsa_agent_t *consumers,
	uint64_t attributes, hsa_signal_t *signal);

	hsa_status_t hsa_signal_destroy(hsa_signal_t signal);

	void hsa_signal_store_relaxed(hsa_signal_t signal, hsa_signal_value_t value);

	void hsa_signal_store_screlease(hsa_signal_t signal, hsa_signal_value_t value);

	hsa_signal_value_t hsa_signal_load_scacquire(hsa_signal_t signal);

	void hsa_signal_subtract_screlease(hsa_signal_t signal,
	hsa_signal_value_t value);

	typedef enum {
	HSA_SIGNAL_CONDITION_EQ = 0,
	HSA_SIGNAL_CONDITION_NE = 1,
	} hsa_signal_condition_t;

	typedef enum {
	HSA_WAIT_STATE_BLOCKED = 0,
	HSA_WAIT_STATE_ACTIVE = 1
	} hsa_wait_state_t;

	hsa_signal_value_t hsa_signal_wait_scacquire(hsa_signal_t signal,
	hsa_signal_condition_t condition,
	hsa_signal_value_t compare_value,
	uint64_t timeout_hint,
	hsa_wait_state_t wait_state_hint);

	typedef enum {
	HSA_QUEUE_TYPE_MULTI = 0,
	HSA_QUEUE_TYPE_SINGLE = 1,
	} hsa_queue_type_t;

	typedef enum {
	HSA_QUEUE_FEATURE_KERNEL_DISPATCH = 1,
	HSA_QUEUE_FEATURE_AGENT_DISPATCH = 2
	} hsa_queue_feature_t;

	typedef uint32_t hsa_queue_type32_t;

	typedef struct hsa_queue_s {
	hsa_queue_type32_t type;
	uint32_t features;

	#ifdef HSA_LARGE_MODEL
	void *base_address;
	#elif defined HSA_LITTLE_ENDIAN
	void *base_address;
	uint32_t reserved0;
	#else
	uint32_t reserved0;
	void *base_address;
	#endif
	hsa_signal_t doorbell_signal;
	uint32_t size;
	uint32_t reserved1;
	uint64_t id;
	} hsa_queue_t;

	hsa_status_t hsa_queue_create(hsa_agent_t agent, uint32_t size,
	hsa_queue_type32_t type,
	void (*callback)(hsa_status_t status,
	hsa_queue_t source, void data),
	void *data, uint32_t private_segment_size,
	uint32_t group_segment_size, hsa_queue_t **queue);

	hsa_status_t hsa_queue_destroy(hsa_queue_t *queue);

	uint64_t hsa_queue_load_read_index_scacquire(const hsa_queue_t *queue);

	uint64_t hsa_queue_add_write_index_relaxed(const hsa_queue_t *queue,
	uint64_t value);

	typedef enum {
	HSA_PACKET_TYPE_KERNEL_DISPATCH = 2,
	HSA_PACKET_TYPE_BARRIER_AND = 3,
	} hsa_packet_type_t;

	typedef enum { HSA_FENCE_SCOPE_SYSTEM = 2 } hsa_fence_scope_t;

	typedef enum {
	HSA_PACKET_HEADER_TYPE = 0,
	HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE = 9,
	HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE = 11
	} hsa_packet_header_t;

	typedef enum {
	HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS = 0
	} hsa_kernel_dispatch_packet_setup_t;

	typedef enum {
	HSA_KERNEL_DISPATCH_PACKET_SETUP_WIDTH_DIMENSIONS = 2
	} hsa_kernel_dispatch_packet_setup_width_t;

	typedef struct hsa_kernel_dispatch_packet_s {
	uint16_t header;
	uint16_t setup;
	uint16_t workgroup_size_x;
	uint16_t workgroup_size_y;
	uint16_t workgroup_size_z;
	uint16_t reserved0;
	uint32_t grid_size_x;
	uint32_t grid_size_y;
	uint32_t grid_size_z;
	uint32_t private_segment_size;
	uint32_t group_segment_size;
	uint64_t kernel_object;
	#ifdef HSA_LARGE_MODEL
	void *kernarg_address;
	#elif defined HSA_LITTLE_ENDIAN
	void *kernarg_address;
	uint32_t reserved1;
	#else
	uint32_t reserved1;
	void *kernarg_address;
	#endif
	uint64_t reserved2;
	hsa_signal_t completion_signal;
	} hsa_kernel_dispatch_packet_t;

	typedef struct hsa_barrier_and_packet_s {
	uint16_t header;
	uint16_t reserved0;
	uint32_t reserved1;
	hsa_signal_t dep_signal[5];
	uint64_t reserved2;
	hsa_signal_t completion_signal;
	} hsa_barrier_and_packet_t;

	typedef enum { HSA_PROFILE_BASE = 0, HSA_PROFILE_FULL = 1 } hsa_profile_t;

	typedef enum {
	HSA_EXECUTABLE_STATE_UNFROZEN = 0,
	HSA_EXECUTABLE_STATE_FROZEN = 1
	} hsa_executable_state_t;

	typedef struct hsa_executable_s {
	uint64_t handle;
	} hsa_executable_t;

	typedef struct hsa_executable_symbol_s {
	uint64_t handle;
	} hsa_executable_symbol_t;

	typedef enum {
	HSA_EXECUTABLE_SYMBOL_INFO_TYPE = 0,
	HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH = 1,
	HSA_EXECUTABLE_SYMBOL_INFO_NAME = 2,
	HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS = 21,
	HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE = 9,
	HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT = 22,
	HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE = 11,
	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_t;

	typedef struct hsa_code_object_s {
	uint64_t handle;
	} hsa_code_object_t;

	typedef enum {
	HSA_SYMBOL_KIND_VARIABLE = 0,
	HSA_SYMBOL_KIND_KERNEL = 1,
	HSA_SYMBOL_KIND_INDIRECT_FUNCTION = 2
	} hsa_symbol_kind_t;

	typedef enum {
	HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT = 0,
	HSA_DEFAULT_FLOAT_ROUNDING_MODE_ZERO = 1,
	HSA_DEFAULT_FLOAT_ROUNDING_MODE_NEAR = 2,
	} hsa_default_float_rounding_mode_t;

	hsa_status_t hsa_memory_copy(void dst, const void src, size_t size);

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

	hsa_status_t 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);

	hsa_status_t hsa_executable_destroy(hsa_executable_t executable);

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

	hsa_status_t hsa_executable_validate(hsa_executable_t executable,
	uint32_t *result);

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

	hsa_status_t 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);

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

	hsa_status_t hsa_code_object_deserialize(void *serialized_code_object,
	size_t serialized_code_object_size,
	const char *options,
	hsa_code_object_t *code_object);

	hsa_status_t hsa_executable_load_code_object(hsa_executable_t executable,
	hsa_agent_t agent,
	hsa_code_object_t code_object,
	const char *options);

	hsa_status_t hsa_code_object_destroy(hsa_code_object_t code_object);

	typedef bool (hsa_amd_signal_handler)(hsa_signal_value_t value, void arg);

	hsa_status_t hsa_amd_signal_async_handler(hsa_signal_t signal,
	hsa_signal_condition_t cond,
	hsa_signal_value_t value,
	hsa_amd_signal_handler handler,
	void *arg);

	#ifdef __cplusplus
	}
	#endif

	#endif