//===--- 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_ERROR_EXCEPTION = 0x1016,
} 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
Codebase Browser
llvm