1//===----------------------------------------------------------------------===//
2//
3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4// See https://llvm.org/LICENSE.txt for license information.
5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6//
7//===----------------------------------------------------------------------===//
8//
9// Minimal HSA declarations without a ROCm dependency.
10//
11//===----------------------------------------------------------------------===//
12
13#ifndef SANITIZER_OFFLOAD_HSA_H
14#define SANITIZER_OFFLOAD_HSA_H
15
16#include <stddef.h>
17#include <stdint.h>
18
19#define SANITIZER_HSA_LIBRARY "libhsa-runtime64"
20
21#define SANITIZER_HSA_FUNCTIONS(X) \
22 X(hsa_iterate_agents) \
23 X(hsa_agent_get_info) \
24 X(hsa_executable_get_symbol_by_name) \
25 X(hsa_executable_symbol_get_info) \
26 X(hsa_memory_copy) \
27 X(hsa_amd_memory_pool_allocate) \
28 X(hsa_amd_memory_pool_free) \
29 X(hsa_amd_agents_allow_access) \
30 X(hsa_amd_memory_pool_get_info) \
31 X(hsa_amd_agent_iterate_memory_pools) \
32 X(hsa_system_get_major_extension_table) \
33 X(hsa_amd_signal_create) \
34 X(hsa_signal_destroy) \
35 X(hsa_signal_store_screlease) \
36 X(hsa_signal_wait_scacquire)
37
38extern "C" {
39
40typedef struct hsa_agent_s {
41 uint64_t handle;
42} hsa_agent_t;
43typedef struct hsa_executable_s {
44 uint64_t handle;
45} hsa_executable_t;
46typedef struct hsa_executable_symbol_s {
47 uint64_t handle;
48} hsa_executable_symbol_t;
49typedef struct hsa_amd_memory_pool_s {
50 uint64_t handle;
51} hsa_amd_memory_pool_t;
52typedef struct hsa_signal_s {
53 uint64_t handle;
54} hsa_signal_t;
55typedef struct hsa_loaded_code_object_s {
56 uint64_t handle;
57} hsa_loaded_code_object_t;
58
59typedef int64_t hsa_signal_value_t;
60
61typedef enum {
62 HSA_STATUS_SUCCESS = 0,
63 HSA_STATUS_ERROR = 0x1001,
64} hsa_status_t;
65
66typedef enum {
67 HSA_DEVICE_TYPE_CPU = 0,
68 HSA_DEVICE_TYPE_GPU = 1,
69} hsa_device_type_t;
70
71typedef enum {
72 HSA_AGENT_INFO_WAVEFRONT_SIZE = 6,
73 HSA_AGENT_INFO_DEVICE = 17,
74} hsa_agent_info_t;
75
76typedef enum { HSA_AMD_SEGMENT_GLOBAL = 0 } hsa_amd_segment_t;
77
78typedef enum {
79 HSA_AMD_MEMORY_POOL_INFO_SEGMENT = 0,
80 HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS = 1,
81} hsa_amd_memory_pool_info_t;
82
83typedef enum {
84 HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_FINE_GRAINED = 2,
85} hsa_amd_memory_pool_global_flag_t;
86
87typedef enum {
88 HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT = 0xA002,
89 HSA_AMD_AGENT_INFO_MAX_WAVES_PER_CU = 0xA00A,
90} hsa_amd_agent_info_t;
91
92typedef enum {
93 HSA_SIGNAL_CONDITION_NE = 1,
94} hsa_signal_condition_t;
95
96typedef enum {
97 HSA_WAIT_STATE_BLOCKED = 0,
98} hsa_wait_state_t;
99
100typedef enum {
101 HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS = 21,
102} hsa_executable_symbol_info_t;
103
104typedef enum { HSA_EXTENSION_AMD_LOADER = 0x201 } hsa_extension_t;
105
106typedef enum {
107 HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_KIND = 2,
108 HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_AGENT = 3,
109 HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_TYPE = 4,
110 HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_MEMORY_BASE =
111 5,
112 HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_MEMORY_SIZE =
113 6,
114 HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_BASE = 9,
115 HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_SIZE = 10,
116} hsa_ven_amd_loader_loaded_code_object_info_t;
117
118typedef enum {
119 HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_KIND_AGENT = 2,
120} hsa_ven_amd_loader_loaded_code_object_kind_t;
121
122typedef enum {
123 HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_MEMORY = 2,
124} hsa_ven_amd_loader_code_object_storage_type_t;
125
126// Version 1.01 of HSA_EXTENSION_AMD_LOADER. Slot order is the extension ABI.
127struct LoaderApi {
128 hsa_status_t (*QueryHostAddress)(const void*, const void**);
129 void* UnusedQuerySegmentDescriptors;
130 void* UnusedQueryExecutable;
131 hsa_status_t (*IterateLoadedCodeObjects)(
132 hsa_executable_t,
133 hsa_status_t (*)(hsa_executable_t, hsa_loaded_code_object_t, void*),
134 void*);
135 hsa_status_t (*GetCodeObjectInfo)(
136 hsa_loaded_code_object_t, hsa_ven_amd_loader_loaded_code_object_info_t,
137 void*);
138};
139static_assert(sizeof(LoaderApi) == 5 * sizeof(void*), "layout drift");
140
141hsa_status_t hsa_init(void);
142hsa_status_t hsa_shut_down(void);
143hsa_status_t hsa_iterate_agents(hsa_status_t (*callback)(hsa_agent_t, void*),
144 void* data);
145hsa_status_t hsa_agent_get_info(hsa_agent_t agent, hsa_agent_info_t attribute,
146 void* value);
147hsa_status_t hsa_executable_destroy(hsa_executable_t executable);
148hsa_status_t hsa_executable_freeze(hsa_executable_t executable,
149 const char* options);
150hsa_status_t hsa_executable_get_symbol_by_name(hsa_executable_t executable,
151 const char* symbol_name,
152 const hsa_agent_t* agent,
153 hsa_executable_symbol_t* symbol);
154hsa_status_t hsa_executable_symbol_get_info(
155 hsa_executable_symbol_t symbol, hsa_executable_symbol_info_t attribute,
156 void* value);
157hsa_status_t hsa_amd_memory_pool_allocate(hsa_amd_memory_pool_t memory_pool,
158 size_t size, uint32_t flags,
159 void** ptr);
160hsa_status_t hsa_amd_memory_pool_free(void* ptr);
161hsa_status_t hsa_amd_memory_pool_get_info(hsa_amd_memory_pool_t memory_pool,
162 hsa_amd_memory_pool_info_t attribute,
163 void* value);
164hsa_status_t hsa_amd_agent_iterate_memory_pools(
165 hsa_agent_t agent, hsa_status_t (*callback)(hsa_amd_memory_pool_t, void*),
166 void* data);
167hsa_status_t hsa_amd_agents_allow_access(uint32_t num_agents,
168 const hsa_agent_t* agents,
169 const uint32_t* flags,
170 const void* ptr);
171hsa_status_t hsa_system_get_major_extension_table(uint16_t extension,
172 uint16_t version_major,
173 size_t table_length,
174 void* table);
175hsa_status_t hsa_memory_copy(void* dst, const void* src, size_t size);
176hsa_status_t hsa_amd_signal_create(hsa_signal_value_t initial_value,
177 uint32_t num_consumers,
178 const hsa_agent_t* consumers,
179 uint64_t attributes, hsa_signal_t* signal);
180hsa_status_t hsa_signal_destroy(hsa_signal_t signal);
181void hsa_signal_store_screlease(hsa_signal_t signal, hsa_signal_value_t value);
182hsa_signal_value_t hsa_signal_wait_scacquire(hsa_signal_t signal,
183 hsa_signal_condition_t condition,
184 hsa_signal_value_t compare_value,
185 uint64_t timeout_hint,
186 hsa_wait_state_t wait_state_hint);
187
188} // extern "C"
189
190#endif // SANITIZER_OFFLOAD_HSA_H
191