rocprofiler-sdk/fwd.h Source File

rocprofiler-sdk/fwd.h Source File#

Rocprofiler SDK Developer API: rocprofiler-sdk/fwd.h Source File
Rocprofiler SDK Developer API 0.6.0
ROCm Profiling API and tools
fwd.h
Go to the documentation of this file.
1// MIT License
2//
3// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
4//
5// Permission is hereby granted, free of charge, to any person obtaining a copy
6// of this software and associated documentation files (the "Software"), to deal
7// in the Software without restriction, including without limitation the rights
8// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
9// copies of the Software, and to permit persons to whom the Software is
10// furnished to do so, subject to the following conditions:
11//
12// The above copyright notice and this permission notice shall be included in all
13// copies or substantial portions of the Software.
14//
15// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
16// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
17// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
18// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
19// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
20// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
21// SOFTWARE.
22
23#pragma once
24
26
27#include <stddef.h>
28#include <stdint.h>
29
30ROCPROFILER_EXTERN_C_INIT
31
32//--------------------------------------------------------------------------------------//
33//
34// ENUMERATIONS
35//
36//--------------------------------------------------------------------------------------//
37
38/**
39 * @defgroup BASIC_DATA_TYPES Basic data types
40 * @brief Basic data types and typedefs
41 *
42 * @{
43 */
44
45// TODO(aelwazir): Do we need to add a null (way) for every handle?
46// TODO(aelwazir): Remove API Data args from the doxygen?
47// TODO(aelwazir): Not everything in bin needs to be installed bin, use libexec or share?
48
49/**
50 * @brief Status codes.
51 */
52typedef enum // NOLINT(performance-enum-size)
53{
54 ROCPROFILER_STATUS_SUCCESS = 0, ///< No error occurred
55 ROCPROFILER_STATUS_ERROR, ///< Generalized error
56 ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND, ///< No valid context for given context id
57 ROCPROFILER_STATUS_ERROR_BUFFER_NOT_FOUND, ///< No valid buffer for given buffer id
58 ROCPROFILER_STATUS_ERROR_KIND_NOT_FOUND, ///< Kind identifier is invalid
59 ROCPROFILER_STATUS_ERROR_OPERATION_NOT_FOUND, ///< Operation identifier is invalid for domain
60 ROCPROFILER_STATUS_ERROR_THREAD_NOT_FOUND, ///< No valid thread for given thread id
61 ROCPROFILER_STATUS_ERROR_AGENT_NOT_FOUND, ///< Agent identifier not found
62 ROCPROFILER_STATUS_ERROR_COUNTER_NOT_FOUND, ///< Counter identifier does not exist
63 ROCPROFILER_STATUS_ERROR_CONTEXT_ERROR, ///< Generalized context error
64 ROCPROFILER_STATUS_ERROR_CONTEXT_INVALID, ///< Context configuration is not valid
65 ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_STARTED, ///< Context was not started (e.g., atomic swap
66 ///< into active array failed)
67 ROCPROFILER_STATUS_ERROR_CONTEXT_CONFLICT, ///< Context operation failed due to a conflict with
68 ///< another context
69 ROCPROFILER_STATUS_ERROR_CONTEXT_ID_NOT_ZERO, ///< Context ID is not initialized to zero
70 ROCPROFILER_STATUS_ERROR_BUFFER_BUSY, ///< buffer operation failed because it currently busy
71 ///< handling another request (e.g. flushing)
72 ROCPROFILER_STATUS_ERROR_SERVICE_ALREADY_CONFIGURED, ///< service has already been configured
73 ///< in context
74 ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED, ///< Function call is not valid outside of
75 ///< rocprofiler configuration (i.e.
76 ///< function called post-initialization)
77 ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED, ///< Function is not implemented
78 ROCPROFILER_STATUS_ERROR_INCOMPATIBLE_ABI, ///< Data structure provided by user is incompatible
79 ///< with current version of rocprofiler
80 ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT, ///< Function invoked with one or more invalid
81 ///< arguments
82 ROCPROFILER_STATUS_ERROR_METRIC_NOT_VALID_FOR_AGENT, ///< Invalid metric supplied to agent.
83 ROCPROFILER_STATUS_ERROR_FINALIZED, ///< invalid because rocprofiler has been finalized
84 ROCPROFILER_STATUS_ERROR_HSA_NOT_LOADED, ///< Call requires HSA to be loaded before performed
85 ROCPROFILER_STATUS_ERROR_DIM_NOT_FOUND, ///< Dimension is not found for counter
86 ROCPROFILER_STATUS_ERROR_PROFILE_COUNTER_NOT_FOUND, ///< Profile could not find counter for GPU
87 ///< agent
88 ROCPROFILER_STATUS_ERROR_AST_GENERATION_FAILED, ///< AST could not be generated correctly
90 ROCPROFILER_STATUS_ERROR_AQL_NO_EVENT_COORD, ///< Event coordinate was not found by AQL profile
91 ROCPROFILER_STATUS_ERROR_INCOMPATIBLE_KERNEL, ///< A service depends on a newer version of KFD
92 ///< (amdgpu kernel driver). Check logs for
93 ///< service that report incompatibility
94 ROCPROFILER_STATUS_ERROR_OUT_OF_RESOURCES, ///< The given resources are
95 ///< insufficient to complete operation
96 ROCPROFILER_STATUS_ERROR_PROFILE_NOT_FOUND, ///< Could not find the counter profile
97 ROCPROFILER_STATUS_ERROR_AGENT_DISPATCH_CONFLICT, ///< Cannot enable both agent and dispatch
98 ///< counting in the same context.
99 ROCPROFILER_STATUS_INTERNAL_NO_AGENT_CONTEXT, ///< No agent context found, may not be an error
101 ROCPROFILER_STATUS_ERROR_NO_PROFILE_QUEUE, ///< Profile queue creation failed
102 ROCPROFILER_STATUS_ERROR_NO_HARDWARE_COUNTERS, ///< No hardware counters were specified
103 ROCPROFILER_STATUS_ERROR_AGENT_MISMATCH, ///< Agent mismatch between profile and context.
104 ROCPROFILER_STATUS_ERROR_NOT_AVAILABLE, ///< The service is not available.
105 ///< Please refer to API functions that return this
106 ///< status code for more information.
107 ROCPROFILER_STATUS_ERROR_EXCEEDS_HW_LIMIT, ///< Exceeds hardware limits for collection.
108 ROCPROFILER_STATUS_ERROR_AGENT_ARCH_NOT_SUPPORTED, ///< Agent HW architecture not supported.
111
112/**
113 * @brief Buffer record categories. This enumeration type is encoded in @ref
114 * rocprofiler_record_header_t category field
115 */
124
125/**
126 * @brief Agent type.
127 */
128typedef enum // NOLINT(performance-enum-size)
129{
130 ROCPROFILER_AGENT_TYPE_NONE = 0, ///< Agent type is unknown
131 ROCPROFILER_AGENT_TYPE_CPU, ///< Agent type is a CPU
132 ROCPROFILER_AGENT_TYPE_GPU, ///< Agent type is a GPU
135
136/**
137 * @brief Service Callback Phase.
138 */
139typedef enum // NOLINT(performance-enum-size)
140{
141 ROCPROFILER_CALLBACK_PHASE_NONE = 0, ///< Callback has no phase
142 ROCPROFILER_CALLBACK_PHASE_ENTER, ///< Callback invoked prior to function execution
144 ROCPROFILER_CALLBACK_PHASE_ENTER, ///< Callback invoked prior to code object loading
145 ROCPROFILER_CALLBACK_PHASE_EXIT, ///< Callback invoked after to function execution
147 ROCPROFILER_CALLBACK_PHASE_EXIT, ///< Callback invoked prior to code object unloading
150
151/**
152 * @brief Service Callback Tracing Kind. @see rocprofiler_configure_callback_tracing_service.
153 */
154typedef enum // NOLINT(performance-enum-size)
155{
157 ROCPROFILER_CALLBACK_TRACING_HSA_CORE_API, ///< @see ::rocprofiler_hsa_core_api_id_t
158 ROCPROFILER_CALLBACK_TRACING_HSA_AMD_EXT_API, ///< @see ::rocprofiler_hsa_amd_ext_api_id_t
159 ROCPROFILER_CALLBACK_TRACING_HSA_IMAGE_EXT_API, ///< @see ::rocprofiler_hsa_image_ext_api_id_t
161 ///< ::rocprofiler_hsa_finalize_ext_api_id_t
162 ROCPROFILER_CALLBACK_TRACING_HIP_RUNTIME_API, ///< @see ::rocprofiler_hip_runtime_api_id_t
163 ROCPROFILER_CALLBACK_TRACING_HIP_COMPILER_API, ///< @see ::rocprofiler_hip_compiler_api_id_t
164 ROCPROFILER_CALLBACK_TRACING_MARKER_CORE_API, ///< @see ::rocprofiler_marker_core_api_id_t
166 ///< ::rocprofiler_marker_control_api_id_t
167 ROCPROFILER_CALLBACK_TRACING_MARKER_NAME_API, ///< @see ::rocprofiler_marker_name_api_id_t
168 ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT, ///< @see ::rocprofiler_code_object_operation_t
169 ROCPROFILER_CALLBACK_TRACING_SCRATCH_MEMORY, ///< @see ::rocprofiler_scratch_memory_operation_t
170 ROCPROFILER_CALLBACK_TRACING_KERNEL_DISPATCH, ///< Callbacks for kernel dispatches
171 ROCPROFILER_CALLBACK_TRACING_MEMORY_COPY, ///< @see ::rocprofiler_memory_copy_operation_t
173 ROCPROFILER_CALLBACK_TRACING_OMPT, ///< @see ::rocprofiler_ompt_operation_t
175 ///< ::rocprofiler_memory_allocation_operation_t
176 ROCPROFILER_CALLBACK_TRACING_RUNTIME_INITIALIZATION, ///< Callback notifying that a runtime
177 ///< library has been initialized
180
181/**
182 * @brief Service Buffer Tracing Kind. @see rocprofiler_configure_buffer_tracing_service.
183 */
184typedef enum // NOLINT(performance-enum-size)
185{
187 ROCPROFILER_BUFFER_TRACING_HSA_CORE_API, ///< @see ::rocprofiler_hsa_core_api_id_t
188 ROCPROFILER_BUFFER_TRACING_HSA_AMD_EXT_API, ///< @see ::rocprofiler_hsa_amd_ext_api_id_t
189 ROCPROFILER_BUFFER_TRACING_HSA_IMAGE_EXT_API, ///< @see ::rocprofiler_hsa_image_ext_api_id_t
191 ///< ::rocprofiler_hsa_finalize_ext_api_id_t
192 ROCPROFILER_BUFFER_TRACING_HIP_RUNTIME_API, ///< @see ::rocprofiler_hip_runtime_api_id_t
193 ROCPROFILER_BUFFER_TRACING_HIP_COMPILER_API, ///< @see ::rocprofiler_hip_compiler_api_id_t
194 ROCPROFILER_BUFFER_TRACING_MARKER_CORE_API, ///< @see ::rocprofiler_marker_core_api_id_t
195 ROCPROFILER_BUFFER_TRACING_MARKER_CONTROL_API, ///< @see ::rocprofiler_marker_control_api_id_t
196 ROCPROFILER_BUFFER_TRACING_MARKER_NAME_API, ///< @see ::rocprofiler_marker_name_api_id_t
197 ROCPROFILER_BUFFER_TRACING_MEMORY_COPY, ///< @see ::rocprofiler_memory_copy_operation_t
198 ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH, ///< Buffer kernel dispatch info
199 ROCPROFILER_BUFFER_TRACING_PAGE_MIGRATION, ///< Buffer page migration info
200 ROCPROFILER_BUFFER_TRACING_SCRATCH_MEMORY, ///< Buffer scratch memory reclaimation info
201 ROCPROFILER_BUFFER_TRACING_CORRELATION_ID_RETIREMENT, ///< Correlation ID in no longer in use
203 ROCPROFILER_BUFFER_TRACING_OMPT, ///< @see ::rocprofiler_ompt_operation_t
205 ///< ::rocprofiler_memory_allocation_operation_t
206 ROCPROFILER_BUFFER_TRACING_RUNTIME_INITIALIZATION, ///< Record indicating a runtime library has
207 ///< been initialized. @see
208 ///< ::rocprofiler_runtime_initialization_operation_t
211
212/**
213 * @brief ROCProfiler Code Object Tracer Operations.
214 */
215typedef enum // NOLINT(performance-enum-size)
216{
217 ROCPROFILER_CODE_OBJECT_NONE = 0, ///< Unknown code object operation
218 ROCPROFILER_CODE_OBJECT_LOAD, ///< Code object containing kernel symbols
223
224/**
225 * @brief Memory Copy Operations.
226 */
227typedef enum // NOLINT(performance-enum-size)
228{
229 ROCPROFILER_MEMORY_COPY_NONE = 0, ///< Unknown memory copy direction
230 ROCPROFILER_MEMORY_COPY_HOST_TO_HOST, ///< Memory copy from host to host
231 ROCPROFILER_MEMORY_COPY_HOST_TO_DEVICE, ///< Memory copy from host to device
232 ROCPROFILER_MEMORY_COPY_DEVICE_TO_HOST, ///< Memory copy from device to host
233 ROCPROFILER_MEMORY_COPY_DEVICE_TO_DEVICE, ///< Memory copy from device to device
236
237/**
238 * @brief Memory Allocation Operation.
239 */
240typedef enum // NOLINT(performance-enum-size)
241{
242 ROCPROFILER_MEMORY_ALLOCATION_NONE = 0, ///< Unknown memory allocation function
243 ROCPROFILER_MEMORY_ALLOCATION_ALLOCATE, ///< Allocate memory function
244 ROCPROFILER_MEMORY_ALLOCATION_VMEM_ALLOCATE, ///< Allocate vmem memory handle
245 ROCPROFILER_MEMORY_ALLOCATION_FREE, ///< Free memory function
246 ROCPROFILER_MEMORY_ALLOCATION_VMEM_FREE, ///< Release vmem memory handle
249
250/**
251 * @brief ROCProfiler Kernel Dispatch Tracing Operation Types.
252 */
253typedef enum // NOLINT(performance-enum-size)
254{
255 ROCPROFILER_KERNEL_DISPATCH_NONE = 0, ///< Unknown kernel dispatch operation
259
260 /// @var ROCPROFILER_KERNEL_DISPATCH_ENQUEUE
261 /// @brief Invoke callback prior to a kernel being enqueued and after the kernel has been
262 /// enqueued. When the phase is ::ROCPROFILER_CALLBACK_PHASE_ENTER, this is an opportunity to
263 /// push an external correlation id and/or modify the active contexts before a kernel is
264 /// launched. Any active contexts containing services related to a kernel dispatch (kernel
265 /// tracing, counter collection, etc.) will be captured after this callback and attached to the
266 /// kernel. These captured contexts will be considered "active" when the kernel completes even
267 /// if the context was stopped before the kernel completes -- this contract is designed to
268 /// ensure that tools do not have to delay stopping a context because of an async operation in
269 /// order to get the data they requested when the async operation was started. When the phase is
270 /// ::ROCPROFILER_CALLBACK_PHASE_EXIT, the active contexts for the kernel dispatch have been
271 /// captured and it is safe to disable those contexts without affecting the delivery of the
272 /// requested data when the kernel completes. It is important to note that, even if the context
273 /// associated with the kernel dispatch callback tracing service is disabled in between the
274 /// enter and exit phase, the exit phase callback is still delievered but that context will not
275 /// be captured when the kernel is enqueued and therefore will not provide a
276 /// ::ROCPROFILER_KERNEL_DISPATCH_COMPLETE callback. Furthermore, it should be
277 /// noted that if a tool encodes information into the `::rocprofiler_user_data_t` output
278 /// parameter in ::rocprofiler_callback_tracing_cb_t, that same value will be delivered in the
279 /// exit phase and in the ::ROCPROFILER_KERNEL_DISPATCH_COMPLETE callback. In
280 /// other words, any modifications to that user data value in the exit phase will not be
281 /// reflected in the ::ROCPROFILER_KERNEL_DISPATCH_COMPLETE callback because a
282 /// copy of that user data struct is attached to the kernel, not a reference to the user data
283 /// struct.
284 ///
285 /// @var ROCPROFILER_KERNEL_DISPATCH_COMPLETE
286 /// @brief Invoke callback after a kernel has completed and the HSA runtime has processed the
287 /// signal indicating that the kernel has completed. The latter half of this statement is
288 /// important. There is no guarantee that these callbacks are invoked in any order related to
289 /// when the kernels were dispatched, i.e. even if kernel A is launched and fully executed
290 /// before kernel B is launched, it is entirely possible that the HSA runtime ends up processing
291 /// the signal associated with kernel B before processing the signal associated with kernel A --
292 /// resulting in rocprofiler-sdk invoking this operation callback for kernel B before invoking
293 /// the callback for kernel A.
295
296/**
297 * @brief PC Sampling Method.
298 */
299typedef enum // NOLINT(performance-enum-size)
300{
301 ROCPROFILER_PC_SAMPLING_METHOD_NONE = 0, ///< Unknown sampling type
302 ROCPROFILER_PC_SAMPLING_METHOD_STOCHASTIC, ///< Stochastic sampling (MI300+)
303 ROCPROFILER_PC_SAMPLING_METHOD_HOST_TRAP, ///< Interval sampling (MI200+)
306
307/**
308 * @brief PC Sampling Unit.
309 */
310typedef enum // NOLINT(performance-enum-size)
311{
312 ROCPROFILER_PC_SAMPLING_UNIT_NONE = 0, ///< Sample interval has unspecified units
313 ROCPROFILER_PC_SAMPLING_UNIT_INSTRUCTIONS, ///< Sample interval is in instructions
314 ROCPROFILER_PC_SAMPLING_UNIT_CYCLES, ///< Sample interval is in cycles
315 ROCPROFILER_PC_SAMPLING_UNIT_TIME, ///< Sample internval is in nanoseconds
318
319/**
320 * @brief Actions when Buffer is full.
321 */
322typedef enum // NOLINT(performance-enum-size)
323{
324 ROCPROFILER_BUFFER_POLICY_NONE = 0, ///< No policy has been set
325 ROCPROFILER_BUFFER_POLICY_DISCARD, ///< Drop records when buffer is full
326 ROCPROFILER_BUFFER_POLICY_LOSSLESS, ///< Block when buffer is full
329
330/**
331 * @brief Page migration event.
332 */
346
347/**
348 * @brief Scratch event kind
349 */
350typedef enum
351{
352 ROCPROFILER_SCRATCH_MEMORY_NONE = 0, ///< Unknown scratch operation
353 ROCPROFILER_SCRATCH_MEMORY_ALLOC, ///< Scratch memory allocation event
354 ROCPROFILER_SCRATCH_MEMORY_FREE, ///< Scratch memory free event
355 ROCPROFILER_SCRATCH_MEMORY_ASYNC_RECLAIM, ///< Scratch memory asynchronously reclaimed
358
359/**
360 * @brief Enumeration for specifying runtime libraries supported by rocprofiler. This enumeration is
361 * used for thread creation callbacks. @see INTERNAL_THREADING.
362 */
372
373/**
374 * @brief Enumeration for specifying intercept tables supported by rocprofiler. This enumeration is
375 * used for intercept tables. @see INTERCEPT_TABLE.
376 */
388
389/**
390 * @brief ROCProfiler Runtime Initialization Tracer Operations.
391 */
392typedef enum // NOLINT(performance-enum-size)
393{
394 ROCPROFILER_RUNTIME_INITIALIZATION_NONE = 0, ///< Unknown runtime initialization
395 ROCPROFILER_RUNTIME_INITIALIZATION_HSA, ///< Application loaded HSA runtime
396 ROCPROFILER_RUNTIME_INITIALIZATION_HIP, ///< Application loaded HIP runtime
397 ROCPROFILER_RUNTIME_INITIALIZATION_MARKER, ///< Application loaded Marker (ROCTx) runtime
398 ROCPROFILER_RUNTIME_INITIALIZATION_RCCL, ///< Application loaded RCCL runtime
401
402/**
403 * @brief Enumeration for specifying the counter info struct version you want.
404 */
411
412/**
413 * @brief Enumeration for distinguishing different buffer record kinds within the
414 * ::ROCPROFILER_BUFFER_CATEGORY_COUNTERS category
415 */
416typedef enum
417{
419 ROCPROFILER_COUNTER_RECORD_PROFILE_COUNTING_DISPATCH_HEADER, ///< ::rocprofiler_dispatch_counting_service_record_t
422
423 /// @var ROCPROFILER_COUNTER_RECORD_KIND_DISPATCH_PROFILE_HEADER
424 /// @brief Indicates the payload type is of type
425 /// ::rocprofiler_dispatch_counting_service_record_t
427
428/**
429 * @brief Enumeration of flags that can be used with some counter api calls
430 */
431typedef enum
432{
434 ROCPROFILER_COUNTER_FLAG_ASYNC, ///< Do not wait for completion before returning.
435 ROCPROFILER_COUNTER_FLAG_APPEND_DEFINITION, ///< Append the counter definition to the system
436 ///< provided counter definition file.
439
440/**
441 * @brief Enumeration for distinguishing different buffer record kinds within the
442 * ::ROCPROFILER_BUFFER_CATEGORY_PC_SAMPLING category
443 */
451
452//--------------------------------------------------------------------------------------//
453//
454// ALIASES
455//
456//--------------------------------------------------------------------------------------//
457
458/**
459 * @brief ROCProfiler Timestamp.
460 */
461typedef uint64_t rocprofiler_timestamp_t;
462
463/**
464 * @brief Thread ID. Value will be equivalent to `syscall(__NR_gettid)`
465 */
466typedef uint64_t rocprofiler_thread_id_t;
467
468/**
469 * @brief Tracing Operation ID. Depending on the kind, operations can be determined.
470 * If the value is equal to zero that means all operations will be considered
471 * for tracing. Detailed API tracing operations can be found at associated header file
472 * for that partiular operation. i.e: For ROCProfiler enumeration of HSA AMD Extended API tracing
473 * operations, look at source/include/rocprofiler-sdk/hsa/amd_ext_api_id.h
474 */
476
477/**
478 * @brief Kernel identifier type
479 *
480 */
481typedef uint64_t rocprofiler_kernel_id_t;
482
483// /**
484// * @brief Sequence identifier type
485// *
486// */
488
489/**
490 * @brief Unique record id encoding both the counter
491 * and dimensional values (positions) for the record.
492 */
494
495/**
496 * @brief A dimension for counter instances. Some example
497 * dimensions include XCC, SM (Shader), etc. This
498 * value represents the dimension beind described
499 * or queried about.
500 */
502
503//--------------------------------------------------------------------------------------//
504//
505// UNIONS
506//
507//--------------------------------------------------------------------------------------//
508
509/**
510 * @brief User-assignable data type
511 *
512 */
514{
515 uint64_t value; ///< usage example: set to process id, thread id, etc.
516 void* ptr; ///< usage example: set to address of data allocation
518
519/**
520 * @brief Stores memory address for profiling
521 *
522 */
524{
525 uint64_t value; ///< usage example: store address in uint64_t format
526 void* ptr; ///< usage example: generic form of address
528
529//--------------------------------------------------------------------------------------//
530//
531// STRUCTS
532//
533//--------------------------------------------------------------------------------------//
534
535/**
536 * @brief Context ID.
537 */
538typedef struct
539{
540 uint64_t handle;
542
543/**
544 * @brief Queue ID.
545 */
546typedef struct
547{
548 uint64_t handle;
550
551/**
552 * @brief ROCProfiler Record Correlation ID.
553 */
559
560/**
561 * @brief The NULL value of an internal correlation ID.
562 */
563#define ROCPROFILER_CORRELATION_ID_INTERNAL_NONE ROCPROFILER_UINT64_C(0)
564
565/**
566 * @struct rocprofiler_buffer_id_t
567 * @brief Buffer ID.
568 */
569typedef struct
570{
571 uint64_t handle;
573
574/**
575 * @brief Agent Identifier
576 */
577typedef struct
578{
579 uint64_t handle;
581
582/**
583 * @brief Counter ID.
584 */
585typedef struct
586{
587 uint64_t handle;
589
590/**
591 * @brief Profile Configurations
592 * @see rocprofiler_create_profile_config for how to create.
593 */
594typedef struct
595{
596 uint64_t handle; // Opaque handle
598
599/**
600 * @brief Multi-dimensional struct of data used to describe GPU workgroup and grid sizes
601 */
602typedef struct rocprofiler_dim3_t
603{
604 uint32_t x;
605 uint32_t y;
606 uint32_t z;
608
609/**
610 * @brief Tracing record
611 *
612 */
623
624/**
625 * @brief Generic record with type identifier(s) and a pointer to data. This data type is used with
626 * buffered data.
627 *
628 * @code{.cpp}
629 * void
630 * tool_tracing_callback(rocprofiler_record_header_t** headers,
631 * size_t num_headers)
632 * {
633 * for(size_t i = 0; i < num_headers; ++i)
634 * {
635 * rocprofiler_record_header_t* header = headers[i];
636 *
637 * if(header->category == ROCPROFILER_BUFFER_CATEGORY_TRACING &&
638 * header->kind == ROCPROFILER_BUFFER_TRACING_HSA_API)
639 * {
640 * // cast to rocprofiler_buffer_tracing_hsa_api_record_t which
641 * // is type associated with this category + kind
642 * auto* record =
643 * static_cast<rocprofiler_buffer_tracing_hsa_api_record_t*>(header->payload);
644 *
645 * // trivial test
646 * assert(record->start_timestamp <= record->end_timestamp);
647 * }
648 * }
649 * }
650 *
651 * @endcode
652 */
653typedef struct
654{
655 union
656 {
657 struct
658 {
659 uint32_t category; ///< rocprofiler_buffer_category_t
660 uint32_t kind; ///< domain
661 };
662 uint64_t hash; ///< generic identifier. You can compute this via: `uint64_t hash = category
663 ///< | ((uint64_t)(kind) << 32)`, e.g.
664 };
665 void* payload;
667
668/**
669 * @brief Function for computing the unsigned 64-bit hash value in @ref rocprofiler_record_header_t
670 * from a category and kind (two unsigned 32-bit values)
671 *
672 * @param [in] category a value from @ref rocprofiler_buffer_category_t
673 * @param [in] kind depending on the category, this is the domain value, e.g., @ref
674 * rocprofiler_buffer_tracing_kind_t value
675 * @return uint64_t hash value of category and kind
676 */
677static inline uint64_t
678rocprofiler_record_header_compute_hash(uint32_t category, uint32_t kind)
679{
680 uint64_t value = category;
681 value |= ((uint64_t)(kind)) << 32;
682 return value;
683}
684
685/**
686 * @brief ROCProfiler kernel dispatch information
687 *
688 */
690{
691 uint64_t size; ///< Size of this struct (minus reserved padding)
692 rocprofiler_agent_id_t agent_id; ///< Agent ID where kernel is launched
693 rocprofiler_queue_id_t queue_id; ///< Queue ID where kernel packet is enqueued
694 rocprofiler_kernel_id_t kernel_id; ///< Kernel identifier
695 rocprofiler_dispatch_id_t dispatch_id; ///< unique id for each dispatch
696 uint32_t private_segment_size; ///< runtime private memory segment size
697 uint32_t group_segment_size; ///< runtime group memory segment size
698 rocprofiler_dim3_t workgroup_size; ///< runtime workgroup size (grid * threads)
699 rocprofiler_dim3_t grid_size; ///< runtime grid size
700 uint8_t reserved_padding[56]; // reserved for extensions w/o ABI break
702
703/**
704 * @brief Details for the dimension, including its size, for a counter record.
705 */
706typedef struct
707{
708 const char* name;
711
712 /// @var id
713 /// @brief Id for this dimension used by @ref rocprofiler_query_record_dimension_position
715
716/**
717 * @brief ROCProfiler Profile Counting Counter Record per instance.
718 */
719typedef struct
720{
721 rocprofiler_counter_instance_id_t id; ///< counter identifier
722 double counter_value; ///< counter value
726
727 /// @var dispatch_id
728 /// @brief A value greater than zero indicates that this counter record is associated with a
729 /// specific dispatch.
730 ///
731 /// This value can be mapped to a dispatch via the `dispatch_info` field (@see
732 /// ::rocprofiler_kernel_dispatch_info_t) of a ::rocprofiler_dispatch_counting_service_data_t
733 /// instance (provided during callback for profile config) or a
734 /// ::rocprofiler_dispatch_counting_service_record_t records (which will be insert into the
735 /// buffer prior to the associated ::rocprofiler_record_counter_t records).
737
738/**
739 * @brief Counter info struct version 0
740 */
741typedef struct
742{
743 rocprofiler_counter_id_t id; ///< Id of this counter
744 const char* name; ///< Name of the counter
745 const char* description; ///< Description of the counter
746 const char* block; ///< Block of the counter (non-derived only)
747 const char* expression; ///< Counter expression (derived counters only)
748 uint8_t is_constant : 1; ///< If this counter is HW constant
749 uint8_t is_derived : 1; ///< If this counter is a derived counter
751
752/**
753 * @brief ROCProfiler SPM Record.
754 *
755 */
756typedef struct
757{
758 /**
759 * Counters, including identifiers to get counter information and Counters
760 * values
761 */
765
766/** @} */
767
768ROCPROFILER_EXTERN_C_FINI
769
771 static_assert(sizeof(rocprofiler_kernel_dispatch_info_t) == 128,
772 "Increasing the size of the kernel dispatch info is not permitted");)
#define ROCPROFILER_CXX_CODE(...)
Definition defines.h:132
void * ptr
usage example: set to address of data allocation
Definition fwd.h:516
const char * name
Name of the counter.
Definition fwd.h:744
const char * block
Block of the counter (non-derived only)
Definition fwd.h:746
rocprofiler_dim3_t grid_size
runtime grid size
Definition fwd.h:699
rocprofiler_user_data_t external
Definition fwd.h:557
rocprofiler_counter_dimension_id_t id
Id for this dimension used by rocprofiler_query_record_dimension_position.
Definition fwd.h:710
uint64_t value
usage example: set to process id, thread id, etc.
Definition fwd.h:515
uint32_t z
Definition fwd.h:606
uint32_t group_segment_size
runtime group memory segment size
Definition fwd.h:697
void * ptr
usage example: generic form of address
Definition fwd.h:526
rocprofiler_user_data_t user_data
Definition fwd.h:724
rocprofiler_agent_id_t agent_id
Agent ID where kernel is launched.
Definition fwd.h:692
uint32_t y
Definition fwd.h:605
uint32_t private_segment_size
runtime private memory segment size
Definition fwd.h:696
rocprofiler_callback_phase_t phase
Definition fwd.h:620
uint64_t handle
Definition fwd.h:579
rocprofiler_counter_id_t id
Id of this counter.
Definition fwd.h:743
rocprofiler_record_counter_t * counters
Definition fwd.h:762
rocprofiler_dispatch_id_t dispatch_id
A value greater than zero indicates that this counter record is associated with a specific dispatch.
Definition fwd.h:723
rocprofiler_context_id_t context_id
Definition fwd.h:615
const char * description
Description of the counter.
Definition fwd.h:745
rocprofiler_callback_tracing_kind_t kind
Definition fwd.h:618
uint8_t is_constant
If this counter is HW constant.
Definition fwd.h:748
uint64_t handle
Definition fwd.h:548
rocprofiler_dispatch_id_t dispatch_id
unique id for each dispatch
Definition fwd.h:695
rocprofiler_thread_id_t thread_id
Definition fwd.h:616
rocprofiler_kernel_id_t kernel_id
Kernel identifier.
Definition fwd.h:694
rocprofiler_tracing_operation_t operation
Definition fwd.h:619
rocprofiler_counter_instance_id_t id
counter identifier
Definition fwd.h:721
const char * expression
Counter expression (derived counters only)
Definition fwd.h:747
uint64_t size
Size of this struct (minus reserved padding)
Definition fwd.h:691
double counter_value
counter value
Definition fwd.h:722
uint64_t counters_count
Definition fwd.h:763
uint64_t value
usage example: store address in uint64_t format
Definition fwd.h:525
rocprofiler_correlation_id_t correlation_id
Definition fwd.h:617
uint32_t x
Definition fwd.h:604
rocprofiler_dim3_t workgroup_size
runtime workgroup size (grid * threads)
Definition fwd.h:698
uint8_t is_derived
If this counter is a derived counter.
Definition fwd.h:749
rocprofiler_queue_id_t queue_id
Queue ID where kernel packet is enqueued.
Definition fwd.h:693
rocprofiler_agent_id_t agent_id
Definition fwd.h:725
rocprofiler_pc_sampling_method_t
PC Sampling Method.
Definition fwd.h:300
rocprofiler_code_object_operation_t
ROCProfiler Code Object Tracer Operations.
Definition fwd.h:216
int32_t rocprofiler_tracing_operation_t
Tracing Operation ID. Depending on the kind, operations can be determined. If the value is equal to z...
Definition fwd.h:475
rocprofiler_buffer_policy_t
Actions when Buffer is full.
Definition fwd.h:323
rocprofiler_scratch_memory_operation_t
Scratch event kind.
Definition fwd.h:351
uint64_t rocprofiler_counter_dimension_id_t
A dimension for counter instances. Some example dimensions include XCC, SM (Shader),...
Definition fwd.h:501
rocprofiler_page_migration_operation_t
Page migration event.
Definition fwd.h:334
rocprofiler_memory_copy_operation_t
Memory Copy Operations.
Definition fwd.h:228
uint64_t rocprofiler_counter_instance_id_t
Unique record id encoding both the counter and dimensional values (positions) for the record.
Definition fwd.h:493
rocprofiler_runtime_initialization_operation_t
ROCProfiler Runtime Initialization Tracer Operations.
Definition fwd.h:393
rocprofiler_memory_allocation_operation_t
Memory Allocation Operation.
Definition fwd.h:241
rocprofiler_counter_record_kind_t
Enumeration for distinguishing different buffer record kinds within the ROCPROFILER_BUFFER_CATEGORY_C...
Definition fwd.h:417
rocprofiler_pc_sampling_unit_t
PC Sampling Unit.
Definition fwd.h:311
rocprofiler_intercept_table_t
Enumeration for specifying intercept tables supported by rocprofiler. This enumeration is used for in...
Definition fwd.h:378
rocprofiler_status_t
Status codes.
Definition fwd.h:53
uint64_t rocprofiler_dispatch_id_t
Definition fwd.h:487
rocprofiler_callback_phase_t
Service Callback Phase.
Definition fwd.h:140
uint64_t rocprofiler_kernel_id_t
Kernel identifier type.
Definition fwd.h:481
uint64_t rocprofiler_thread_id_t
Thread ID. Value will be equivalent to syscall(__NR_gettid)
Definition fwd.h:466
rocprofiler_buffer_tracing_kind_t
Service Buffer Tracing Kind.
Definition fwd.h:185
rocprofiler_pc_sampling_record_kind_t
Enumeration for distinguishing different buffer record kinds within the ROCPROFILER_BUFFER_CATEGORY_P...
Definition fwd.h:445
rocprofiler_kernel_dispatch_operation_t
ROCProfiler Kernel Dispatch Tracing Operation Types.
Definition fwd.h:254
rocprofiler_runtime_library_t
Enumeration for specifying runtime libraries supported by rocprofiler. This enumeration is used for t...
Definition fwd.h:364
uint64_t rocprofiler_timestamp_t
ROCProfiler Timestamp.
Definition fwd.h:461
rocprofiler_buffer_category_t
Buffer record categories. This enumeration type is encoded in rocprofiler_record_header_t category fi...
Definition fwd.h:117
rocprofiler_callback_tracing_kind_t
Service Callback Tracing Kind.
Definition fwd.h:155
rocprofiler_agent_type_t
Agent type.
Definition fwd.h:129
rocprofiler_counter_info_version_id_t
Enumeration for specifying the counter info struct version you want.
Definition fwd.h:406
rocprofiler_counter_flag_t
Enumeration of flags that can be used with some counter api calls.
Definition fwd.h:432
@ ROCPROFILER_PC_SAMPLING_METHOD_STOCHASTIC
Stochastic sampling (MI300+)
Definition fwd.h:302
@ ROCPROFILER_PC_SAMPLING_METHOD_HOST_TRAP
Interval sampling (MI200+)
Definition fwd.h:303
@ ROCPROFILER_PC_SAMPLING_METHOD_LAST
Definition fwd.h:304
@ ROCPROFILER_PC_SAMPLING_METHOD_NONE
Unknown sampling type.
Definition fwd.h:301
@ ROCPROFILER_CODE_OBJECT_HOST_KERNEL_SYMBOL_REGISTER
Kernel symbols - Host.
Definition fwd.h:220
@ ROCPROFILER_CODE_OBJECT_DEVICE_KERNEL_SYMBOL_REGISTER
Kernel symbols - Device.
Definition fwd.h:219
@ ROCPROFILER_CODE_OBJECT_LAST
Definition fwd.h:221
@ ROCPROFILER_CODE_OBJECT_NONE
Unknown code object operation.
Definition fwd.h:217
@ ROCPROFILER_CODE_OBJECT_LOAD
Code object containing kernel symbols.
Definition fwd.h:218
@ ROCPROFILER_BUFFER_POLICY_DISCARD
Drop records when buffer is full.
Definition fwd.h:325
@ ROCPROFILER_BUFFER_POLICY_NONE
No policy has been set.
Definition fwd.h:324
@ ROCPROFILER_BUFFER_POLICY_LOSSLESS
Block when buffer is full.
Definition fwd.h:326
@ ROCPROFILER_BUFFER_POLICY_LAST
Definition fwd.h:327
@ ROCPROFILER_SCRATCH_MEMORY_ALLOC
Scratch memory allocation event.
Definition fwd.h:353
@ ROCPROFILER_SCRATCH_MEMORY_FREE
Scratch memory free event.
Definition fwd.h:354
@ ROCPROFILER_SCRATCH_MEMORY_ASYNC_RECLAIM
Scratch memory asynchronously reclaimed.
Definition fwd.h:355
@ ROCPROFILER_SCRATCH_MEMORY_LAST
Definition fwd.h:356
@ ROCPROFILER_SCRATCH_MEMORY_NONE
Unknown scratch operation.
Definition fwd.h:352
@ ROCPROFILER_PAGE_MIGRATION_NONE
Unknown event.
Definition fwd.h:335
@ ROCPROFILER_PAGE_MIGRATION_LAST
Definition fwd.h:344
@ ROCPROFILER_PAGE_MIGRATION_PAGE_FAULT_START
Definition fwd.h:338
@ ROCPROFILER_PAGE_MIGRATION_PAGE_FAULT_END
Definition fwd.h:339
@ ROCPROFILER_PAGE_MIGRATION_UNMAP_FROM_GPU
Definition fwd.h:342
@ ROCPROFILER_PAGE_MIGRATION_QUEUE_EVICTION
Definition fwd.h:340
@ ROCPROFILER_PAGE_MIGRATION_PAGE_MIGRATE_END
Definition fwd.h:337
@ ROCPROFILER_PAGE_MIGRATION_PAGE_MIGRATE_START
Definition fwd.h:336
@ ROCPROFILER_PAGE_MIGRATION_DROPPED_EVENT
Definition fwd.h:343
@ ROCPROFILER_PAGE_MIGRATION_QUEUE_RESTORE
Definition fwd.h:341
@ ROCPROFILER_MEMORY_COPY_LAST
Definition fwd.h:234
@ ROCPROFILER_MEMORY_COPY_DEVICE_TO_DEVICE
Memory copy from device to device.
Definition fwd.h:233
@ ROCPROFILER_MEMORY_COPY_HOST_TO_HOST
Memory copy from host to host.
Definition fwd.h:230
@ ROCPROFILER_MEMORY_COPY_DEVICE_TO_HOST
Memory copy from device to host.
Definition fwd.h:232
@ ROCPROFILER_MEMORY_COPY_NONE
Unknown memory copy direction.
Definition fwd.h:229
@ ROCPROFILER_MEMORY_COPY_HOST_TO_DEVICE
Memory copy from host to device.
Definition fwd.h:231
@ ROCPROFILER_RUNTIME_INITIALIZATION_LAST
Definition fwd.h:399
@ ROCPROFILER_RUNTIME_INITIALIZATION_NONE
Unknown runtime initialization.
Definition fwd.h:394
@ ROCPROFILER_RUNTIME_INITIALIZATION_HIP
Application loaded HIP runtime.
Definition fwd.h:396
@ ROCPROFILER_RUNTIME_INITIALIZATION_MARKER
Application loaded Marker (ROCTx) runtime.
Definition fwd.h:397
@ ROCPROFILER_RUNTIME_INITIALIZATION_RCCL
Application loaded RCCL runtime.
Definition fwd.h:398
@ ROCPROFILER_RUNTIME_INITIALIZATION_HSA
Application loaded HSA runtime.
Definition fwd.h:395
@ ROCPROFILER_MEMORY_ALLOCATION_NONE
Unknown memory allocation function.
Definition fwd.h:242
@ ROCPROFILER_MEMORY_ALLOCATION_ALLOCATE
Allocate memory function.
Definition fwd.h:243
@ ROCPROFILER_MEMORY_ALLOCATION_FREE
Free memory function.
Definition fwd.h:245
@ ROCPROFILER_MEMORY_ALLOCATION_VMEM_FREE
Release vmem memory handle.
Definition fwd.h:246
@ ROCPROFILER_MEMORY_ALLOCATION_VMEM_ALLOCATE
Allocate vmem memory handle.
Definition fwd.h:244
@ ROCPROFILER_MEMORY_ALLOCATION_LAST
Definition fwd.h:247
@ ROCPROFILER_COUNTER_RECORD_PROFILE_COUNTING_DISPATCH_HEADER
rocprofiler_dispatch_counting_service_record_t
Definition fwd.h:419
@ ROCPROFILER_COUNTER_RECORD_VALUE
Definition fwd.h:420
@ ROCPROFILER_COUNTER_RECORD_LAST
Definition fwd.h:421
@ ROCPROFILER_COUNTER_RECORD_NONE
Definition fwd.h:418
@ ROCPROFILER_PC_SAMPLING_UNIT_LAST
Definition fwd.h:316
@ ROCPROFILER_PC_SAMPLING_UNIT_TIME
Sample internval is in nanoseconds.
Definition fwd.h:315
@ ROCPROFILER_PC_SAMPLING_UNIT_NONE
Sample interval has unspecified units.
Definition fwd.h:312
@ ROCPROFILER_PC_SAMPLING_UNIT_INSTRUCTIONS
Sample interval is in instructions.
Definition fwd.h:313
@ ROCPROFILER_PC_SAMPLING_UNIT_CYCLES
Sample interval is in cycles.
Definition fwd.h:314
@ ROCPROFILER_MARKER_NAME_TABLE
Definition fwd.h:384
@ ROCPROFILER_MARKER_CONTROL_TABLE
Definition fwd.h:383
@ ROCPROFILER_RCCL_TABLE
Definition fwd.h:385
@ ROCPROFILER_HIP_COMPILER_TABLE
Definition fwd.h:381
@ ROCPROFILER_MARKER_CORE_TABLE
Definition fwd.h:382
@ ROCPROFILER_TABLE_LAST
Definition fwd.h:386
@ ROCPROFILER_HSA_TABLE
Definition fwd.h:379
@ ROCPROFILER_HIP_RUNTIME_TABLE
Definition fwd.h:380
@ ROCPROFILER_STATUS_ERROR_FINALIZED
invalid because rocprofiler has been finalized
Definition fwd.h:83
@ ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_STARTED
Context was not started (e.g., atomic swap into active array failed)
Definition fwd.h:65
@ ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT
Function invoked with one or more invalid arguments.
Definition fwd.h:80
@ ROCPROFILER_STATUS_ERROR_NO_PROFILE_QUEUE
Profile queue creation failed.
Definition fwd.h:101
@ ROCPROFILER_STATUS_ERROR_CONTEXT_ERROR
Generalized context error.
Definition fwd.h:63
@ ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED
Function call is not valid outside of rocprofiler configuration (i.e. function called post-initializa...
Definition fwd.h:74
@ ROCPROFILER_STATUS_ERROR_BUFFER_NOT_FOUND
No valid buffer for given buffer id.
Definition fwd.h:57
@ ROCPROFILER_STATUS_ERROR_PROFILE_NOT_FOUND
Could not find the counter profile.
Definition fwd.h:96
@ ROCPROFILER_STATUS_ERROR_COUNTER_NOT_FOUND
Counter identifier does not exist.
Definition fwd.h:62
@ ROCPROFILER_STATUS_ERROR_THREAD_NOT_FOUND
No valid thread for given thread id.
Definition fwd.h:60
@ ROCPROFILER_STATUS_ERROR_CONTEXT_CONFLICT
Context operation failed due to a conflict with another context.
Definition fwd.h:67
@ ROCPROFILER_STATUS_ERROR_SERVICE_ALREADY_CONFIGURED
service has already been configured in context
Definition fwd.h:72
@ ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED
Function is not implemented.
Definition fwd.h:77
@ ROCPROFILER_STATUS_ERROR_AGENT_MISMATCH
Agent mismatch between profile and context.
Definition fwd.h:103
@ ROCPROFILER_STATUS_ERROR_AGENT_DISPATCH_CONFLICT
Cannot enable both agent and dispatch counting in the same context.
Definition fwd.h:97
@ ROCPROFILER_STATUS_ERROR_OPERATION_NOT_FOUND
Operation identifier is invalid for domain.
Definition fwd.h:59
@ ROCPROFILER_STATUS_ERROR_AQL_NO_EVENT_COORD
Event coordinate was not found by AQL profile.
Definition fwd.h:90
@ ROCPROFILER_STATUS_ERROR_NOT_AVAILABLE
The service is not available. Please refer to API functions that return this status code for more inf...
Definition fwd.h:104
@ ROCPROFILER_STATUS_ERROR_SAMPLE_RATE_EXCEEDED
Sample rate exceeded.
Definition fwd.h:100
@ ROCPROFILER_STATUS_ERROR_CONTEXT_INVALID
Context configuration is not valid.
Definition fwd.h:64
@ ROCPROFILER_STATUS_ERROR_OUT_OF_RESOURCES
The given resources are insufficient to complete operation.
Definition fwd.h:94
@ ROCPROFILER_STATUS_LAST
Definition fwd.h:109
@ ROCPROFILER_STATUS_ERROR_AGENT_NOT_FOUND
Agent identifier not found.
Definition fwd.h:61
@ ROCPROFILER_STATUS_ERROR
Generalized error.
Definition fwd.h:55
@ ROCPROFILER_STATUS_ERROR_EXCEEDS_HW_LIMIT
Exceeds hardware limits for collection.
Definition fwd.h:107
@ ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND
No valid context for given context id.
Definition fwd.h:56
@ ROCPROFILER_STATUS_ERROR_KIND_NOT_FOUND
Kind identifier is invalid.
Definition fwd.h:58
@ ROCPROFILER_STATUS_INTERNAL_NO_AGENT_CONTEXT
No agent context found, may not be an error.
Definition fwd.h:99
@ ROCPROFILER_STATUS_ERROR_AST_GENERATION_FAILED
AST could not be generated correctly.
Definition fwd.h:88
@ ROCPROFILER_STATUS_ERROR_INCOMPATIBLE_KERNEL
A service depends on a newer version of KFD (amdgpu kernel driver). Check logs for service that repor...
Definition fwd.h:91
@ ROCPROFILER_STATUS_ERROR_PROFILE_COUNTER_NOT_FOUND
Profile could not find counter for GPU agent.
Definition fwd.h:86
@ ROCPROFILER_STATUS_ERROR_AGENT_ARCH_NOT_SUPPORTED
Agent HW architecture not supported.
Definition fwd.h:108
@ ROCPROFILER_STATUS_ERROR_DIM_NOT_FOUND
Dimension is not found for counter.
Definition fwd.h:85
@ ROCPROFILER_STATUS_ERROR_METRIC_NOT_VALID_FOR_AGENT
Invalid metric supplied to agent.
Definition fwd.h:82
@ ROCPROFILER_STATUS_ERROR_BUFFER_BUSY
buffer operation failed because it currently busy handling another request (e.g. flushing)
Definition fwd.h:70
@ ROCPROFILER_STATUS_SUCCESS
No error occurred.
Definition fwd.h:54
@ ROCPROFILER_STATUS_ERROR_INCOMPATIBLE_ABI
Data structure provided by user is incompatible with current version of rocprofiler.
Definition fwd.h:78
@ ROCPROFILER_STATUS_ERROR_CONTEXT_ID_NOT_ZERO
Context ID is not initialized to zero.
Definition fwd.h:69
@ ROCPROFILER_STATUS_ERROR_NO_HARDWARE_COUNTERS
No hardware counters were specified.
Definition fwd.h:102
@ ROCPROFILER_STATUS_ERROR_AST_NOT_FOUND
AST was not found.
Definition fwd.h:89
@ ROCPROFILER_STATUS_ERROR_HSA_NOT_LOADED
Call requires HSA to be loaded before performed.
Definition fwd.h:84
@ ROCPROFILER_CALLBACK_PHASE_NONE
Callback has no phase.
Definition fwd.h:141
@ ROCPROFILER_CALLBACK_PHASE_UNLOAD
Callback invoked prior to code object unloading.
Definition fwd.h:146
@ ROCPROFILER_CALLBACK_PHASE_EXIT
Callback invoked after to function execution.
Definition fwd.h:145
@ ROCPROFILER_CALLBACK_PHASE_ENTER
Callback invoked prior to function execution.
Definition fwd.h:142
@ ROCPROFILER_CALLBACK_PHASE_LAST
Definition fwd.h:148
@ ROCPROFILER_CALLBACK_PHASE_LOAD
Callback invoked prior to code object loading.
Definition fwd.h:143
@ ROCPROFILER_BUFFER_TRACING_PAGE_MIGRATION
Buffer page migration info.
Definition fwd.h:199
@ ROCPROFILER_BUFFER_TRACING_MARKER_CORE_API
Definition fwd.h:194
@ ROCPROFILER_BUFFER_TRACING_RUNTIME_INITIALIZATION
Record indicating a runtime library has been initialized.
Definition fwd.h:206
@ ROCPROFILER_BUFFER_TRACING_MEMORY_ALLOCATION
Definition fwd.h:204
@ ROCPROFILER_BUFFER_TRACING_MARKER_NAME_API
Definition fwd.h:196
@ ROCPROFILER_BUFFER_TRACING_LAST
Definition fwd.h:209
@ ROCPROFILER_BUFFER_TRACING_CORRELATION_ID_RETIREMENT
Correlation ID in no longer in use.
Definition fwd.h:201
@ ROCPROFILER_BUFFER_TRACING_HSA_CORE_API
Definition fwd.h:187
@ ROCPROFILER_BUFFER_TRACING_MEMORY_COPY
Definition fwd.h:197
@ ROCPROFILER_BUFFER_TRACING_RCCL_API
RCCL tracing.
Definition fwd.h:202
@ ROCPROFILER_BUFFER_TRACING_SCRATCH_MEMORY
Buffer scratch memory reclaimation info.
Definition fwd.h:200
@ ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH
Buffer kernel dispatch info.
Definition fwd.h:198
@ ROCPROFILER_BUFFER_TRACING_HIP_RUNTIME_API
Definition fwd.h:192
@ ROCPROFILER_BUFFER_TRACING_MARKER_CONTROL_API
Definition fwd.h:195
@ ROCPROFILER_BUFFER_TRACING_OMPT
Definition fwd.h:203
@ ROCPROFILER_BUFFER_TRACING_NONE
Definition fwd.h:186
@ ROCPROFILER_BUFFER_TRACING_HSA_IMAGE_EXT_API
Definition fwd.h:189
@ ROCPROFILER_BUFFER_TRACING_HIP_COMPILER_API
Definition fwd.h:193
@ ROCPROFILER_BUFFER_TRACING_HSA_AMD_EXT_API
Definition fwd.h:188
@ ROCPROFILER_BUFFER_TRACING_HSA_FINALIZE_EXT_API
Definition fwd.h:190
@ ROCPROFILER_PC_SAMPLING_RECORD_LAST
Definition fwd.h:449
@ ROCPROFILER_PC_SAMPLING_RECORD_HOST_TRAP_V0_SAMPLE
rocprofiler_pc_sampling_record_host_trap_v0_t
Definition fwd.h:447
@ ROCPROFILER_PC_SAMPLING_RECORD_NONE
Definition fwd.h:446
@ ROCPROFILER_PC_SAMPLING_RECORD_STOCHASTIC_V0_SAMPLE
for the future use
Definition fwd.h:448
@ ROCPROFILER_KERNEL_DISPATCH_ENQUEUE
Definition fwd.h:256
@ ROCPROFILER_KERNEL_DISPATCH_COMPLETE
Definition fwd.h:257
@ ROCPROFILER_KERNEL_DISPATCH_NONE
Unknown kernel dispatch operation.
Definition fwd.h:255
@ ROCPROFILER_KERNEL_DISPATCH_LAST
Definition fwd.h:258
@ ROCPROFILER_RCCL_LIBRARY
Definition fwd.h:369
@ ROCPROFILER_LIBRARY_LAST
Definition fwd.h:370
@ ROCPROFILER_HIP_LIBRARY
Definition fwd.h:367
@ ROCPROFILER_LIBRARY
Definition fwd.h:365
@ ROCPROFILER_HSA_LIBRARY
Definition fwd.h:366
@ ROCPROFILER_MARKER_LIBRARY
Definition fwd.h:368
@ ROCPROFILER_BUFFER_CATEGORY_TRACING
Definition fwd.h:119
@ ROCPROFILER_BUFFER_CATEGORY_NONE
Definition fwd.h:118
@ ROCPROFILER_BUFFER_CATEGORY_PC_SAMPLING
Definition fwd.h:120
@ ROCPROFILER_BUFFER_CATEGORY_LAST
Definition fwd.h:122
@ ROCPROFILER_BUFFER_CATEGORY_COUNTERS
Definition fwd.h:121
@ ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT
Definition fwd.h:168
@ ROCPROFILER_CALLBACK_TRACING_MEMORY_ALLOCATION
Definition fwd.h:174
@ ROCPROFILER_CALLBACK_TRACING_OMPT
Definition fwd.h:173
@ ROCPROFILER_CALLBACK_TRACING_MARKER_CONTROL_API
Definition fwd.h:165
@ ROCPROFILER_CALLBACK_TRACING_KERNEL_DISPATCH
Callbacks for kernel dispatches.
Definition fwd.h:170
@ ROCPROFILER_CALLBACK_TRACING_RCCL_API
RCCL tracing.
Definition fwd.h:172
@ ROCPROFILER_CALLBACK_TRACING_RUNTIME_INITIALIZATION
Callback notifying that a runtime library has been initialized.
Definition fwd.h:176
@ ROCPROFILER_CALLBACK_TRACING_HIP_COMPILER_API
Definition fwd.h:163
@ ROCPROFILER_CALLBACK_TRACING_HSA_AMD_EXT_API
Definition fwd.h:158
@ ROCPROFILER_CALLBACK_TRACING_MARKER_NAME_API
Definition fwd.h:167
@ ROCPROFILER_CALLBACK_TRACING_HSA_CORE_API
Definition fwd.h:157
@ ROCPROFILER_CALLBACK_TRACING_LAST
Definition fwd.h:178
@ ROCPROFILER_CALLBACK_TRACING_MEMORY_COPY
Definition fwd.h:171
@ ROCPROFILER_CALLBACK_TRACING_NONE
Definition fwd.h:156
@ ROCPROFILER_CALLBACK_TRACING_SCRATCH_MEMORY
Definition fwd.h:169
@ ROCPROFILER_CALLBACK_TRACING_HIP_RUNTIME_API
Definition fwd.h:162
@ ROCPROFILER_CALLBACK_TRACING_HSA_IMAGE_EXT_API
Definition fwd.h:159
@ ROCPROFILER_CALLBACK_TRACING_MARKER_CORE_API
Definition fwd.h:164
@ ROCPROFILER_CALLBACK_TRACING_HSA_FINALIZE_EXT_API
Definition fwd.h:160
@ ROCPROFILER_AGENT_TYPE_NONE
Agent type is unknown.
Definition fwd.h:130
@ ROCPROFILER_AGENT_TYPE_GPU
Agent type is a GPU.
Definition fwd.h:132
@ ROCPROFILER_AGENT_TYPE_LAST
Definition fwd.h:133
@ ROCPROFILER_AGENT_TYPE_CPU
Agent type is a CPU.
Definition fwd.h:131
@ ROCPROFILER_COUNTER_INFO_VERSION_0
Definition fwd.h:408
@ ROCPROFILER_COUNTER_INFO_VERSION_LAST
Definition fwd.h:409
@ ROCPROFILER_COUNTER_INFO_VERSION_NONE
Definition fwd.h:407
@ ROCPROFILER_COUNTER_FLAG_LAST
Definition fwd.h:437
@ ROCPROFILER_COUNTER_FLAG_NONE
Definition fwd.h:433
@ ROCPROFILER_COUNTER_FLAG_APPEND_DEFINITION
Append the counter definition to the system provided counter definition file.
Definition fwd.h:435
@ ROCPROFILER_COUNTER_FLAG_ASYNC
Do not wait for completion before returning.
Definition fwd.h:434
Agent Identifier.
Definition fwd.h:578
Context ID.
Definition fwd.h:539
ROCProfiler Record Correlation ID.
Definition fwd.h:555
Counter ID.
Definition fwd.h:586
Counter info struct version 0.
Definition fwd.h:742
Multi-dimensional struct of data used to describe GPU workgroup and grid sizes.
Definition fwd.h:603
ROCProfiler kernel dispatch information.
Definition fwd.h:690
Profile Configurations.
Definition fwd.h:595
ROCProfiler Profile Counting Counter Record per instance.
Definition fwd.h:720
Details for the dimension, including its size, for a counter record.
Definition fwd.h:707
Generic record with type identifier(s) and a pointer to data. This data type is used with buffered da...
Definition fwd.h:654
ROCProfiler SPM Record.
Definition fwd.h:757
Stores memory address for profiling.
Definition fwd.h:524
User-assignable data type.
Definition fwd.h:514