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_RUNTIME_ALLOC_ALLOWED = 5,
82} hsa_amd_memory_pool_info_t;
83
84typedef enum {
85 HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_FINE_GRAINED = 2,
86 HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED = 4,
87} hsa_amd_memory_pool_global_flag_t;
88
89typedef enum {
90 HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT = 0xA002,
91 HSA_AMD_AGENT_INFO_MAX_WAVES_PER_CU = 0xA00A,
92} hsa_amd_agent_info_t;
93
94typedef enum {
95 HSA_SIGNAL_CONDITION_NE = 1,
96} hsa_signal_condition_t;
97
98typedef enum {
99 HSA_WAIT_STATE_BLOCKED = 0,
100} hsa_wait_state_t;
101
102typedef enum {
103 HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS = 21,
104} hsa_executable_symbol_info_t;
105
106typedef enum { HSA_EXTENSION_AMD_LOADER = 0x201 } hsa_extension_t;
107
108typedef enum {
109 HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_KIND = 2,
110 HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_AGENT = 3,
111 HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_TYPE = 4,
112 HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_MEMORY_BASE =
113 5,
114 HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_MEMORY_SIZE =
115 6,
116 HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_BASE = 9,
117 HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_SIZE = 10,
118} hsa_ven_amd_loader_loaded_code_object_info_t;
119
120typedef enum {
121 HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_KIND_AGENT = 2,
122} hsa_ven_amd_loader_loaded_code_object_kind_t;
123
124typedef enum {
125 HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_MEMORY = 2,
126} hsa_ven_amd_loader_code_object_storage_type_t;
127
128// Version 1.01 of HSA_EXTENSION_AMD_LOADER. Slot order is the extension ABI.
129struct LoaderApi {
130 hsa_status_t (*QueryHostAddress)(const void*, const void**);
131 void* UnusedQuerySegmentDescriptors;
132 void* UnusedQueryExecutable;
133 hsa_status_t (*IterateLoadedCodeObjects)(
134 hsa_executable_t,
135 hsa_status_t (*)(hsa_executable_t, hsa_loaded_code_object_t, void*),
136 void*);
137 hsa_status_t (*GetCodeObjectInfo)(
138 hsa_loaded_code_object_t, hsa_ven_amd_loader_loaded_code_object_info_t,
139 void*);
140};
141static_assert(sizeof(LoaderApi) == 5 * sizeof(void*), "layout drift");
142
143hsa_status_t hsa_init(void);
144hsa_status_t hsa_shut_down(void);
145hsa_status_t hsa_iterate_agents(hsa_status_t (*callback)(hsa_agent_t, void*),
146 void* data);
147hsa_status_t hsa_agent_get_info(hsa_agent_t agent, hsa_agent_info_t attribute,
148 void* value);
149hsa_status_t hsa_executable_destroy(hsa_executable_t executable);
150hsa_status_t hsa_executable_freeze(hsa_executable_t executable,
151 const char* options);
152hsa_status_t hsa_executable_get_symbol_by_name(hsa_executable_t executable,
153 const char* symbol_name,
154 const hsa_agent_t* agent,
155 hsa_executable_symbol_t* symbol);
156hsa_status_t hsa_executable_symbol_get_info(
157 hsa_executable_symbol_t symbol, hsa_executable_symbol_info_t attribute,
158 void* value);
159hsa_status_t hsa_amd_memory_pool_allocate(hsa_amd_memory_pool_t memory_pool,
160 size_t size, uint32_t flags,
161 void** ptr);
162hsa_status_t hsa_amd_memory_pool_free(void* ptr);
163hsa_status_t hsa_amd_memory_pool_get_info(hsa_amd_memory_pool_t memory_pool,
164 hsa_amd_memory_pool_info_t attribute,
165 void* value);
166hsa_status_t hsa_amd_agent_iterate_memory_pools(
167 hsa_agent_t agent, hsa_status_t (*callback)(hsa_amd_memory_pool_t, void*),
168 void* data);
169hsa_status_t hsa_amd_agents_allow_access(uint32_t num_agents,
170 const hsa_agent_t* agents,
171 const uint32_t* flags,
172 const void* ptr);
173hsa_status_t hsa_system_get_major_extension_table(uint16_t extension,
174 uint16_t version_major,
175 size_t table_length,
176 void* table);
177hsa_status_t hsa_memory_copy(void* dst, const void* src, size_t size);
178hsa_status_t hsa_amd_signal_create(hsa_signal_value_t initial_value,
179 uint32_t num_consumers,
180 const hsa_agent_t* consumers,
181 uint64_t attributes, hsa_signal_t* signal);
182hsa_status_t hsa_signal_destroy(hsa_signal_t signal);
183void hsa_signal_store_screlease(hsa_signal_t signal, hsa_signal_value_t value);
184hsa_signal_value_t hsa_signal_wait_scacquire(hsa_signal_t signal,
185 hsa_signal_condition_t condition,
186 hsa_signal_value_t compare_value,
187 uint64_t timeout_hint,
188 hsa_wait_state_t wait_state_hint);
189
190} // extern "C"
191
192#endif // SANITIZER_OFFLOAD_HSA_H
193