1 ////////////////////////////////////////////////////////////////////////////////
2 //
3 // Copyright (C) 2014-2020 Advanced Micro Devices Inc. All rights reserved.
4 //
5 // Permission is hereby granted, free of charge, to any person or organization
6 // obtaining a copy of the software and accompanying documentation covered by
7 // this license (the "Software") to use, reproduce, display, distribute,
8 // execute, and transmit the Software, and to prepare derivative works of the
9 // Software, and to permit third-parties to whom the Software is furnished to
10 // do so, all subject to the following:
11 //
12 // The copyright notices in the Software and this entire statement, including
13 // the above license grant, this restriction and the following disclaimer,
14 // must be included in all copies of the Software, in whole or in part, and
15 // all derivative works of the Software, unless such copies or derivative
16 // works are solely in the form of machine-executable object code generated by
17 // a source language processor.
18 //
19 // THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
20 // IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
21 // FITNESS FOR A PARTICULAR PURPOSE, TITLE AND NON-INFRINGEMENT. IN NO EVENT
22 // SHALL THE COPYRIGHT HOLDERS OR ANYONE DISTRIBUTING THE SOFTWARE BE LIABLE
23 // FOR ANY DAMAGES OR OTHER LIABILITY, WHETHER IN CONTRACT, TORT OR OTHERWISE,
24 // ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
25 // DEALINGS IN THE SOFTWARE.
26 //
27 ////////////////////////////////////////////////////////////////////////////////
28
29 #ifndef HSA_RUNTIME_INC_HSA_H_
30 #define HSA_RUNTIME_INC_HSA_H_
31
32 #include <stddef.h> /* size_t */
33 #include <stdint.h> /* uintXX_t */
34
35 #ifndef __cplusplus
36 #include <stdbool.h> /* bool */
37 #endif /* __cplusplus */
38
39 // Placeholder for calling convention and import/export macros
40 #ifndef HSA_CALL
41 #define HSA_CALL
42 #endif
43
44 #ifndef HSA_EXPORT_DECORATOR
45 #ifdef __GNUC__
46 #define HSA_EXPORT_DECORATOR __attribute__ ((visibility ("default")))
47 #else
48 #define HSA_EXPORT_DECORATOR
49 #endif
50 #endif
51 #define HSA_API_EXPORT HSA_EXPORT_DECORATOR HSA_CALL
52 #define HSA_API_IMPORT HSA_CALL
53
54 #if !defined(HSA_API) && defined(HSA_EXPORT)
55 #define HSA_API HSA_API_EXPORT
56 #else
57 #define HSA_API HSA_API_IMPORT
58 #endif
59
60 // Detect and set large model builds.
61 #undef HSA_LARGE_MODEL
62 #if defined(__LP64__) || defined(_M_X64)
63 #define HSA_LARGE_MODEL
64 #endif
65
66 // Try to detect CPU endianness
67 #if !defined(LITTLEENDIAN_CPU) && !defined(BIGENDIAN_CPU)
68 #if defined(__i386__) || defined(__x86_64__) || defined(_M_IX86) || \
69 defined(_M_X64)
70 #define LITTLEENDIAN_CPU
71 #endif
72 #endif
73
74 #undef HSA_LITTLE_ENDIAN
75 #if defined(LITTLEENDIAN_CPU)
76 #define HSA_LITTLE_ENDIAN
77 #elif defined(BIGENDIAN_CPU)
78 #else
79 #error "BIGENDIAN_CPU or LITTLEENDIAN_CPU must be defined"
80 #endif
81
82 #ifndef HSA_DEPRECATED
83 #define HSA_DEPRECATED
84 //#ifdef __GNUC__
85 //#define HSA_DEPRECATED __attribute__((deprecated))
86 //#else
87 //#define HSA_DEPRECATED __declspec(deprecated)
88 //#endif
89 #endif
90
91 #define HSA_VERSION_1_0 1
92
93 #ifdef __cplusplus
94 extern "C" {
95 #endif /* __cplusplus */
96
97 /** \defgroup status Runtime Notifications
98 * @{
99 */
100
101 /**
102 * @brief Status codes.
103 */
104 typedef enum {
105 /**
106 * The function has been executed successfully.
107 */
108 HSA_STATUS_SUCCESS = 0x0,
109 /**
110 * A traversal over a list of elements has been interrupted by the
111 * application before completing.
112 */
113 HSA_STATUS_INFO_BREAK = 0x1,
114 /**
115 * A generic error has occurred.
116 */
117 HSA_STATUS_ERROR = 0x1000,
118 /**
119 * One of the actual arguments does not meet a precondition stated in the
120 * documentation of the corresponding formal argument.
121 */
122 HSA_STATUS_ERROR_INVALID_ARGUMENT = 0x1001,
123 /**
124 * The requested queue creation is not valid.
125 */
126 HSA_STATUS_ERROR_INVALID_QUEUE_CREATION = 0x1002,
127 /**
128 * The requested allocation is not valid.
129 */
130 HSA_STATUS_ERROR_INVALID_ALLOCATION = 0x1003,
131 /**
132 * The agent is invalid.
133 */
134 HSA_STATUS_ERROR_INVALID_AGENT = 0x1004,
135 /**
136 * The memory region is invalid.
137 */
138 HSA_STATUS_ERROR_INVALID_REGION = 0x1005,
139 /**
140 * The signal is invalid.
141 */
142 HSA_STATUS_ERROR_INVALID_SIGNAL = 0x1006,
143 /**
144 * The queue is invalid.
145 */
146 HSA_STATUS_ERROR_INVALID_QUEUE = 0x1007,
147 /**
148 * The HSA runtime failed to allocate the necessary resources. This error
149 * may also occur when the HSA runtime needs to spawn threads or create
150 * internal OS-specific events.
151 */
152 HSA_STATUS_ERROR_OUT_OF_RESOURCES = 0x1008,
153 /**
154 * The AQL packet is malformed.
155 */
156 HSA_STATUS_ERROR_INVALID_PACKET_FORMAT = 0x1009,
157 /**
158 * An error has been detected while releasing a resource.
159 */
160 HSA_STATUS_ERROR_RESOURCE_FREE = 0x100A,
161 /**
162 * An API other than ::hsa_init has been invoked while the reference count
163 * of the HSA runtime is 0.
164 */
165 HSA_STATUS_ERROR_NOT_INITIALIZED = 0x100B,
166 /**
167 * The maximum reference count for the object has been reached.
168 */
169 HSA_STATUS_ERROR_REFCOUNT_OVERFLOW = 0x100C,
170 /**
171 * The arguments passed to a functions are not compatible.
172 */
173 HSA_STATUS_ERROR_INCOMPATIBLE_ARGUMENTS = 0x100D,
174 /**
175 * The index is invalid.
176 */
177 HSA_STATUS_ERROR_INVALID_INDEX = 0x100E,
178 /**
179 * The instruction set architecture is invalid.
180 */
181 HSA_STATUS_ERROR_INVALID_ISA = 0x100F,
182 /**
183 * The instruction set architecture name is invalid.
184 */
185 HSA_STATUS_ERROR_INVALID_ISA_NAME = 0x1017,
186 /**
187 * The code object is invalid.
188 */
189 HSA_STATUS_ERROR_INVALID_CODE_OBJECT = 0x1010,
190 /**
191 * The executable is invalid.
192 */
193 HSA_STATUS_ERROR_INVALID_EXECUTABLE = 0x1011,
194 /**
195 * The executable is frozen.
196 */
197 HSA_STATUS_ERROR_FROZEN_EXECUTABLE = 0x1012,
198 /**
199 * There is no symbol with the given name.
200 */
201 HSA_STATUS_ERROR_INVALID_SYMBOL_NAME = 0x1013,
202 /**
203 * The variable is already defined.
204 */
205 HSA_STATUS_ERROR_VARIABLE_ALREADY_DEFINED = 0x1014,
206 /**
207 * The variable is undefined.
208 */
209 HSA_STATUS_ERROR_VARIABLE_UNDEFINED = 0x1015,
210 /**
211 * An HSAIL operation resulted in a hardware exception.
212 */
213 HSA_STATUS_ERROR_EXCEPTION = 0x1016,
214 /**
215 * The code object symbol is invalid.
216 */
217 HSA_STATUS_ERROR_INVALID_CODE_SYMBOL = 0x1018,
218 /**
219 * The executable symbol is invalid.
220 */
221 HSA_STATUS_ERROR_INVALID_EXECUTABLE_SYMBOL = 0x1019,
222 /**
223 * The file descriptor is invalid.
224 */
225 HSA_STATUS_ERROR_INVALID_FILE = 0x1020,
226 /**
227 * The code object reader is invalid.
228 */
229 HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER = 0x1021,
230 /**
231 * The cache is invalid.
232 */
233 HSA_STATUS_ERROR_INVALID_CACHE = 0x1022,
234 /**
235 * The wavefront is invalid.
236 */
237 HSA_STATUS_ERROR_INVALID_WAVEFRONT = 0x1023,
238 /**
239 * The signal group is invalid.
240 */
241 HSA_STATUS_ERROR_INVALID_SIGNAL_GROUP = 0x1024,
242 /**
243 * The HSA runtime is not in the configuration state.
244 */
245 HSA_STATUS_ERROR_INVALID_RUNTIME_STATE = 0x1025,
246 /**
247 * The queue received an error that may require process termination.
248 */
249 HSA_STATUS_ERROR_FATAL = 0x1026
250 } hsa_status_t;
251
252 /**
253 * @brief Query additional information about a status code.
254 *
255 * @param[in] status Status code.
256 *
257 * @param[out] status_string A NUL-terminated string that describes the error
258 * status.
259 *
260 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
261 *
262 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
263 * initialized.
264 *
265 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p status is an invalid
266 * status code, or @p status_string is NULL.
267 */
268 hsa_status_t HSA_API hsa_status_string(
269 hsa_status_t status,
270 const char ** status_string);
271
272 /** @} */
273
274 /** \defgroup common Common Definitions
275 * @{
276 */
277
278 /**
279 * @brief Three-dimensional coordinate.
280 */
281 typedef struct hsa_dim3_s {
282 /**
283 * X dimension.
284 */
285 uint32_t x;
286
287 /**
288 * Y dimension.
289 */
290 uint32_t y;
291
292 /**
293 * Z dimension.
294 */
295 uint32_t z;
296 } hsa_dim3_t;
297
298 /**
299 * @brief Access permissions.
300 */
301 typedef enum {
302 /**
303 * Read-only access.
304 */
305 HSA_ACCESS_PERMISSION_RO = 1,
306 /**
307 * Write-only access.
308 */
309 HSA_ACCESS_PERMISSION_WO = 2,
310 /**
311 * Read and write access.
312 */
313 HSA_ACCESS_PERMISSION_RW = 3
314 } hsa_access_permission_t;
315
316 /**
317 * @brief POSIX file descriptor.
318 */
319 typedef int hsa_file_t;
320
321 /** @} **/
322
323
324 /** \defgroup initshutdown Initialization and Shut Down
325 * @{
326 */
327
328 /**
329 * @brief Initialize the HSA runtime.
330 *
331 * @details Initializes the HSA runtime if it is not already initialized, and
332 * increases the reference counter associated with the HSA runtime for the
333 * current process. Invocation of any HSA function other than ::hsa_init results
334 * in undefined behavior if the current HSA runtime reference counter is less
335 * than one.
336 *
337 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
338 *
339 * @retval ::HSA_STATUS_ERROR_OUT_OF_RESOURCES The HSA runtime failed to allocate
340 * the required resources.
341 *
342 * @retval ::HSA_STATUS_ERROR_REFCOUNT_OVERFLOW The HSA runtime reference
343 * count reaches INT32_MAX.
344 */
345 hsa_status_t HSA_API hsa_init();
346
347 /**
348 * @brief Shut down the HSA runtime.
349 *
350 * @details Decreases the reference count of the HSA runtime instance. When the
351 * reference count reaches 0, the HSA runtime is no longer considered valid
352 * but the application might call ::hsa_init to initialize the HSA runtime
353 * again.
354 *
355 * Once the reference count of the HSA runtime reaches 0, all the resources
356 * associated with it (queues, signals, agent information, etc.) are
357 * considered invalid and any attempt to reference them in subsequent API calls
358 * results in undefined behavior. When the reference count reaches 0, the HSA
359 * runtime may release resources associated with it.
360 *
361 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
362 *
363 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
364 * initialized.
365 *
366 */
367 hsa_status_t HSA_API hsa_shut_down();
368
369 /** @} **/
370
371 /** \defgroup agentinfo System and Agent Information
372 * @{
373 */
374
375 /**
376 * @brief Endianness. A convention used to interpret the bytes making up a data
377 * word.
378 */
379 typedef enum {
380 /**
381 * The least significant byte is stored in the smallest address.
382 */
383 HSA_ENDIANNESS_LITTLE = 0,
384 /**
385 * The most significant byte is stored in the smallest address.
386 */
387 HSA_ENDIANNESS_BIG = 1
388 } hsa_endianness_t;
389
390 /**
391 * @brief Machine model. A machine model determines the size of certain data
392 * types in HSA runtime and an agent.
393 */
394 typedef enum {
395 /**
396 * Small machine model. Addresses use 32 bits.
397 */
398 HSA_MACHINE_MODEL_SMALL = 0,
399 /**
400 * Large machine model. Addresses use 64 bits.
401 */
402 HSA_MACHINE_MODEL_LARGE = 1
403 } hsa_machine_model_t;
404
405 /**
406 * @brief Profile. A profile indicates a particular level of feature
407 * support. For example, in the base profile the application must use the HSA
408 * runtime allocator to reserve shared virtual memory, while in the full profile
409 * any host pointer can be shared across all the agents.
410 */
411 typedef enum {
412 /**
413 * Base profile.
414 */
415 HSA_PROFILE_BASE = 0,
416 /**
417 * Full profile.
418 */
419 HSA_PROFILE_FULL = 1
420 } hsa_profile_t;
421
422 /**
423 * @brief System attributes.
424 */
425 typedef enum {
426 /**
427 * Major version of the HSA runtime specification supported by the
428 * implementation. The type of this attribute is uint16_t.
429 */
430 HSA_SYSTEM_INFO_VERSION_MAJOR = 0,
431 /**
432 * Minor version of the HSA runtime specification supported by the
433 * implementation. The type of this attribute is uint16_t.
434 */
435 HSA_SYSTEM_INFO_VERSION_MINOR = 1,
436 /**
437 * Current timestamp. The value of this attribute monotonically increases at a
438 * constant rate. The type of this attribute is uint64_t.
439 */
440 HSA_SYSTEM_INFO_TIMESTAMP = 2,
441 /**
442 * Timestamp value increase rate, in Hz. The timestamp (clock) frequency is
443 * in the range 1-400MHz. The type of this attribute is uint64_t.
444 */
445 HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY = 3,
446 /**
447 * Maximum duration of a signal wait operation. Expressed as a count based on
448 * the timestamp frequency. The type of this attribute is uint64_t.
449 */
450 HSA_SYSTEM_INFO_SIGNAL_MAX_WAIT = 4,
451 /**
452 * Endianness of the system. The type of this attribute is ::hsa_endianness_t.
453 */
454 HSA_SYSTEM_INFO_ENDIANNESS = 5,
455 /**
456 * Machine model supported by the HSA runtime. The type of this attribute is
457 * ::hsa_machine_model_t.
458 */
459 HSA_SYSTEM_INFO_MACHINE_MODEL = 6,
460 /**
461 * Bit-mask indicating which extensions are supported by the
462 * implementation. An extension with an ID of @p i is supported if the bit at
463 * position @p i is set. The type of this attribute is uint8_t[128].
464 */
465 HSA_SYSTEM_INFO_EXTENSIONS = 7,
466 /**
467 * String containing the ROCr build identifier.
468 */
469 HSA_AMD_SYSTEM_INFO_BUILD_VERSION = 0x200
470 } hsa_system_info_t;
471
472 /**
473 * @brief Get the current value of a system attribute.
474 *
475 * @param[in] attribute Attribute to query.
476 *
477 * @param[out] value Pointer to an application-allocated buffer where to store
478 * the value of the attribute. If the buffer passed by the application is not
479 * large enough to hold the value of @p attribute, the behavior is undefined.
480 *
481 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
482 *
483 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
484 * initialized.
485 *
486 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p attribute is an invalid
487 * system attribute, or @p value is NULL.
488 */
489 hsa_status_t HSA_API hsa_system_get_info(
490 hsa_system_info_t attribute,
491 void* value);
492
493 /**
494 * @brief HSA extensions.
495 */
496 typedef enum {
497 /**
498 * Finalizer extension.
499 */
500 HSA_EXTENSION_FINALIZER = 0,
501 /**
502 * Images extension.
503 */
504 HSA_EXTENSION_IMAGES = 1,
505
506 /**
507 * Performance counter extension.
508 */
509 HSA_EXTENSION_PERFORMANCE_COUNTERS = 2,
510
511 /**
512 * Profiling events extension.
513 */
514 HSA_EXTENSION_PROFILING_EVENTS = 3,
515 /**
516 * Extension count.
517 */
518 HSA_EXTENSION_STD_LAST = 3,
519 /**
520 * First AMD extension number.
521 */
522 HSA_AMD_FIRST_EXTENSION = 0x200,
523 /**
524 * Profiler extension.
525 */
526 HSA_EXTENSION_AMD_PROFILER = 0x200,
527 /**
528 * Loader extension.
529 */
530 HSA_EXTENSION_AMD_LOADER = 0x201,
531 /**
532 * AqlProfile extension.
533 */
534 HSA_EXTENSION_AMD_AQLPROFILE = 0x202,
535 /**
536 * Last AMD extension.
537 */
538 HSA_AMD_LAST_EXTENSION = 0x202
539 } hsa_extension_t;
540
541 /**
542 * @brief Query the name of a given extension.
543 *
544 * @param[in] extension Extension identifier. If the extension is not supported
545 * by the implementation (see ::HSA_SYSTEM_INFO_EXTENSIONS), the behavior
546 * is undefined.
547 *
548 * @param[out] name Pointer to a memory location where the HSA runtime stores
549 * the extension name. The extension name is a NUL-terminated string.
550 *
551 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
552 *
553 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
554 * initialized.
555 *
556 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p extension is not a valid
557 * extension, or @p name is NULL.
558 */
559 hsa_status_t HSA_API hsa_extension_get_name(
560 uint16_t extension,
561 const char **name);
562
563 /**
564 * @deprecated
565 *
566 * @brief Query if a given version of an extension is supported by the HSA
567 * implementation.
568 *
569 * @param[in] extension Extension identifier.
570 *
571 * @param[in] version_major Major version number.
572 *
573 * @param[in] version_minor Minor version number.
574 *
575 * @param[out] result Pointer to a memory location where the HSA runtime stores
576 * the result of the check. The result is true if the specified version of the
577 * extension is supported, and false otherwise.
578 *
579 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
580 *
581 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
582 * initialized.
583 *
584 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p extension is not a valid
585 * extension, or @p result is NULL.
586 */
587 hsa_status_t HSA_API HSA_DEPRECATED hsa_system_extension_supported(
588 uint16_t extension,
589 uint16_t version_major,
590 uint16_t version_minor,
591 bool* result);
592
593 /**
594 * @brief Query if a given version of an extension is supported by the HSA
595 * implementation. All minor versions from 0 up to the returned @p version_minor
596 * must be supported by the implementation.
597 *
598 * @param[in] extension Extension identifier.
599 *
600 * @param[in] version_major Major version number.
601 *
602 * @param[out] version_minor Minor version number.
603 *
604 * @param[out] result Pointer to a memory location where the HSA runtime stores
605 * the result of the check. The result is true if the specified version of the
606 * extension is supported, and false otherwise.
607 *
608 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
609 *
610 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
611 * initialized.
612 *
613 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p extension is not a valid
614 * extension, or @p version_minor is NULL, or @p result is NULL.
615 */
616 hsa_status_t HSA_API hsa_system_major_extension_supported(
617 uint16_t extension,
618 uint16_t version_major,
619 uint16_t *version_minor,
620 bool* result);
621
622
623 /**
624 * @deprecated
625 *
626 * @brief Retrieve the function pointers corresponding to a given version of an
627 * extension. Portable applications are expected to invoke the extension API
628 * using the returned function pointers
629 *
630 * @details The application is responsible for verifying that the given version
631 * of the extension is supported by the HSA implementation (see
632 * ::hsa_system_extension_supported). If the given combination of extension,
633 * major version, and minor version is not supported by the implementation, the
634 * behavior is undefined.
635 *
636 * @param[in] extension Extension identifier.
637 *
638 * @param[in] version_major Major version number for which to retrieve the
639 * function pointer table.
640 *
641 * @param[in] version_minor Minor version number for which to retrieve the
642 * function pointer table.
643 *
644 * @param[out] table Pointer to an application-allocated function pointer table
645 * that is populated by the HSA runtime. Must not be NULL. The memory associated
646 * with table can be reused or freed after the function returns.
647 *
648 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
649 *
650 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
651 * initialized.
652 *
653 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p extension is not a valid
654 * extension, or @p table is NULL.
655 */
656 hsa_status_t HSA_API HSA_DEPRECATED hsa_system_get_extension_table(
657 uint16_t extension,
658 uint16_t version_major,
659 uint16_t version_minor,
660 void *table);
661
662 /**
663 * @brief Retrieve the function pointers corresponding to a given major version
664 * of an extension. Portable applications are expected to invoke the extension
665 * API using the returned function pointers.
666 *
667 * @details The application is responsible for verifying that the given major
668 * version of the extension is supported by the HSA implementation (see
669 * ::hsa_system_major_extension_supported). If the given combination of extension
670 * and major version is not supported by the implementation, the behavior is
671 * undefined. Additionally if the length doesn't allow space for a full minor
672 * version, it is implementation defined if only some of the function pointers for
673 * that minor version get written.
674 *
675 * @param[in] extension Extension identifier.
676 *
677 * @param[in] version_major Major version number for which to retrieve the
678 * function pointer table.
679 *
680 * @param[in] table_length Size in bytes of the function pointer table to be
681 * populated. The implementation will not write more than this many bytes to the
682 * table.
683 *
684 * @param[out] table Pointer to an application-allocated function pointer table
685 * that is populated by the HSA runtime. Must not be NULL. The memory associated
686 * with table can be reused or freed after the function returns.
687 *
688 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
689 *
690 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
691 * initialized.
692 *
693 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p extension is not a valid
694 * extension, or @p table is NULL.
695 */
696 hsa_status_t HSA_API hsa_system_get_major_extension_table(
697 uint16_t extension,
698 uint16_t version_major,
699 size_t table_length,
700 void *table);
701
702 /**
703 * @brief Struct containing an opaque handle to an agent, a device that participates in
704 * the HSA memory model. An agent can submit AQL packets for execution, and
705 * may also accept AQL packets for execution (agent dispatch packets or kernel
706 * dispatch packets launching HSAIL-derived binaries).
707 */
708 typedef struct hsa_agent_s {
709 /**
710 * Opaque handle. Two handles reference the same object of the enclosing type
711 * if and only if they are equal.
712 */
713 uint64_t handle;
714 } hsa_agent_t;
715
716 /**
717 * @brief Agent features.
718 */
719 typedef enum {
720 /**
721 * The agent supports AQL packets of kernel dispatch type. If this
722 * feature is enabled, the agent is also a kernel agent.
723 */
724 HSA_AGENT_FEATURE_KERNEL_DISPATCH = 1,
725 /**
726 * The agent supports AQL packets of agent dispatch type.
727 */
728 HSA_AGENT_FEATURE_AGENT_DISPATCH = 2
729 } hsa_agent_feature_t;
730
731 /**
732 * @brief Hardware device type.
733 */
734 typedef enum {
735 /**
736 * CPU device.
737 */
738 HSA_DEVICE_TYPE_CPU = 0,
739 /**
740 * GPU device.
741 */
742 HSA_DEVICE_TYPE_GPU = 1,
743 /**
744 * DSP device.
745 */
746 HSA_DEVICE_TYPE_DSP = 2
747 } hsa_device_type_t;
748
749 /**
750 * @brief Default floating-point rounding mode.
751 */
752 typedef enum {
753 /**
754 * Use a default floating-point rounding mode specified elsewhere.
755 */
756 HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT = 0,
757 /**
758 * Operations that specify the default floating-point mode are rounded to zero
759 * by default.
760 */
761 HSA_DEFAULT_FLOAT_ROUNDING_MODE_ZERO = 1,
762 /**
763 * Operations that specify the default floating-point mode are rounded to the
764 * nearest representable number and that ties should be broken by selecting
765 * the value with an even least significant bit.
766 */
767 HSA_DEFAULT_FLOAT_ROUNDING_MODE_NEAR = 2
768 } hsa_default_float_rounding_mode_t;
769
770 /**
771 * @brief Agent attributes.
772 */
773 typedef enum {
774 /**
775 * Agent name. The type of this attribute is a NUL-terminated char[64]. The
776 * name must be at most 63 characters long (not including the NUL terminator)
777 * and all array elements not used for the name must be NUL.
778 */
779 HSA_AGENT_INFO_NAME = 0,
780 /**
781 * Name of vendor. The type of this attribute is a NUL-terminated char[64].
782 * The name must be at most 63 characters long (not including the NUL
783 * terminator) and all array elements not used for the name must be NUL.
784 */
785 HSA_AGENT_INFO_VENDOR_NAME = 1,
786 /**
787 * Agent capability. The type of this attribute is ::hsa_agent_feature_t.
788 */
789 HSA_AGENT_INFO_FEATURE = 2,
790 /**
791 * @deprecated Query ::HSA_ISA_INFO_MACHINE_MODELS for a given intruction set
792 * architecture supported by the agent instead. If more than one ISA is
793 * supported by the agent, the returned value corresponds to the first ISA
794 * enumerated by ::hsa_agent_iterate_isas.
795 *
796 * Machine model supported by the agent. The type of this attribute is
797 * ::hsa_machine_model_t.
798 */
799 HSA_AGENT_INFO_MACHINE_MODEL = 3,
800 /**
801 * @deprecated Query ::HSA_ISA_INFO_PROFILES for a given intruction set
802 * architecture supported by the agent instead. If more than one ISA is
803 * supported by the agent, the returned value corresponds to the first ISA
804 * enumerated by ::hsa_agent_iterate_isas.
805 *
806 * Profile supported by the agent. The type of this attribute is
807 * ::hsa_profile_t.
808 */
809 HSA_AGENT_INFO_PROFILE = 4,
810 /**
811 * @deprecated Query ::HSA_ISA_INFO_DEFAULT_FLOAT_ROUNDING_MODES for a given
812 * intruction set architecture supported by the agent instead. If more than
813 * one ISA is supported by the agent, the returned value corresponds to the
814 * first ISA enumerated by ::hsa_agent_iterate_isas.
815 *
816 * Default floating-point rounding mode. The type of this attribute is
817 * ::hsa_default_float_rounding_mode_t, but the value
818 * ::HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT is not allowed.
819 */
820 HSA_AGENT_INFO_DEFAULT_FLOAT_ROUNDING_MODE = 5,
821 /**
822 * @deprecated Query ::HSA_ISA_INFO_BASE_PROFILE_DEFAULT_FLOAT_ROUNDING_MODES
823 * for a given intruction set architecture supported by the agent instead. If
824 * more than one ISA is supported by the agent, the returned value corresponds
825 * to the first ISA enumerated by ::hsa_agent_iterate_isas.
826 *
827 * A bit-mask of ::hsa_default_float_rounding_mode_t values, representing the
828 * default floating-point rounding modes supported by the agent in the Base
829 * profile. The type of this attribute is uint32_t. The default floating-point
830 * rounding mode (::HSA_AGENT_INFO_DEFAULT_FLOAT_ROUNDING_MODE) bit must not
831 * be set.
832 */
833 HSA_AGENT_INFO_BASE_PROFILE_DEFAULT_FLOAT_ROUNDING_MODES = 23,
834 /**
835 * @deprecated Query ::HSA_ISA_INFO_FAST_F16_OPERATION for a given intruction
836 * set architecture supported by the agent instead. If more than one ISA is
837 * supported by the agent, the returned value corresponds to the first ISA
838 * enumerated by ::hsa_agent_iterate_isas.
839 *
840 * Flag indicating that the f16 HSAIL operation is at least as fast as the
841 * f32 operation in the current agent. The value of this attribute is
842 * undefined if the agent is not a kernel agent. The type of this
843 * attribute is bool.
844 */
845 HSA_AGENT_INFO_FAST_F16_OPERATION = 24,
846 /**
847 * @deprecated Query ::HSA_WAVEFRONT_INFO_SIZE for a given wavefront and
848 * intruction set architecture supported by the agent instead. If more than
849 * one ISA is supported by the agent, the returned value corresponds to the
850 * first ISA enumerated by ::hsa_agent_iterate_isas and the first wavefront
851 * enumerated by ::hsa_isa_iterate_wavefronts for that ISA.
852 *
853 * Number of work-items in a wavefront. Must be a power of 2 in the range
854 * [1,256]. The value of this attribute is undefined if the agent is not
855 * a kernel agent. The type of this attribute is uint32_t.
856 */
857 HSA_AGENT_INFO_WAVEFRONT_SIZE = 6,
858 /**
859 * @deprecated Query ::HSA_ISA_INFO_WORKGROUP_MAX_DIM for a given intruction
860 * set architecture supported by the agent instead. If more than one ISA is
861 * supported by the agent, the returned value corresponds to the first ISA
862 * enumerated by ::hsa_agent_iterate_isas.
863 *
864 * Maximum number of work-items of each dimension of a work-group. Each
865 * maximum must be greater than 0. No maximum can exceed the value of
866 * ::HSA_AGENT_INFO_WORKGROUP_MAX_SIZE. The value of this attribute is
867 * undefined if the agent is not a kernel agent. The type of this
868 * attribute is uint16_t[3].
869 */
870 HSA_AGENT_INFO_WORKGROUP_MAX_DIM = 7,
871 /**
872 * @deprecated Query ::HSA_ISA_INFO_WORKGROUP_MAX_SIZE for a given intruction
873 * set architecture supported by the agent instead. If more than one ISA is
874 * supported by the agent, the returned value corresponds to the first ISA
875 * enumerated by ::hsa_agent_iterate_isas.
876 *
877 * Maximum total number of work-items in a work-group. The value of this
878 * attribute is undefined if the agent is not a kernel agent. The type
879 * of this attribute is uint32_t.
880 */
881 HSA_AGENT_INFO_WORKGROUP_MAX_SIZE = 8,
882 /**
883 * @deprecated Query ::HSA_ISA_INFO_GRID_MAX_DIM for a given intruction set
884 * architecture supported by the agent instead.
885 *
886 * Maximum number of work-items of each dimension of a grid. Each maximum must
887 * be greater than 0, and must not be smaller than the corresponding value in
888 * ::HSA_AGENT_INFO_WORKGROUP_MAX_DIM. No maximum can exceed the value of
889 * ::HSA_AGENT_INFO_GRID_MAX_SIZE. The value of this attribute is undefined
890 * if the agent is not a kernel agent. The type of this attribute is
891 * ::hsa_dim3_t.
892 */
893 HSA_AGENT_INFO_GRID_MAX_DIM = 9,
894 /**
895 * @deprecated Query ::HSA_ISA_INFO_GRID_MAX_SIZE for a given intruction set
896 * architecture supported by the agent instead. If more than one ISA is
897 * supported by the agent, the returned value corresponds to the first ISA
898 * enumerated by ::hsa_agent_iterate_isas.
899 *
900 * Maximum total number of work-items in a grid. The value of this attribute
901 * is undefined if the agent is not a kernel agent. The type of this
902 * attribute is uint32_t.
903 */
904 HSA_AGENT_INFO_GRID_MAX_SIZE = 10,
905 /**
906 * @deprecated Query ::HSA_ISA_INFO_FBARRIER_MAX_SIZE for a given intruction
907 * set architecture supported by the agent instead. If more than one ISA is
908 * supported by the agent, the returned value corresponds to the first ISA
909 * enumerated by ::hsa_agent_iterate_isas.
910 *
911 * Maximum number of fbarriers per work-group. Must be at least 32. The value
912 * of this attribute is undefined if the agent is not a kernel agent. The
913 * type of this attribute is uint32_t.
914 */
915 HSA_AGENT_INFO_FBARRIER_MAX_SIZE = 11,
916 /**
917 * @deprecated The maximum number of queues is not statically determined.
918 *
919 * Maximum number of queues that can be active (created but not destroyed) at
920 * one time in the agent. The type of this attribute is uint32_t.
921 */
922 HSA_AGENT_INFO_QUEUES_MAX = 12,
923 /**
924 * Minimum number of packets that a queue created in the agent
925 * can hold. Must be a power of 2 greater than 0. Must not exceed
926 * the value of ::HSA_AGENT_INFO_QUEUE_MAX_SIZE. The type of this
927 * attribute is uint32_t.
928 */
929 HSA_AGENT_INFO_QUEUE_MIN_SIZE = 13,
930 /**
931 * Maximum number of packets that a queue created in the agent can
932 * hold. Must be a power of 2 greater than 0. The type of this attribute
933 * is uint32_t.
934 */
935 HSA_AGENT_INFO_QUEUE_MAX_SIZE = 14,
936 /**
937 * Type of a queue created in the agent. The type of this attribute is
938 * ::hsa_queue_type32_t.
939 */
940 HSA_AGENT_INFO_QUEUE_TYPE = 15,
941 /**
942 * @deprecated NUMA information is not exposed anywhere else in the API.
943 *
944 * Identifier of the NUMA node associated with the agent. The type of this
945 * attribute is uint32_t.
946 */
947 HSA_AGENT_INFO_NODE = 16,
948 /**
949 * Type of hardware device associated with the agent. The type of this
950 * attribute is ::hsa_device_type_t.
951 */
952 HSA_AGENT_INFO_DEVICE = 17,
953 /**
954 * @deprecated Query ::hsa_agent_iterate_caches to retrieve information about
955 * the caches present in a given agent.
956 *
957 * Array of data cache sizes (L1..L4). Each size is expressed in bytes. A size
958 * of 0 for a particular level indicates that there is no cache information
959 * for that level. The type of this attribute is uint32_t[4].
960 */
961 HSA_AGENT_INFO_CACHE_SIZE = 18,
962 /**
963 * @deprecated An agent may support multiple instruction set
964 * architectures. See ::hsa_agent_iterate_isas. If more than one ISA is
965 * supported by the agent, the returned value corresponds to the first ISA
966 * enumerated by ::hsa_agent_iterate_isas.
967 *
968 * Instruction set architecture of the agent. The type of this attribute
969 * is ::hsa_isa_t.
970 */
971 HSA_AGENT_INFO_ISA = 19,
972 /**
973 * Bit-mask indicating which extensions are supported by the agent. An
974 * extension with an ID of @p i is supported if the bit at position @p i is
975 * set. The type of this attribute is uint8_t[128].
976 */
977 HSA_AGENT_INFO_EXTENSIONS = 20,
978 /**
979 * Major version of the HSA runtime specification supported by the
980 * agent. The type of this attribute is uint16_t.
981 */
982 HSA_AGENT_INFO_VERSION_MAJOR = 21,
983 /**
984 * Minor version of the HSA runtime specification supported by the
985 * agent. The type of this attribute is uint16_t.
986 */
987 HSA_AGENT_INFO_VERSION_MINOR = 22
988
989 } hsa_agent_info_t;
990
991 /**
992 * @brief Get the current value of an attribute for a given agent.
993 *
994 * @param[in] agent A valid agent.
995 *
996 * @param[in] attribute Attribute to query.
997 *
998 * @param[out] value Pointer to an application-allocated buffer where to store
999 * the value of the attribute. If the buffer passed by the application is not
1000 * large enough to hold the value of @p attribute, the behavior is undefined.
1001 *
1002 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
1003 *
1004 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
1005 * initialized.
1006 *
1007 * @retval ::HSA_STATUS_ERROR_INVALID_AGENT The agent is invalid.
1008 *
1009 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p attribute is an invalid
1010 * agent attribute, or @p value is NULL.
1011 */
1012 hsa_status_t HSA_API hsa_agent_get_info(
1013 hsa_agent_t agent,
1014 hsa_agent_info_t attribute,
1015 void* value);
1016
1017 /**
1018 * @brief Iterate over the available agents, and invoke an
1019 * application-defined callback on every iteration.
1020 *
1021 * @param[in] callback Callback to be invoked once per agent. The HSA
1022 * runtime passes two arguments to the callback: the agent and the
1023 * application data. If @p callback returns a status other than
1024 * ::HSA_STATUS_SUCCESS for a particular iteration, the traversal stops and
1025 * ::hsa_iterate_agents returns that status value.
1026 *
1027 * @param[in] data Application data that is passed to @p callback on every
1028 * iteration. May be NULL.
1029 *
1030 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
1031 *
1032 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
1033 * initialized.
1034 *
1035 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p callback is NULL.
1036 */
1037 hsa_status_t HSA_API hsa_iterate_agents(
1038 hsa_status_t (*callback)(hsa_agent_t agent, void* data),
1039 void* data);
1040
1041 /*
1042
1043 // If we do not know the size of an attribute, we need to query it first
1044 // Note: this API will not be in the spec unless needed
1045 hsa_status_t HSA_API hsa_agent_get_info_size(
1046 hsa_agent_t agent,
1047 hsa_agent_info_t attribute,
1048 size_t* size);
1049
1050 // Set the value of an agents attribute
1051 // Note: this API will not be in the spec unless needed
1052 hsa_status_t HSA_API hsa_agent_set_info(
1053 hsa_agent_t agent,
1054 hsa_agent_info_t attribute,
1055 void* value);
1056
1057 */
1058
1059 /**
1060 * @brief Exception policies applied in the presence of hardware exceptions.
1061 */
1062 typedef enum {
1063 /**
1064 * If a hardware exception is detected, a work-item signals an exception.
1065 */
1066 HSA_EXCEPTION_POLICY_BREAK = 1,
1067 /**
1068 * If a hardware exception is detected, a hardware status bit is set.
1069 */
1070 HSA_EXCEPTION_POLICY_DETECT = 2
1071 } hsa_exception_policy_t;
1072
1073 /**
1074 * @deprecated Use ::hsa_isa_get_exception_policies for a given intruction set
1075 * architecture supported by the agent instead. If more than one ISA is
1076 * supported by the agent, this function uses the first value returned by
1077 * ::hsa_agent_iterate_isas.
1078 *
1079 * @brief Retrieve the exception policy support for a given combination of
1080 * agent and profile
1081 *
1082 * @param[in] agent Agent.
1083 *
1084 * @param[in] profile Profile.
1085 *
1086 * @param[out] mask Pointer to a memory location where the HSA runtime stores a
1087 * mask of ::hsa_exception_policy_t values. Must not be NULL.
1088 *
1089 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
1090 *
1091 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
1092 * initialized.
1093 *
1094 * @retval ::HSA_STATUS_ERROR_INVALID_AGENT The agent is invalid.
1095 *
1096 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p profile is not a valid
1097 * profile, or @p mask is NULL.
1098 *
1099 */
1100 hsa_status_t HSA_API HSA_DEPRECATED hsa_agent_get_exception_policies(
1101 hsa_agent_t agent,
1102 hsa_profile_t profile,
1103 uint16_t *mask);
1104
1105 /**
1106 * @brief Cache handle.
1107 */
1108 typedef struct hsa_cache_s {
1109 /**
1110 * Opaque handle. Two handles reference the same object of the enclosing type
1111 * if and only if they are equal.
1112 */
1113 uint64_t handle;
1114 } hsa_cache_t;
1115
1116 /**
1117 * @brief Cache attributes.
1118 */
1119 typedef enum {
1120 /**
1121 * The length of the cache name in bytes, not including the NUL terminator.
1122 * The type of this attribute is uint32_t.
1123 */
1124 HSA_CACHE_INFO_NAME_LENGTH = 0,
1125 /**
1126 * Human-readable description. The type of this attribute is a NUL-terminated
1127 * character array with the length equal to the value of
1128 * ::HSA_CACHE_INFO_NAME_LENGTH attribute.
1129 */
1130 HSA_CACHE_INFO_NAME = 1,
1131 /**
1132 * Cache level. A L1 cache must return a value of 1, a L2 must return a value
1133 * of 2, and so on. The type of this attribute is uint8_t.
1134 */
1135 HSA_CACHE_INFO_LEVEL = 2,
1136 /**
1137 * Cache size, in bytes. A value of 0 indicates that there is no size
1138 * information available. The type of this attribute is uint32_t.
1139 */
1140 HSA_CACHE_INFO_SIZE = 3
1141 } hsa_cache_info_t;
1142
1143 /**
1144 * @brief Get the current value of an attribute for a given cache object.
1145 *
1146 * @param[in] cache Cache.
1147 *
1148 * @param[in] attribute Attribute to query.
1149 *
1150 * @param[out] value Pointer to an application-allocated buffer where to store
1151 * the value of the attribute. If the buffer passed by the application is not
1152 * large enough to hold the value of @p attribute, the behavior is undefined.
1153 *
1154 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
1155 *
1156 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
1157 * initialized.
1158 *
1159 * @retval ::HSA_STATUS_ERROR_INVALID_CACHE The cache is invalid.
1160 *
1161 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p attribute is an invalid
1162 * instruction set architecture attribute, or @p value is
1163 * NULL.
1164 */
1165 hsa_status_t HSA_API hsa_cache_get_info(
1166 hsa_cache_t cache,
1167 hsa_cache_info_t attribute,
1168 void* value);
1169
1170 /**
1171 * @brief Iterate over the memory caches of a given agent, and
1172 * invoke an application-defined callback on every iteration.
1173 *
1174 * @details Caches are visited in ascending order according to the value of the
1175 * ::HSA_CACHE_INFO_LEVEL attribute.
1176 *
1177 * @param[in] agent A valid agent.
1178 *
1179 * @param[in] callback Callback to be invoked once per cache that is present in
1180 * the agent. The HSA runtime passes two arguments to the callback: the cache
1181 * and the application data. If @p callback returns a status other than
1182 * ::HSA_STATUS_SUCCESS for a particular iteration, the traversal stops and
1183 * that value is returned.
1184 *
1185 * @param[in] data Application data that is passed to @p callback on every
1186 * iteration. May be NULL.
1187 *
1188 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
1189 *
1190 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
1191 * initialized.
1192 *
1193 * @retval ::HSA_STATUS_ERROR_INVALID_AGENT The agent is invalid.
1194 *
1195 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p callback is NULL.
1196 */
1197 hsa_status_t HSA_API hsa_agent_iterate_caches(
1198 hsa_agent_t agent,
1199 hsa_status_t (*callback)(hsa_cache_t cache, void* data),
1200 void* data);
1201
1202 /**
1203 * @deprecated
1204 *
1205 * @brief Query if a given version of an extension is supported by an agent
1206 *
1207 * @param[in] extension Extension identifier.
1208 *
1209 * @param[in] agent Agent.
1210 *
1211 * @param[in] version_major Major version number.
1212 *
1213 * @param[in] version_minor Minor version number.
1214 *
1215 * @param[out] result Pointer to a memory location where the HSA runtime stores
1216 * the result of the check. The result is true if the specified version of the
1217 * extension is supported, and false otherwise. The result must be false if
1218 * ::hsa_system_extension_supported returns false for the same extension
1219 * version.
1220 *
1221 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
1222 *
1223 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
1224 * initialized.
1225 *
1226 * @retval ::HSA_STATUS_ERROR_INVALID_AGENT The agent is invalid.
1227 *
1228 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p extension is not a valid
1229 * extension, or @p result is NULL.
1230 */
1231 hsa_status_t HSA_API HSA_DEPRECATED hsa_agent_extension_supported(
1232 uint16_t extension,
1233 hsa_agent_t agent,
1234 uint16_t version_major,
1235 uint16_t version_minor,
1236 bool* result);
1237
1238 /**
1239 * @brief Query if a given version of an extension is supported by an agent. All
1240 * minor versions from 0 up to the returned @p version_minor must be supported.
1241 *
1242 * @param[in] extension Extension identifier.
1243 *
1244 * @param[in] agent Agent.
1245 *
1246 * @param[in] version_major Major version number.
1247 *
1248 * @param[out] version_minor Minor version number.
1249 *
1250 * @param[out] result Pointer to a memory location where the HSA runtime stores
1251 * the result of the check. The result is true if the specified version of the
1252 * extension is supported, and false otherwise. The result must be false if
1253 * ::hsa_system_extension_supported returns false for the same extension
1254 * version.
1255 *
1256 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
1257 *
1258 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
1259 * initialized.
1260 *
1261 * @retval ::HSA_STATUS_ERROR_INVALID_AGENT The agent is invalid.
1262 *
1263 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p extension is not a valid
1264 * extension, or @p version_minor is NULL, or @p result is NULL.
1265 */
1266 hsa_status_t HSA_API hsa_agent_major_extension_supported(
1267 uint16_t extension,
1268 hsa_agent_t agent,
1269 uint16_t version_major,
1270 uint16_t *version_minor,
1271 bool* result);
1272
1273
1274 /** @} */
1275
1276
1277 /** \defgroup signals Signals
1278 * @{
1279 */
1280
1281 /**
1282 * @brief Signal handle.
1283 */
1284 typedef struct hsa_signal_s {
1285 /**
1286 * Opaque handle. Two handles reference the same object of the enclosing type
1287 * if and only if they are equal. The value 0 is reserved.
1288 */
1289 uint64_t handle;
1290 } hsa_signal_t;
1291
1292 /**
1293 * @brief Signal value. The value occupies 32 bits in small machine mode, and 64
1294 * bits in large machine mode.
1295 */
1296 #ifdef HSA_LARGE_MODEL
1297 typedef int64_t hsa_signal_value_t;
1298 #else
1299 typedef int32_t hsa_signal_value_t;
1300 #endif
1301
1302 /**
1303 * @brief Create a signal.
1304 *
1305 * @param[in] initial_value Initial value of the signal.
1306 *
1307 * @param[in] num_consumers Size of @p consumers. A value of 0 indicates that
1308 * any agent might wait on the signal.
1309 *
1310 * @param[in] consumers List of agents that might consume (wait on) the
1311 * signal. If @p num_consumers is 0, this argument is ignored; otherwise, the
1312 * HSA runtime might use the list to optimize the handling of the signal
1313 * object. If an agent not listed in @p consumers waits on the returned
1314 * signal, the behavior is undefined. The memory associated with @p consumers
1315 * can be reused or freed after the function returns.
1316 *
1317 * @param[out] signal Pointer to a memory location where the HSA runtime will
1318 * store the newly created signal handle. Must not be NULL.
1319 *
1320 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
1321 *
1322 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
1323 * initialized.
1324 *
1325 * @retval ::HSA_STATUS_ERROR_OUT_OF_RESOURCES The HSA runtime failed to allocate
1326 * the required resources.
1327 *
1328 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p signal is NULL, @p
1329 * num_consumers is greater than 0 but @p consumers is NULL, or @p consumers
1330 * contains duplicates.
1331 */
1332 hsa_status_t HSA_API hsa_signal_create(
1333 hsa_signal_value_t initial_value,
1334 uint32_t num_consumers,
1335 const hsa_agent_t *consumers,
1336 hsa_signal_t *signal);
1337
1338 /**
1339 * @brief Destroy a signal previous created by ::hsa_signal_create.
1340 *
1341 * @param[in] signal Signal.
1342 *
1343 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
1344 *
1345 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
1346 * initialized.
1347 *
1348 * @retval ::HSA_STATUS_ERROR_INVALID_SIGNAL @p signal is invalid.
1349 *
1350 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT The handle in @p signal is 0.
1351 */
1352 hsa_status_t HSA_API hsa_signal_destroy(
1353 hsa_signal_t signal);
1354
1355 /**
1356 * @brief Atomically read the current value of a signal.
1357 *
1358 * @param[in] signal Signal.
1359 *
1360 * @return Value of the signal.
1361 */
1362 hsa_signal_value_t HSA_API hsa_signal_load_scacquire(
1363 hsa_signal_t signal);
1364
1365 /**
1366 * @copydoc hsa_signal_load_scacquire
1367 */
1368 hsa_signal_value_t HSA_API hsa_signal_load_relaxed(
1369 hsa_signal_t signal);
1370
1371 /**
1372 * @deprecated Renamed as ::hsa_signal_load_scacquire.
1373 *
1374 * @copydoc hsa_signal_load_scacquire
1375 */
1376 hsa_signal_value_t HSA_API HSA_DEPRECATED hsa_signal_load_acquire(
1377 hsa_signal_t signal);
1378
1379 /**
1380 * @brief Atomically set the value of a signal.
1381 *
1382 * @details If the value of the signal is changed, all the agents waiting
1383 * on @p signal for which @p value satisfies their wait condition are awakened.
1384 *
1385 * @param[in] signal Signal.
1386 *
1387 * @param[in] value New signal value.
1388 */
1389 void HSA_API hsa_signal_store_relaxed(
1390 hsa_signal_t signal,
1391 hsa_signal_value_t value);
1392
1393 /**
1394 * @copydoc hsa_signal_store_relaxed
1395 */
1396 void HSA_API hsa_signal_store_screlease(
1397 hsa_signal_t signal,
1398 hsa_signal_value_t value);
1399
1400 /**
1401 * @deprecated Renamed as ::hsa_signal_store_screlease.
1402 *
1403 * @copydoc hsa_signal_store_screlease
1404 */
1405 void HSA_API HSA_DEPRECATED hsa_signal_store_release(
1406 hsa_signal_t signal,
1407 hsa_signal_value_t value);
1408
1409 /**
1410 * @brief Atomically set the value of a signal without necessarily notifying the
1411 * the agents waiting on it.
1412 *
1413 * @details The agents waiting on @p signal may not wake up even when the new
1414 * value satisfies their wait condition. If the application wants to update the
1415 * signal and there is no need to notify any agent, invoking this function can
1416 * be more efficient than calling the non-silent counterpart.
1417 *
1418 * @param[in] signal Signal.
1419 *
1420 * @param[in] value New signal value.
1421 */
1422 void HSA_API hsa_signal_silent_store_relaxed(
1423 hsa_signal_t signal,
1424 hsa_signal_value_t value);
1425
1426 /**
1427 * @copydoc hsa_signal_silent_store_relaxed
1428 */
1429 void HSA_API hsa_signal_silent_store_screlease(
1430 hsa_signal_t signal,
1431 hsa_signal_value_t value);
1432
1433 /**
1434 * @brief Atomically set the value of a signal and return its previous value.
1435 *
1436 * @details If the value of the signal is changed, all the agents waiting
1437 * on @p signal for which @p value satisfies their wait condition are awakened.
1438 *
1439 * @param[in] signal Signal. If @p signal is a queue doorbell signal, the
1440 * behavior is undefined.
1441 *
1442 * @param[in] value New value.
1443 *
1444 * @return Value of the signal prior to the exchange.
1445 *
1446 */
1447 hsa_signal_value_t HSA_API hsa_signal_exchange_scacq_screl(
1448 hsa_signal_t signal,
1449 hsa_signal_value_t value);
1450
1451 /**
1452 * @deprecated Renamed as ::hsa_signal_exchange_scacq_screl.
1453 *
1454 * @copydoc hsa_signal_exchange_scacq_screl
1455 */
1456 hsa_signal_value_t HSA_API HSA_DEPRECATED hsa_signal_exchange_acq_rel(
1457 hsa_signal_t signal,
1458 hsa_signal_value_t value);
1459
1460 /**
1461 * @copydoc hsa_signal_exchange_scacq_screl
1462 */
1463 hsa_signal_value_t HSA_API hsa_signal_exchange_scacquire(
1464 hsa_signal_t signal,
1465 hsa_signal_value_t value);
1466
1467 /**
1468 * @deprecated Renamed as ::hsa_signal_exchange_scacquire.
1469 *
1470 * @copydoc hsa_signal_exchange_scacquire
1471 */
1472 hsa_signal_value_t HSA_API HSA_DEPRECATED hsa_signal_exchange_acquire(
1473 hsa_signal_t signal,
1474 hsa_signal_value_t value);
1475
1476 /**
1477 * @copydoc hsa_signal_exchange_scacq_screl
1478 */
1479 hsa_signal_value_t HSA_API hsa_signal_exchange_relaxed(
1480 hsa_signal_t signal,
1481 hsa_signal_value_t value);
1482 /**
1483 * @copydoc hsa_signal_exchange_scacq_screl
1484 */
1485 hsa_signal_value_t HSA_API hsa_signal_exchange_screlease(
1486 hsa_signal_t signal,
1487 hsa_signal_value_t value);
1488
1489 /**
1490 * @deprecated Renamed as ::hsa_signal_exchange_screlease.
1491 *
1492 * @copydoc hsa_signal_exchange_screlease
1493 */
1494 hsa_signal_value_t HSA_API HSA_DEPRECATED hsa_signal_exchange_release(
1495 hsa_signal_t signal,
1496 hsa_signal_value_t value);
1497
1498 /**
1499 * @brief Atomically set the value of a signal if the observed value is equal to
1500 * the expected value. The observed value is returned regardless of whether the
1501 * replacement was done.
1502 *
1503 * @details If the value of the signal is changed, all the agents waiting
1504 * on @p signal for which @p value satisfies their wait condition are awakened.
1505 *
1506 * @param[in] signal Signal. If @p signal is a queue
1507 * doorbell signal, the behavior is undefined.
1508 *
1509 * @param[in] expected Value to compare with.
1510 *
1511 * @param[in] value New value.
1512 *
1513 * @return Observed value of the signal.
1514 *
1515 */
1516 hsa_signal_value_t HSA_API hsa_signal_cas_scacq_screl(
1517 hsa_signal_t signal,
1518 hsa_signal_value_t expected,
1519 hsa_signal_value_t value);
1520
1521
1522 /**
1523 * @deprecated Renamed as ::hsa_signal_cas_scacq_screl.
1524 *
1525 * @copydoc hsa_signal_cas_scacq_screl
1526 */
1527 hsa_signal_value_t HSA_API HSA_DEPRECATED hsa_signal_cas_acq_rel(
1528 hsa_signal_t signal,
1529 hsa_signal_value_t expected,
1530 hsa_signal_value_t value);
1531
1532 /**
1533 * @copydoc hsa_signal_cas_scacq_screl
1534 */
1535 hsa_signal_value_t HSA_API hsa_signal_cas_scacquire(
1536 hsa_signal_t signal,
1537 hsa_signal_value_t expected,
1538 hsa_signal_value_t value);
1539
1540 /**
1541 * @deprecated Renamed as ::hsa_signal_cas_scacquire.
1542 *
1543 * @copydoc hsa_signal_cas_scacquire
1544 */
1545 hsa_signal_value_t HSA_API HSA_DEPRECATED hsa_signal_cas_acquire(
1546 hsa_signal_t signal,
1547 hsa_signal_value_t expected,
1548 hsa_signal_value_t value);
1549
1550 /**
1551 * @copydoc hsa_signal_cas_scacq_screl
1552 */
1553 hsa_signal_value_t HSA_API hsa_signal_cas_relaxed(
1554 hsa_signal_t signal,
1555 hsa_signal_value_t expected,
1556 hsa_signal_value_t value);
1557
1558 /**
1559 * @copydoc hsa_signal_cas_scacq_screl
1560 */
1561 hsa_signal_value_t HSA_API hsa_signal_cas_screlease(
1562 hsa_signal_t signal,
1563 hsa_signal_value_t expected,
1564 hsa_signal_value_t value);
1565
1566 /**
1567 * @deprecated Renamed as ::hsa_signal_cas_screlease.
1568 *
1569 * @copydoc hsa_signal_cas_screlease
1570 */
1571 hsa_signal_value_t HSA_API HSA_DEPRECATED hsa_signal_cas_release(
1572 hsa_signal_t signal,
1573 hsa_signal_value_t expected,
1574 hsa_signal_value_t value);
1575
1576 /**
1577 * @brief Atomically increment the value of a signal by a given amount.
1578 *
1579 * @details If the value of the signal is changed, all the agents waiting on
1580 * @p signal for which @p value satisfies their wait condition are awakened.
1581 *
1582 * @param[in] signal Signal. If @p signal is a queue doorbell signal, the
1583 * behavior is undefined.
1584 *
1585 * @param[in] value Value to add to the value of the signal.
1586 *
1587 */
1588 void HSA_API hsa_signal_add_scacq_screl(
1589 hsa_signal_t signal,
1590 hsa_signal_value_t value);
1591
1592 /**
1593 * @deprecated Renamed as ::hsa_signal_add_scacq_screl.
1594 *
1595 * @copydoc hsa_signal_add_scacq_screl
1596 */
1597 void HSA_API HSA_DEPRECATED hsa_signal_add_acq_rel(
1598 hsa_signal_t signal,
1599 hsa_signal_value_t value);
1600
1601 /**
1602 * @copydoc hsa_signal_add_scacq_screl
1603 */
1604 void HSA_API hsa_signal_add_scacquire(
1605 hsa_signal_t signal,
1606 hsa_signal_value_t value);
1607
1608 /**
1609 * @deprecated Renamed as ::hsa_signal_add_scacquire.
1610 *
1611 * @copydoc hsa_signal_add_scacquire
1612 */
1613 void HSA_API HSA_DEPRECATED hsa_signal_add_acquire(
1614 hsa_signal_t signal,
1615 hsa_signal_value_t value);
1616
1617 /**
1618 * @copydoc hsa_signal_add_scacq_screl
1619 */
1620 void HSA_API hsa_signal_add_relaxed(
1621 hsa_signal_t signal,
1622 hsa_signal_value_t value);
1623
1624 /**
1625 * @copydoc hsa_signal_add_scacq_screl
1626 */
1627 void HSA_API hsa_signal_add_screlease(
1628 hsa_signal_t signal,
1629 hsa_signal_value_t value);
1630
1631
1632 /**
1633 * @deprecated Renamed as ::hsa_signal_add_screlease.
1634 *
1635 * @copydoc hsa_signal_add_screlease
1636 */
1637 void HSA_API HSA_DEPRECATED hsa_signal_add_release(
1638 hsa_signal_t signal,
1639 hsa_signal_value_t value);
1640
1641 /**
1642 * @brief Atomically decrement the value of a signal by a given amount.
1643 *
1644 * @details If the value of the signal is changed, all the agents waiting on
1645 * @p signal for which @p value satisfies their wait condition are awakened.
1646 *
1647 * @param[in] signal Signal. If @p signal is a queue doorbell signal, the
1648 * behavior is undefined.
1649 *
1650 * @param[in] value Value to subtract from the value of the signal.
1651 *
1652 */
1653 void HSA_API hsa_signal_subtract_scacq_screl(
1654 hsa_signal_t signal,
1655 hsa_signal_value_t value);
1656
1657
1658 /**
1659 * @deprecated Renamed as ::hsa_signal_subtract_scacq_screl.
1660 *
1661 * @copydoc hsa_signal_subtract_scacq_screl
1662 */
1663 void HSA_API HSA_DEPRECATED hsa_signal_subtract_acq_rel(
1664 hsa_signal_t signal,
1665 hsa_signal_value_t value);
1666
1667 /**
1668 * @copydoc hsa_signal_subtract_scacq_screl
1669 */
1670 void HSA_API hsa_signal_subtract_scacquire(
1671 hsa_signal_t signal,
1672 hsa_signal_value_t value);
1673
1674 /**
1675 * @deprecated Renamed as ::hsa_signal_subtract_scacquire.
1676 *
1677 * @copydoc hsa_signal_subtract_scacquire
1678 */
1679 void HSA_API HSA_DEPRECATED hsa_signal_subtract_acquire(
1680 hsa_signal_t signal,
1681 hsa_signal_value_t value);
1682
1683 /**
1684 * @copydoc hsa_signal_subtract_scacq_screl
1685 */
1686 void HSA_API hsa_signal_subtract_relaxed(
1687 hsa_signal_t signal,
1688 hsa_signal_value_t value);
1689
1690 /**
1691 * @copydoc hsa_signal_subtract_scacq_screl
1692 */
1693 void HSA_API hsa_signal_subtract_screlease(
1694 hsa_signal_t signal,
1695 hsa_signal_value_t value);
1696
1697
1698 /**
1699 * @deprecated Renamed as ::hsa_signal_subtract_screlease.
1700 *
1701 * @copydoc hsa_signal_subtract_screlease
1702 */
1703 void HSA_API HSA_DEPRECATED hsa_signal_subtract_release(
1704 hsa_signal_t signal,
1705 hsa_signal_value_t value);
1706
1707 /**
1708 * @brief Atomically perform a bitwise AND operation between the value of a
1709 * signal and a given value.
1710 *
1711 * @details If the value of the signal is changed, all the agents waiting on
1712 * @p signal for which @p value satisfies their wait condition are awakened.
1713 *
1714 * @param[in] signal Signal. If @p signal is a queue doorbell signal, the
1715 * behavior is undefined.
1716 *
1717 * @param[in] value Value to AND with the value of the signal.
1718 *
1719 */
1720 void HSA_API hsa_signal_and_scacq_screl(
1721 hsa_signal_t signal,
1722 hsa_signal_value_t value);
1723
1724 /**
1725 * @deprecated Renamed as ::hsa_signal_and_scacq_screl.
1726 *
1727 * @copydoc hsa_signal_and_scacq_screl
1728 */
1729 void HSA_API HSA_DEPRECATED hsa_signal_and_acq_rel(
1730 hsa_signal_t signal,
1731 hsa_signal_value_t value);
1732
1733 /**
1734 * @copydoc hsa_signal_and_scacq_screl
1735 */
1736 void HSA_API hsa_signal_and_scacquire(
1737 hsa_signal_t signal,
1738 hsa_signal_value_t value);
1739
1740 /**
1741 * @deprecated Renamed as ::hsa_signal_and_scacquire.
1742 *
1743 * @copydoc hsa_signal_and_scacquire
1744 */
1745 void HSA_API HSA_DEPRECATED hsa_signal_and_acquire(
1746 hsa_signal_t signal,
1747 hsa_signal_value_t value);
1748
1749 /**
1750 * @copydoc hsa_signal_and_scacq_screl
1751 */
1752 void HSA_API hsa_signal_and_relaxed(
1753 hsa_signal_t signal,
1754 hsa_signal_value_t value);
1755
1756 /**
1757 * @copydoc hsa_signal_and_scacq_screl
1758 */
1759 void HSA_API hsa_signal_and_screlease(
1760 hsa_signal_t signal,
1761 hsa_signal_value_t value);
1762
1763
1764 /**
1765 * @deprecated Renamed as ::hsa_signal_and_screlease.
1766 *
1767 * @copydoc hsa_signal_and_screlease
1768 */
1769 void HSA_API HSA_DEPRECATED hsa_signal_and_release(
1770 hsa_signal_t signal,
1771 hsa_signal_value_t value);
1772
1773 /**
1774 * @brief Atomically perform a bitwise OR operation between the value of a
1775 * signal and a given value.
1776 *
1777 * @details If the value of the signal is changed, all the agents waiting on
1778 * @p signal for which @p value satisfies their wait condition are awakened.
1779 *
1780 * @param[in] signal Signal. If @p signal is a queue doorbell signal, the
1781 * behavior is undefined.
1782 *
1783 * @param[in] value Value to OR with the value of the signal.
1784 */
1785 void HSA_API hsa_signal_or_scacq_screl(
1786 hsa_signal_t signal,
1787 hsa_signal_value_t value);
1788
1789
1790 /**
1791 * @deprecated Renamed as ::hsa_signal_or_scacq_screl.
1792 *
1793 * @copydoc hsa_signal_or_scacq_screl
1794 */
1795 void HSA_API HSA_DEPRECATED hsa_signal_or_acq_rel(
1796 hsa_signal_t signal,
1797 hsa_signal_value_t value);
1798
1799 /**
1800 * @copydoc hsa_signal_or_scacq_screl
1801 */
1802 void HSA_API hsa_signal_or_scacquire(
1803 hsa_signal_t signal,
1804 hsa_signal_value_t value);
1805
1806 /**
1807 * @deprecated Renamed as ::hsa_signal_or_scacquire.
1808 *
1809 * @copydoc hsa_signal_or_scacquire
1810 */
1811 void HSA_API HSA_DEPRECATED hsa_signal_or_acquire(
1812 hsa_signal_t signal,
1813 hsa_signal_value_t value);
1814
1815 /**
1816 * @copydoc hsa_signal_or_scacq_screl
1817 */
1818 void HSA_API hsa_signal_or_relaxed(
1819 hsa_signal_t signal,
1820 hsa_signal_value_t value);
1821
1822 /**
1823 * @copydoc hsa_signal_or_scacq_screl
1824 */
1825 void HSA_API hsa_signal_or_screlease(
1826 hsa_signal_t signal,
1827 hsa_signal_value_t value);
1828
1829 /**
1830 * @deprecated Renamed as ::hsa_signal_or_screlease.
1831 *
1832 * @copydoc hsa_signal_or_screlease
1833 */
1834 void HSA_API HSA_DEPRECATED hsa_signal_or_release(
1835 hsa_signal_t signal,
1836 hsa_signal_value_t value);
1837
1838 /**
1839 * @brief Atomically perform a bitwise XOR operation between the value of a
1840 * signal and a given value.
1841 *
1842 * @details If the value of the signal is changed, all the agents waiting on
1843 * @p signal for which @p value satisfies their wait condition are awakened.
1844 *
1845 * @param[in] signal Signal. If @p signal is a queue doorbell signal, the
1846 * behavior is undefined.
1847 *
1848 * @param[in] value Value to XOR with the value of the signal.
1849 *
1850 */
1851 void HSA_API hsa_signal_xor_scacq_screl(
1852 hsa_signal_t signal,
1853 hsa_signal_value_t value);
1854
1855
1856 /**
1857 * @deprecated Renamed as ::hsa_signal_xor_scacq_screl.
1858 *
1859 * @copydoc hsa_signal_xor_scacq_screl
1860 */
1861 void HSA_API HSA_DEPRECATED hsa_signal_xor_acq_rel(
1862 hsa_signal_t signal,
1863 hsa_signal_value_t value);
1864
1865 /**
1866 * @copydoc hsa_signal_xor_scacq_screl
1867 */
1868 void HSA_API hsa_signal_xor_scacquire(
1869 hsa_signal_t signal,
1870 hsa_signal_value_t value);
1871
1872 /**
1873 * @deprecated Renamed as ::hsa_signal_xor_scacquire.
1874 *
1875 * @copydoc hsa_signal_xor_scacquire
1876 */
1877 void HSA_API HSA_DEPRECATED hsa_signal_xor_acquire(
1878 hsa_signal_t signal,
1879 hsa_signal_value_t value);
1880
1881 /**
1882 * @copydoc hsa_signal_xor_scacq_screl
1883 */
1884 void HSA_API hsa_signal_xor_relaxed(
1885 hsa_signal_t signal,
1886 hsa_signal_value_t value);
1887
1888 /**
1889 * @copydoc hsa_signal_xor_scacq_screl
1890 */
1891 void HSA_API hsa_signal_xor_screlease(
1892 hsa_signal_t signal,
1893 hsa_signal_value_t value);
1894
1895 /**
1896 * @deprecated Renamed as ::hsa_signal_xor_screlease.
1897 *
1898 * @copydoc hsa_signal_xor_screlease
1899 */
1900 void HSA_API HSA_DEPRECATED hsa_signal_xor_release(
1901 hsa_signal_t signal,
1902 hsa_signal_value_t value);
1903
1904 /**
1905 * @brief Wait condition operator.
1906 */
1907 typedef enum {
1908 /**
1909 * The two operands are equal.
1910 */
1911 HSA_SIGNAL_CONDITION_EQ = 0,
1912 /**
1913 * The two operands are not equal.
1914 */
1915 HSA_SIGNAL_CONDITION_NE = 1,
1916 /**
1917 * The first operand is less than the second operand.
1918 */
1919 HSA_SIGNAL_CONDITION_LT = 2,
1920 /**
1921 * The first operand is greater than or equal to the second operand.
1922 */
1923 HSA_SIGNAL_CONDITION_GTE = 3
1924 } hsa_signal_condition_t;
1925
1926 /**
1927 * @brief State of the application thread during a signal wait.
1928 */
1929 typedef enum {
1930 /**
1931 * The application thread may be rescheduled while waiting on the signal.
1932 */
1933 HSA_WAIT_STATE_BLOCKED = 0,
1934 /**
1935 * The application thread stays active while waiting on a signal.
1936 */
1937 HSA_WAIT_STATE_ACTIVE = 1
1938 } hsa_wait_state_t;
1939
1940
1941 /**
1942 * @brief Wait until a signal value satisfies a specified condition, or a
1943 * certain amount of time has elapsed.
1944 *
1945 * @details A wait operation can spuriously resume at any time sooner than the
1946 * timeout (for example, due to system or other external factors) even when the
1947 * condition has not been met.
1948 *
1949 * The function is guaranteed to return if the signal value satisfies the
1950 * condition at some point in time during the wait, but the value returned to
1951 * the application might not satisfy the condition. The application must ensure
1952 * that signals are used in such way that wait wakeup conditions are not
1953 * invalidated before dependent threads have woken up.
1954 *
1955 * When the wait operation internally loads the value of the passed signal, it
1956 * uses the memory order indicated in the function name.
1957 *
1958 * @param[in] signal Signal.
1959 *
1960 * @param[in] condition Condition used to compare the signal value with @p
1961 * compare_value.
1962 *
1963 * @param[in] compare_value Value to compare with.
1964 *
1965 * @param[in] timeout_hint Maximum duration of the wait. Specified in the same
1966 * unit as the system timestamp. The operation might block for a shorter or
1967 * longer time even if the condition is not met. A value of UINT64_MAX indicates
1968 * no maximum.
1969 *
1970 * @param[in] wait_state_hint Hint used by the application to indicate the
1971 * preferred waiting state. The actual waiting state is ultimately decided by
1972 * HSA runtime and may not match the provided hint. A value of
1973 * ::HSA_WAIT_STATE_ACTIVE may improve the latency of response to a signal
1974 * update by avoiding rescheduling overhead.
1975 *
1976 * @return Observed value of the signal, which might not satisfy the specified
1977 * condition.
1978 *
1979 */
1980 hsa_signal_value_t HSA_API hsa_signal_wait_scacquire(
1981 hsa_signal_t signal,
1982 hsa_signal_condition_t condition,
1983 hsa_signal_value_t compare_value,
1984 uint64_t timeout_hint,
1985 hsa_wait_state_t wait_state_hint);
1986
1987 /**
1988 * @copydoc hsa_signal_wait_scacquire
1989 */
1990 hsa_signal_value_t HSA_API hsa_signal_wait_relaxed(
1991 hsa_signal_t signal,
1992 hsa_signal_condition_t condition,
1993 hsa_signal_value_t compare_value,
1994 uint64_t timeout_hint,
1995 hsa_wait_state_t wait_state_hint);
1996
1997 /**
1998 * @deprecated Renamed as ::hsa_signal_wait_scacquire.
1999 *
2000 * @copydoc hsa_signal_wait_scacquire
2001 */
2002 hsa_signal_value_t HSA_API HSA_DEPRECATED hsa_signal_wait_acquire(
2003 hsa_signal_t signal,
2004 hsa_signal_condition_t condition,
2005 hsa_signal_value_t compare_value,
2006 uint64_t timeout_hint,
2007 hsa_wait_state_t wait_state_hint);
2008
2009 /**
2010 * @brief Group of signals.
2011 */
2012 typedef struct hsa_signal_group_s {
2013 /**
2014 * Opaque handle. Two handles reference the same object of the enclosing type
2015 * if and only if they are equal.
2016 */
2017 uint64_t handle;
2018 } hsa_signal_group_t;
2019
2020 /**
2021 * @brief Create a signal group.
2022 *
2023 * @param[in] num_signals Number of elements in @p signals. Must not be 0.
2024 *
2025 * @param[in] signals List of signals in the group. The list must not contain
2026 * any repeated elements. Must not be NULL.
2027 *
2028 * @param[in] num_consumers Number of elements in @p consumers. Must not be 0.
2029 *
2030 * @param[in] consumers List of agents that might consume (wait on) the signal
2031 * group. The list must not contain repeated elements, and must be a subset of
2032 * the set of agents that are allowed to wait on all the signals in the
2033 * group. If an agent not listed in @p consumers waits on the returned group,
2034 * the behavior is undefined. The memory associated with @p consumers can be
2035 * reused or freed after the function returns. Must not be NULL.
2036 *
2037 * @param[out] signal_group Pointer to newly created signal group. Must not be
2038 * NULL.
2039 *
2040 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
2041 *
2042 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
2043 * initialized.
2044 *
2045 * @retval ::HSA_STATUS_ERROR_OUT_OF_RESOURCES The HSA runtime failed to allocate
2046 * the required resources.
2047 *
2048 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p num_signals is 0, @p signals
2049 * is NULL, @p num_consumers is 0, @p consumers is NULL, or @p signal_group is
2050 * NULL.
2051 */
2052 hsa_status_t HSA_API hsa_signal_group_create(
2053 uint32_t num_signals,
2054 const hsa_signal_t *signals,
2055 uint32_t num_consumers,
2056 const hsa_agent_t *consumers,
2057 hsa_signal_group_t *signal_group);
2058
2059 /**
2060 * @brief Destroy a signal group previous created by ::hsa_signal_group_create.
2061 *
2062 * @param[in] signal_group Signal group.
2063 *
2064 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
2065 *
2066 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
2067 * initialized.
2068 *
2069 * @retval ::HSA_STATUS_ERROR_INVALID_SIGNAL_GROUP @p signal_group is invalid.
2070 */
2071 hsa_status_t HSA_API hsa_signal_group_destroy(
2072 hsa_signal_group_t signal_group);
2073
2074 /**
2075 * @brief Wait until the value of at least one of the signals in a signal group
2076 * satisfies its associated condition.
2077 *
2078 * @details The function is guaranteed to return if the value of at least one of
2079 * the signals in the group satisfies its associated condition at some point in
2080 * time during the wait, but the signal value returned to the application may no
2081 * longer satisfy the condition. The application must ensure that signals in the
2082 * group are used in such way that wait wakeup conditions are not invalidated
2083 * before dependent threads have woken up.
2084 *
2085 * When this operation internally loads the value of the passed signal, it uses
2086 * the memory order indicated in the function name.
2087 *
2088 * @param[in] signal_group Signal group.
2089 *
2090 * @param[in] conditions List of conditions. Each condition, and the value at
2091 * the same index in @p compare_values, is used to compare the value of the
2092 * signal at that index in @p signal_group (the signal passed by the application
2093 * to ::hsa_signal_group_create at that particular index). The size of @p
2094 * conditions must not be smaller than the number of signals in @p signal_group;
2095 * any extra elements are ignored. Must not be NULL.
2096 *
2097 * @param[in] compare_values List of comparison values. The size of @p
2098 * compare_values must not be smaller than the number of signals in @p
2099 * signal_group; any extra elements are ignored. Must not be NULL.
2100 *
2101 * @param[in] wait_state_hint Hint used by the application to indicate the
2102 * preferred waiting state. The actual waiting state is decided by the HSA runtime
2103 * and may not match the provided hint. A value of ::HSA_WAIT_STATE_ACTIVE may
2104 * improve the latency of response to a signal update by avoiding rescheduling
2105 * overhead.
2106 *
2107 * @param[out] signal Signal in the group that satisfied the associated
2108 * condition. If several signals satisfied their condition, the function can
2109 * return any of those signals. Must not be NULL.
2110 *
2111 * @param[out] value Observed value for @p signal, which might no longer satisfy
2112 * the specified condition. Must not be NULL.
2113 *
2114 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
2115 *
2116 * @retval ::HSA_STATUS_ERROR_INVALID_SIGNAL_GROUP @p signal_group is invalid.
2117 *
2118 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p conditions is NULL, @p
2119 * compare_values is NULL, @p signal is NULL, or @p value is NULL.
2120 */
2121 hsa_status_t HSA_API hsa_signal_group_wait_any_scacquire(
2122 hsa_signal_group_t signal_group,
2123 const hsa_signal_condition_t *conditions,
2124 const hsa_signal_value_t *compare_values,
2125 hsa_wait_state_t wait_state_hint,
2126 hsa_signal_t *signal,
2127 hsa_signal_value_t *value);
2128
2129 /**
2130 * @copydoc hsa_signal_group_wait_any_scacquire
2131 */
2132 hsa_status_t HSA_API hsa_signal_group_wait_any_relaxed(
2133 hsa_signal_group_t signal_group,
2134 const hsa_signal_condition_t *conditions,
2135 const hsa_signal_value_t *compare_values,
2136 hsa_wait_state_t wait_state_hint,
2137 hsa_signal_t *signal,
2138 hsa_signal_value_t *value);
2139
2140 /** @} */
2141
2142 /** \defgroup memory Memory
2143 * @{
2144 */
2145
2146 /**
2147 * @brief A memory region represents a block of virtual memory with certain
2148 * properties. For example, the HSA runtime represents fine-grained memory in
2149 * the global segment using a region. A region might be associated with more
2150 * than one agent.
2151 */
2152 typedef struct hsa_region_s {
2153 /**
2154 * Opaque handle. Two handles reference the same object of the enclosing type
2155 * if and only if they are equal.
2156 */
2157 uint64_t handle;
2158 } hsa_region_t;
2159
2160 /** @} */
2161
2162
2163 /** \defgroup queue Queues
2164 * @{
2165 */
2166
2167 /**
2168 * @brief Queue type. Intended to be used for dynamic queue protocol
2169 * determination.
2170 */
2171 typedef enum {
2172 /**
2173 * Queue supports multiple producers. Use of multiproducer queue mechanics is
2174 * required.
2175 */
2176 HSA_QUEUE_TYPE_MULTI = 0,
2177 /**
2178 * Queue only supports a single producer. In some scenarios, the application
2179 * may want to limit the submission of AQL packets to a single agent. Queues
2180 * that support a single producer may be more efficient than queues supporting
2181 * multiple producers. Use of multiproducer queue mechanics is not supported.
2182 */
2183 HSA_QUEUE_TYPE_SINGLE = 1,
2184 /**
2185 * Queue supports multiple producers and cooperative dispatches. Cooperative
2186 * dispatches are able to use GWS synchronization. Queues of this type may be
2187 * limited in number. The runtime may return the same queue to serve multiple
2188 * ::hsa_queue_create calls when this type is given. Callers must inspect the
2189 * returned queue to discover queue size. Queues of this type are reference
2190 * counted and require a matching number of ::hsa_queue_destroy calls to
2191 * release. Use of multiproducer queue mechanics is required. See
2192 * ::HSA_AMD_AGENT_INFO_COOPERATIVE_QUEUES to query agent support for this
2193 * type.
2194 */
2195 HSA_QUEUE_TYPE_COOPERATIVE = 2
2196 } hsa_queue_type_t;
2197
2198 /**
2199 * @brief A fixed-size type used to represent ::hsa_queue_type_t constants.
2200 */
2201 typedef uint32_t hsa_queue_type32_t;
2202
2203 /**
2204 * @brief Queue features.
2205 */
2206 typedef enum {
2207 /**
2208 * Queue supports kernel dispatch packets.
2209 */
2210 HSA_QUEUE_FEATURE_KERNEL_DISPATCH = 1,
2211
2212 /**
2213 * Queue supports agent dispatch packets.
2214 */
2215 HSA_QUEUE_FEATURE_AGENT_DISPATCH = 2
2216 } hsa_queue_feature_t;
2217
2218 /**
2219 * @brief User mode queue.
2220 *
2221 * @details The queue structure is read-only and allocated by the HSA runtime,
2222 * but agents can directly modify the contents of the buffer pointed by @a
2223 * base_address, or use HSA runtime APIs to access the doorbell signal.
2224 *
2225 */
2226 typedef struct hsa_queue_s {
2227 /**
2228 * Queue type.
2229 */
2230 hsa_queue_type32_t type;
2231
2232 /**
2233 * Queue features mask. This is a bit-field of ::hsa_queue_feature_t
2234 * values. Applications should ignore any unknown set bits.
2235 */
2236 uint32_t features;
2237
2238 #ifdef HSA_LARGE_MODEL
2239 void* base_address;
2240 #elif defined HSA_LITTLE_ENDIAN
2241 /**
2242 * Starting address of the HSA runtime-allocated buffer used to store the AQL
2243 * packets. Must be aligned to the size of an AQL packet.
2244 */
2245 void* base_address;
2246 /**
2247 * Reserved. Must be 0.
2248 */
2249 uint32_t reserved0;
2250 #else
2251 uint32_t reserved0;
2252 void* base_address;
2253 #endif
2254
2255 /**
2256 * Signal object used by the application to indicate the ID of a packet that
2257 * is ready to be processed. The HSA runtime manages the doorbell signal. If
2258 * the application tries to replace or destroy this signal, the behavior is
2259 * undefined.
2260 *
2261 * If @a type is ::HSA_QUEUE_TYPE_SINGLE, the doorbell signal value must be
2262 * updated in a monotonically increasing fashion. If @a type is
2263 * ::HSA_QUEUE_TYPE_MULTI, the doorbell signal value can be updated with any
2264 * value.
2265 */
2266 hsa_signal_t doorbell_signal;
2267
2268 /**
2269 * Maximum number of packets the queue can hold. Must be a power of 2.
2270 */
2271 uint32_t size;
2272 /**
2273 * Reserved. Must be 0.
2274 */
2275 uint32_t reserved1;
2276 /**
2277 * Queue identifier, which is unique over the lifetime of the application.
2278 */
2279 uint64_t id;
2280
2281 } hsa_queue_t;
2282
2283 /**
2284 * @brief Create a user mode queue.
2285 *
2286 * @details The HSA runtime creates the queue structure, the underlying packet
2287 * buffer, the completion signal, and the write and read indexes. The initial
2288 * value of the write and read indexes is 0. The type of every packet in the
2289 * buffer is initialized to ::HSA_PACKET_TYPE_INVALID.
2290 *
2291 * The application should only rely on the error code returned to determine if
2292 * the queue is valid.
2293 *
2294 * @param[in] agent Agent where to create the queue.
2295 *
2296 * @param[in] size Number of packets the queue is expected to
2297 * hold. Must be a power of 2 between 1 and the value of
2298 * ::HSA_AGENT_INFO_QUEUE_MAX_SIZE in @p agent. The size of the newly
2299 * created queue is the maximum of @p size and the value of
2300 * ::HSA_AGENT_INFO_QUEUE_MIN_SIZE in @p agent.
2301 *
2302 * @param[in] type Type of the queue, a bitwise OR of hsa_queue_type_t values.
2303 * If the value of ::HSA_AGENT_INFO_QUEUE_TYPE in @p agent is ::HSA_QUEUE_TYPE_SINGLE,
2304 * then @p type must also be ::HSA_QUEUE_TYPE_SINGLE.
2305 *
2306 * @param[in] callback Callback invoked by the HSA runtime for every
2307 * asynchronous event related to the newly created queue. May be NULL. The HSA
2308 * runtime passes three arguments to the callback: a code identifying the event
2309 * that triggered the invocation, a pointer to the queue where the event
2310 * originated, and the application data.
2311 *
2312 * @param[in] data Application data that is passed to @p callback on every
2313 * iteration. May be NULL.
2314 *
2315 * @param[in] private_segment_size Hint indicating the maximum
2316 * expected private segment usage per work-item, in bytes. There may
2317 * be performance degradation if the application places a kernel
2318 * dispatch packet in the queue and the corresponding private segment
2319 * usage exceeds @p private_segment_size. If the application does not
2320 * want to specify any particular value for this argument, @p
2321 * private_segment_size must be UINT32_MAX. If the queue does not
2322 * support kernel dispatch packets, this argument is ignored.
2323 *
2324 * @param[in] group_segment_size Hint indicating the maximum expected
2325 * group segment usage per work-group, in bytes. There may be
2326 * performance degradation if the application places a kernel dispatch
2327 * packet in the queue and the corresponding group segment usage
2328 * exceeds @p group_segment_size. If the application does not want to
2329 * specify any particular value for this argument, @p
2330 * group_segment_size must be UINT32_MAX. If the queue does not
2331 * support kernel dispatch packets, this argument is ignored.
2332 *
2333 * @param[out] queue Memory location where the HSA runtime stores a pointer to
2334 * the newly created queue.
2335 *
2336 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
2337 *
2338 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
2339 * initialized.
2340 *
2341 * @retval ::HSA_STATUS_ERROR_OUT_OF_RESOURCES The HSA runtime failed to allocate
2342 * the required resources.
2343 *
2344 * @retval ::HSA_STATUS_ERROR_INVALID_AGENT The agent is invalid.
2345 *
2346 * @retval ::HSA_STATUS_ERROR_INVALID_QUEUE_CREATION @p agent does not
2347 * support queues of the given type.
2348 *
2349 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p size is not a power of two,
2350 * @p size is 0, @p type is an invalid queue type, or @p queue is NULL.
2351 *
2352 */
2353 hsa_status_t HSA_API hsa_queue_create(
2354 hsa_agent_t agent,
2355 uint32_t size,
2356 hsa_queue_type32_t type,
2357 void (*callback)(hsa_status_t status, hsa_queue_t *source, void *data),
2358 void *data,
2359 uint32_t private_segment_size,
2360 uint32_t group_segment_size,
2361 hsa_queue_t **queue);
2362
2363 /**
2364 * @brief Create a queue for which the application or a kernel is responsible
2365 * for processing the AQL packets.
2366 *
2367 * @details The application can use this function to create queues where AQL
2368 * packets are not parsed by the packet processor associated with an agent,
2369 * but rather by a unit of execution running on that agent (for example, a
2370 * thread in the host application).
2371 *
2372 * The application is responsible for ensuring that all the producers and
2373 * consumers of the resulting queue can access the provided doorbell signal
2374 * and memory region. The application is also responsible for ensuring that the
2375 * unit of execution processing the queue packets supports the indicated
2376 * features (AQL packet types).
2377 *
2378 * When the queue is created, the HSA runtime allocates the packet buffer using
2379 * @p region, and the write and read indexes. The initial value of the write and
2380 * read indexes is 0, and the type of every packet in the buffer is initialized
2381 * to ::HSA_PACKET_TYPE_INVALID. The value of the @e size, @e type, @e features,
2382 * and @e doorbell_signal fields in the returned queue match the values passed
2383 * by the application.
2384 *
2385 * @param[in] region Memory region that the HSA runtime should use to allocate
2386 * the AQL packet buffer and any other queue metadata.
2387 *
2388 * @param[in] size Number of packets the queue is expected to hold. Must be a
2389 * power of 2 greater than 0.
2390 *
2391 * @param[in] type Queue type.
2392 *
2393 * @param[in] features Supported queue features. This is a bit-field of
2394 * ::hsa_queue_feature_t values.
2395 *
2396 * @param[in] doorbell_signal Doorbell signal that the HSA runtime must
2397 * associate with the returned queue. The signal handle must not be 0.
2398 *
2399 * @param[out] queue Memory location where the HSA runtime stores a pointer to
2400 * the newly created queue. The application should not rely on the value
2401 * returned for this argument but only in the status code to determine if the
2402 * queue is valid. Must not be NULL.
2403 *
2404 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
2405 *
2406 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
2407 * initialized.
2408 *
2409 * @retval ::HSA_STATUS_ERROR_OUT_OF_RESOURCES The HSA runtime failed to allocate
2410 * the required resources.
2411 *
2412 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p size is not a power of two, @p
2413 * size is 0, @p type is an invalid queue type, the doorbell signal handle is
2414 * 0, or @p queue is NULL.
2415 *
2416 */
2417 hsa_status_t HSA_API hsa_soft_queue_create(
2418 hsa_region_t region,
2419 uint32_t size,
2420 hsa_queue_type32_t type,
2421 uint32_t features,
2422 hsa_signal_t doorbell_signal,
2423 hsa_queue_t **queue);
2424
2425 /**
2426 * @brief Destroy a user mode queue.
2427 *
2428 * @details When a queue is destroyed, the state of the AQL packets that have
2429 * not been yet fully processed (their completion phase has not finished)
2430 * becomes undefined. It is the responsibility of the application to ensure that
2431 * all pending queue operations are finished if their results are required.
2432 *
2433 * The resources allocated by the HSA runtime during queue creation (queue
2434 * structure, ring buffer, doorbell signal) are released. The queue should not
2435 * be accessed after being destroyed.
2436 *
2437 * @param[in] queue Pointer to a queue created using ::hsa_queue_create.
2438 *
2439 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
2440 *
2441 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
2442 * initialized.
2443 *
2444 * @retval ::HSA_STATUS_ERROR_INVALID_QUEUE The queue is invalid.
2445 *
2446 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p queue is NULL.
2447 */
2448 hsa_status_t HSA_API hsa_queue_destroy(
2449 hsa_queue_t *queue);
2450
2451 /**
2452 * @brief Inactivate a queue.
2453 *
2454 * @details Inactivating the queue aborts any pending executions and prevent any
2455 * new packets from being processed. Any more packets written to the queue once
2456 * it is inactivated will be ignored by the packet processor.
2457 *
2458 * @param[in] queue Pointer to a queue.
2459 *
2460 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
2461 *
2462 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
2463 * initialized.
2464 *
2465 * @retval ::HSA_STATUS_ERROR_INVALID_QUEUE The queue is invalid.
2466 *
2467 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p queue is NULL.
2468 */
2469 hsa_status_t HSA_API hsa_queue_inactivate(
2470 hsa_queue_t *queue);
2471
2472 /**
2473 * @deprecated Renamed as ::hsa_queue_load_read_index_scacquire.
2474 *
2475 * @copydoc hsa_queue_load_read_index_scacquire
2476 */
2477 uint64_t HSA_API HSA_DEPRECATED hsa_queue_load_read_index_acquire(
2478 const hsa_queue_t *queue);
2479
2480 /**
2481 * @brief Atomically load the read index of a queue.
2482 *
2483 * @param[in] queue Pointer to a queue.
2484 *
2485 * @return Read index of the queue pointed by @p queue.
2486 */
2487 uint64_t HSA_API hsa_queue_load_read_index_scacquire(
2488 const hsa_queue_t *queue);
2489
2490 /**
2491 * @copydoc hsa_queue_load_read_index_scacquire
2492 */
2493 uint64_t HSA_API hsa_queue_load_read_index_relaxed(
2494 const hsa_queue_t *queue);
2495
2496 /**
2497 * @deprecated Renamed as ::hsa_queue_load_write_index_scacquire.
2498 *
2499 * @copydoc hsa_queue_load_write_index_scacquire
2500 */
2501 uint64_t HSA_API HSA_DEPRECATED hsa_queue_load_write_index_acquire(
2502 const hsa_queue_t *queue);
2503
2504 /**
2505 * @brief Atomically load the write index of a queue.
2506 *
2507 * @param[in] queue Pointer to a queue.
2508 *
2509 * @return Write index of the queue pointed by @p queue.
2510 */
2511 uint64_t HSA_API hsa_queue_load_write_index_scacquire(
2512 const hsa_queue_t *queue);
2513
2514 /**
2515 * @copydoc hsa_queue_load_write_index_scacquire
2516 */
2517 uint64_t HSA_API hsa_queue_load_write_index_relaxed(
2518 const hsa_queue_t *queue);
2519
2520 /**
2521 * @brief Atomically set the write index of a queue.
2522 *
2523 * @details It is recommended that the application uses this function to update
2524 * the write index when there is a single agent submitting work to the queue
2525 * (the queue type is ::HSA_QUEUE_TYPE_SINGLE).
2526 *
2527 * @param[in] queue Pointer to a queue.
2528 *
2529 * @param[in] value Value to assign to the write index.
2530 *
2531 */
2532 void HSA_API hsa_queue_store_write_index_relaxed(
2533 const hsa_queue_t *queue,
2534 uint64_t value);
2535
2536 /**
2537 * @deprecated Renamed as ::hsa_queue_store_write_index_screlease.
2538 *
2539 * @copydoc hsa_queue_store_write_index_screlease
2540 */
2541 void HSA_API HSA_DEPRECATED hsa_queue_store_write_index_release(
2542 const hsa_queue_t *queue,
2543 uint64_t value);
2544
2545 /**
2546 * @copydoc hsa_queue_store_write_index_relaxed
2547 */
2548 void HSA_API hsa_queue_store_write_index_screlease(
2549 const hsa_queue_t *queue,
2550 uint64_t value);
2551
2552 /**
2553 * @deprecated Renamed as ::hsa_queue_cas_write_index_scacq_screl.
2554 *
2555 * @copydoc hsa_queue_cas_write_index_scacq_screl
2556 */
2557 uint64_t HSA_API HSA_DEPRECATED hsa_queue_cas_write_index_acq_rel(
2558 const hsa_queue_t *queue,
2559 uint64_t expected,
2560 uint64_t value);
2561
2562 /**
2563 * @brief Atomically set the write index of a queue if the observed value is
2564 * equal to the expected value. The application can inspect the returned value
2565 * to determine if the replacement was done.
2566 *
2567 * @param[in] queue Pointer to a queue.
2568 *
2569 * @param[in] expected Expected value.
2570 *
2571 * @param[in] value Value to assign to the write index if @p expected matches
2572 * the observed write index. Must be greater than @p expected.
2573 *
2574 * @return Previous value of the write index.
2575 */
2576 uint64_t HSA_API hsa_queue_cas_write_index_scacq_screl(
2577 const hsa_queue_t *queue,
2578 uint64_t expected,
2579 uint64_t value);
2580
2581 /**
2582 * @deprecated Renamed as ::hsa_queue_cas_write_index_scacquire.
2583 *
2584 * @copydoc hsa_queue_cas_write_index_scacquire
2585 */
2586 uint64_t HSA_API HSA_DEPRECATED hsa_queue_cas_write_index_acquire(
2587 const hsa_queue_t *queue,
2588 uint64_t expected,
2589 uint64_t value);
2590
2591 /**
2592 * @copydoc hsa_queue_cas_write_index_scacq_screl
2593 */
2594 uint64_t HSA_API hsa_queue_cas_write_index_scacquire(
2595 const hsa_queue_t *queue,
2596 uint64_t expected,
2597 uint64_t value);
2598
2599 /**
2600 * @copydoc hsa_queue_cas_write_index_scacq_screl
2601 */
2602 uint64_t HSA_API hsa_queue_cas_write_index_relaxed(
2603 const hsa_queue_t *queue,
2604 uint64_t expected,
2605 uint64_t value);
2606
2607 /**
2608 * @deprecated Renamed as ::hsa_queue_cas_write_index_screlease.
2609 *
2610 * @copydoc hsa_queue_cas_write_index_screlease
2611 */
2612 uint64_t HSA_API HSA_DEPRECATED hsa_queue_cas_write_index_release(
2613 const hsa_queue_t *queue,
2614 uint64_t expected,
2615 uint64_t value);
2616
2617 /**
2618 * @copydoc hsa_queue_cas_write_index_scacq_screl
2619 */
2620 uint64_t HSA_API hsa_queue_cas_write_index_screlease(
2621 const hsa_queue_t *queue,
2622 uint64_t expected,
2623 uint64_t value);
2624
2625 /**
2626 * @deprecated Renamed as ::hsa_queue_add_write_index_scacq_screl.
2627 *
2628 * @copydoc hsa_queue_add_write_index_scacq_screl
2629 */
2630 uint64_t HSA_API HSA_DEPRECATED hsa_queue_add_write_index_acq_rel(
2631 const hsa_queue_t *queue,
2632 uint64_t value);
2633
2634 /**
2635 * @brief Atomically increment the write index of a queue by an offset.
2636 *
2637 * @param[in] queue Pointer to a queue.
2638 *
2639 * @param[in] value Value to add to the write index.
2640 *
2641 * @return Previous value of the write index.
2642 */
2643 uint64_t HSA_API hsa_queue_add_write_index_scacq_screl(
2644 const hsa_queue_t *queue,
2645 uint64_t value);
2646
2647 /**
2648 * @deprecated Renamed as ::hsa_queue_add_write_index_scacquire.
2649 *
2650 * @copydoc hsa_queue_add_write_index_scacquire
2651 */
2652 uint64_t HSA_API HSA_DEPRECATED hsa_queue_add_write_index_acquire(
2653 const hsa_queue_t *queue,
2654 uint64_t value);
2655
2656 /**
2657 * @copydoc hsa_queue_add_write_index_scacq_screl
2658 */
2659 uint64_t HSA_API hsa_queue_add_write_index_scacquire(
2660 const hsa_queue_t *queue,
2661 uint64_t value);
2662
2663 /**
2664 * @copydoc hsa_queue_add_write_index_scacq_screl
2665 */
2666 uint64_t HSA_API hsa_queue_add_write_index_relaxed(
2667 const hsa_queue_t *queue,
2668 uint64_t value);
2669
2670 /**
2671 * @deprecated Renamed as ::hsa_queue_add_write_index_screlease.
2672 *
2673 * @copydoc hsa_queue_add_write_index_screlease
2674 */
2675 uint64_t HSA_API HSA_DEPRECATED hsa_queue_add_write_index_release(
2676 const hsa_queue_t *queue,
2677 uint64_t value);
2678
2679 /**
2680 * @copydoc hsa_queue_add_write_index_scacq_screl
2681 */
2682 uint64_t HSA_API hsa_queue_add_write_index_screlease(
2683 const hsa_queue_t *queue,
2684 uint64_t value);
2685
2686 /**
2687 * @brief Atomically set the read index of a queue.
2688 *
2689 * @details Modifications of the read index are not allowed and result in
2690 * undefined behavior if the queue is associated with an agent for which
2691 * only the corresponding packet processor is permitted to update the read
2692 * index.
2693 *
2694 * @param[in] queue Pointer to a queue.
2695 *
2696 * @param[in] value Value to assign to the read index.
2697 *
2698 */
2699 void HSA_API hsa_queue_store_read_index_relaxed(
2700 const hsa_queue_t *queue,
2701 uint64_t value);
2702
2703 /**
2704 * @deprecated Renamed as ::hsa_queue_store_read_index_screlease.
2705 *
2706 * @copydoc hsa_queue_store_read_index_screlease
2707 */
2708 void HSA_API HSA_DEPRECATED hsa_queue_store_read_index_release(
2709 const hsa_queue_t *queue,
2710 uint64_t value);
2711
2712 /**
2713 * @copydoc hsa_queue_store_read_index_relaxed
2714 */
2715 void HSA_API hsa_queue_store_read_index_screlease(
2716 const hsa_queue_t *queue,
2717 uint64_t value);
2718 /** @} */
2719
2720
2721 /** \defgroup aql Architected Queuing Language
2722 * @{
2723 */
2724
2725 /**
2726 * @brief Packet type.
2727 */
2728 typedef enum {
2729 /**
2730 * Vendor-specific packet.
2731 */
2732 HSA_PACKET_TYPE_VENDOR_SPECIFIC = 0,
2733 /**
2734 * The packet has been processed in the past, but has not been reassigned to
2735 * the packet processor. A packet processor must not process a packet of this
2736 * type. All queues support this packet type.
2737 */
2738 HSA_PACKET_TYPE_INVALID = 1,
2739 /**
2740 * Packet used by agents for dispatching jobs to kernel agents. Not all
2741 * queues support packets of this type (see ::hsa_queue_feature_t).
2742 */
2743 HSA_PACKET_TYPE_KERNEL_DISPATCH = 2,
2744 /**
2745 * Packet used by agents to delay processing of subsequent packets, and to
2746 * express complex dependencies between multiple packets. All queues support
2747 * this packet type.
2748 */
2749 HSA_PACKET_TYPE_BARRIER_AND = 3,
2750 /**
2751 * Packet used by agents for dispatching jobs to agents. Not all
2752 * queues support packets of this type (see ::hsa_queue_feature_t).
2753 */
2754 HSA_PACKET_TYPE_AGENT_DISPATCH = 4,
2755 /**
2756 * Packet used by agents to delay processing of subsequent packets, and to
2757 * express complex dependencies between multiple packets. All queues support
2758 * this packet type.
2759 */
2760 HSA_PACKET_TYPE_BARRIER_OR = 5
2761 } hsa_packet_type_t;
2762
2763 /**
2764 * @brief Scope of the memory fence operation associated with a packet.
2765 */
2766 typedef enum {
2767 /**
2768 * No scope (no fence is applied). The packet relies on external fences to
2769 * ensure visibility of memory updates.
2770 */
2771 HSA_FENCE_SCOPE_NONE = 0,
2772 /**
2773 * The fence is applied with agent scope for the global segment.
2774 */
2775 HSA_FENCE_SCOPE_AGENT = 1,
2776 /**
2777 * The fence is applied across both agent and system scope for the global
2778 * segment.
2779 */
2780 HSA_FENCE_SCOPE_SYSTEM = 2
2781 } hsa_fence_scope_t;
2782
2783 /**
2784 * @brief Sub-fields of the @a header field that is present in any AQL
2785 * packet. The offset (with respect to the address of @a header) of a sub-field
2786 * is identical to its enumeration constant. The width of each sub-field is
2787 * determined by the corresponding value in ::hsa_packet_header_width_t. The
2788 * offset and the width are expressed in bits.
2789 */
2790 typedef enum {
2791 /**
2792 * Packet type. The value of this sub-field must be one of
2793 * ::hsa_packet_type_t. If the type is ::HSA_PACKET_TYPE_VENDOR_SPECIFIC, the
2794 * packet layout is vendor-specific.
2795 */
2796 HSA_PACKET_HEADER_TYPE = 0,
2797 /**
2798 * Barrier bit. If the barrier bit is set, the processing of the current
2799 * packet only launches when all preceding packets (within the same queue) are
2800 * complete.
2801 */
2802 HSA_PACKET_HEADER_BARRIER = 8,
2803 /**
2804 * Acquire fence scope. The value of this sub-field determines the scope and
2805 * type of the memory fence operation applied before the packet enters the
2806 * active phase. An acquire fence ensures that any subsequent global segment
2807 * or image loads by any unit of execution that belongs to a dispatch that has
2808 * not yet entered the active phase on any queue of the same kernel agent,
2809 * sees any data previously released at the scopes specified by the acquire
2810 * fence. The value of this sub-field must be one of ::hsa_fence_scope_t.
2811 */
2812 HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE = 9,
2813 /**
2814 * @deprecated Renamed as ::HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE.
2815 */
2816 HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE = 9,
2817 /**
2818 * Release fence scope, The value of this sub-field determines the scope and
2819 * type of the memory fence operation applied after kernel completion but
2820 * before the packet is completed. A release fence makes any global segment or
2821 * image data that was stored by any unit of execution that belonged to a
2822 * dispatch that has completed the active phase on any queue of the same
2823 * kernel agent visible in all the scopes specified by the release fence. The
2824 * value of this sub-field must be one of ::hsa_fence_scope_t.
2825 */
2826 HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE = 11,
2827 /**
2828 * @deprecated Renamed as ::HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE.
2829 */
2830 HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE = 11
2831 } hsa_packet_header_t;
2832
2833 /**
2834 * @brief Width (in bits) of the sub-fields in ::hsa_packet_header_t.
2835 */
2836 typedef enum {
2837 HSA_PACKET_HEADER_WIDTH_TYPE = 8,
2838 HSA_PACKET_HEADER_WIDTH_BARRIER = 1,
2839 HSA_PACKET_HEADER_WIDTH_SCACQUIRE_FENCE_SCOPE = 2,
2840 /**
2841 * @deprecated Use HSA_PACKET_HEADER_WIDTH_SCACQUIRE_FENCE_SCOPE.
2842 */
2843 HSA_PACKET_HEADER_WIDTH_ACQUIRE_FENCE_SCOPE = 2,
2844 HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE = 2,
2845 /**
2846 * @deprecated Use HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE.
2847 */
2848 HSA_PACKET_HEADER_WIDTH_RELEASE_FENCE_SCOPE = 2
2849 } hsa_packet_header_width_t;
2850
2851 /**
2852 * @brief Sub-fields of the kernel dispatch packet @a setup field. The offset
2853 * (with respect to the address of @a setup) of a sub-field is identical to its
2854 * enumeration constant. The width of each sub-field is determined by the
2855 * corresponding value in ::hsa_kernel_dispatch_packet_setup_width_t. The
2856 * offset and the width are expressed in bits.
2857 */
2858 typedef enum {
2859 /**
2860 * Number of dimensions of the grid. Valid values are 1, 2, or 3.
2861 *
2862 */
2863 HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS = 0
2864 } hsa_kernel_dispatch_packet_setup_t;
2865
2866 /**
2867 * @brief Width (in bits) of the sub-fields in
2868 * ::hsa_kernel_dispatch_packet_setup_t.
2869 */
2870 typedef enum {
2871 HSA_KERNEL_DISPATCH_PACKET_SETUP_WIDTH_DIMENSIONS = 2
2872 } hsa_kernel_dispatch_packet_setup_width_t;
2873
2874 /**
2875 * @brief AQL kernel dispatch packet
2876 */
2877 typedef struct hsa_kernel_dispatch_packet_s {
2878 /**
2879 * Packet header. Used to configure multiple packet parameters such as the
2880 * packet type. The parameters are described by ::hsa_packet_header_t.
2881 */
2882 uint16_t header;
2883
2884 /**
2885 * Dispatch setup parameters. Used to configure kernel dispatch parameters
2886 * such as the number of dimensions in the grid. The parameters are described
2887 * by ::hsa_kernel_dispatch_packet_setup_t.
2888 */
2889 uint16_t setup;
2890
2891 /**
2892 * X dimension of work-group, in work-items. Must be greater than 0.
2893 */
2894 uint16_t workgroup_size_x;
2895
2896 /**
2897 * Y dimension of work-group, in work-items. Must be greater than
2898 * 0. If the grid has 1 dimension, the only valid value is 1.
2899 */
2900 uint16_t workgroup_size_y;
2901
2902 /**
2903 * Z dimension of work-group, in work-items. Must be greater than
2904 * 0. If the grid has 1 or 2 dimensions, the only valid value is 1.
2905 */
2906 uint16_t workgroup_size_z;
2907
2908 /**
2909 * Reserved. Must be 0.
2910 */
2911 uint16_t reserved0;
2912
2913 /**
2914 * X dimension of grid, in work-items. Must be greater than 0. Must
2915 * not be smaller than @a workgroup_size_x.
2916 */
2917 uint32_t grid_size_x;
2918
2919 /**
2920 * Y dimension of grid, in work-items. Must be greater than 0. If the grid has
2921 * 1 dimension, the only valid value is 1. Must not be smaller than @a
2922 * workgroup_size_y.
2923 */
2924 uint32_t grid_size_y;
2925
2926 /**
2927 * Z dimension of grid, in work-items. Must be greater than 0. If the grid has
2928 * 1 or 2 dimensions, the only valid value is 1. Must not be smaller than @a
2929 * workgroup_size_z.
2930 */
2931 uint32_t grid_size_z;
2932
2933 /**
2934 * Size in bytes of private memory allocation request (per work-item).
2935 */
2936 uint32_t private_segment_size;
2937
2938 /**
2939 * Size in bytes of group memory allocation request (per work-group). Must not
2940 * be less than the sum of the group memory used by the kernel (and the
2941 * functions it calls directly or indirectly) and the dynamically allocated
2942 * group segment variables.
2943 */
2944 uint32_t group_segment_size;
2945
2946 /**
2947 * Opaque handle to a code object that includes an implementation-defined
2948 * executable code for the kernel.
2949 */
2950 uint64_t kernel_object;
2951
2952 #ifdef HSA_LARGE_MODEL
2953 void* kernarg_address;
2954 #elif defined HSA_LITTLE_ENDIAN
2955 /**
2956 * Pointer to a buffer containing the kernel arguments. May be NULL.
2957 *
2958 * The buffer must be allocated using ::hsa_memory_allocate, and must not be
2959 * modified once the kernel dispatch packet is enqueued until the dispatch has
2960 * completed execution.
2961 */
2962 void* kernarg_address;
2963 /**
2964 * Reserved. Must be 0.
2965 */
2966 uint32_t reserved1;
2967 #else
2968 uint32_t reserved1;
2969 void* kernarg_address;
2970 #endif
2971
2972 /**
2973 * Reserved. Must be 0.
2974 */
2975 uint64_t reserved2;
2976
2977 /**
2978 * Signal used to indicate completion of the job. The application can use the
2979 * special signal handle 0 to indicate that no signal is used.
2980 */
2981 hsa_signal_t completion_signal;
2982
2983 } hsa_kernel_dispatch_packet_t;
2984
2985 /**
2986 * @brief Agent dispatch packet.
2987 */
2988 typedef struct hsa_agent_dispatch_packet_s {
2989 /**
2990 * Packet header. Used to configure multiple packet parameters such as the
2991 * packet type. The parameters are described by ::hsa_packet_header_t.
2992 */
2993 uint16_t header;
2994
2995 /**
2996 * Application-defined function to be performed by the destination agent.
2997 */
2998 uint16_t type;
2999
3000 /**
3001 * Reserved. Must be 0.
3002 */
3003 uint32_t reserved0;
3004
3005 #ifdef HSA_LARGE_MODEL
3006 void* return_address;
3007 #elif defined HSA_LITTLE_ENDIAN
3008 /**
3009 * Address where to store the function return values, if any.
3010 */
3011 void* return_address;
3012 /**
3013 * Reserved. Must be 0.
3014 */
3015 uint32_t reserved1;
3016 #else
3017 uint32_t reserved1;
3018 void* return_address;
3019 #endif
3020
3021 /**
3022 * Function arguments.
3023 */
3024 uint64_t arg[4];
3025
3026 /**
3027 * Reserved. Must be 0.
3028 */
3029 uint64_t reserved2;
3030
3031 /**
3032 * Signal used to indicate completion of the job. The application can use the
3033 * special signal handle 0 to indicate that no signal is used.
3034 */
3035 hsa_signal_t completion_signal;
3036
3037 } hsa_agent_dispatch_packet_t;
3038
3039 /**
3040 * @brief Barrier-AND packet.
3041 */
3042 typedef struct hsa_barrier_and_packet_s {
3043 /**
3044 * Packet header. Used to configure multiple packet parameters such as the
3045 * packet type. The parameters are described by ::hsa_packet_header_t.
3046 */
3047 uint16_t header;
3048
3049 /**
3050 * Reserved. Must be 0.
3051 */
3052 uint16_t reserved0;
3053
3054 /**
3055 * Reserved. Must be 0.
3056 */
3057 uint32_t reserved1;
3058
3059 /**
3060 * Array of dependent signal objects. Signals with a handle value of 0 are
3061 * allowed and are interpreted by the packet processor as satisfied
3062 * dependencies.
3063 */
3064 hsa_signal_t dep_signal[5];
3065
3066 /**
3067 * Reserved. Must be 0.
3068 */
3069 uint64_t reserved2;
3070
3071 /**
3072 * Signal used to indicate completion of the job. The application can use the
3073 * special signal handle 0 to indicate that no signal is used.
3074 */
3075 hsa_signal_t completion_signal;
3076
3077 } hsa_barrier_and_packet_t;
3078
3079 /**
3080 * @brief Barrier-OR packet.
3081 */
3082 typedef struct hsa_barrier_or_packet_s {
3083 /**
3084 * Packet header. Used to configure multiple packet parameters such as the
3085 * packet type. The parameters are described by ::hsa_packet_header_t.
3086 */
3087 uint16_t header;
3088
3089 /**
3090 * Reserved. Must be 0.
3091 */
3092 uint16_t reserved0;
3093
3094 /**
3095 * Reserved. Must be 0.
3096 */
3097 uint32_t reserved1;
3098
3099 /**
3100 * Array of dependent signal objects. Signals with a handle value of 0 are
3101 * allowed and are interpreted by the packet processor as dependencies not
3102 * satisfied.
3103 */
3104 hsa_signal_t dep_signal[5];
3105
3106 /**
3107 * Reserved. Must be 0.
3108 */
3109 uint64_t reserved2;
3110
3111 /**
3112 * Signal used to indicate completion of the job. The application can use the
3113 * special signal handle 0 to indicate that no signal is used.
3114 */
3115 hsa_signal_t completion_signal;
3116
3117 } hsa_barrier_or_packet_t;
3118
3119 /** @} */
3120
3121 /** \addtogroup memory Memory
3122 * @{
3123 */
3124
3125 /**
3126 * @brief Memory segments associated with a region.
3127 */
3128 typedef enum {
3129 /**
3130 * Global segment. Used to hold data that is shared by all agents.
3131 */
3132 HSA_REGION_SEGMENT_GLOBAL = 0,
3133 /**
3134 * Read-only segment. Used to hold data that remains constant during the
3135 * execution of a kernel.
3136 */
3137 HSA_REGION_SEGMENT_READONLY = 1,
3138 /**
3139 * Private segment. Used to hold data that is local to a single work-item.
3140 */
3141 HSA_REGION_SEGMENT_PRIVATE = 2,
3142 /**
3143 * Group segment. Used to hold data that is shared by the work-items of a
3144 * work-group.
3145 */
3146 HSA_REGION_SEGMENT_GROUP = 3,
3147 /**
3148 * Kernarg segment. Used to store kernel arguments.
3149 */
3150 HSA_REGION_SEGMENT_KERNARG = 4
3151 } hsa_region_segment_t;
3152
3153 /**
3154 * @brief Global region flags.
3155 */
3156 typedef enum {
3157 /**
3158 * The application can use memory in the region to store kernel arguments, and
3159 * provide the values for the kernarg segment of a kernel dispatch. If this
3160 * flag is set, then ::HSA_REGION_GLOBAL_FLAG_FINE_GRAINED must be set.
3161 */
3162 HSA_REGION_GLOBAL_FLAG_KERNARG = 1,
3163 /**
3164 * Updates to memory in this region are immediately visible to all the
3165 * agents under the terms of the HSA memory model. If this
3166 * flag is set, then ::HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED must not be set.
3167 */
3168 HSA_REGION_GLOBAL_FLAG_FINE_GRAINED = 2,
3169 /**
3170 * Updates to memory in this region can be performed by a single agent at
3171 * a time. If a different agent in the system is allowed to access the
3172 * region, the application must explicitely invoke ::hsa_memory_assign_agent
3173 * in order to transfer ownership to that agent for a particular buffer.
3174 */
3175 HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED = 4
3176 } hsa_region_global_flag_t;
3177
3178 /**
3179 * @brief Attributes of a memory region.
3180 */
3181 typedef enum {
3182 /**
3183 * Segment where memory in the region can be used. The type of this
3184 * attribute is ::hsa_region_segment_t.
3185 */
3186 HSA_REGION_INFO_SEGMENT = 0,
3187 /**
3188 * Flag mask. The value of this attribute is undefined if the value of
3189 * ::HSA_REGION_INFO_SEGMENT is not ::HSA_REGION_SEGMENT_GLOBAL. The type of
3190 * this attribute is uint32_t, a bit-field of ::hsa_region_global_flag_t
3191 * values.
3192 */
3193 HSA_REGION_INFO_GLOBAL_FLAGS = 1,
3194 /**
3195 * Size of this region, in bytes. The type of this attribute is size_t.
3196 */
3197 HSA_REGION_INFO_SIZE = 2,
3198 /**
3199 * Maximum allocation size in this region, in bytes. Must not exceed the value
3200 * of ::HSA_REGION_INFO_SIZE. The type of this attribute is size_t.
3201 *
3202 * If the region is in the global or readonly segments, this is the maximum
3203 * size that the application can pass to ::hsa_memory_allocate.
3204 *
3205 * If the region is in the group segment, this is the maximum size (per
3206 * work-group) that can be requested for a given kernel dispatch. If the
3207 * region is in the private segment, this is the maximum size (per work-item)
3208 * that can be requested for a specific kernel dispatch, and must be at least
3209 * 256 bytes.
3210 */
3211 HSA_REGION_INFO_ALLOC_MAX_SIZE = 4,
3212 /**
3213 * Maximum size (per work-group) of private memory that can be requested for a
3214 * specific kernel dispatch. Must be at least 65536 bytes. The type of this
3215 * attribute is uint32_t. The value of this attribute is undefined if the
3216 * region is not in the private segment.
3217 */
3218 HSA_REGION_INFO_ALLOC_MAX_PRIVATE_WORKGROUP_SIZE = 8,
3219 /**
3220 * Indicates whether memory in this region can be allocated using
3221 * ::hsa_memory_allocate. The type of this attribute is bool.
3222 *
3223 * The value of this flag is always false for regions in the group and private
3224 * segments.
3225 */
3226 HSA_REGION_INFO_RUNTIME_ALLOC_ALLOWED = 5,
3227 /**
3228 * Allocation granularity of buffers allocated by ::hsa_memory_allocate in
3229 * this region. The size of a buffer allocated in this region is a multiple of
3230 * the value of this attribute. The value of this attribute is only defined if
3231 * ::HSA_REGION_INFO_RUNTIME_ALLOC_ALLOWED is true for this region. The type
3232 * of this attribute is size_t.
3233 */
3234 HSA_REGION_INFO_RUNTIME_ALLOC_GRANULE = 6,
3235 /**
3236 * Alignment of buffers allocated by ::hsa_memory_allocate in this region. The
3237 * value of this attribute is only defined if
3238 * ::HSA_REGION_INFO_RUNTIME_ALLOC_ALLOWED is true for this region, and must be
3239 * a power of 2. The type of this attribute is size_t.
3240 */
3241 HSA_REGION_INFO_RUNTIME_ALLOC_ALIGNMENT = 7
3242 } hsa_region_info_t;
3243
3244 /**
3245 * @brief Get the current value of an attribute of a region.
3246 *
3247 * @param[in] region A valid region.
3248 *
3249 * @param[in] attribute Attribute to query.
3250 *
3251 * @param[out] value Pointer to a application-allocated buffer where to store
3252 * the value of the attribute. If the buffer passed by the application is not
3253 * large enough to hold the value of @p attribute, the behavior is undefined.
3254 *
3255 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
3256 *
3257 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
3258 * initialized.
3259 *
3260 * @retval ::HSA_STATUS_ERROR_INVALID_REGION The region is invalid.
3261 *
3262 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p attribute is an invalid
3263 * region attribute, or @p value is NULL.
3264 */
3265 hsa_status_t HSA_API hsa_region_get_info(
3266 hsa_region_t region,
3267 hsa_region_info_t attribute,
3268 void* value);
3269
3270 /**
3271 * @brief Iterate over the memory regions associated with a given agent, and
3272 * invoke an application-defined callback on every iteration.
3273 *
3274 * @param[in] agent A valid agent.
3275 *
3276 * @param[in] callback Callback to be invoked once per region that is
3277 * accessible from the agent. The HSA runtime passes two arguments to the
3278 * callback, the region and the application data. If @p callback returns a
3279 * status other than ::HSA_STATUS_SUCCESS for a particular iteration, the
3280 * traversal stops and ::hsa_agent_iterate_regions returns that status value.
3281 *
3282 * @param[in] data Application data that is passed to @p callback on every
3283 * iteration. May be NULL.
3284 *
3285 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
3286 *
3287 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
3288 * initialized.
3289 *
3290 * @retval ::HSA_STATUS_ERROR_INVALID_AGENT The agent is invalid.
3291 *
3292 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p callback is NULL.
3293 */
3294 hsa_status_t HSA_API hsa_agent_iterate_regions(
3295 hsa_agent_t agent,
3296 hsa_status_t (*callback)(hsa_region_t region, void* data),
3297 void* data);
3298
3299 /**
3300 * @brief Allocate a block of memory in a given region.
3301 *
3302 * @param[in] region Region where to allocate memory from. The region must have
3303 * the ::HSA_REGION_INFO_RUNTIME_ALLOC_ALLOWED flag set.
3304 *
3305 * @param[in] size Allocation size, in bytes. Must not be zero. This value is
3306 * rounded up to the nearest multiple of ::HSA_REGION_INFO_RUNTIME_ALLOC_GRANULE
3307 * in @p region.
3308 *
3309 * @param[out] ptr Pointer to the location where to store the base address of
3310 * the allocated block. The returned base address is aligned to the value of
3311 * ::HSA_REGION_INFO_RUNTIME_ALLOC_ALIGNMENT in @p region. If the allocation
3312 * fails, the returned value is undefined.
3313 *
3314 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
3315 *
3316 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
3317 * initialized.
3318 *
3319 * @retval ::HSA_STATUS_ERROR_OUT_OF_RESOURCES The HSA runtime failed to allocate
3320 * the required resources.
3321 *
3322 * @retval ::HSA_STATUS_ERROR_INVALID_REGION The region is invalid.
3323 *
3324 * @retval ::HSA_STATUS_ERROR_INVALID_ALLOCATION The host is not allowed to
3325 * allocate memory in @p region, or @p size is greater than the value of
3326 * HSA_REGION_INFO_ALLOC_MAX_SIZE in @p region.
3327 *
3328 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p ptr is NULL, or @p size is 0.
3329 */
3330 hsa_status_t HSA_API hsa_memory_allocate(hsa_region_t region,
3331 size_t size,
3332 void** ptr);
3333
3334 /**
3335 * @brief Deallocate a block of memory previously allocated using
3336 * ::hsa_memory_allocate.
3337 *
3338 * @param[in] ptr Pointer to a memory block. If @p ptr does not match a value
3339 * previously returned by ::hsa_memory_allocate, the behavior is undefined.
3340 *
3341 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
3342 *
3343 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
3344 * initialized.
3345 */
3346 hsa_status_t HSA_API hsa_memory_free(void* ptr);
3347
3348 /**
3349 * @brief Copy a block of memory from the location pointed to by @p src to the
3350 * memory block pointed to by @p dst.
3351 *
3352 * @param[out] dst Buffer where the content is to be copied. If @p dst is in
3353 * coarse-grained memory, the copied data is only visible to the agent currently
3354 * assigned (::hsa_memory_assign_agent) to @p dst.
3355 *
3356 * @param[in] src A valid pointer to the source of data to be copied. The source
3357 * buffer must not overlap with the destination buffer. If the source buffer is
3358 * in coarse-grained memory then it must be assigned to an agent, from which the
3359 * data will be retrieved.
3360 *
3361 * @param[in] size Number of bytes to copy. If @p size is 0, no copy is
3362 * performed and the function returns success. Copying a number of bytes larger
3363 * than the size of the buffers pointed by @p dst or @p src results in undefined
3364 * behavior.
3365 *
3366 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
3367 *
3368 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
3369 * initialized.
3370 *
3371 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT The source or destination
3372 * pointers are NULL.
3373 */
3374 hsa_status_t HSA_API hsa_memory_copy(
3375 void *dst,
3376 const void *src,
3377 size_t size);
3378
3379 /**
3380 * @brief Change the ownership of a global, coarse-grained buffer.
3381 *
3382 * @details The contents of a coarse-grained buffer are visible to an agent
3383 * only after ownership has been explicitely transferred to that agent. Once the
3384 * operation completes, the previous owner cannot longer access the data in the
3385 * buffer.
3386 *
3387 * An implementation of the HSA runtime is allowed, but not required, to change
3388 * the physical location of the buffer when ownership is transferred to a
3389 * different agent. In general the application must not assume this
3390 * behavior. The virtual location (address) of the passed buffer is never
3391 * modified.
3392 *
3393 * @param[in] ptr Base address of a global buffer. The pointer must match an
3394 * address previously returned by ::hsa_memory_allocate. The size of the buffer
3395 * affected by the ownership change is identical to the size of that previous
3396 * allocation. If @p ptr points to a fine-grained global buffer, no operation is
3397 * performed and the function returns success. If @p ptr does not point to
3398 * global memory, the behavior is undefined.
3399 *
3400 * @param[in] agent Agent that becomes the owner of the buffer. The
3401 * application is responsible for ensuring that @p agent has access to the
3402 * region that contains the buffer. It is allowed to change ownership to an
3403 * agent that is already the owner of the buffer, with the same or different
3404 * access permissions.
3405 *
3406 * @param[in] access Access permissions requested for the new owner.
3407 *
3408 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
3409 *
3410 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
3411 * initialized.
3412 *
3413 * @retval ::HSA_STATUS_ERROR_INVALID_AGENT The agent is invalid.
3414 *
3415 * @retval ::HSA_STATUS_ERROR_OUT_OF_RESOURCES The HSA runtime failed to allocate
3416 * the required resources.
3417 *
3418 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p ptr is NULL, or @p access is
3419 * not a valid access value.
3420 */
3421 hsa_status_t HSA_API hsa_memory_assign_agent(
3422 void *ptr,
3423 hsa_agent_t agent,
3424 hsa_access_permission_t access);
3425
3426 /**
3427 *
3428 * @brief Register a global, fine-grained buffer.
3429 *
3430 * @details Registering a buffer serves as an indication to the HSA runtime that
3431 * the memory might be accessed from a kernel agent other than the
3432 * host. Registration is a performance hint that allows the HSA runtime
3433 * implementation to know which buffers will be accessed by some of the kernel
3434 * agents ahead of time.
3435 *
3436 * Registration is only recommended for buffers in the global segment that have
3437 * not been allocated using the HSA allocator (::hsa_memory_allocate), but an OS
3438 * allocator instead. Registering an OS-allocated buffer in the base profile is
3439 * equivalent to a no-op.
3440 *
3441 * Registrations should not overlap.
3442 *
3443 * @param[in] ptr A buffer in global, fine-grained memory. If a NULL pointer is
3444 * passed, no operation is performed. If the buffer has been allocated using
3445 * ::hsa_memory_allocate, or has already been registered, no operation is
3446 * performed.
3447 *
3448 * @param[in] size Requested registration size in bytes. A size of 0 is
3449 * only allowed if @p ptr is NULL.
3450 *
3451 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
3452 *
3453 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
3454 * initialized.
3455 *
3456 * @retval ::HSA_STATUS_ERROR_OUT_OF_RESOURCES The HSA runtime failed to allocate
3457 * the required resources.
3458 *
3459 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p size is 0 but @p ptr
3460 * is not NULL.
3461 */
3462 hsa_status_t HSA_API hsa_memory_register(
3463 void *ptr,
3464 size_t size);
3465
3466 /**
3467 *
3468 * @brief Deregister memory previously registered using ::hsa_memory_register.
3469 *
3470 * @details If the memory interval being deregistered does not match a previous
3471 * registration (start and end addresses), the behavior is undefined.
3472 *
3473 * @param[in] ptr A pointer to the base of the buffer to be deregistered. If
3474 * a NULL pointer is passed, no operation is performed.
3475 *
3476 * @param[in] size Size of the buffer to be deregistered.
3477 *
3478 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
3479 *
3480 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
3481 * initialized.
3482 *
3483 */
3484 hsa_status_t HSA_API hsa_memory_deregister(
3485 void *ptr,
3486 size_t size);
3487
3488 /** @} */
3489
3490
3491 /** \defgroup instruction-set-architecture Instruction Set Architecture.
3492 * @{
3493 */
3494
3495 /**
3496 * @brief Instruction set architecture.
3497 */
3498 typedef struct hsa_isa_s {
3499 /**
3500 * Opaque handle. Two handles reference the same object of the enclosing type
3501 * if and only if they are equal.
3502 */
3503 uint64_t handle;
3504 } hsa_isa_t;
3505
3506 /**
3507 * @brief Retrieve a reference to an instruction set architecture handle out of
3508 * a symbolic name.
3509 *
3510 * @param[in] name Vendor-specific name associated with a a particular
3511 * instruction set architecture. @p name must start with the vendor name and a
3512 * colon (for example, "AMD:"). The rest of the name is vendor-specific. Must be
3513 * a NUL-terminated string.
3514 *
3515 * @param[out] isa Memory location where the HSA runtime stores the ISA handle
3516 * corresponding to the given name. Must not be NULL.
3517 *
3518 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
3519 *
3520 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
3521 * initialized.
3522 *
3523 * @retval ::HSA_STATUS_ERROR_INVALID_ISA_NAME The given name does not
3524 * correspond to any instruction set architecture.
3525 *
3526 * @retval ::HSA_STATUS_ERROR_OUT_OF_RESOURCES The HSA runtime failed to
3527 * allocate the required resources.
3528 *
3529 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p name is NULL, or @p isa is
3530 * NULL.
3531 */
3532 hsa_status_t HSA_API hsa_isa_from_name(
3533 const char *name,
3534 hsa_isa_t *isa);
3535
3536 /**
3537 * @brief Iterate over the instruction sets supported by the given agent, and
3538 * invoke an application-defined callback on every iteration. The iterator is
3539 * deterministic: if an agent supports several instruction set architectures,
3540 * they are traversed in the same order in every invocation of this function.
3541 *
3542 * @param[in] agent A valid agent.
3543 *
3544 * @param[in] callback Callback to be invoked once per instruction set
3545 * architecture. The HSA runtime passes two arguments to the callback: the
3546 * ISA and the application data. If @p callback returns a status other than
3547 * ::HSA_STATUS_SUCCESS for a particular iteration, the traversal stops and
3548 * that status value is returned.
3549 *
3550 * @param[in] data Application data that is passed to @p callback on every
3551 * iteration. May be NULL.
3552 *
3553 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
3554 *
3555 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
3556 * initialized.
3557 *
3558 * @retval ::HSA_STATUS_ERROR_INVALID_AGENT The agent is invalid.
3559 *
3560 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p callback is NULL.
3561 */
3562 hsa_status_t HSA_API hsa_agent_iterate_isas(
3563 hsa_agent_t agent,
3564 hsa_status_t (*callback)(hsa_isa_t isa, void *data),
3565 void *data);
3566
3567 /**
3568 * @brief Instruction set architecture attributes.
3569 */
3570 typedef enum {
3571 /**
3572 * The length of the ISA name in bytes, not including the NUL terminator. The
3573 * type of this attribute is uint32_t.
3574 */
3575 HSA_ISA_INFO_NAME_LENGTH = 0,
3576 /**
3577 * Human-readable description. The type of this attribute is character array
3578 * with the length equal to the value of ::HSA_ISA_INFO_NAME_LENGTH attribute.
3579 */
3580 HSA_ISA_INFO_NAME = 1,
3581 /**
3582 * @deprecated
3583 *
3584 * Number of call conventions supported by the instruction set architecture.
3585 * Must be greater than zero. The type of this attribute is uint32_t.
3586 */
3587 HSA_ISA_INFO_CALL_CONVENTION_COUNT = 2,
3588 /**
3589 * @deprecated
3590 *
3591 * Number of work-items in a wavefront for a given call convention. Must be a
3592 * power of 2 in the range [1,256]. The type of this attribute is uint32_t.
3593 */
3594 HSA_ISA_INFO_CALL_CONVENTION_INFO_WAVEFRONT_SIZE = 3,
3595 /**
3596 * @deprecated
3597 *
3598 * Number of wavefronts per compute unit for a given call convention. In
3599 * practice, other factors (for example, the amount of group memory used by a
3600 * work-group) may further limit the number of wavefronts per compute
3601 * unit. The type of this attribute is uint32_t.
3602 */
3603 HSA_ISA_INFO_CALL_CONVENTION_INFO_WAVEFRONTS_PER_COMPUTE_UNIT = 4,
3604 /**
3605 * Machine models supported by the instruction set architecture. The type of
3606 * this attribute is a bool[2]. If the ISA supports the small machine model,
3607 * the element at index ::HSA_MACHINE_MODEL_SMALL is true. If the ISA supports
3608 * the large model, the element at index ::HSA_MACHINE_MODEL_LARGE is true.
3609 */
3610 HSA_ISA_INFO_MACHINE_MODELS = 5,
3611 /**
3612 * Profiles supported by the instruction set architecture. The type of this
3613 * attribute is a bool[2]. If the ISA supports the base profile, the element
3614 * at index ::HSA_PROFILE_BASE is true. If the ISA supports the full profile,
3615 * the element at index ::HSA_PROFILE_FULL is true.
3616 */
3617 HSA_ISA_INFO_PROFILES = 6,
3618 /**
3619 * Default floating-point rounding modes supported by the instruction set
3620 * architecture. The type of this attribute is a bool[3]. The value at a given
3621 * index is true if the corresponding rounding mode in
3622 * ::hsa_default_float_rounding_mode_t is supported. At least one default mode
3623 * has to be supported.
3624 *
3625 * If the default mode is supported, then
3626 * ::HSA_ISA_INFO_BASE_PROFILE_DEFAULT_FLOAT_ROUNDING_MODES must report that
3627 * both the zero and the near roundings modes are supported.
3628 */
3629 HSA_ISA_INFO_DEFAULT_FLOAT_ROUNDING_MODES = 7,
3630 /**
3631 * Default floating-point rounding modes supported by the instruction set
3632 * architecture in the Base profile. The type of this attribute is a
3633 * bool[3]. The value at a given index is true if the corresponding rounding
3634 * mode in ::hsa_default_float_rounding_mode_t is supported. The value at
3635 * index HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT must be false. At least one
3636 * of the values at indexes ::HSA_DEFAULT_FLOAT_ROUNDING_MODE_ZERO or
3637 * HSA_DEFAULT_FLOAT_ROUNDING_MODE_NEAR must be true.
3638 */
3639 HSA_ISA_INFO_BASE_PROFILE_DEFAULT_FLOAT_ROUNDING_MODES = 8,
3640 /**
3641 * Flag indicating that the f16 HSAIL operation is at least as fast as the
3642 * f32 operation in the instruction set architecture. The type of this
3643 * attribute is bool.
3644 */
3645 HSA_ISA_INFO_FAST_F16_OPERATION = 9,
3646 /**
3647 * Maximum number of work-items of each dimension of a work-group. Each
3648 * maximum must be greater than 0. No maximum can exceed the value of
3649 * ::HSA_ISA_INFO_WORKGROUP_MAX_SIZE. The type of this attribute is
3650 * uint16_t[3].
3651 */
3652 HSA_ISA_INFO_WORKGROUP_MAX_DIM = 12,
3653 /**
3654 * Maximum total number of work-items in a work-group. The type
3655 * of this attribute is uint32_t.
3656 */
3657 HSA_ISA_INFO_WORKGROUP_MAX_SIZE = 13,
3658 /**
3659 * Maximum number of work-items of each dimension of a grid. Each maximum must
3660 * be greater than 0, and must not be smaller than the corresponding value in
3661 * ::HSA_ISA_INFO_WORKGROUP_MAX_DIM. No maximum can exceed the value of
3662 * ::HSA_ISA_INFO_GRID_MAX_SIZE. The type of this attribute is
3663 * ::hsa_dim3_t.
3664 */
3665 HSA_ISA_INFO_GRID_MAX_DIM = 14,
3666 /**
3667 * Maximum total number of work-items in a grid. The type of this
3668 * attribute is uint64_t.
3669 */
3670 HSA_ISA_INFO_GRID_MAX_SIZE = 16,
3671 /**
3672 * Maximum number of fbarriers per work-group. Must be at least 32. The
3673 * type of this attribute is uint32_t.
3674 */
3675 HSA_ISA_INFO_FBARRIER_MAX_SIZE = 17
3676 } hsa_isa_info_t;
3677
3678 /**
3679 * @deprecated The concept of call convention has been deprecated. If the
3680 * application wants to query the value of an attribute for a given instruction
3681 * set architecture, use ::hsa_isa_get_info_alt instead. If the application
3682 * wants to query an attribute that is specific to a given combination of ISA
3683 * and wavefront, use ::hsa_wavefront_get_info.
3684 *
3685 * @brief Get the current value of an attribute for a given instruction set
3686 * architecture (ISA).
3687 *
3688 * @param[in] isa A valid instruction set architecture.
3689 *
3690 * @param[in] attribute Attribute to query.
3691 *
3692 * @param[in] index Call convention index. Used only for call convention
3693 * attributes, otherwise ignored. Must have a value between 0 (inclusive) and
3694 * the value of the attribute ::HSA_ISA_INFO_CALL_CONVENTION_COUNT (not
3695 * inclusive) in @p isa.
3696 *
3697 * @param[out] value Pointer to an application-allocated buffer where to store
3698 * the value of the attribute. If the buffer passed by the application is not
3699 * large enough to hold the value of @p attribute, the behavior is undefined.
3700 *
3701 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
3702 *
3703 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
3704 * initialized.
3705 *
3706 * @retval ::HSA_STATUS_ERROR_INVALID_ISA The instruction set architecture is
3707 * invalid.
3708 *
3709 * @retval ::HSA_STATUS_ERROR_INVALID_INDEX The index is out of range.
3710 *
3711 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p attribute is an invalid
3712 * instruction set architecture attribute, or @p value is
3713 * NULL.
3714 */
3715 hsa_status_t HSA_API HSA_DEPRECATED hsa_isa_get_info(
3716 hsa_isa_t isa,
3717 hsa_isa_info_t attribute,
3718 uint32_t index,
3719 void *value);
3720
3721 /**
3722 * @brief Get the current value of an attribute for a given instruction set
3723 * architecture (ISA).
3724 *
3725 * @param[in] isa A valid instruction set architecture.
3726 *
3727 * @param[in] attribute Attribute to query.
3728 *
3729 * @param[out] value Pointer to an application-allocated buffer where to store
3730 * the value of the attribute. If the buffer passed by the application is not
3731 * large enough to hold the value of @p attribute, the behavior is undefined.
3732 *
3733 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
3734 *
3735 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
3736 * initialized.
3737 *
3738 * @retval ::HSA_STATUS_ERROR_INVALID_ISA The instruction set architecture is
3739 * invalid.
3740 *
3741 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p attribute is an invalid
3742 * instruction set architecture attribute, or @p value is
3743 * NULL.
3744 */
3745 hsa_status_t HSA_API hsa_isa_get_info_alt(
3746 hsa_isa_t isa,
3747 hsa_isa_info_t attribute,
3748 void *value);
3749
3750 /**
3751 * @brief Retrieve the exception policy support for a given combination of
3752 * instruction set architecture and profile.
3753 *
3754 * @param[in] isa A valid instruction set architecture.
3755 *
3756 * @param[in] profile Profile.
3757 *
3758 * @param[out] mask Pointer to a memory location where the HSA runtime stores a
3759 * mask of ::hsa_exception_policy_t values. Must not be NULL.
3760 *
3761 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
3762 *
3763 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
3764 * initialized.
3765 *
3766 * @retval ::HSA_STATUS_ERROR_INVALID_ISA The instruction set architecture is
3767 * invalid.
3768 *
3769 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p profile is not a valid
3770 * profile, or @p mask is NULL.
3771 */
3772 hsa_status_t HSA_API hsa_isa_get_exception_policies(
3773 hsa_isa_t isa,
3774 hsa_profile_t profile,
3775 uint16_t *mask);
3776
3777 /**
3778 * @brief Floating-point types.
3779 */
3780 typedef enum {
3781 /**
3782 * 16-bit floating-point type.
3783 */
3784 HSA_FP_TYPE_16 = 1,
3785 /**
3786 * 32-bit floating-point type.
3787 */
3788 HSA_FP_TYPE_32 = 2,
3789 /**
3790 * 64-bit floating-point type.
3791 */
3792 HSA_FP_TYPE_64 = 4
3793 } hsa_fp_type_t;
3794
3795 /**
3796 * @brief Flush to zero modes.
3797 */
3798 typedef enum {
3799 /**
3800 * Flush to zero.
3801 */
3802 HSA_FLUSH_MODE_FTZ = 1,
3803 /**
3804 * Do not flush to zero.
3805 */
3806 HSA_FLUSH_MODE_NON_FTZ = 2
3807 } hsa_flush_mode_t;
3808
3809 /**
3810 * @brief Round methods.
3811 */
3812 typedef enum {
3813 /**
3814 * Single round method.
3815 */
3816 HSA_ROUND_METHOD_SINGLE = 1,
3817 /**
3818 * Double round method.
3819 */
3820 HSA_ROUND_METHOD_DOUBLE = 2
3821 } hsa_round_method_t;
3822
3823 /**
3824 * @brief Retrieve the round method (single or double) used to implement the
3825 * floating-point multiply add instruction (mad) for a given combination of
3826 * instruction set architecture, floating-point type, and flush to zero
3827 * modifier.
3828 *
3829 * @param[in] isa Instruction set architecture.
3830 *
3831 * @param[in] fp_type Floating-point type.
3832 *
3833 * @param[in] flush_mode Flush to zero modifier.
3834 *
3835 * @param[out] round_method Pointer to a memory location where the HSA
3836 * runtime stores the round method used by the implementation. Must not be NULL.
3837 *
3838 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
3839 *
3840 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
3841 * initialized.
3842 *
3843 * @retval ::HSA_STATUS_ERROR_INVALID_ISA The instruction set architecture is
3844 * invalid.
3845 *
3846 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p fp_type is not a valid
3847 * floating-point type, or @p flush_mode is not a valid flush to zero modifier,
3848 * or @p round_method is NULL.
3849 */
3850 hsa_status_t HSA_API hsa_isa_get_round_method(
3851 hsa_isa_t isa,
3852 hsa_fp_type_t fp_type,
3853 hsa_flush_mode_t flush_mode,
3854 hsa_round_method_t *round_method);
3855
3856 /**
3857 * @brief Wavefront handle
3858 */
3859 typedef struct hsa_wavefront_s {
3860 /**
3861 * Opaque handle. Two handles reference the same object of the enclosing type
3862 * if and only if they are equal.
3863 */
3864 uint64_t handle;
3865 } hsa_wavefront_t;
3866
3867 /**
3868 * @brief Wavefront attributes.
3869 */
3870 typedef enum {
3871 /**
3872 * Number of work-items in the wavefront. Must be a power of 2 in the range
3873 * [1,256]. The type of this attribute is uint32_t.
3874 */
3875 HSA_WAVEFRONT_INFO_SIZE = 0
3876 } hsa_wavefront_info_t;
3877
3878 /**
3879 * @brief Get the current value of a wavefront attribute.
3880 *
3881 * @param[in] wavefront A wavefront.
3882 *
3883 * @param[in] attribute Attribute to query.
3884 *
3885 * @param[out] value Pointer to an application-allocated buffer where to store
3886 * the value of the attribute. If the buffer passed by the application is not
3887 * large enough to hold the value of @p attribute, the behavior is undefined.
3888 *
3889 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
3890 *
3891 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
3892 * initialized.
3893 *
3894 * @retval ::HSA_STATUS_ERROR_INVALID_WAVEFRONT The wavefront is invalid.
3895 *
3896 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p attribute is an invalid
3897 * wavefront attribute, or @p value is NULL.
3898 */
3899 hsa_status_t HSA_API hsa_wavefront_get_info(
3900 hsa_wavefront_t wavefront,
3901 hsa_wavefront_info_t attribute,
3902 void *value);
3903
3904 /**
3905 * @brief Iterate over the different wavefronts supported by an instruction set
3906 * architecture, and invoke an application-defined callback on every iteration.
3907 *
3908 * @param[in] isa Instruction set architecture.
3909 *
3910 * @param[in] callback Callback to be invoked once per wavefront that is
3911 * supported by the agent. The HSA runtime passes two arguments to the callback:
3912 * the wavefront handle and the application data. If @p callback returns a
3913 * status other than ::HSA_STATUS_SUCCESS for a particular iteration, the
3914 * traversal stops and that value is returned.
3915 *
3916 * @param[in] data Application data that is passed to @p callback on every
3917 * iteration. May be NULL.
3918 *
3919 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
3920 *
3921 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
3922 * initialized.
3923 *
3924 * @retval ::HSA_STATUS_ERROR_INVALID_ISA The instruction set architecture is
3925 * invalid.
3926 *
3927 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p callback is NULL.
3928 */
3929 hsa_status_t HSA_API hsa_isa_iterate_wavefronts(
3930 hsa_isa_t isa,
3931 hsa_status_t (*callback)(hsa_wavefront_t wavefront, void *data),
3932 void *data);
3933
3934 /**
3935 * @deprecated Use ::hsa_agent_iterate_isas to query which instructions set
3936 * architectures are supported by a given agent.
3937 *
3938 * @brief Check if the instruction set architecture of a code object can be
3939 * executed on an agent associated with another architecture.
3940 *
3941 * @param[in] code_object_isa Instruction set architecture associated with a
3942 * code object.
3943 *
3944 * @param[in] agent_isa Instruction set architecture associated with an agent.
3945 *
3946 * @param[out] result Pointer to a memory location where the HSA runtime stores
3947 * the result of the check. If the two architectures are compatible, the result
3948 * is true; if they are incompatible, the result is false.
3949 *
3950 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
3951 *
3952 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
3953 * initialized.
3954 *
3955 * @retval ::HSA_STATUS_ERROR_INVALID_ISA @p code_object_isa or @p agent_isa are
3956 * invalid.
3957 *
3958 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p result is NULL.
3959 */
3960 hsa_status_t HSA_API HSA_DEPRECATED hsa_isa_compatible(
3961 hsa_isa_t code_object_isa,
3962 hsa_isa_t agent_isa,
3963 bool *result);
3964
3965 /** @} */
3966
3967
3968 /** \defgroup executable Executable
3969 * @{
3970 */
3971
3972 /**
3973 * @brief Code object reader handle. A code object reader is used to
3974 * load a code object from file (when created using
3975 * ::hsa_code_object_reader_create_from_file), or from memory (if created using
3976 * ::hsa_code_object_reader_create_from_memory).
3977 */
3978 typedef struct hsa_code_object_reader_s {
3979 /**
3980 * Opaque handle. Two handles reference the same object of the enclosing type
3981 * if and only if they are equal.
3982 */
3983 uint64_t handle;
3984 } hsa_code_object_reader_t;
3985
3986 /**
3987 * @brief Create a code object reader to operate on a file.
3988 *
3989 * @param[in] file File descriptor. The file must have been opened by
3990 * application with at least read permissions prior calling this function. The
3991 * file must contain a vendor-specific code object.
3992 *
3993 * The file is owned and managed by the application; the lifetime of the file
3994 * descriptor must exceed that of any associated code object reader.
3995 *
3996 * @param[out] code_object_reader Memory location to store the newly created
3997 * code object reader handle. Must not be NULL.
3998 *
3999 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
4000 *
4001 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
4002 * initialized.
4003 *
4004 * @retval ::HSA_STATUS_ERROR_INVALID_FILE @p file is invalid.
4005 *
4006 * @retval ::HSA_STATUS_ERROR_OUT_OF_RESOURCES The HSA runtime failed to
4007 * allocate the required resources.
4008 *
4009 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p code_object_reader is NULL.
4010 */
4011 hsa_status_t HSA_API hsa_code_object_reader_create_from_file(
4012 hsa_file_t file,
4013 hsa_code_object_reader_t *code_object_reader);
4014
4015 /**
4016 * @brief Create a code object reader to operate on memory.
4017 *
4018 * @param[in] code_object Memory buffer that contains a vendor-specific code
4019 * object. The buffer is owned and managed by the application; the lifetime of
4020 * the buffer must exceed that of any associated code object reader.
4021 *
4022 * @param[in] size Size of the buffer pointed to by @p code_object. Must not be
4023 * 0.
4024 *
4025 * @param[out] code_object_reader Memory location to store newly created code
4026 * object reader handle. Must not be NULL.
4027 *
4028 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
4029 *
4030 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
4031 * initialized.
4032 *
4033 * @retval ::HSA_STATUS_ERROR_OUT_OF_RESOURCES The HSA runtime failed to
4034 * allocate the required resources.
4035 *
4036 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p code_object is NULL, @p size
4037 * is zero, or @p code_object_reader is NULL.
4038 */
4039 hsa_status_t HSA_API hsa_code_object_reader_create_from_memory(
4040 const void *code_object,
4041 size_t size,
4042 hsa_code_object_reader_t *code_object_reader);
4043
4044 /**
4045 * @brief Destroy a code object reader.
4046 *
4047 * @details The code object reader handle becomes invalid after completion of
4048 * this function. Any file or memory used to create the code object read is not
4049 * closed, removed, or deallocated by this function.
4050 *
4051 * @param[in] code_object_reader Code object reader to destroy.
4052 *
4053 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
4054 *
4055 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
4056 * initialized.
4057 *
4058 * @retval ::HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER @p code_object_reader
4059 * is invalid.
4060 */
4061 hsa_status_t HSA_API hsa_code_object_reader_destroy(
4062 hsa_code_object_reader_t code_object_reader);
4063
4064 /**
4065 * @brief Struct containing an opaque handle to an executable, which contains
4066 * ISA for finalized kernels and indirect functions together with the allocated
4067 * global or readonly segment variables they reference.
4068 */
4069 typedef struct hsa_executable_s {
4070 /**
4071 * Opaque handle. Two handles reference the same object of the enclosing type
4072 * if and only if they are equal.
4073 */
4074 uint64_t handle;
4075 } hsa_executable_t;
4076
4077 /**
4078 * @brief Executable state.
4079 */
4080 typedef enum {
4081 /**
4082 * Executable state, which allows the user to load code objects and define
4083 * external variables. Variable addresses, kernel code handles, and
4084 * indirect function code handles are not available in query operations until
4085 * the executable is frozen (zero always returned).
4086 */
4087 HSA_EXECUTABLE_STATE_UNFROZEN = 0,
4088 /**
4089 * Executable state, which allows the user to query variable addresses,
4090 * kernel code handles, and indirect function code handles using query
4091 * operations. Loading new code objects, as well as defining external
4092 * variables, is not allowed in this state.
4093 */
4094 HSA_EXECUTABLE_STATE_FROZEN = 1
4095 } hsa_executable_state_t;
4096
4097 /**
4098 * @deprecated Use ::hsa_executable_create_alt instead, which allows the
4099 * application to specify the default floating-point rounding mode of the
4100 * executable and assumes an unfrozen initial state.
4101 *
4102 * @brief Create an empty executable.
4103 *
4104 * @param[in] profile Profile used in the executable.
4105 *
4106 * @param[in] executable_state Executable state. If the state is
4107 * ::HSA_EXECUTABLE_STATE_FROZEN, the resulting executable is useless because no
4108 * code objects can be loaded, and no variables can be defined.
4109 *
4110 * @param[in] options Standard and vendor-specific options. Unknown options are
4111 * ignored. A standard option begins with the "-hsa_" prefix. Options beginning
4112 * with the "-hsa_ext_<extension_name>_" prefix are reserved for extensions. A
4113 * vendor-specific option begins with the "-<vendor_name>_" prefix. Must be a
4114 * NUL-terminated string. May be NULL.
4115 *
4116 * @param[out] executable Memory location where the HSA runtime stores the newly
4117 * created executable handle.
4118 *
4119 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
4120 *
4121 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
4122 * initialized.
4123 *
4124 * @retval ::HSA_STATUS_ERROR_OUT_OF_RESOURCES The HSA runtime failed to
4125 * allocate the required resources.
4126 *
4127 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p profile is invalid, or
4128 * @p executable is NULL.
4129 */
4130 hsa_status_t HSA_API HSA_DEPRECATED hsa_executable_create(
4131 hsa_profile_t profile,
4132 hsa_executable_state_t executable_state,
4133 const char *options,
4134 hsa_executable_t *executable);
4135
4136 /**
4137 * @brief Create an empty executable.
4138 *
4139 * @param[in] profile Profile used in the executable.
4140 *
4141 * @param[in] default_float_rounding_mode Default floating-point rounding mode
4142 * used in the executable. Allowed rounding modes are near and zero (default is
4143 * not allowed).
4144 *
4145 * @param[in] options Standard and vendor-specific options. Unknown options are
4146 * ignored. A standard option begins with the "-hsa_" prefix. Options beginning
4147 * with the "-hsa_ext_<extension_name>_" prefix are reserved for extensions. A
4148 * vendor-specific option begins with the "-<vendor_name>_" prefix. Must be a
4149 * NUL-terminated string. May be NULL.
4150 *
4151 * @param[out] executable Memory location where the HSA runtime stores newly
4152 * created executable handle. The initial state of the executable is
4153 * ::HSA_EXECUTABLE_STATE_UNFROZEN.
4154 *
4155 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
4156 *
4157 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
4158 * initialized.
4159 *
4160 * @retval ::HSA_STATUS_ERROR_OUT_OF_RESOURCES The HSA runtime failed to
4161 * allocate the required resources.
4162 *
4163 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p profile is invalid, or
4164 * @p executable is NULL.
4165 */
4166 hsa_status_t HSA_API hsa_executable_create_alt(
4167 hsa_profile_t profile,
4168 hsa_default_float_rounding_mode_t default_float_rounding_mode,
4169 const char *options,
4170 hsa_executable_t *executable);
4171
4172 /**
4173 * @brief Destroy an executable.
4174 *
4175 * @details An executable handle becomes invalid after the executable has been
4176 * destroyed. Code object handles that were loaded into this executable are
4177 * still valid after the executable has been destroyed, and can be used as
4178 * intended. Resources allocated outside and associated with this executable
4179 * (such as external global or readonly variables) can be released after the
4180 * executable has been destroyed.
4181 *
4182 * Executable should not be destroyed while kernels are in flight.
4183 *
4184 * @param[in] executable Executable.
4185 *
4186 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
4187 *
4188 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
4189 * initialized.
4190 *
4191 * @retval ::HSA_STATUS_ERROR_INVALID_EXECUTABLE The executable is invalid.
4192 */
4193 hsa_status_t HSA_API hsa_executable_destroy(
4194 hsa_executable_t executable);
4195
4196 /**
4197 * @brief Loaded code object handle.
4198 */
4199 typedef struct hsa_loaded_code_object_s {
4200 /**
4201 * Opaque handle. Two handles reference the same object of the enclosing type
4202 * if and only if they are equal.
4203 */
4204 uint64_t handle;
4205 } hsa_loaded_code_object_t;
4206
4207 /**
4208 * @brief Load a program code object into an executable.
4209 *
4210 * @details A program code object contains information about resources that are
4211 * accessible by all kernel agents that run the executable, and can be loaded
4212 * at most once into an executable.
4213 *
4214 * If the program code object uses extensions, the implementation must support
4215 * them for this operation to return successfully.
4216 *
4217 * @param[in] executable Executable.
4218 *
4219 * @param[in] code_object_reader A code object reader that holds the program
4220 * code object to load. If a code object reader is destroyed before all the
4221 * associated executables are destroyed, the behavior is undefined.
4222 *
4223 * @param[in] options Standard and vendor-specific options. Unknown options are
4224 * ignored. A standard option begins with the "-hsa_" prefix. Options beginning
4225 * with the "-hsa_ext_<extension_name>_" prefix are reserved for extensions. A
4226 * vendor-specific option begins with the "-<vendor_name>_" prefix. Must be a
4227 * NUL-terminated string. May be NULL.
4228 *
4229 * @param[out] loaded_code_object Pointer to a memory location where the HSA
4230 * runtime stores the loaded code object handle. May be NULL.
4231 *
4232 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
4233 *
4234 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
4235 * initialized.
4236 *
4237 * @retval ::HSA_STATUS_ERROR_OUT_OF_RESOURCES The HSA runtime failed to
4238 * allocate the required resources.
4239 *
4240 * @retval ::HSA_STATUS_ERROR_INVALID_EXECUTABLE The executable is invalid.
4241 *
4242 * @retval ::HSA_STATUS_ERROR_FROZEN_EXECUTABLE The executable is frozen.
4243 *
4244 * @retval ::HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER @p code_object_reader
4245 * is invalid.
4246 *
4247 * @retval ::HSA_STATUS_ERROR_INCOMPATIBLE_ARGUMENTS The program code object is
4248 * not compatible with the executable or the implementation (for example, the
4249 * code object uses an extension that is not supported by the implementation).
4250 */
4251 hsa_status_t HSA_API hsa_executable_load_program_code_object(
4252 hsa_executable_t executable,
4253 hsa_code_object_reader_t code_object_reader,
4254 const char *options,
4255 hsa_loaded_code_object_t *loaded_code_object);
4256
4257 /**
4258 * @brief Load an agent code object into an executable.
4259 *
4260 * @details The agent code object contains all defined agent
4261 * allocation variables, functions, indirect functions, and kernels in a given
4262 * program for a given instruction set architecture.
4263 *
4264 * Any module linkage declaration must have been defined either by a define
4265 * variable or by loading a code object that has a symbol with module linkage
4266 * definition.
4267 *
4268 * The default floating-point rounding mode of the code object associated with
4269 * @p code_object_reader must match that of the executable
4270 * (::HSA_EXECUTABLE_INFO_DEFAULT_FLOAT_ROUNDING_MODE), or be default (in which
4271 * case the value of ::HSA_EXECUTABLE_INFO_DEFAULT_FLOAT_ROUNDING_MODE is used).
4272 * If the agent code object uses extensions, the implementation and the agent
4273 * must support them for this operation to return successfully.
4274 *
4275 * @param[in] executable Executable.
4276 *
4277 * @param[in] agent Agent to load code object for. A code object can be loaded
4278 * into an executable at most once for a given agent. The instruction set
4279 * architecture of the code object must be supported by the agent.
4280 *
4281 * @param[in] code_object_reader A code object reader that holds the code object
4282 * to load. If a code object reader is destroyed before all the associated
4283 * executables are destroyed, the behavior is undefined.
4284 *
4285 * @param[in] options Standard and vendor-specific options. Unknown options are
4286 * ignored. A standard option begins with the "-hsa_" prefix. Options beginning
4287 * with the "-hsa_ext_<extension_name>_" prefix are reserved for extensions. A
4288 * vendor-specific option begins with the "-<vendor_name>_" prefix. Must be a
4289 * NUL-terminated string. May be NULL.
4290 *
4291 * @param[out] loaded_code_object Pointer to a memory location where the HSA
4292 * runtime stores the loaded code object handle. May be NULL.
4293 *
4294 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
4295 *
4296 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
4297 * initialized.
4298 *
4299 * @retval ::HSA_STATUS_ERROR_OUT_OF_RESOURCES The HSA runtime failed to
4300 * allocate the required resources.
4301 *
4302 * @retval ::HSA_STATUS_ERROR_INVALID_EXECUTABLE The executable is invalid.
4303 *
4304 * @retval ::HSA_STATUS_ERROR_FROZEN_EXECUTABLE The executable is frozen.
4305 *
4306 * @retval ::HSA_STATUS_ERROR_INVALID_AGENT The agent is invalid.
4307 *
4308 * @retval ::HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER @p code_object_reader
4309 * is invalid.
4310 *
4311 * @retval ::HSA_STATUS_ERROR_INCOMPATIBLE_ARGUMENTS The code object read by @p
4312 * code_object_reader is not compatible with the agent (for example, the agent
4313 * does not support the instruction set architecture of the code object), the
4314 * executable (for example, there is a default floating-point mode mismatch
4315 * between the two), or the implementation.
4316 */
4317 hsa_status_t HSA_API hsa_executable_load_agent_code_object(
4318 hsa_executable_t executable,
4319 hsa_agent_t agent,
4320 hsa_code_object_reader_t code_object_reader,
4321 const char *options,
4322 hsa_loaded_code_object_t *loaded_code_object);
4323
4324 /**
4325 * @brief Freeze the executable.
4326 *
4327 * @details No modifications to executable can be made after freezing: no code
4328 * objects can be loaded to the executable, and no external variables can be
4329 * defined. Freezing the executable does not prevent querying the executable's
4330 * attributes. The application must define all the external variables in an
4331 * executable before freezing it.
4332 *
4333 * @param[in] executable Executable.
4334 *
4335 * @param[in] options Standard and vendor-specific options. Unknown options are
4336 * ignored. A standard option begins with the "-hsa_" prefix. Options beginning
4337 * with the "-hsa_ext_<extension_name>_" prefix are reserved for extensions. A
4338 * vendor-specific option begins with the "-<vendor_name>_" prefix. Must be a
4339 * NUL-terminated string. May be NULL.
4340 *
4341 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
4342 *
4343 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
4344 * initialized.
4345 *
4346 * @retval ::HSA_STATUS_ERROR_INVALID_EXECUTABLE The executable is invalid.
4347 *
4348 * @retval ::HSA_STATUS_ERROR_VARIABLE_UNDEFINED One or more variables are
4349 * undefined in the executable.
4350 *
4351 * @retval ::HSA_STATUS_ERROR_FROZEN_EXECUTABLE @p executable is already frozen.
4352 */
4353 hsa_status_t HSA_API hsa_executable_freeze(
4354 hsa_executable_t executable,
4355 const char *options);
4356
4357 /**
4358 * @brief Executable attributes.
4359 */
4360 typedef enum {
4361 /**
4362 * Profile this executable is created for. The type of this attribute is
4363 * ::hsa_profile_t.
4364 */
4365 HSA_EXECUTABLE_INFO_PROFILE = 1,
4366 /**
4367 * Executable state. The type of this attribute is ::hsa_executable_state_t.
4368 */
4369 HSA_EXECUTABLE_INFO_STATE = 2,
4370 /**
4371 * Default floating-point rounding mode specified when executable was created.
4372 * The type of this attribute is ::hsa_default_float_rounding_mode_t.
4373 */
4374 HSA_EXECUTABLE_INFO_DEFAULT_FLOAT_ROUNDING_MODE = 3
4375 } hsa_executable_info_t;
4376
4377 /**
4378 * @brief Get the current value of an attribute for a given executable.
4379 *
4380 * @param[in] executable Executable.
4381 *
4382 * @param[in] attribute Attribute to query.
4383 *
4384 * @param[out] value Pointer to an application-allocated buffer where to store
4385 * the value of the attribute. If the buffer passed by the application is not
4386 * large enough to hold the value of @p attribute, the behavior is undefined.
4387 *
4388 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
4389 *
4390 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
4391 * initialized.
4392 *
4393 * @retval ::HSA_STATUS_ERROR_INVALID_EXECUTABLE The executable is invalid.
4394 *
4395 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p attribute is an invalid
4396 * executable attribute, or @p value is NULL.
4397 */
4398 hsa_status_t HSA_API hsa_executable_get_info(
4399 hsa_executable_t executable,
4400 hsa_executable_info_t attribute,
4401 void *value);
4402
4403 /**
4404 * @brief Define an external global variable with program allocation.
4405 *
4406 * @details This function allows the application to provide the definition
4407 * of a variable in the global segment memory with program allocation. The
4408 * variable must be defined before loading a code object into an executable.
4409 * In addition, code objects loaded must not define the variable.
4410 *
4411 * @param[in] executable Executable. Must not be in frozen state.
4412 *
4413 * @param[in] variable_name Name of the variable. The Programmer's Reference
4414 * Manual describes the standard name mangling scheme.
4415 *
4416 * @param[in] address Address where the variable is defined. This address must
4417 * be in global memory and can be read and written by any agent in the
4418 * system. The application cannot deallocate the buffer pointed by @p address
4419 * before @p executable is destroyed.
4420 *
4421 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
4422 *
4423 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
4424 * initialized.
4425 *
4426 * @retval ::HSA_STATUS_ERROR_OUT_OF_RESOURCES The HSA runtime failed to
4427 * allocate the required resources.
4428 *
4429 * @retval ::HSA_STATUS_ERROR_INVALID_EXECUTABLE The executable is invalid.
4430 *
4431 * @retval ::HSA_STATUS_ERROR_VARIABLE_ALREADY_DEFINED The variable is
4432 * already defined.
4433 *
4434 * @retval ::HSA_STATUS_ERROR_INVALID_SYMBOL_NAME There is no variable with the
4435 * @p variable_name.
4436 *
4437 * @retval ::HSA_STATUS_ERROR_FROZEN_EXECUTABLE @p executable is frozen.
4438 *
4439 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p variable_name is NULL.
4440 */
4441 hsa_status_t HSA_API hsa_executable_global_variable_define(
4442 hsa_executable_t executable,
4443 const char *variable_name,
4444 void *address);
4445
4446 /**
4447 * @brief Define an external global variable with agent allocation.
4448 *
4449 * @details This function allows the application to provide the definition
4450 * of a variable in the global segment memory with agent allocation. The
4451 * variable must be defined before loading a code object into an executable.
4452 * In addition, code objects loaded must not define the variable.
4453 *
4454 * @param[in] executable Executable. Must not be in frozen state.
4455 *
4456 * @param[in] agent Agent for which the variable is being defined.
4457 *
4458 * @param[in] variable_name Name of the variable. The Programmer's Reference
4459 * Manual describes the standard name mangling scheme.
4460 *
4461 * @param[in] address Address where the variable is defined. This address must
4462 * have been previously allocated using ::hsa_memory_allocate in a global region
4463 * that is only visible to @p agent. The application cannot deallocate the
4464 * buffer pointed by @p address before @p executable is destroyed.
4465 *
4466 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
4467 *
4468 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
4469 * initialized.
4470 *
4471 * @retval ::HSA_STATUS_ERROR_OUT_OF_RESOURCES The HSA runtime failed to
4472 * allocate the required resources.
4473 *
4474 * @retval ::HSA_STATUS_ERROR_INVALID_EXECUTABLE The executable is invalid.
4475 *
4476 * @retval ::HSA_STATUS_ERROR_INVALID_AGENT @p agent is invalid.
4477 *
4478 * @retval ::HSA_STATUS_ERROR_VARIABLE_ALREADY_DEFINED The variable is
4479 * already defined.
4480 *
4481 * @retval ::HSA_STATUS_ERROR_INVALID_SYMBOL_NAME There is no variable with the
4482 * @p variable_name.
4483 *
4484 * @retval ::HSA_STATUS_ERROR_FROZEN_EXECUTABLE @p executable is frozen.
4485 *
4486 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p variable_name is NULL.
4487 */
4488 hsa_status_t HSA_API hsa_executable_agent_global_variable_define(
4489 hsa_executable_t executable,
4490 hsa_agent_t agent,
4491 const char *variable_name,
4492 void *address);
4493
4494 /**
4495 * @brief Define an external readonly variable.
4496 *
4497 * @details This function allows the application to provide the definition
4498 * of a variable in the readonly segment memory. The variable must be defined
4499 * before loading a code object into an executable. In addition, code objects
4500 * loaded must not define the variable.
4501 *
4502 * @param[in] executable Executable. Must not be in frozen state.
4503 *
4504 * @param[in] agent Agent for which the variable is being defined.
4505 *
4506 * @param[in] variable_name Name of the variable. The Programmer's Reference
4507 * Manual describes the standard name mangling scheme.
4508 *
4509 * @param[in] address Address where the variable is defined. This address must
4510 * have been previously allocated using ::hsa_memory_allocate in a readonly
4511 * region associated with @p agent. The application cannot deallocate the buffer
4512 * pointed by @p address before @p executable is destroyed.
4513 *
4514 * @param[in] address Address where the variable is defined. The buffer pointed
4515 * by @p address is owned by the application, and cannot be deallocated before
4516 * @p executable is destroyed.
4517 *
4518 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
4519 *
4520 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
4521 * initialized.
4522 *
4523 * @retval ::HSA_STATUS_ERROR_OUT_OF_RESOURCES The HSA runtime failed to
4524 * allocate the required resources.
4525 *
4526 * @retval ::HSA_STATUS_ERROR_INVALID_EXECUTABLE Executable is invalid.
4527 *
4528 * @retval ::HSA_STATUS_ERROR_INVALID_AGENT @p agent is invalid.
4529 *
4530 * @retval ::HSA_STATUS_ERROR_VARIABLE_ALREADY_DEFINED The variable is
4531 * already defined.
4532 *
4533 * @retval ::HSA_STATUS_ERROR_INVALID_SYMBOL_NAME There is no variable with the
4534 * @p variable_name.
4535 *
4536 * @retval ::HSA_STATUS_ERROR_FROZEN_EXECUTABLE @p executable is frozen.
4537 *
4538 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p variable_name is NULL.
4539 */
4540 hsa_status_t HSA_API hsa_executable_readonly_variable_define(
4541 hsa_executable_t executable,
4542 hsa_agent_t agent,
4543 const char *variable_name,
4544 void *address);
4545
4546 /**
4547 * @brief Validate an executable. Checks that all code objects have matching
4548 * machine model, profile, and default floating-point rounding mode. Checks that
4549 * all declarations have definitions. Checks declaration-definition
4550 * compatibility (see the HSA Programming Reference Manual for compatibility
4551 * rules). Invoking this function is equivalent to invoking
4552 * ::hsa_executable_validate_alt with no options.
4553 *
4554 * @param[in] executable Executable. Must be in frozen state.
4555 *
4556 * @param[out] result Memory location where the HSA runtime stores the
4557 * validation result. If the executable passes validation, the result is 0.
4558 *
4559 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
4560 *
4561 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
4562 * initialized.
4563 *
4564 * @retval ::HSA_STATUS_ERROR_INVALID_EXECUTABLE @p executable is invalid.
4565 *
4566 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p result is NULL.
4567 */
4568 hsa_status_t HSA_API hsa_executable_validate(
4569 hsa_executable_t executable,
4570 uint32_t *result);
4571
4572 /**
4573 * @brief Validate an executable. Checks that all code objects have matching
4574 * machine model, profile, and default floating-point rounding mode. Checks that
4575 * all declarations have definitions. Checks declaration-definition
4576 * compatibility (see the HSA Programming Reference Manual for compatibility
4577 * rules).
4578 *
4579 * @param[in] executable Executable. Must be in frozen state.
4580 *
4581 * @param[in] options Standard and vendor-specific options. Unknown options are
4582 * ignored. A standard option begins with the "-hsa_" prefix. Options beginning
4583 * with the "-hsa_ext_<extension_name>_" prefix are reserved for extensions. A
4584 * vendor-specific option begins with the "-<vendor_name>_" prefix. Must be a
4585 * NUL-terminated string. May be NULL.
4586 *
4587 * @param[out] result Memory location where the HSA runtime stores the
4588 * validation result. If the executable passes validation, the result is 0.
4589 *
4590 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
4591 *
4592 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
4593 * initialized.
4594 *
4595 * @retval ::HSA_STATUS_ERROR_INVALID_EXECUTABLE @p executable is invalid.
4596 *
4597 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p result is NULL.
4598 */
4599 hsa_status_t HSA_API hsa_executable_validate_alt(
4600 hsa_executable_t executable,
4601 const char *options,
4602 uint32_t *result);
4603
4604 /**
4605 * @brief Executable symbol handle.
4606 *
4607 * The lifetime of an executable object symbol matches that of the executable
4608 * associated with it. An operation on a symbol whose associated executable has
4609 * been destroyed results in undefined behavior.
4610 */
4611 typedef struct hsa_executable_symbol_s {
4612 /**
4613 * Opaque handle. Two handles reference the same object of the enclosing type
4614 * if and only if they are equal.
4615 */
4616 uint64_t handle;
4617 } hsa_executable_symbol_t;
4618
4619 /**
4620 * @deprecated Use ::hsa_executable_get_symbol_by_name instead.
4621 *
4622 * @brief Get the symbol handle for a given a symbol name.
4623 *
4624 * @param[in] executable Executable.
4625 *
4626 * @param[in] module_name Module name. Must be NULL if the symbol has
4627 * program linkage.
4628 *
4629 * @param[in] symbol_name Symbol name.
4630 *
4631 * @param[in] agent Agent associated with the symbol. If the symbol is
4632 * independent of any agent (for example, a variable with program
4633 * allocation), this argument is ignored.
4634 *
4635 * @param[in] call_convention Call convention associated with the symbol. If the
4636 * symbol does not correspond to an indirect function, this argument is ignored.
4637 *
4638 * @param[out] symbol Memory location where the HSA runtime stores the symbol
4639 * handle.
4640 *
4641 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
4642 *
4643 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
4644 * initialized.
4645 *
4646 * @retval ::HSA_STATUS_ERROR_INVALID_EXECUTABLE The executable is invalid.
4647 *
4648 * @retval ::HSA_STATUS_ERROR_INVALID_SYMBOL_NAME There is no symbol with a name
4649 * that matches @p symbol_name.
4650 *
4651 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p symbol_name is NULL, or
4652 * @p symbol is NULL.
4653 */
4654 hsa_status_t HSA_API HSA_DEPRECATED hsa_executable_get_symbol(
4655 hsa_executable_t executable,
4656 const char *module_name,
4657 const char *symbol_name,
4658 hsa_agent_t agent,
4659 int32_t call_convention,
4660 hsa_executable_symbol_t *symbol);
4661
4662 /**
4663 * @brief Retrieve the symbol handle corresponding to a given a symbol name.
4664 *
4665 * @param[in] executable Executable.
4666 *
4667 * @param[in] symbol_name Symbol name. Must be a NUL-terminated character
4668 * array. The Programmer's Reference Manual describes the standard name mangling
4669 * scheme.
4670 *
4671 * @param[in] agent Pointer to the agent for which the symbol with the given
4672 * name is defined. If the symbol corresponding to the given name has program
4673 * allocation, @p agent must be NULL.
4674 *
4675 * @param[out] symbol Memory location where the HSA runtime stores the symbol
4676 * handle. Must not be NULL.
4677 *
4678 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
4679 *
4680 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
4681 * initialized.
4682 *
4683 * @retval ::HSA_STATUS_ERROR_INVALID_EXECUTABLE The executable is invalid.
4684 *
4685 * @retval ::HSA_STATUS_ERROR_INVALID_SYMBOL_NAME There is no symbol with a name
4686 * that matches @p symbol_name.
4687 *
4688 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p symbol_name is NULL, or @p
4689 * symbol is NULL.
4690 */
4691 hsa_status_t HSA_API hsa_executable_get_symbol_by_name(
4692 hsa_executable_t executable,
4693 const char *symbol_name,
4694 const hsa_agent_t *agent,
4695 hsa_executable_symbol_t *symbol);
4696
4697 /**
4698 * @brief Symbol type.
4699 */
4700 typedef enum {
4701 /**
4702 * Variable.
4703 */
4704 HSA_SYMBOL_KIND_VARIABLE = 0,
4705 /**
4706 * Kernel.
4707 */
4708 HSA_SYMBOL_KIND_KERNEL = 1,
4709 /**
4710 * Indirect function.
4711 */
4712 HSA_SYMBOL_KIND_INDIRECT_FUNCTION = 2
4713 } hsa_symbol_kind_t;
4714
4715 /**
4716 * @brief Linkage type of a symbol.
4717 */
4718 typedef enum {
4719 /**
4720 * Module linkage.
4721 */
4722 HSA_SYMBOL_LINKAGE_MODULE = 0,
4723 /**
4724 * Program linkage.
4725 */
4726 HSA_SYMBOL_LINKAGE_PROGRAM = 1
4727 } hsa_symbol_linkage_t;
4728
4729 /**
4730 * @brief Allocation type of a variable.
4731 */
4732 typedef enum {
4733 /**
4734 * Agent allocation.
4735 */
4736 HSA_VARIABLE_ALLOCATION_AGENT = 0,
4737 /**
4738 * Program allocation.
4739 */
4740 HSA_VARIABLE_ALLOCATION_PROGRAM = 1
4741 } hsa_variable_allocation_t;
4742
4743 /**
4744 * @brief Memory segment associated with a variable.
4745 */
4746 typedef enum {
4747 /**
4748 * Global memory segment.
4749 */
4750 HSA_VARIABLE_SEGMENT_GLOBAL = 0,
4751 /**
4752 * Readonly memory segment.
4753 */
4754 HSA_VARIABLE_SEGMENT_READONLY = 1
4755 } hsa_variable_segment_t;
4756
4757 /**
4758 * @brief Executable symbol attributes.
4759 */
4760 typedef enum {
4761 /**
4762 * The kind of the symbol. The type of this attribute is ::hsa_symbol_kind_t.
4763 */
4764 HSA_EXECUTABLE_SYMBOL_INFO_TYPE = 0,
4765 /**
4766 * The length of the symbol name in bytes, not including the NUL terminator.
4767 * The type of this attribute is uint32_t.
4768 */
4769 HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH = 1,
4770 /**
4771 * The name of the symbol. The type of this attribute is character array with
4772 * the length equal to the value of ::HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH
4773 * attribute.
4774 */
4775 HSA_EXECUTABLE_SYMBOL_INFO_NAME = 2,
4776 /**
4777 * @deprecated
4778 *
4779 * The length of the module name in bytes (not including the NUL terminator)
4780 * to which this symbol belongs if this symbol has module linkage, otherwise 0
4781 * is returned. The type of this attribute is uint32_t.
4782 */
4783 HSA_EXECUTABLE_SYMBOL_INFO_MODULE_NAME_LENGTH = 3,
4784 /**
4785 * @deprecated
4786 *
4787 * The module name to which this symbol belongs if this symbol has module
4788 * linkage, otherwise an empty string is returned. The type of this attribute
4789 * is character array with the length equal to the value of
4790 * ::HSA_EXECUTABLE_SYMBOL_INFO_MODULE_NAME_LENGTH attribute.
4791 */
4792 HSA_EXECUTABLE_SYMBOL_INFO_MODULE_NAME = 4,
4793 /**
4794 * @deprecated
4795 *
4796 * Agent associated with this symbol. If the symbol is a variable, the
4797 * value of this attribute is only defined if
4798 * ::HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ALLOCATION is
4799 * ::HSA_VARIABLE_ALLOCATION_AGENT. The type of this attribute is hsa_agent_t.
4800 */
4801 HSA_EXECUTABLE_SYMBOL_INFO_AGENT = 20,
4802 /**
4803 * The address of the variable. The value of this attribute is undefined if
4804 * the symbol is not a variable. The type of this attribute is uint64_t.
4805 *
4806 * If executable's state is ::HSA_EXECUTABLE_STATE_UNFROZEN, then 0 is
4807 * returned.
4808 */
4809 HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS = 21,
4810 /**
4811 * The linkage kind of the symbol. The type of this attribute is
4812 * ::hsa_symbol_linkage_t.
4813 */
4814 HSA_EXECUTABLE_SYMBOL_INFO_LINKAGE = 5,
4815 /**
4816 * Indicates whether the symbol corresponds to a definition. The type of this
4817 * attribute is bool.
4818 */
4819 HSA_EXECUTABLE_SYMBOL_INFO_IS_DEFINITION = 17,
4820 /**
4821 * @deprecated
4822 *
4823 * The allocation kind of the variable. The value of this attribute is
4824 * undefined if the symbol is not a variable. The type of this attribute is
4825 * ::hsa_variable_allocation_t.
4826 */
4827 HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ALLOCATION = 6,
4828 /**
4829 * @deprecated
4830 *
4831 * The segment kind of the variable. The value of this attribute is undefined
4832 * if the symbol is not a variable. The type of this attribute is
4833 * ::hsa_variable_segment_t.
4834 */
4835 HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SEGMENT = 7,
4836 /**
4837 * @deprecated
4838 *
4839 * Alignment of the symbol in memory. The value of this attribute is undefined
4840 * if the symbol is not a variable. The type of this attribute is uint32_t.
4841 *
4842 * The current alignment of the variable in memory may be greater than the
4843 * value specified in the source program variable declaration.
4844 */
4845 HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ALIGNMENT = 8,
4846 /**
4847 * @deprecated
4848 *
4849 * Size of the variable. The value of this attribute is undefined if
4850 * the symbol is not a variable. The type of this attribute is uint32_t.
4851 *
4852 * A value of 0 is returned if the variable is an external variable and has an
4853 * unknown dimension.
4854 */
4855 HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE = 9,
4856 /**
4857 * @deprecated
4858 *
4859 * Indicates whether the variable is constant. The value of this attribute is
4860 * undefined if the symbol is not a variable. The type of this attribute is
4861 * bool.
4862 */
4863 HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_IS_CONST = 10,
4864 /**
4865 * Kernel object handle, used in the kernel dispatch packet. The value of this
4866 * attribute is undefined if the symbol is not a kernel. The type of this
4867 * attribute is uint64_t.
4868 *
4869 * If the state of the executable is ::HSA_EXECUTABLE_STATE_UNFROZEN, then 0
4870 * is returned.
4871 */
4872 HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT = 22,
4873 /**
4874 * Size of kernarg segment memory that is required to hold the values of the
4875 * kernel arguments, in bytes. Must be a multiple of 16. The value of this
4876 * attribute is undefined if the symbol is not a kernel. The type of this
4877 * attribute is uint32_t.
4878 */
4879 HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE = 11,
4880 /**
4881 * Alignment (in bytes) of the buffer used to pass arguments to the kernel,
4882 * which is the maximum of 16 and the maximum alignment of any of the kernel
4883 * arguments. The value of this attribute is undefined if the symbol is not a
4884 * kernel. The type of this attribute is uint32_t.
4885 */
4886 HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_ALIGNMENT = 12,
4887 /**
4888 * Size of static group segment memory required by the kernel (per
4889 * work-group), in bytes. The value of this attribute is undefined
4890 * if the symbol is not a kernel. The type of this attribute is uint32_t.
4891 *
4892 * The reported amount does not include any dynamically allocated group
4893 * segment memory that may be requested by the application when a kernel is
4894 * dispatched.
4895 */
4896 HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE = 13,
4897 /**
4898 * Size of static private, spill, and arg segment memory required by
4899 * this kernel (per work-item), in bytes. The value of this attribute is
4900 * undefined if the symbol is not a kernel. The type of this attribute is
4901 * uint32_t.
4902 *
4903 * If the value of ::HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_DYNAMIC_CALLSTACK is
4904 * true, the kernel may use more private memory than the reported value, and
4905 * the application must add the dynamic call stack usage to @a
4906 * private_segment_size when populating a kernel dispatch packet.
4907 */
4908 HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE = 14,
4909 /**
4910 * Dynamic callstack flag. The value of this attribute is undefined if the
4911 * symbol is not a kernel. The type of this attribute is bool.
4912 *
4913 * If this flag is set (the value is true), the kernel uses a dynamically
4914 * sized call stack. This can happen if recursive calls, calls to indirect
4915 * functions, or the HSAIL alloca instruction are present in the kernel.
4916 */
4917 HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_DYNAMIC_CALLSTACK = 15,
4918 /**
4919 * @deprecated
4920 *
4921 * Call convention of the kernel. The value of this attribute is undefined if
4922 * the symbol is not a kernel. The type of this attribute is uint32_t.
4923 */
4924 HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_CALL_CONVENTION = 18,
4925 /**
4926 * Indirect function object handle. The value of this attribute is undefined
4927 * if the symbol is not an indirect function, or the associated agent does
4928 * not support the Full Profile. The type of this attribute depends on the
4929 * machine model: the type is uint32_t for small machine model, and uint64_t
4930 * for large model.
4931 *
4932 * If the state of the executable is ::HSA_EXECUTABLE_STATE_UNFROZEN, then 0
4933 * is returned.
4934 */
4935 HSA_EXECUTABLE_SYMBOL_INFO_INDIRECT_FUNCTION_OBJECT = 23,
4936 /**
4937 * @deprecated
4938 *
4939 * Call convention of the indirect function. The value of this attribute is
4940 * undefined if the symbol is not an indirect function, or the associated
4941 * agent does not support the Full Profile. The type of this attribute is
4942 * uint32_t.
4943 */
4944 HSA_EXECUTABLE_SYMBOL_INFO_INDIRECT_FUNCTION_CALL_CONVENTION = 16
4945 } hsa_executable_symbol_info_t;
4946
4947 /**
4948 * @brief Get the current value of an attribute for a given executable symbol.
4949 *
4950 * @param[in] executable_symbol Executable symbol.
4951 *
4952 * @param[in] attribute Attribute to query.
4953 *
4954 * @param[out] value Pointer to an application-allocated buffer where to store
4955 * the value of the attribute. If the buffer passed by the application is not
4956 * large enough to hold the value of @p attribute, the behavior is undefined.
4957 *
4958 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
4959 *
4960 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
4961 * initialized.
4962 *
4963 * @retval ::HSA_STATUS_ERROR_INVALID_EXECUTABLE_SYMBOL The executable symbol is
4964 * invalid.
4965 *
4966 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p attribute is an invalid
4967 * executable symbol attribute, or @p value is NULL.
4968 */
4969 hsa_status_t HSA_API hsa_executable_symbol_get_info(
4970 hsa_executable_symbol_t executable_symbol,
4971 hsa_executable_symbol_info_t attribute,
4972 void *value);
4973
4974 /**
4975 * @deprecated
4976 *
4977 * @brief Iterate over the symbols in a executable, and invoke an
4978 * application-defined callback on every iteration.
4979 *
4980 * @param[in] executable Executable.
4981 *
4982 * @param[in] callback Callback to be invoked once per executable symbol. The
4983 * HSA runtime passes three arguments to the callback: the executable, a symbol,
4984 * and the application data. If @p callback returns a status other than
4985 * ::HSA_STATUS_SUCCESS for a particular iteration, the traversal stops and
4986 * ::hsa_executable_iterate_symbols returns that status value.
4987 *
4988 * @param[in] data Application data that is passed to @p callback on every
4989 * iteration. May be NULL.
4990 *
4991 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
4992 *
4993 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
4994 * initialized.
4995 *
4996 * @retval ::HSA_STATUS_ERROR_INVALID_EXECUTABLE The executable is invalid.
4997 *
4998 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p callback is NULL.
4999 */
5000 hsa_status_t HSA_API HSA_DEPRECATED hsa_executable_iterate_symbols(
5001 hsa_executable_t executable,
5002 hsa_status_t (*callback)(hsa_executable_t exec,
5003 hsa_executable_symbol_t symbol,
5004 void *data),
5005 void *data);
5006
5007 /**
5008 * @brief Iterate over the kernels, indirect functions, and agent allocation
5009 * variables in an executable for a given agent, and invoke an application-
5010 * defined callback on every iteration.
5011 *
5012 * @param[in] executable Executable.
5013 *
5014 * @param[in] agent Agent.
5015 *
5016 * @param[in] callback Callback to be invoked once per executable symbol. The
5017 * HSA runtime passes three arguments to the callback: the executable, a symbol,
5018 * and the application data. If @p callback returns a status other than
5019 * ::HSA_STATUS_SUCCESS for a particular iteration, the traversal stops and
5020 * ::hsa_executable_iterate_symbols returns that status value.
5021 *
5022 * @param[in] data Application data that is passed to @p callback on every
5023 * iteration. May be NULL.
5024 *
5025 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
5026 *
5027 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
5028 * initialized.
5029 *
5030 * @retval ::HSA_STATUS_ERROR_INVALID_EXECUTABLE The executable is invalid.
5031 *
5032 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p callback is NULL.
5033 */
5034 hsa_status_t HSA_API hsa_executable_iterate_agent_symbols(
5035 hsa_executable_t executable,
5036 hsa_agent_t agent,
5037 hsa_status_t (*callback)(hsa_executable_t exec,
5038 hsa_agent_t agent,
5039 hsa_executable_symbol_t symbol,
5040 void *data),
5041 void *data);
5042
5043 /**
5044 * @brief Iterate over the program allocation variables in an executable, and
5045 * invoke an application-defined callback on every iteration.
5046 *
5047 * @param[in] executable Executable.
5048 *
5049 * @param[in] callback Callback to be invoked once per executable symbol. The
5050 * HSA runtime passes three arguments to the callback: the executable, a symbol,
5051 * and the application data. If @p callback returns a status other than
5052 * ::HSA_STATUS_SUCCESS for a particular iteration, the traversal stops and
5053 * ::hsa_executable_iterate_symbols returns that status value.
5054 *
5055 * @param[in] data Application data that is passed to @p callback on every
5056 * iteration. May be NULL.
5057 *
5058 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
5059 *
5060 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
5061 * initialized.
5062 *
5063 * @retval ::HSA_STATUS_ERROR_INVALID_EXECUTABLE The executable is invalid.
5064 *
5065 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p callback is NULL.
5066 */
5067 hsa_status_t HSA_API hsa_executable_iterate_program_symbols(
5068 hsa_executable_t executable,
5069 hsa_status_t (*callback)(hsa_executable_t exec,
5070 hsa_executable_symbol_t symbol,
5071 void *data),
5072 void *data);
5073
5074 /** @} */
5075
5076
5077 /** \defgroup code-object Code Objects (deprecated).
5078 * @{
5079 */
5080
5081 /**
5082 * @deprecated
5083 *
5084 * @brief Struct containing an opaque handle to a code object, which contains
5085 * ISA for finalized kernels and indirect functions together with information
5086 * about the global or readonly segment variables they reference.
5087 */
5088 typedef struct hsa_code_object_s {
5089 /**
5090 * Opaque handle. Two handles reference the same object of the enclosing type
5091 * if and only if they are equal.
5092 */
5093 uint64_t handle;
5094 } hsa_code_object_t;
5095
5096 /**
5097 * @deprecated
5098 *
5099 * @brief Application data handle that is passed to the serialization
5100 * and deserialization functions.
5101 */
5102 typedef struct hsa_callback_data_s {
5103 /**
5104 * Opaque handle.
5105 */
5106 uint64_t handle;
5107 } hsa_callback_data_t;
5108
5109 /**
5110 * @deprecated
5111 *
5112 * @brief Serialize a code object. Can be used for offline finalization,
5113 * install-time finalization, disk code caching, etc.
5114 *
5115 * @param[in] code_object Code object.
5116 *
5117 * @param[in] alloc_callback Callback function for memory allocation. Must not
5118 * be NULL. The HSA runtime passes three arguments to the callback: the
5119 * allocation size, the application data, and a pointer to a memory location
5120 * where the application stores the allocation result. The HSA runtime invokes
5121 * @p alloc_callback once to allocate a buffer that contains the serialized
5122 * version of @p code_object. If the callback returns a status code other than
5123 * ::HSA_STATUS_SUCCESS, this function returns the same code.
5124 *
5125 * @param[in] callback_data Application data that is passed to @p
5126 * alloc_callback. May be NULL.
5127 *
5128 * @param[in] options Standard and vendor-specific options. Unknown options are
5129 * ignored. A standard option begins with the "-hsa_" prefix. Options beginning
5130 * with the "-hsa_ext_<extension_name>_" prefix are reserved for extensions. A
5131 * vendor-specific option begins with the "-<vendor_name>_" prefix. Must be a
5132 * NUL-terminated string. May be NULL.
5133 *
5134 * @param[out] serialized_code_object Memory location where the HSA runtime
5135 * stores a pointer to the serialized code object. Must not be NULL.
5136 *
5137 * @param[out] serialized_code_object_size Memory location where the HSA runtime
5138 * stores the size (in bytes) of @p serialized_code_object. The returned value
5139 * matches the allocation size passed by the HSA runtime to @p
5140 * alloc_callback. Must not be NULL.
5141 *
5142 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
5143 *
5144 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
5145 * initialized.
5146 *
5147 * @retval ::HSA_STATUS_ERROR_OUT_OF_RESOURCES The HSA runtime failed to
5148 * allocate the required resources.
5149 *
5150 * @retval ::HSA_STATUS_ERROR_INVALID_CODE_OBJECT @p code_object is invalid.
5151 *
5152 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p alloc_callback, @p
5153 * serialized_code_object, or @p serialized_code_object_size are NULL.
5154 */
5155 hsa_status_t HSA_API HSA_DEPRECATED hsa_code_object_serialize(
5156 hsa_code_object_t code_object,
5157 hsa_status_t (*alloc_callback)(size_t size,
5158 hsa_callback_data_t data,
5159 void **address),
5160 hsa_callback_data_t callback_data,
5161 const char *options,
5162 void **serialized_code_object,
5163 size_t *serialized_code_object_size);
5164
5165 /**
5166 * @deprecated
5167 *
5168 * @brief Deserialize a code object.
5169 *
5170 * @param[in] serialized_code_object A serialized code object. Must not be NULL.
5171 *
5172 * @param[in] serialized_code_object_size The size (in bytes) of @p
5173 * serialized_code_object. Must not be 0.
5174 *
5175 * @param[in] options Standard and vendor-specific options. Unknown options are
5176 * ignored. A standard option begins with the "-hsa_" prefix. Options beginning
5177 * with the "-hsa_ext_<extension_name>_" prefix are reserved for extensions. A
5178 * vendor-specific option begins with the "-<vendor_name>_" prefix. Must be a
5179 * NUL-terminated string. May be NULL.
5180 *
5181 * @param[out] code_object Memory location where the HSA runtime stores the
5182 * deserialized code object.
5183 *
5184 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
5185 *
5186 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
5187 * initialized.
5188 *
5189 * @retval ::HSA_STATUS_ERROR_OUT_OF_RESOURCES The HSA runtime failed to
5190 * allocate the required resources.
5191 *
5192 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p serialized_code_object, or @p
5193 * code_object are NULL, or @p serialized_code_object_size is 0.
5194 */
5195 hsa_status_t HSA_API HSA_DEPRECATED hsa_code_object_deserialize(
5196 void *serialized_code_object,
5197 size_t serialized_code_object_size,
5198 const char *options,
5199 hsa_code_object_t *code_object);
5200
5201 /**
5202 * @deprecated
5203 *
5204 * @brief Destroy a code object.
5205 *
5206 * @details The lifetime of a code object must exceed that of any executable
5207 * where it has been loaded. If an executable that loaded @p code_object has not
5208 * been destroyed, the behavior is undefined.
5209 *
5210 * @param[in] code_object Code object. The handle becomes invalid after it has
5211 * been destroyed.
5212 *
5213 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
5214 *
5215 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
5216 * initialized.
5217 *
5218 * @retval ::HSA_STATUS_ERROR_INVALID_CODE_OBJECT @p code_object is invalid.
5219 */
5220 hsa_status_t HSA_API HSA_DEPRECATED hsa_code_object_destroy(
5221 hsa_code_object_t code_object);
5222
5223 /**
5224 * @deprecated
5225 *
5226 * @brief Code object type.
5227 */
5228 typedef enum {
5229 /**
5230 * Produces code object that contains ISA for all kernels and indirect
5231 * functions in HSA source.
5232 */
5233 HSA_CODE_OBJECT_TYPE_PROGRAM = 0
5234 } hsa_code_object_type_t;
5235
5236 /**
5237 * @deprecated
5238 *
5239 * @brief Code object attributes.
5240 */
5241 typedef enum {
5242 /**
5243 * The version of the code object. The type of this attribute is a
5244 * NUL-terminated char[64]. The name must be at most 63 characters long (not
5245 * including the NUL terminator) and all array elements not used for the name
5246 * must be NUL.
5247 */
5248 HSA_CODE_OBJECT_INFO_VERSION = 0,
5249 /**
5250 * Type of code object. The type of this attribute is
5251 * ::hsa_code_object_type_t.
5252 */
5253 HSA_CODE_OBJECT_INFO_TYPE = 1,
5254 /**
5255 * Instruction set architecture this code object is produced for. The type of
5256 * this attribute is ::hsa_isa_t.
5257 */
5258 HSA_CODE_OBJECT_INFO_ISA = 2,
5259 /**
5260 * Machine model this code object is produced for. The type of this attribute
5261 * is ::hsa_machine_model_t.
5262 */
5263 HSA_CODE_OBJECT_INFO_MACHINE_MODEL = 3,
5264 /**
5265 * Profile this code object is produced for. The type of this attribute is
5266 * ::hsa_profile_t.
5267 */
5268 HSA_CODE_OBJECT_INFO_PROFILE = 4,
5269 /**
5270 * Default floating-point rounding mode used when the code object is
5271 * produced. The type of this attribute is
5272 * ::hsa_default_float_rounding_mode_t.
5273 */
5274 HSA_CODE_OBJECT_INFO_DEFAULT_FLOAT_ROUNDING_MODE = 5
5275 } hsa_code_object_info_t;
5276
5277 /**
5278 * @deprecated
5279 *
5280 * @brief Get the current value of an attribute for a given code object.
5281 *
5282 * @param[in] code_object Code object.
5283 *
5284 * @param[in] attribute Attribute to query.
5285 *
5286 * @param[out] value Pointer to an application-allocated buffer where to store
5287 * the value of the attribute. If the buffer passed by the application is not
5288 * large enough to hold the value of @p attribute, the behavior is undefined.
5289 *
5290 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
5291 *
5292 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
5293 * initialized.
5294 *
5295 * @retval ::HSA_STATUS_ERROR_INVALID_CODE_OBJECT @p code_object is invalid.
5296 *
5297 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p attribute is an invalid
5298 * code object attribute, or @p value is NULL.
5299 */
5300 hsa_status_t HSA_API HSA_DEPRECATED hsa_code_object_get_info(
5301 hsa_code_object_t code_object,
5302 hsa_code_object_info_t attribute,
5303 void *value);
5304
5305 /**
5306 * @deprecated
5307 *
5308 * @brief Load code object into the executable.
5309 *
5310 * @details Every global or readonly variable that is external must be defined
5311 * before loading the code object. An internal global or readonly variable is
5312 * allocated once the code object, that is being loaded, references this
5313 * variable and this variable is not allocated.
5314 *
5315 * Any module linkage declaration must have been defined either by a define
5316 * variable or by loading a code object that has a symbol with module linkage
5317 * definition.
5318 *
5319 * @param[in] executable Executable.
5320 *
5321 * @param[in] agent Agent to load code object for. The agent must support the
5322 * default floating-point rounding mode used by @p code_object.
5323 *
5324 * @param[in] code_object Code object to load. The lifetime of the code object
5325 * must exceed that of the executable: if @p code_object is destroyed before @p
5326 * executable, the behavior is undefined.
5327 *
5328 * @param[in] options Standard and vendor-specific options. Unknown options are
5329 * ignored. A standard option begins with the "-hsa_" prefix. Options beginning
5330 * with the "-hsa_ext_<extension_name>_" prefix are reserved for extensions. A
5331 * vendor-specific option begins with the "-<vendor_name>_" prefix. Must be a
5332 * NUL-terminated string. May be NULL.
5333 *
5334 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
5335 *
5336 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
5337 * initialized.
5338 *
5339 * @retval ::HSA_STATUS_ERROR_OUT_OF_RESOURCES The HSA runtime failed to
5340 * allocate the required resources.
5341 *
5342 * @retval ::HSA_STATUS_ERROR_INVALID_EXECUTABLE The executable is invalid.
5343 *
5344 * @retval ::HSA_STATUS_ERROR_INVALID_AGENT The agent is invalid.
5345 *
5346 * @retval ::HSA_STATUS_ERROR_INVALID_CODE_OBJECT @p code_object is invalid.
5347 *
5348 * @retval ::HSA_STATUS_ERROR_INCOMPATIBLE_ARGUMENTS @p agent is not compatible
5349 * with @p code_object (for example, @p agent does not support the default
5350 * floating-point rounding mode specified by @p code_object), or @p code_object
5351 * is not compatible with @p executable (for example, @p code_object and @p
5352 * executable have different machine models or profiles).
5353 *
5354 * @retval ::HSA_STATUS_ERROR_FROZEN_EXECUTABLE @p executable is frozen.
5355 */
5356 hsa_status_t HSA_API HSA_DEPRECATED hsa_executable_load_code_object(
5357 hsa_executable_t executable,
5358 hsa_agent_t agent,
5359 hsa_code_object_t code_object,
5360 const char *options);
5361
5362 /**
5363 * @deprecated
5364 *
5365 * @brief Code object symbol handle.
5366 *
5367 * The lifetime of a code object symbol matches that of the code object
5368 * associated with it. An operation on a symbol whose associated code object has
5369 * been destroyed results in undefined behavior.
5370 */
5371 typedef struct hsa_code_symbol_s {
5372 /**
5373 * Opaque handle. Two handles reference the same object of the enclosing type
5374 * if and only if they are equal.
5375 */
5376 uint64_t handle;
5377 } hsa_code_symbol_t;
5378
5379 /**
5380 * @deprecated
5381 *
5382 * @brief Get the symbol handle within a code object for a given a symbol name.
5383 *
5384 * @param[in] code_object Code object.
5385 *
5386 * @param[in] symbol_name Symbol name.
5387 *
5388 * @param[out] symbol Memory location where the HSA runtime stores the symbol
5389 * handle.
5390 *
5391 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
5392 *
5393 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
5394 * initialized.
5395 *
5396 * @retval ::HSA_STATUS_ERROR_INVALID_CODE_OBJECT @p code_object is invalid.
5397 *
5398 * @retval ::HSA_STATUS_ERROR_INVALID_SYMBOL_NAME There is no symbol with a name
5399 * that matches @p symbol_name.
5400 *
5401 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p symbol_name is NULL, or
5402 * @p symbol is NULL.
5403 */
5404 hsa_status_t HSA_API HSA_DEPRECATED hsa_code_object_get_symbol(
5405 hsa_code_object_t code_object,
5406 const char *symbol_name,
5407 hsa_code_symbol_t *symbol);
5408
5409 /**
5410 * @deprecated
5411 *
5412 * @brief Get the symbol handle within a code object for a given a symbol name.
5413 *
5414 * @param[in] code_object Code object.
5415 *
5416 * @param[in] module_name Module name. Must be NULL if the symbol has
5417 * program linkage.
5418 *
5419 * @param[in] symbol_name Symbol name.
5420 *
5421 * @param[out] symbol Memory location where the HSA runtime stores the symbol
5422 * handle.
5423 *
5424 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
5425 *
5426 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
5427 * initialized.
5428 *
5429 * @retval ::HSA_STATUS_ERROR_INVALID_CODE_OBJECT @p code_object is invalid.
5430 *
5431 * @retval ::HSA_STATUS_ERROR_INVALID_SYMBOL_NAME There is no symbol with a name
5432 * that matches @p symbol_name.
5433 *
5434 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p symbol_name is NULL, or
5435 * @p symbol is NULL.
5436 */
5437 hsa_status_t HSA_API HSA_DEPRECATED hsa_code_object_get_symbol_from_name(
5438 hsa_code_object_t code_object,
5439 const char *module_name,
5440 const char *symbol_name,
5441 hsa_code_symbol_t *symbol);
5442
5443 /**
5444 * @deprecated
5445 *
5446 * @brief Code object symbol attributes.
5447 */
5448 typedef enum {
5449 /**
5450 * The type of the symbol. The type of this attribute is ::hsa_symbol_kind_t.
5451 */
5452 HSA_CODE_SYMBOL_INFO_TYPE = 0,
5453 /**
5454 * The length of the symbol name in bytes, not including the NUL terminator.
5455 * The type of this attribute is uint32_t.
5456 */
5457 HSA_CODE_SYMBOL_INFO_NAME_LENGTH = 1,
5458 /**
5459 * The name of the symbol. The type of this attribute is character array with
5460 * the length equal to the value of ::HSA_CODE_SYMBOL_INFO_NAME_LENGTH
5461 * attribute.
5462 */
5463 HSA_CODE_SYMBOL_INFO_NAME = 2,
5464 /**
5465 * The length of the module name in bytes (not including the NUL terminator)
5466 * to which this symbol belongs if this symbol has module linkage, otherwise 0
5467 * is returned. The type of this attribute is uint32_t.
5468 */
5469 HSA_CODE_SYMBOL_INFO_MODULE_NAME_LENGTH = 3,
5470 /**
5471 * The module name to which this symbol belongs if this symbol has module
5472 * linkage, otherwise an empty string is returned. The type of this attribute
5473 * is character array with the length equal to the value of
5474 * ::HSA_CODE_SYMBOL_INFO_MODULE_NAME_LENGTH attribute.
5475 */
5476 HSA_CODE_SYMBOL_INFO_MODULE_NAME = 4,
5477 /**
5478 * The linkage kind of the symbol. The type of this attribute is
5479 * ::hsa_symbol_linkage_t.
5480 */
5481 HSA_CODE_SYMBOL_INFO_LINKAGE = 5,
5482 /**
5483 * Indicates whether the symbol corresponds to a definition. The type of this
5484 * attribute is bool.
5485 */
5486 HSA_CODE_SYMBOL_INFO_IS_DEFINITION = 17,
5487 /**
5488 * The allocation kind of the variable. The value of this attribute is
5489 * undefined if the symbol is not a variable. The type of this attribute is
5490 * ::hsa_variable_allocation_t.
5491 */
5492 HSA_CODE_SYMBOL_INFO_VARIABLE_ALLOCATION = 6,
5493 /**
5494 * The segment kind of the variable. The value of this attribute is
5495 * undefined if the symbol is not a variable. The type of this attribute is
5496 * ::hsa_variable_segment_t.
5497 */
5498 HSA_CODE_SYMBOL_INFO_VARIABLE_SEGMENT = 7,
5499 /**
5500 * Alignment of the symbol in memory. The value of this attribute is undefined
5501 * if the symbol is not a variable. The type of this attribute is uint32_t.
5502 *
5503 * The current alignment of the variable in memory may be greater than the
5504 * value specified in the source program variable declaration.
5505 */
5506 HSA_CODE_SYMBOL_INFO_VARIABLE_ALIGNMENT = 8,
5507 /**
5508 * Size of the variable. The value of this attribute is undefined if the
5509 * symbol is not a variable. The type of this attribute is uint32_t.
5510 *
5511 * A size of 0 is returned if the variable is an external variable and has an
5512 * unknown dimension.
5513 */
5514 HSA_CODE_SYMBOL_INFO_VARIABLE_SIZE = 9,
5515 /**
5516 * Indicates whether the variable is constant. The value of this attribute is
5517 * undefined if the symbol is not a variable. The type of this attribute is
5518 * bool.
5519 */
5520 HSA_CODE_SYMBOL_INFO_VARIABLE_IS_CONST = 10,
5521 /**
5522 * Size of kernarg segment memory that is required to hold the values of the
5523 * kernel arguments, in bytes. Must be a multiple of 16. The value of this
5524 * attribute is undefined if the symbol is not a kernel. The type of this
5525 * attribute is uint32_t.
5526 */
5527 HSA_CODE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE = 11,
5528 /**
5529 * Alignment (in bytes) of the buffer used to pass arguments to the kernel,
5530 * which is the maximum of 16 and the maximum alignment of any of the kernel
5531 * arguments. The value of this attribute is undefined if the symbol is not a
5532 * kernel. The type of this attribute is uint32_t.
5533 */
5534 HSA_CODE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_ALIGNMENT = 12,
5535 /**
5536 * Size of static group segment memory required by the kernel (per
5537 * work-group), in bytes. The value of this attribute is undefined
5538 * if the symbol is not a kernel. The type of this attribute is uint32_t.
5539 *
5540 * The reported amount does not include any dynamically allocated group
5541 * segment memory that may be requested by the application when a kernel is
5542 * dispatched.
5543 */
5544 HSA_CODE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE = 13,
5545 /**
5546 * Size of static private, spill, and arg segment memory required by
5547 * this kernel (per work-item), in bytes. The value of this attribute is
5548 * undefined if the symbol is not a kernel. The type of this attribute is
5549 * uint32_t.
5550 *
5551 * If the value of ::HSA_CODE_SYMBOL_INFO_KERNEL_DYNAMIC_CALLSTACK is true,
5552 * the kernel may use more private memory than the reported value, and the
5553 * application must add the dynamic call stack usage to @a
5554 * private_segment_size when populating a kernel dispatch packet.
5555 */
5556 HSA_CODE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE = 14,
5557 /**
5558 * Dynamic callstack flag. The value of this attribute is undefined if the
5559 * symbol is not a kernel. The type of this attribute is bool.
5560 *
5561 * If this flag is set (the value is true), the kernel uses a dynamically
5562 * sized call stack. This can happen if recursive calls, calls to indirect
5563 * functions, or the HSAIL alloca instruction are present in the kernel.
5564 */
5565 HSA_CODE_SYMBOL_INFO_KERNEL_DYNAMIC_CALLSTACK = 15,
5566 /**
5567 * Call convention of the kernel. The value of this attribute is undefined if
5568 * the symbol is not a kernel. The type of this attribute is uint32_t.
5569 */
5570 HSA_CODE_SYMBOL_INFO_KERNEL_CALL_CONVENTION = 18,
5571 /**
5572 * Call convention of the indirect function. The value of this attribute is
5573 * undefined if the symbol is not an indirect function. The type of this
5574 * attribute is uint32_t.
5575 */
5576 HSA_CODE_SYMBOL_INFO_INDIRECT_FUNCTION_CALL_CONVENTION = 16
5577 } hsa_code_symbol_info_t;
5578
5579 /**
5580 * @deprecated
5581 *
5582 * @brief Get the current value of an attribute for a given code symbol.
5583 *
5584 * @param[in] code_symbol Code symbol.
5585 *
5586 * @param[in] attribute Attribute to query.
5587 *
5588 * @param[out] value Pointer to an application-allocated buffer where to store
5589 * the value of the attribute. If the buffer passed by the application is not
5590 * large enough to hold the value of @p attribute, the behavior is undefined.
5591 *
5592 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
5593 *
5594 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
5595 * initialized.
5596 *
5597 * @retval ::HSA_STATUS_ERROR_INVALID_CODE_SYMBOL The code symbol is invalid.
5598 *
5599 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p attribute is an invalid
5600 * code symbol attribute, or @p value is NULL.
5601 */
5602 hsa_status_t HSA_API HSA_DEPRECATED hsa_code_symbol_get_info(
5603 hsa_code_symbol_t code_symbol,
5604 hsa_code_symbol_info_t attribute,
5605 void *value);
5606
5607 /**
5608 * @deprecated
5609 *
5610 * @brief Iterate over the symbols in a code object, and invoke an
5611 * application-defined callback on every iteration.
5612 *
5613 * @param[in] code_object Code object.
5614 *
5615 * @param[in] callback Callback to be invoked once per code object symbol. The
5616 * HSA runtime passes three arguments to the callback: the code object, a
5617 * symbol, and the application data. If @p callback returns a status other than
5618 * ::HSA_STATUS_SUCCESS for a particular iteration, the traversal stops and
5619 * ::hsa_code_object_iterate_symbols returns that status value.
5620 *
5621 * @param[in] data Application data that is passed to @p callback on every
5622 * iteration. May be NULL.
5623 *
5624 * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
5625 *
5626 * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
5627 * initialized.
5628 *
5629 * @retval ::HSA_STATUS_ERROR_INVALID_CODE_OBJECT @p code_object is invalid.
5630 *
5631 * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p callback is NULL.
5632 */
5633 hsa_status_t HSA_API HSA_DEPRECATED hsa_code_object_iterate_symbols(
5634 hsa_code_object_t code_object,
5635 hsa_status_t (*callback)(hsa_code_object_t code_object,
5636 hsa_code_symbol_t symbol,
5637 void *data),
5638 void *data);
5639
5640 /** @} */
5641
5642 #ifdef __cplusplus
5643 } // end extern "C" block
5644 #endif
5645
5646 #endif // header guard