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 5647