hsa.h

Go to the documentation of this file.

//
// The University of Illinois/NCSA
// Open Source License (NCSA)
// 
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
// 
// Developed by:
// 
//                 AMD Research and AMD HSA Software Development
// 
//                 Advanced Micro Devices, Inc.
// 
//                 www.amd.com
// 
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with the Software without restriction, including without limitation
// the rights to use, copy, modify, merge, publish, distribute, sublicense,
// and/or sell copies of the Software, and to permit persons to whom the
// Software is furnished to do so, subject to the following conditions:
// 
//  - Redistributions of source code must retain the above copyright notice,
//    this list of conditions and the following disclaimers.
//  - Redistributions in binary form must reproduce the above copyright
//    notice, this list of conditions and the following disclaimers in
//    the documentation and/or other materials provided with the distribution.
//  - Neither the names of Advanced Micro Devices, Inc,
//    nor the names of its contributors may be used to endorse or promote
//    products derived from this Software without specific prior written
//    permission.
// 
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
// THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR
// OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE,
// ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
// DEALINGS WITH THE SOFTWARE.
//

#ifndef HSA_RUNTIME_INC_HSA_H_
#define HSA_RUNTIME_INC_HSA_H_

#include <stddef.h>   /* size_t */
#include <stdint.h>   /* uintXX_t */

#ifndef __cplusplus
#include <stdbool.h>  /* bool */
#endif /* __cplusplus */

// Placeholder for calling convention and import/export macros
#ifndef HSA_CALL
#define HSA_CALL
#endif

#ifndef HSA_EXPORT_DECORATOR
#ifdef __GNUC__
#define HSA_EXPORT_DECORATOR __attribute__ ((visibility ("default")))
#else
#define HSA_EXPORT_DECORATOR
#endif
#endif
#define HSA_API_EXPORT HSA_EXPORT_DECORATOR HSA_CALL
#define HSA_API_IMPORT HSA_CALL

#if !defined(HSA_API) && defined(HSA_EXPORT)
#define HSA_API HSA_API_EXPORT
#else
#define HSA_API HSA_API_IMPORT
#endif

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

// Try to detect CPU endianness
#if !defined(LITTLEENDIAN_CPU) && !defined(BIGENDIAN_CPU)
#if defined(__i386__) || defined(__x86_64__) || defined(_M_IX86) || \
    defined(_M_X64)
#define LITTLEENDIAN_CPU
#endif
#endif

#undef HSA_LITTLE_ENDIAN
#if defined(LITTLEENDIAN_CPU)
#define HSA_LITTLE_ENDIAN
#elif defined(BIGENDIAN_CPU)
#else
#error "BIGENDIAN_CPU or LITTLEENDIAN_CPU must be defined"
#endif

#ifndef HSA_DEPRECATED
#define HSA_DEPRECATED
//#ifdef __GNUC__
//#define HSA_DEPRECATED __attribute__((deprecated))
//#else
//#define HSA_DEPRECATED __declspec(deprecated)
//#endif
#endif

#define HSA_VERSION_1_0                              1

#ifdef __cplusplus
extern "C" {
#endif  /* __cplusplus */

typedef enum {
    HSA_STATUS_SUCCESS = 0x0,
    HSA_STATUS_INFO_BREAK = 0x1,
    HSA_STATUS_ERROR = 0x1000,
    HSA_STATUS_ERROR_INVALID_ARGUMENT = 0x1001,
    HSA_STATUS_ERROR_INVALID_QUEUE_CREATION = 0x1002,
    HSA_STATUS_ERROR_INVALID_ALLOCATION = 0x1003,
    HSA_STATUS_ERROR_INVALID_AGENT = 0x1004,
    HSA_STATUS_ERROR_INVALID_REGION = 0x1005,
    HSA_STATUS_ERROR_INVALID_SIGNAL = 0x1006,
    HSA_STATUS_ERROR_INVALID_QUEUE = 0x1007,
    HSA_STATUS_ERROR_OUT_OF_RESOURCES = 0x1008,
    HSA_STATUS_ERROR_INVALID_PACKET_FORMAT = 0x1009,
    HSA_STATUS_ERROR_RESOURCE_FREE = 0x100A,
    HSA_STATUS_ERROR_NOT_INITIALIZED = 0x100B,
    HSA_STATUS_ERROR_REFCOUNT_OVERFLOW = 0x100C,
    HSA_STATUS_ERROR_INCOMPATIBLE_ARGUMENTS = 0x100D,
    HSA_STATUS_ERROR_INVALID_INDEX = 0x100E,
    HSA_STATUS_ERROR_INVALID_ISA = 0x100F,
    HSA_STATUS_ERROR_INVALID_ISA_NAME = 0x1017,
    HSA_STATUS_ERROR_INVALID_CODE_OBJECT = 0x1010,
    HSA_STATUS_ERROR_INVALID_EXECUTABLE = 0x1011,
    HSA_STATUS_ERROR_FROZEN_EXECUTABLE = 0x1012,
    HSA_STATUS_ERROR_INVALID_SYMBOL_NAME = 0x1013,
    HSA_STATUS_ERROR_VARIABLE_ALREADY_DEFINED = 0x1014,
    HSA_STATUS_ERROR_VARIABLE_UNDEFINED = 0x1015,
    HSA_STATUS_ERROR_EXCEPTION = 0x1016,
    HSA_STATUS_ERROR_INVALID_CODE_SYMBOL = 0x1018,
    HSA_STATUS_ERROR_INVALID_EXECUTABLE_SYMBOL = 0x1019,
    HSA_STATUS_ERROR_INVALID_FILE = 0x1020,
    HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER = 0x1021,
    HSA_STATUS_ERROR_INVALID_CACHE = 0x1022,
    HSA_STATUS_ERROR_INVALID_WAVEFRONT = 0x1023,
    HSA_STATUS_ERROR_INVALID_SIGNAL_GROUP = 0x1024,
    HSA_STATUS_ERROR_INVALID_RUNTIME_STATE = 0x1025

} hsa_status_t;

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

typedef enum {
  HSA_ACCESS_PERMISSION_RO = 1,
  HSA_ACCESS_PERMISSION_WO = 2,
  HSA_ACCESS_PERMISSION_RW = 3
} hsa_access_permission_t;

typedef int hsa_file_t;

hsa_status_t HSA_API hsa_init();

hsa_status_t HSA_API hsa_shut_down();

typedef enum {
    HSA_ENDIANNESS_LITTLE = 0,
    HSA_ENDIANNESS_BIG = 1
} hsa_endianness_t;

typedef enum {
    HSA_MACHINE_MODEL_SMALL = 0,
    HSA_MACHINE_MODEL_LARGE = 1
} hsa_machine_model_t;

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

typedef enum {
  HSA_SYSTEM_INFO_VERSION_MAJOR = 0,
  HSA_SYSTEM_INFO_VERSION_MINOR = 1,
  HSA_SYSTEM_INFO_TIMESTAMP = 2,
  HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY = 3,
  HSA_SYSTEM_INFO_SIGNAL_MAX_WAIT = 4,
  HSA_SYSTEM_INFO_ENDIANNESS = 5,
  HSA_SYSTEM_INFO_MACHINE_MODEL = 6,
  HSA_SYSTEM_INFO_EXTENSIONS = 7
} hsa_system_info_t;

hsa_status_t HSA_API hsa_system_get_info(
    hsa_system_info_t attribute,
    void* value);

typedef enum {
  HSA_EXTENSION_FINALIZER = 0,
  HSA_EXTENSION_IMAGES = 1,

  HSA_EXTENSION_PERFORMANCE_COUNTERS = 2,

  HSA_EXTENSION_PROFILING_EVENTS = 3,
  HSA_EXTENSION_STD_LAST = 3,
  HSA_AMD_FIRST_EXTENSION = 0x200,
  HSA_EXTENSION_AMD_PROFILER = 0x200,
  HSA_EXTENSION_AMD_LOADER = 0x201,
  HSA_AMD_LAST_EXTENSION = 0x201
} hsa_extension_t;

hsa_status_t HSA_API hsa_extension_get_name(
    uint16_t extension,
    const char **name);

hsa_status_t HSA_API HSA_DEPRECATED hsa_system_extension_supported(
    uint16_t extension,
    uint16_t version_major,
    uint16_t version_minor,
    bool* result);

hsa_status_t HSA_API hsa_system_major_extension_supported(
    uint16_t extension,
    uint16_t version_major,
    uint16_t *version_minor,
    bool* result);
    

hsa_status_t HSA_API HSA_DEPRECATED hsa_system_get_extension_table(
    uint16_t extension,
    uint16_t version_major,
    uint16_t version_minor,
    void *table);

hsa_status_t HSA_API hsa_system_get_major_extension_table(
    uint16_t extension,
    uint16_t version_major,
    size_t table_length,
    void *table);    

typedef struct hsa_agent_s {
  uint64_t handle;
} hsa_agent_t;

typedef enum {
    HSA_AGENT_FEATURE_KERNEL_DISPATCH = 1,
    HSA_AGENT_FEATURE_AGENT_DISPATCH = 2
} hsa_agent_feature_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_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;

typedef enum {
  HSA_AGENT_INFO_NAME = 0,
  HSA_AGENT_INFO_VENDOR_NAME = 1,
  HSA_AGENT_INFO_FEATURE = 2,
  HSA_AGENT_INFO_MACHINE_MODEL = 3,
  HSA_AGENT_INFO_PROFILE = 4,
  HSA_AGENT_INFO_DEFAULT_FLOAT_ROUNDING_MODE = 5,
  HSA_AGENT_INFO_BASE_PROFILE_DEFAULT_FLOAT_ROUNDING_MODES = 23,
  HSA_AGENT_INFO_FAST_F16_OPERATION = 24,
  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_QUEUE_TYPE = 15,
  HSA_AGENT_INFO_NODE = 16,
  HSA_AGENT_INFO_DEVICE = 17,
  HSA_AGENT_INFO_CACHE_SIZE = 18,
  HSA_AGENT_INFO_ISA = 19,
  HSA_AGENT_INFO_EXTENSIONS = 20,
  HSA_AGENT_INFO_VERSION_MAJOR = 21,
  HSA_AGENT_INFO_VERSION_MINOR = 22

} hsa_agent_info_t;

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

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

/*

// If we do not know the size of an attribute, we need to query it first
// Note: this API will not be in the spec unless needed
hsa_status_t HSA_API hsa_agent_get_info_size(
    hsa_agent_t agent,
    hsa_agent_info_t attribute,
    size_t* size);

// Set the value of an agents attribute
// Note: this API will not be in the spec unless needed
hsa_status_t HSA_API hsa_agent_set_info(
    hsa_agent_t agent,
    hsa_agent_info_t attribute,
    void* value);

*/

typedef enum {
    HSA_EXCEPTION_POLICY_BREAK = 1,
    HSA_EXCEPTION_POLICY_DETECT = 2
} hsa_exception_policy_t;

hsa_status_t HSA_API HSA_DEPRECATED hsa_agent_get_exception_policies(
    hsa_agent_t agent,
    hsa_profile_t profile,
    uint16_t *mask);

typedef struct hsa_cache_s {
  uint64_t handle;
} hsa_cache_t;

typedef enum {
  HSA_CACHE_INFO_NAME_LENGTH = 0,
  HSA_CACHE_INFO_NAME = 1,
  HSA_CACHE_INFO_LEVEL = 2,
  HSA_CACHE_INFO_SIZE = 3
} hsa_cache_info_t;

hsa_status_t HSA_API hsa_cache_get_info(
    hsa_cache_t cache,
    hsa_cache_info_t attribute,
    void* value);

hsa_status_t HSA_API hsa_agent_iterate_caches(
    hsa_agent_t agent,
    hsa_status_t (*callback)(hsa_cache_t cache, void* data),
    void* data);

hsa_status_t HSA_API HSA_DEPRECATED hsa_agent_extension_supported(
    uint16_t extension,
    hsa_agent_t agent,
    uint16_t version_major,
    uint16_t version_minor,
    bool* result);

hsa_status_t HSA_API hsa_agent_major_extension_supported(
    uint16_t extension,
    hsa_agent_t agent,
    uint16_t version_major,
    uint16_t *version_minor,
    bool* result);
    

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_API 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_API hsa_signal_destroy(
    hsa_signal_t signal);

hsa_signal_value_t HSA_API hsa_signal_load_scacquire(
    hsa_signal_t signal);

hsa_signal_value_t HSA_API hsa_signal_load_relaxed(
    hsa_signal_t signal);

hsa_signal_value_t HSA_API HSA_DEPRECATED hsa_signal_load_acquire(
    hsa_signal_t signal);

void HSA_API hsa_signal_store_relaxed(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API hsa_signal_store_screlease(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API HSA_DEPRECATED hsa_signal_store_release(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API hsa_signal_silent_store_relaxed(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API hsa_signal_silent_store_screlease(
    hsa_signal_t signal,
    hsa_signal_value_t value);

hsa_signal_value_t HSA_API hsa_signal_exchange_scacq_screl(
    hsa_signal_t signal,
    hsa_signal_value_t value);

hsa_signal_value_t HSA_API HSA_DEPRECATED hsa_signal_exchange_acq_rel(
    hsa_signal_t signal,
    hsa_signal_value_t value);

hsa_signal_value_t HSA_API hsa_signal_exchange_scacquire(
    hsa_signal_t signal,
    hsa_signal_value_t value);

hsa_signal_value_t HSA_API HSA_DEPRECATED hsa_signal_exchange_acquire(
    hsa_signal_t signal,
    hsa_signal_value_t value);

hsa_signal_value_t HSA_API hsa_signal_exchange_relaxed(
    hsa_signal_t signal,
    hsa_signal_value_t value);
hsa_signal_value_t HSA_API hsa_signal_exchange_screlease(
    hsa_signal_t signal,
    hsa_signal_value_t value);

hsa_signal_value_t HSA_API HSA_DEPRECATED hsa_signal_exchange_release(
    hsa_signal_t signal,
    hsa_signal_value_t value);

hsa_signal_value_t HSA_API hsa_signal_cas_scacq_screl(
    hsa_signal_t signal,
    hsa_signal_value_t expected,
    hsa_signal_value_t value);


hsa_signal_value_t HSA_API HSA_DEPRECATED hsa_signal_cas_acq_rel(
    hsa_signal_t signal,
    hsa_signal_value_t expected,
    hsa_signal_value_t value);

hsa_signal_value_t HSA_API hsa_signal_cas_scacquire(
    hsa_signal_t signal,
    hsa_signal_value_t expected,
    hsa_signal_value_t value);

hsa_signal_value_t HSA_API HSA_DEPRECATED hsa_signal_cas_acquire(
    hsa_signal_t signal,
    hsa_signal_value_t expected,
    hsa_signal_value_t value);

hsa_signal_value_t HSA_API hsa_signal_cas_relaxed(
    hsa_signal_t signal,
    hsa_signal_value_t expected,
    hsa_signal_value_t value);

hsa_signal_value_t HSA_API hsa_signal_cas_screlease(
    hsa_signal_t signal,
    hsa_signal_value_t expected,
    hsa_signal_value_t value);

hsa_signal_value_t HSA_API HSA_DEPRECATED hsa_signal_cas_release(
    hsa_signal_t signal,
    hsa_signal_value_t expected,
    hsa_signal_value_t value);

void HSA_API hsa_signal_add_scacq_screl(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API HSA_DEPRECATED hsa_signal_add_acq_rel(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API hsa_signal_add_scacquire(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API HSA_DEPRECATED hsa_signal_add_acquire(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API hsa_signal_add_relaxed(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API hsa_signal_add_screlease(
    hsa_signal_t signal,
    hsa_signal_value_t value);


void HSA_API HSA_DEPRECATED hsa_signal_add_release(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API hsa_signal_subtract_scacq_screl(
    hsa_signal_t signal,
    hsa_signal_value_t value);


void HSA_API HSA_DEPRECATED hsa_signal_subtract_acq_rel(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API hsa_signal_subtract_scacquire(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API HSA_DEPRECATED hsa_signal_subtract_acquire(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API hsa_signal_subtract_relaxed(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API hsa_signal_subtract_screlease(
    hsa_signal_t signal,
    hsa_signal_value_t value);


void HSA_API HSA_DEPRECATED hsa_signal_subtract_release(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API hsa_signal_and_scacq_screl(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API HSA_DEPRECATED hsa_signal_and_acq_rel(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API hsa_signal_and_scacquire(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API HSA_DEPRECATED hsa_signal_and_acquire(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API hsa_signal_and_relaxed(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API hsa_signal_and_screlease(
    hsa_signal_t signal,
    hsa_signal_value_t value);


void HSA_API HSA_DEPRECATED hsa_signal_and_release(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API hsa_signal_or_scacq_screl(
    hsa_signal_t signal,
    hsa_signal_value_t value);


void HSA_API HSA_DEPRECATED hsa_signal_or_acq_rel(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API hsa_signal_or_scacquire(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API HSA_DEPRECATED hsa_signal_or_acquire(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API hsa_signal_or_relaxed(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API hsa_signal_or_screlease(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API HSA_DEPRECATED hsa_signal_or_release(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API hsa_signal_xor_scacq_screl(
    hsa_signal_t signal,
    hsa_signal_value_t value);


void HSA_API HSA_DEPRECATED hsa_signal_xor_acq_rel(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API hsa_signal_xor_scacquire(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API HSA_DEPRECATED hsa_signal_xor_acquire(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API hsa_signal_xor_relaxed(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API hsa_signal_xor_screlease(
    hsa_signal_t signal,
    hsa_signal_value_t value);

void HSA_API HSA_DEPRECATED hsa_signal_xor_release(
    hsa_signal_t signal,
    hsa_signal_value_t value);

typedef enum {
    HSA_SIGNAL_CONDITION_EQ = 0,
    HSA_SIGNAL_CONDITION_NE = 1,
    HSA_SIGNAL_CONDITION_LT = 2,
    HSA_SIGNAL_CONDITION_GTE = 3
} 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_API 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);

hsa_signal_value_t HSA_API hsa_signal_wait_relaxed(
    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);

hsa_signal_value_t HSA_API HSA_DEPRECATED hsa_signal_wait_acquire(
    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 struct hsa_signal_group_s {
  uint64_t handle;
} hsa_signal_group_t;

hsa_status_t HSA_API hsa_signal_group_create(
    uint32_t num_signals,
    const hsa_signal_t *signals,
    uint32_t num_consumers,
    const hsa_agent_t *consumers,
    hsa_signal_group_t *signal_group);

hsa_status_t HSA_API hsa_signal_group_destroy(
    hsa_signal_group_t signal_group);

hsa_status_t HSA_API hsa_signal_group_wait_any_scacquire(
    hsa_signal_group_t signal_group,
    const hsa_signal_condition_t *conditions,
    const hsa_signal_value_t *compare_values,
    hsa_wait_state_t wait_state_hint,
    hsa_signal_t *signal,
    hsa_signal_value_t *value);

hsa_status_t HSA_API hsa_signal_group_wait_any_relaxed(
    hsa_signal_group_t signal_group,
    const hsa_signal_condition_t *conditions,
    const hsa_signal_value_t *compare_values,
    hsa_wait_state_t wait_state_hint,
    hsa_signal_t *signal,
    hsa_signal_value_t *value);

typedef struct hsa_region_s {
  uint64_t handle;
} hsa_region_t;

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

typedef uint32_t hsa_queue_type32_t;

typedef enum {
  HSA_QUEUE_FEATURE_KERNEL_DISPATCH = 1,

  HSA_QUEUE_FEATURE_AGENT_DISPATCH = 2
} hsa_queue_feature_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_API 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_API hsa_soft_queue_create(
    hsa_region_t region,
    uint32_t size,
    hsa_queue_type32_t type,
    uint32_t features,
    hsa_signal_t doorbell_signal,
    hsa_queue_t **queue);

hsa_status_t HSA_API hsa_queue_destroy(
    hsa_queue_t *queue);

hsa_status_t HSA_API hsa_queue_inactivate(
    hsa_queue_t *queue);

uint64_t HSA_API HSA_DEPRECATED hsa_queue_load_read_index_acquire(
    const hsa_queue_t *queue);

uint64_t HSA_API hsa_queue_load_read_index_scacquire(
    const hsa_queue_t *queue);

uint64_t HSA_API hsa_queue_load_read_index_relaxed(
    const hsa_queue_t *queue);

uint64_t HSA_API HSA_DEPRECATED hsa_queue_load_write_index_acquire(
    const hsa_queue_t *queue);

uint64_t HSA_API hsa_queue_load_write_index_scacquire(
    const hsa_queue_t *queue);

uint64_t HSA_API hsa_queue_load_write_index_relaxed(
    const hsa_queue_t *queue);

void HSA_API hsa_queue_store_write_index_relaxed(
    const hsa_queue_t *queue,
    uint64_t value);

void HSA_API HSA_DEPRECATED hsa_queue_store_write_index_release(
    const hsa_queue_t *queue,
    uint64_t value);

void HSA_API hsa_queue_store_write_index_screlease(
    const hsa_queue_t *queue,
    uint64_t value);

uint64_t HSA_API HSA_DEPRECATED hsa_queue_cas_write_index_acq_rel(
    const hsa_queue_t *queue,
    uint64_t expected,
    uint64_t value);

uint64_t HSA_API hsa_queue_cas_write_index_scacq_screl(
    const hsa_queue_t *queue,
    uint64_t expected,
    uint64_t value);

uint64_t HSA_API HSA_DEPRECATED hsa_queue_cas_write_index_acquire(
    const hsa_queue_t *queue,
    uint64_t expected,
    uint64_t value);

uint64_t HSA_API hsa_queue_cas_write_index_scacquire(
    const hsa_queue_t *queue,
    uint64_t expected,
    uint64_t value);

uint64_t HSA_API hsa_queue_cas_write_index_relaxed(
    const hsa_queue_t *queue,
    uint64_t expected,
    uint64_t value);

uint64_t HSA_API HSA_DEPRECATED hsa_queue_cas_write_index_release(
    const hsa_queue_t *queue,
    uint64_t expected,
    uint64_t value);

uint64_t HSA_API hsa_queue_cas_write_index_screlease(
    const hsa_queue_t *queue,
    uint64_t expected,
    uint64_t value);

uint64_t HSA_API HSA_DEPRECATED hsa_queue_add_write_index_acq_rel(
    const hsa_queue_t *queue,
    uint64_t value);

uint64_t HSA_API hsa_queue_add_write_index_scacq_screl(
    const hsa_queue_t *queue,
    uint64_t value);

uint64_t HSA_API HSA_DEPRECATED hsa_queue_add_write_index_acquire(
    const hsa_queue_t *queue,
    uint64_t value);

uint64_t HSA_API hsa_queue_add_write_index_scacquire(
    const hsa_queue_t *queue,
    uint64_t value);

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

uint64_t HSA_API HSA_DEPRECATED hsa_queue_add_write_index_release(
    const hsa_queue_t *queue,
    uint64_t value);

uint64_t HSA_API hsa_queue_add_write_index_screlease(
    const hsa_queue_t *queue,
    uint64_t value);

void HSA_API hsa_queue_store_read_index_relaxed(
    const hsa_queue_t *queue,
    uint64_t value);

void HSA_API HSA_DEPRECATED hsa_queue_store_read_index_release(
    const hsa_queue_t *queue,
    uint64_t value);

void HSA_API hsa_queue_store_read_index_screlease(
   const hsa_queue_t *queue,
   uint64_t value);
typedef enum {
  HSA_PACKET_TYPE_VENDOR_SPECIFIC = 0,
  HSA_PACKET_TYPE_INVALID = 1,
  HSA_PACKET_TYPE_KERNEL_DISPATCH = 2,
  HSA_PACKET_TYPE_BARRIER_AND = 3,
  HSA_PACKET_TYPE_AGENT_DISPATCH = 4,
  HSA_PACKET_TYPE_BARRIER_OR = 5
} hsa_packet_type_t;

typedef enum {
  HSA_FENCE_SCOPE_NONE = 0,
  HSA_FENCE_SCOPE_AGENT = 1,
  HSA_FENCE_SCOPE_SYSTEM = 2
} hsa_fence_scope_t;

 typedef enum {
   HSA_PACKET_HEADER_TYPE = 0,
   HSA_PACKET_HEADER_BARRIER = 8,
   HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE = 9,
   HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE = 9,
   HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE = 11,
   HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE = 11
 } hsa_packet_header_t;

 typedef enum {
   HSA_PACKET_HEADER_WIDTH_TYPE = 8,
   HSA_PACKET_HEADER_WIDTH_BARRIER = 1,
   HSA_PACKET_HEADER_WIDTH_SCACQUIRE_FENCE_SCOPE = 2,
   HSA_PACKET_HEADER_WIDTH_ACQUIRE_FENCE_SCOPE = 2,
   HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE = 2,
   HSA_PACKET_HEADER_WIDTH_RELEASE_FENCE_SCOPE = 2
 } hsa_packet_header_width_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_agent_dispatch_packet_s {
  uint16_t header;

  uint16_t type;

  uint32_t reserved0;

#ifdef HSA_LARGE_MODEL
  void* return_address;
#elif defined HSA_LITTLE_ENDIAN

  void* return_address;
  uint32_t reserved1;
#else
  uint32_t reserved1;
  void* return_address;
#endif

  uint64_t arg[4];

  uint64_t reserved2;

  hsa_signal_t completion_signal;

} hsa_agent_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 struct hsa_barrier_or_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_or_packet_t;

typedef enum {
  HSA_REGION_SEGMENT_GLOBAL = 0,
  HSA_REGION_SEGMENT_READONLY = 1,
  HSA_REGION_SEGMENT_PRIVATE = 2,
  HSA_REGION_SEGMENT_GROUP = 3,
  HSA_REGION_SEGMENT_KERNARG = 4
} hsa_region_segment_t;

typedef enum {
  HSA_REGION_GLOBAL_FLAG_KERNARG = 1,
  HSA_REGION_GLOBAL_FLAG_FINE_GRAINED = 2,
  HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED = 4
} hsa_region_global_flag_t;

typedef enum {
  HSA_REGION_INFO_SEGMENT = 0,
  HSA_REGION_INFO_GLOBAL_FLAGS = 1,
  HSA_REGION_INFO_SIZE = 2,
  HSA_REGION_INFO_ALLOC_MAX_SIZE = 4,
  HSA_REGION_INFO_ALLOC_MAX_PRIVATE_WORKGROUP_SIZE = 8,
  HSA_REGION_INFO_RUNTIME_ALLOC_ALLOWED = 5,
  HSA_REGION_INFO_RUNTIME_ALLOC_GRANULE = 6,
  HSA_REGION_INFO_RUNTIME_ALLOC_ALIGNMENT = 7
} hsa_region_info_t;

hsa_status_t HSA_API hsa_region_get_info(
    hsa_region_t region,
    hsa_region_info_t attribute,
    void* value);

hsa_status_t HSA_API hsa_agent_iterate_regions(
    hsa_agent_t agent,
    hsa_status_t (*callback)(hsa_region_t region, void* data),
    void* data);

hsa_status_t HSA_API hsa_memory_allocate(hsa_region_t region,
    size_t size,
    void** ptr);

hsa_status_t HSA_API hsa_memory_free(void* ptr);

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

hsa_status_t HSA_API hsa_memory_assign_agent(
    void *ptr,
    hsa_agent_t agent,
    hsa_access_permission_t access);

hsa_status_t HSA_API hsa_memory_register(
    void *ptr,
    size_t size);

hsa_status_t HSA_API hsa_memory_deregister(
    void *ptr,
    size_t size);

typedef struct hsa_isa_s {
  uint64_t handle;
} hsa_isa_t;

hsa_status_t HSA_API hsa_isa_from_name(
    const char *name,
    hsa_isa_t *isa);

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

typedef enum {
  HSA_ISA_INFO_NAME_LENGTH = 0,
  HSA_ISA_INFO_NAME = 1,
  HSA_ISA_INFO_CALL_CONVENTION_COUNT = 2,
  HSA_ISA_INFO_CALL_CONVENTION_INFO_WAVEFRONT_SIZE = 3,
  HSA_ISA_INFO_CALL_CONVENTION_INFO_WAVEFRONTS_PER_COMPUTE_UNIT = 4,
  HSA_ISA_INFO_MACHINE_MODELS = 5,
  HSA_ISA_INFO_PROFILES = 6,
  HSA_ISA_INFO_DEFAULT_FLOAT_ROUNDING_MODES = 7,
  HSA_ISA_INFO_BASE_PROFILE_DEFAULT_FLOAT_ROUNDING_MODES = 8,
  HSA_ISA_INFO_FAST_F16_OPERATION = 9,
  HSA_ISA_INFO_WORKGROUP_MAX_DIM = 12,
  HSA_ISA_INFO_WORKGROUP_MAX_SIZE = 13,
  HSA_ISA_INFO_GRID_MAX_DIM = 14,
  HSA_ISA_INFO_GRID_MAX_SIZE = 16,
  HSA_ISA_INFO_FBARRIER_MAX_SIZE = 17
} hsa_isa_info_t;

hsa_status_t HSA_API HSA_DEPRECATED hsa_isa_get_info(
    hsa_isa_t isa,
    hsa_isa_info_t attribute,
    uint32_t index,
    void *value);

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

hsa_status_t HSA_API hsa_isa_get_exception_policies(
    hsa_isa_t isa,
    hsa_profile_t profile,
    uint16_t *mask);

typedef enum {
  HSA_FP_TYPE_16 = 1,
  HSA_FP_TYPE_32 = 2,
  HSA_FP_TYPE_64 = 4
} hsa_fp_type_t;

typedef enum {
  HSA_FLUSH_MODE_FTZ = 1,
  HSA_FLUSH_MODE_NON_FTZ = 2
} hsa_flush_mode_t;

typedef enum {
  HSA_ROUND_METHOD_SINGLE = 1,
  HSA_ROUND_METHOD_DOUBLE = 2
} hsa_round_method_t;

hsa_status_t HSA_API hsa_isa_get_round_method(
    hsa_isa_t isa,
    hsa_fp_type_t fp_type,
    hsa_flush_mode_t flush_mode,
    hsa_round_method_t *round_method);

typedef struct hsa_wavefront_s {
  uint64_t handle;
} hsa_wavefront_t;

typedef enum {
  HSA_WAVEFRONT_INFO_SIZE = 0
} hsa_wavefront_info_t;

hsa_status_t HSA_API hsa_wavefront_get_info(
    hsa_wavefront_t wavefront,
    hsa_wavefront_info_t attribute,
    void *value);

hsa_status_t HSA_API hsa_isa_iterate_wavefronts(
    hsa_isa_t isa,
    hsa_status_t (*callback)(hsa_wavefront_t wavefront, void *data),
    void *data);

hsa_status_t HSA_API HSA_DEPRECATED hsa_isa_compatible(
    hsa_isa_t code_object_isa,
    hsa_isa_t agent_isa,
    bool *result);

typedef struct hsa_code_object_reader_s {
  uint64_t handle;
} hsa_code_object_reader_t;

hsa_status_t HSA_API hsa_code_object_reader_create_from_file(
    hsa_file_t file,
    hsa_code_object_reader_t *code_object_reader);

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

hsa_status_t HSA_API hsa_code_object_reader_destroy(
    hsa_code_object_reader_t code_object_reader);

typedef struct hsa_executable_s {
  uint64_t handle;
} hsa_executable_t;

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

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

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

hsa_status_t HSA_API hsa_executable_destroy(
    hsa_executable_t executable);

typedef struct hsa_loaded_code_object_s {
  uint64_t handle;
} hsa_loaded_code_object_t;

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

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

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

typedef enum {
  HSA_EXECUTABLE_INFO_PROFILE = 1,
  HSA_EXECUTABLE_INFO_STATE = 2,
  HSA_EXECUTABLE_INFO_DEFAULT_FLOAT_ROUNDING_MODE = 3
} hsa_executable_info_t;

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

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

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

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

hsa_status_t HSA_API hsa_executable_validate(
    hsa_executable_t executable,
    uint32_t *result);

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

typedef struct hsa_executable_symbol_s {
  uint64_t handle;
} hsa_executable_symbol_t;

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

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

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_SYMBOL_LINKAGE_MODULE = 0,
  HSA_SYMBOL_LINKAGE_PROGRAM = 1
} hsa_symbol_linkage_t;

typedef enum {
  HSA_VARIABLE_ALLOCATION_AGENT = 0,
  HSA_VARIABLE_ALLOCATION_PROGRAM = 1
} hsa_variable_allocation_t;

typedef enum {
  HSA_VARIABLE_SEGMENT_GLOBAL = 0,
  HSA_VARIABLE_SEGMENT_READONLY = 1
} hsa_variable_segment_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_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
} hsa_executable_symbol_info_t;

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

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

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

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

typedef struct hsa_code_object_s {
  uint64_t handle;
} hsa_code_object_t;

typedef struct hsa_callback_data_s {
  uint64_t handle;
} hsa_callback_data_t;

hsa_status_t HSA_API HSA_DEPRECATED hsa_code_object_serialize(
    hsa_code_object_t code_object,
    hsa_status_t (*alloc_callback)(size_t size,
                                   hsa_callback_data_t data,
                                   void **address),
    hsa_callback_data_t callback_data,
    const char *options,
    void **serialized_code_object,
    size_t *serialized_code_object_size);

hsa_status_t HSA_API HSA_DEPRECATED 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_API HSA_DEPRECATED hsa_code_object_destroy(
    hsa_code_object_t code_object);

typedef enum {
  HSA_CODE_OBJECT_TYPE_PROGRAM = 0
} hsa_code_object_type_t;

typedef enum {
  HSA_CODE_OBJECT_INFO_VERSION = 0,
  HSA_CODE_OBJECT_INFO_TYPE = 1,
  HSA_CODE_OBJECT_INFO_ISA = 2,
  HSA_CODE_OBJECT_INFO_MACHINE_MODEL = 3,
  HSA_CODE_OBJECT_INFO_PROFILE = 4,
  HSA_CODE_OBJECT_INFO_DEFAULT_FLOAT_ROUNDING_MODE = 5
} hsa_code_object_info_t;

hsa_status_t HSA_API HSA_DEPRECATED hsa_code_object_get_info(
    hsa_code_object_t code_object,
    hsa_code_object_info_t attribute,
    void *value);

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

typedef struct hsa_code_symbol_s {
  uint64_t handle;
} hsa_code_symbol_t;

hsa_status_t HSA_API HSA_DEPRECATED hsa_code_object_get_symbol(
    hsa_code_object_t code_object,
    const char *symbol_name,
    hsa_code_symbol_t *symbol);

hsa_status_t HSA_API HSA_DEPRECATED hsa_code_object_get_symbol_from_name(
    hsa_code_object_t code_object,
    const char *module_name,
    const char *symbol_name,
    hsa_code_symbol_t *symbol);

typedef enum {
  HSA_CODE_SYMBOL_INFO_TYPE = 0,
  HSA_CODE_SYMBOL_INFO_NAME_LENGTH = 1,
  HSA_CODE_SYMBOL_INFO_NAME = 2,
  HSA_CODE_SYMBOL_INFO_MODULE_NAME_LENGTH = 3,
  HSA_CODE_SYMBOL_INFO_MODULE_NAME = 4,
  HSA_CODE_SYMBOL_INFO_LINKAGE = 5,
  HSA_CODE_SYMBOL_INFO_IS_DEFINITION = 17,
  HSA_CODE_SYMBOL_INFO_VARIABLE_ALLOCATION = 6,
  HSA_CODE_SYMBOL_INFO_VARIABLE_SEGMENT = 7,
  HSA_CODE_SYMBOL_INFO_VARIABLE_ALIGNMENT = 8,
  HSA_CODE_SYMBOL_INFO_VARIABLE_SIZE = 9,
  HSA_CODE_SYMBOL_INFO_VARIABLE_IS_CONST = 10,
  HSA_CODE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE = 11,
  HSA_CODE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_ALIGNMENT = 12,
  HSA_CODE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE = 13,
  HSA_CODE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE = 14,
  HSA_CODE_SYMBOL_INFO_KERNEL_DYNAMIC_CALLSTACK = 15,
  HSA_CODE_SYMBOL_INFO_KERNEL_CALL_CONVENTION = 18,
  HSA_CODE_SYMBOL_INFO_INDIRECT_FUNCTION_CALL_CONVENTION = 16
} hsa_code_symbol_info_t;

hsa_status_t HSA_API HSA_DEPRECATED hsa_code_symbol_get_info(
    hsa_code_symbol_t code_symbol,
    hsa_code_symbol_info_t attribute,
    void *value);

hsa_status_t HSA_API HSA_DEPRECATED hsa_code_object_iterate_symbols(
    hsa_code_object_t code_object,
    hsa_status_t (*callback)(hsa_code_object_t code_object,
                             hsa_code_symbol_t symbol,
                             void *data),
    void *data);

#ifdef __cplusplus
}  // end extern "C" block
#endif

#endif  // header guard