1//===- InstrProfilingPlatformROCmHSA.cpp - ROCm HSA device drain ---------===//
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// HSA resident-image drain (Linux only).
10//
11// This pass walks loaded executables on each GPU agent and drains their
12// __llvm_profile_sections tables. It covers static HIP images and device-linked
13// programs without host shadows, while avoiding lookups that would load unused
14// static images. It shares processDeviceOffloadPrf() and section-bounds dedup
15// with the dynamic-module drain.
16//
17//===----------------------------------------------------------------------===//
18
19#if defined(__linux__)
20
21extern "C" {
22#include "InstrProfiling.h"
23#include "InstrProfilingPort.h"
24}
25
26#include "InstrProfilingPlatformROCmInternal.h"
27#include "interception/interception.h"
28// C (not C++) headers: clang_rt.profile is built -nostdinc++.
29#include <stddef.h>
30#include <stdint.h>
31#include <stdio.h>
32#include <stdlib.h>
33#include <string.h>
34
35using namespace __prof_rocm;
36
37// Mirrored HSA declarations the drain needs (dlopen'd, not linked). See the
38// header for the rationale; the values are HSA's stable C ABI.
39#include "InstrProfilingPlatformROCmHSADefs.h"
40
41#ifdef PROFILE_VERIFY_HSA_ABI
42// When the real ROCm headers are available at build time (developer installs
43// and the downstream GPU CI), check that the mirror above still matches them.
44#include <hsa/hsa.h>
45#include <hsa/hsa_ven_amd_loader.h>
46
47static_assert(PROF_HSA_STATUS_SUCCESS == HSA_STATUS_SUCCESS, "HSA ABI drift");
48static_assert(PROF_HSA_STATUS_INFO_BREAK == HSA_STATUS_INFO_BREAK,
49 "HSA ABI drift");
50static_assert(PROF_HSA_AGENT_INFO_NAME == HSA_AGENT_INFO_NAME, "HSA ABI drift");
51static_assert(PROF_HSA_AGENT_INFO_DEVICE == HSA_AGENT_INFO_DEVICE,
52 "HSA ABI drift");
53static_assert(PROF_HSA_DEVICE_TYPE_GPU == HSA_DEVICE_TYPE_GPU, "HSA ABI drift");
54static_assert(PROF_HSA_SYMBOL_KIND_VARIABLE == HSA_SYMBOL_KIND_VARIABLE,
55 "HSA ABI drift");
56static_assert(PROF_HSA_EXECUTABLE_SYMBOL_INFO_TYPE ==
57 HSA_EXECUTABLE_SYMBOL_INFO_TYPE,
58 "HSA ABI drift");
59static_assert(PROF_HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH ==
60 HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH,
61 "HSA ABI drift");
62static_assert(PROF_HSA_EXECUTABLE_SYMBOL_INFO_NAME ==
63 HSA_EXECUTABLE_SYMBOL_INFO_NAME,
64 "HSA ABI drift");
65static_assert(PROF_HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS ==
66 HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS,
67 "HSA ABI drift");
68static_assert(PROF_HSA_EXTENSION_AMD_LOADER == HSA_EXTENSION_AMD_LOADER,
69 "HSA ABI drift");
70
71static_assert(sizeof(prof_hsa_agent_t) == sizeof(hsa_agent_t), "HSA ABI drift");
72static_assert(sizeof(prof_hsa_executable_t) == sizeof(hsa_executable_t),
73 "HSA ABI drift");
74static_assert(sizeof(prof_hsa_executable_symbol_t) ==
75 sizeof(hsa_executable_symbol_t),
76 "HSA ABI drift");
77
78static_assert(sizeof(prof_hsa_loader_segment_descriptor_t) ==
79 sizeof(hsa_ven_amd_loader_segment_descriptor_t),
80 "HSA ABI drift");
81static_assert(offsetof(prof_hsa_loader_segment_descriptor_t, agent) ==
82 offsetof(hsa_ven_amd_loader_segment_descriptor_t, agent),
83 "HSA ABI drift");
84static_assert(offsetof(prof_hsa_loader_segment_descriptor_t, executable) ==
85 offsetof(hsa_ven_amd_loader_segment_descriptor_t, executable),
86 "HSA ABI drift");
87static_assert(offsetof(prof_hsa_loader_segment_descriptor_t, segment_base) ==
88 offsetof(hsa_ven_amd_loader_segment_descriptor_t,
89 segment_base),
90 "HSA ABI drift");
91static_assert(offsetof(prof_hsa_loader_segment_descriptor_t, segment_size) ==
92 offsetof(hsa_ven_amd_loader_segment_descriptor_t,
93 segment_size),
94 "HSA ABI drift");
95
96// We fetch the loader pfn table by raw layout, so query_segment_descriptors
97// must sit at the same offset as in the real table.
98static_assert(offsetof(prof_hsa_loader_pfn_t, query_segment_descriptors) ==
99 offsetof(hsa_ven_amd_loader_1_00_pfn_t,
100 hsa_ven_amd_loader_query_segment_descriptors),
101 "HSA ABI drift");
102#endif // PROFILE_VERIFY_HSA_ABI
103
104static hsa_iterate_agents_ty pHsaIterateAgents = nullptr;
105static hsa_agent_get_info_ty pHsaAgentGetInfo = nullptr;
106static hsa_executable_iterate_agent_symbols_ty pHsaExecIterAgentSyms = nullptr;
107static hsa_executable_symbol_get_info_ty pHsaSymGetInfo = nullptr;
108static hsa_loader_query_segment_descriptors_ty pQuerySegDescs = nullptr;
109
110/* Status-check shorthands, in the spirit of the thin HIP wrappers in
111 * InstrProfilingPlatformROCm.cpp: every HSA entry point returns
112 * prof_hsa_status_t. hsaOkOrBreak() also accepts INFO_BREAK, which the
113 * iterate_* callbacks use to stop early and is not an error. */
114static inline bool hsaOk(prof_hsa_status_t St) {
115 return St == PROF_HSA_STATUS_SUCCESS;
116}
117static inline bool hsaOkOrBreak(prof_hsa_status_t St) {
118 return St == PROF_HSA_STATUS_SUCCESS || St == PROF_HSA_STATUS_INFO_BREAK;
119}
120
121/* 0 = not attempted, 1 = ready, -1 = unavailable. Acquire/release atomics: a
122 * thread observing HsaRuntimeState==1 also sees the published p* pointers. */
123static int HsaRuntimeState = 0;
124
125static int setHsaRuntimeState(int S) {
126 __atomic_store_n(&HsaRuntimeState, S, __ATOMIC_RELEASE);
127 return S > 0 ? 0 : -1;
128}
129
130/* Resolve HSA entry points and the AMD loader extension once, and confirm HIP's
131 * hipMemcpy is reachable for the device-to-host copies. */
132static int loadHsaRuntimePointers(void) {
133 int State = __atomic_load_n(&HsaRuntimeState, __ATOMIC_ACQUIRE);
134 if (State)
135 return State > 0 ? 0 : -1;
136
137 if (!__interception::DynamicLoaderAvailable()) {
138 if (isVerboseMode())
139 PROF_NOTE("%s", "Dynamic library loading not available - "
140 "HSA device profiling disabled\n");
141 return setHsaRuntimeState(-1);
142 }
143
144 void *Hsa = __interception::OpenLibrary(name: "libhsa-runtime64.so");
145 if (!Hsa)
146 Hsa = __interception::OpenLibrary(name: "libhsa-runtime64.so.1");
147 if (!Hsa) {
148 if (isVerboseMode())
149 PROF_NOTE("%s", "libhsa-runtime64.so not loadable - "
150 "HSA device profiling disabled\n");
151 return setHsaRuntimeState(-1);
152 }
153
154 hsa_init_ty pHsaInit =
155 (hsa_init_ty)__interception::LookupSymbol(handle: Hsa, symbol: "hsa_init");
156 hsa_system_get_major_extension_table_ty pGetExtTable =
157 (hsa_system_get_major_extension_table_ty)__interception::LookupSymbol(
158 handle: Hsa, symbol: "hsa_system_get_major_extension_table");
159 pHsaIterateAgents = (hsa_iterate_agents_ty)__interception::LookupSymbol(
160 handle: Hsa, symbol: "hsa_iterate_agents");
161 pHsaAgentGetInfo = (hsa_agent_get_info_ty)__interception::LookupSymbol(
162 handle: Hsa, symbol: "hsa_agent_get_info");
163 pHsaExecIterAgentSyms =
164 (hsa_executable_iterate_agent_symbols_ty)__interception::LookupSymbol(
165 handle: Hsa, symbol: "hsa_executable_iterate_agent_symbols");
166 pHsaSymGetInfo =
167 (hsa_executable_symbol_get_info_ty)__interception::LookupSymbol(
168 handle: Hsa, symbol: "hsa_executable_symbol_get_info");
169
170 if (!pHsaInit || !pGetExtTable || !pHsaIterateAgents || !pHsaAgentGetInfo ||
171 !pHsaExecIterAgentSyms || !pHsaSymGetInfo) {
172 PROF_WARN("%s",
173 "required HSA symbols missing - HSA device profiling disabled\n");
174 return setHsaRuntimeState(-1);
175 }
176
177 /* Bring HSA up lazily on the first drain (idempotent, refcounted), never from
178 * a library constructor -- see the fork-safety note at end of file. */
179 prof_hsa_status_t St = pHsaInit();
180 if (!hsaOkOrBreak(St)) {
181 if (isVerboseMode())
182 PROF_NOTE("hsa_init failed (0x%x) - HSA device profiling disabled\n", St);
183 return setHsaRuntimeState(-1);
184 }
185
186 prof_hsa_loader_pfn_t LoaderApi;
187 __builtin_memset(&LoaderApi, 0, sizeof(LoaderApi));
188 St = pGetExtTable(PROF_HSA_EXTENSION_AMD_LOADER, 1, sizeof(LoaderApi),
189 &LoaderApi);
190 if (!hsaOk(St) || !LoaderApi.query_segment_descriptors) {
191 PROF_WARN("AMD loader extension unavailable (0x%x) - "
192 "HSA device profiling disabled\n",
193 St);
194 return setHsaRuntimeState(-1);
195 }
196 pQuerySegDescs = LoaderApi.query_segment_descriptors;
197
198 /* The device-to-host copies go through the shared HIP loader. */
199 ensureHipLoaded();
200 if (!hipMemcpyAvailable()) {
201 PROF_WARN("%s", "hipMemcpy unavailable - HSA device profiling disabled\n");
202 return setHsaRuntimeState(-1);
203 }
204
205 if (isVerboseMode())
206 PROF_NOTE("%s", "HSA + HIP runtime resolved for device profiling\n");
207 return setHsaRuntimeState(1);
208}
209
210int __prof_rocm::hsaRuntimeAvailable(void) {
211 return loadHsaRuntimePointers() == 0;
212}
213
214/* The canonical device bounds-table symbol from InstrProfilingPlatformGPU.c. */
215static const char ProfileSectionsSymbol[] = "__llvm_profile_sections";
216
217/* Dedup of drained section-bounds tuples, shared with the host-shadow path
218 * (processDeviceOffloadPrf records here on every successful drain) so each
219 * unique counter set is drained exactly once across both paths.
220 */
221static ProfBoundsSet SeenBounds;
222
223/* Has this bounds tuple already been drained? Pure check, no state mutation. */
224static int profBoundsAlreadyDrained(const void *D, const void *C,
225 const void *N) {
226 return SeenBounds.contains(D, C, N);
227}
228
229/* Record a drained bounds tuple. Idempotent; call only after a successful drain
230 * so a failed attempt stays retryable. */
231void __prof_rocm::profRecordDrainedBounds(const void *D, const void *C,
232 const void *N) {
233 SeenBounds.record(D, C, N);
234}
235
236#define PROF_MAX_GPU_AGENTS 64
237
238/* Buffer size for HSA agent names and symbol names we read back; both the
239 * device arch string and the __llvm_profile_sections symbol are far shorter. */
240#define PROF_HSA_NAME_MAX 64
241
242namespace {
243struct GpuAgent {
244 prof_hsa_agent_t agent;
245 char arch[PROF_HSA_NAME_MAX];
246};
247
248struct WalkState {
249 GpuAgent agents[PROF_MAX_GPU_AGENTS];
250 int num_agents;
251 int total_found;
252 int total_drained;
253};
254
255/* Per (agent, executable) symbol-iteration state. */
256struct SymbolState {
257 const char *arch;
258 int found;
259 int drained;
260};
261} // namespace
262
263/* HSA per-symbol callback: when it finds a __llvm_profile_sections variable,
264 * drain it via processDeviceOffloadPrf() unless the host-shadow path (or an
265 * earlier agent) already handled the same bounds. */
266static prof_hsa_status_t onSymbol(prof_hsa_executable_t, prof_hsa_agent_t,
267 prof_hsa_executable_symbol_t Sym,
268 void *Data) {
269 SymbolState *S = (SymbolState *)Data;
270
271 prof_hsa_symbol_kind_t Kind;
272 if (!hsaOk(
273 St: pHsaSymGetInfo(Sym, PROF_HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &Kind)) ||
274 Kind != PROF_HSA_SYMBOL_KIND_VARIABLE)
275 return PROF_HSA_STATUS_SUCCESS;
276
277 uint32_t NameLen = 0;
278 if (!hsaOk(St: pHsaSymGetInfo(Sym, PROF_HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH,
279 &NameLen)) ||
280 NameLen != sizeof(ProfileSectionsSymbol) - 1)
281 return PROF_HSA_STATUS_SUCCESS;
282
283 char NameBuf[PROF_HSA_NAME_MAX];
284 if (NameLen + 1 > sizeof(NameBuf))
285 return PROF_HSA_STATUS_SUCCESS;
286 if (!hsaOk(
287 St: pHsaSymGetInfo(Sym, PROF_HSA_EXECUTABLE_SYMBOL_INFO_NAME, NameBuf)))
288 return PROF_HSA_STATUS_SUCCESS;
289 NameBuf[NameLen] = '\0';
290
291 if (strcmp(s1: NameBuf, s2: ProfileSectionsSymbol) != 0)
292 return PROF_HSA_STATUS_SUCCESS;
293
294 uint64_t Addr = 0;
295 if (!hsaOk(St: pHsaSymGetInfo(
296 Sym, PROF_HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS, &Addr)) ||
297 Addr == 0) {
298 if (isVerboseMode())
299 PROF_NOTE("%s", "failed to read __llvm_profile_sections address\n");
300 return PROF_HSA_STATUS_SUCCESS;
301 }
302
303 S->found++;
304
305 // Read the bounds table first to dedup (and detect empty sections) before
306 // the full copy/relocate done by processDeviceOffloadPrf.
307 __llvm_profile_gpu_sections Sec;
308 if (memcpyDeviceToHost(Dst: &Sec, Src: (void *)(uintptr_t)Addr, Size: sizeof(Sec)) != 0) {
309 PROF_WARN("%s", "failed to copy device bounds table\n");
310 return PROF_HSA_STATUS_SUCCESS;
311 }
312 if (profBoundsAlreadyDrained(D: Sec.DataStart, C: Sec.CountersStart,
313 N: Sec.NamesStart)) {
314 if (isVerboseMode())
315 PROF_NOTE("%s", "device bounds already drained, skipping\n");
316 return PROF_HSA_STATUS_SUCCESS;
317 }
318
319 size_t DataBytes = (const char *)Sec.DataStop - (const char *)Sec.DataStart;
320 size_t CntsBytes =
321 (const char *)Sec.CountersStop - (const char *)Sec.CountersStart;
322 if (DataBytes == 0 || CntsBytes == 0) {
323 // Empty code object: nothing to write. Mark seen so we don't revisit it.
324 profRecordDrainedBounds(D: Sec.DataStart, C: Sec.CountersStart, N: Sec.NamesStart);
325 return PROF_HSA_STATUS_SUCCESS;
326 }
327
328 // Name HSA-drained objects in their own ".hsaN" suffix space so they never
329 // collide with the host-shadow path's "arch"/"arch.<i>" filenames. The drain
330 // latch (HsaDrainCompleted) already prevents re-draining an object, so a
331 // plain per-drain counter is enough for uniqueness.
332 static int DrainIndex = 0;
333 char Target[96];
334 snprintf(s: Target, maxlen: sizeof(Target), format: "%s.hsa%d", S->arch, DrainIndex);
335
336 // Record the bounds (and advance the index) only on a successful write so a
337 // transient error stays retryable on a later agent or collect call.
338 if (processDeviceOffloadPrf(DeviceOffloadPrf: (void *)(uintptr_t)Addr, Target, Sections: nullptr) == 0) {
339 S->drained++;
340 DrainIndex++;
341 profRecordDrainedBounds(D: Sec.DataStart, C: Sec.CountersStart, N: Sec.NamesStart);
342 }
343
344 return PROF_HSA_STATUS_SUCCESS;
345}
346
347static prof_hsa_status_t collectAgent(prof_hsa_agent_t Agent, void *Data) {
348 prof_hsa_device_type_t DevType;
349 if (!hsaOk(St: pHsaAgentGetInfo(Agent, PROF_HSA_AGENT_INFO_DEVICE, &DevType)) ||
350 DevType != PROF_HSA_DEVICE_TYPE_GPU)
351 return PROF_HSA_STATUS_SUCCESS;
352
353 WalkState *W = (WalkState *)Data;
354 if (W->num_agents >= PROF_MAX_GPU_AGENTS)
355 return PROF_HSA_STATUS_SUCCESS;
356
357 GpuAgent &GA = W->agents[W->num_agents++];
358 GA.agent = Agent;
359 char Name[PROF_HSA_NAME_MAX];
360 __builtin_memset(Name, 0, sizeof(Name));
361 pHsaAgentGetInfo(Agent, PROF_HSA_AGENT_INFO_NAME, Name);
362 size_t N = strnlen(string: Name, maxlen: sizeof(GA.arch) - 1);
363 __builtin_memcpy(GA.arch, Name, N);
364 GA.arch[N] = '\0';
365 if (!GA.arch[0])
366 strncpy(dest: GA.arch, src: "amdgpu", n: sizeof(GA.arch) - 1);
367
368 if (isVerboseMode())
369 PROF_NOTE("GPU agent %d: %s\n", W->num_agents - 1, GA.arch);
370 return PROF_HSA_STATUS_SUCCESS;
371}
372
373/* Reentrancy guard and "drained at least once" latch (both acquire/release). */
374static int HsaDrainInProgress = 0;
375static int HsaDrainCompleted = 0;
376
377int __prof_rocm::drainDevicesViaHsa(void) {
378 if (__atomic_load_n(&HsaDrainCompleted, __ATOMIC_ACQUIRE))
379 return 0;
380
381 int Expected = 0;
382 if (!__atomic_compare_exchange_n(&HsaDrainInProgress, &Expected, 1,
383 /*weak=*/0, __ATOMIC_ACQ_REL,
384 __ATOMIC_ACQUIRE))
385 return 0;
386
387 struct InProgressGuard {
388 ~InProgressGuard() {
389 __atomic_store_n(&HsaDrainInProgress, 0, __ATOMIC_RELEASE);
390 }
391 } _Guard;
392
393 if (loadHsaRuntimePointers() != 0)
394 return 0; /* Runtime unavailable: stay retryable. */
395
396 WalkState W;
397 __builtin_memset(&W, 0, sizeof(W));
398 prof_hsa_status_t St = pHsaIterateAgents(collectAgent, &W);
399 if (!hsaOkOrBreak(St)) {
400 PROF_WARN("hsa_iterate_agents failed (0x%x)\n", St);
401 return -1;
402 }
403 if (W.num_agents == 0) {
404 if (isVerboseMode())
405 PROF_NOTE("%s", "no GPU agents present; nothing to drain (will retry)\n");
406 return 0;
407 }
408
409 /* query_segment_descriptors ships in every loader-extension version, is more
410 * permissive than iterate_executables on ROCm, and yields the loaded
411 * (agent, executable) pairs directly. */
412 size_t NumSegs = 0;
413 St = pQuerySegDescs(nullptr, &NumSegs);
414 if (!hsaOk(St)) {
415 PROF_WARN("query_segment_descriptors(count) failed (0x%x)\n", St);
416 return -1;
417 }
418 if (NumSegs == 0) {
419 if (isVerboseMode())
420 PROF_NOTE("%s", "no loaded segments; nothing to drain (will retry)\n");
421 return 0;
422 }
423
424 prof_hsa_loader_segment_descriptor_t *Segs =
425 (prof_hsa_loader_segment_descriptor_t *)calloc(nmemb: NumSegs, size: sizeof(*Segs));
426 if (!Segs) {
427 PROF_ERR("%s\n", "failed to allocate segment descriptor array");
428 return -1;
429 }
430 UniqueFree SegsOwner(Segs);
431
432 St = pQuerySegDescs(Segs, &NumSegs);
433 if (!hsaOk(St)) {
434 PROF_WARN("query_segment_descriptors(fetch) failed (0x%x)\n", St);
435 return -1;
436 }
437
438 if (isVerboseMode())
439 PROF_NOTE("query_segment_descriptors: %zu segments\n", NumSegs);
440
441 // Walk each unique (agent, executable) pair once.
442 struct SeenPair {
443 uint64_t agent;
444 uint64_t exec;
445 };
446 enum { kSeenPairsInitCap = 64 };
447 SeenPair *Seen = nullptr;
448 int NumPairs = 0;
449 int CapPairs = 0;
450 int IterFailures = 0;
451
452 for (size_t i = 0; i < NumSegs; ++i) {
453 if (Segs[i].executable.handle == 0 || Segs[i].agent.handle == 0)
454 continue;
455
456 bool AlreadySeen = false;
457 for (int j = 0; j < NumPairs; ++j)
458 if (Seen[j].agent == Segs[i].agent.handle &&
459 Seen[j].exec == Segs[i].executable.handle) {
460 AlreadySeen = true;
461 break;
462 }
463 if (AlreadySeen)
464 continue;
465 if (growArray(Arr: (void **)&Seen, Cap: &CapPairs, MinCount: NumPairs + 1, InitCap: kSeenPairsInitCap,
466 ElemSize: sizeof(*Seen)) == 0) {
467 Seen[NumPairs].agent = Segs[i].agent.handle;
468 Seen[NumPairs].exec = Segs[i].executable.handle;
469 NumPairs++;
470 }
471
472 const char *Arch = nullptr;
473 for (int k = 0; k < W.num_agents; ++k)
474 if (W.agents[k].agent.handle == Segs[i].agent.handle) {
475 Arch = W.agents[k].arch;
476 break;
477 }
478 if (!Arch)
479 continue; /* not a GPU agent we collected */
480
481 SymbolState S;
482 __builtin_memset(&S, 0, sizeof(S));
483 S.arch = Arch;
484 if (isVerboseMode())
485 PROF_NOTE("walking executable 0x%llx on %s\n",
486 (unsigned long long)Segs[i].executable.handle, Arch);
487 prof_hsa_status_t IterSt =
488 pHsaExecIterAgentSyms(Segs[i].executable, Segs[i].agent, onSymbol, &S);
489 if (!hsaOkOrBreak(St: IterSt)) {
490 PROF_WARN("hsa_executable_iterate_agent_symbols on executable 0x%llx "
491 "failed (0x%x)\n",
492 (unsigned long long)Segs[i].executable.handle, IterSt);
493 IterFailures++;
494 }
495 W.total_found += S.found;
496 W.total_drained += S.drained;
497 }
498
499 if (isVerboseMode())
500 PROF_NOTE("HSA walk complete: agents=%d pairs=%d found=%d drained=%d "
501 "iter-failures=%d\n",
502 W.num_agents, NumPairs, W.total_found, W.total_drained,
503 IterFailures);
504
505 free(ptr: Seen);
506
507 /* Latch only when we actually drained data. A "found nothing new" walk is
508 * deliberately not latched: an early collect can precede any kernel launch,
509 * and latching it would suppress the real exit-time drain. No-op walks are
510 * cheap to repeat. */
511 if (W.total_drained > 0)
512 __atomic_store_n(&HsaDrainCompleted, 1, __ATOMIC_RELEASE);
513 return (IterFailures > 0) ? -1 : 0;
514}
515
516/* Fork-safety: deliberately no library constructor calling hsa_init(). */
517
518#endif /* defined(__linux__) && !defined(_WIN32) -- HSA drain */
519