rocprofiler-sdk/callback_tracing.h Source File

rocprofiler-sdk/callback_tracing.h Source File#

Rocprofiler SDK Developer API: rocprofiler-sdk/callback_tracing.h Source File
Rocprofiler SDK Developer API 0.6.0
ROCm Profiling API and tools
callback_tracing.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#include <rocprofiler-sdk/fwd.h>
27#include <rocprofiler-sdk/hip.h>
28#include <rocprofiler-sdk/hsa.h>
32
33#include <hsa/hsa.h>
34#include <hsa/hsa_amd_tool.h>
35#include <hsa/hsa_ext_amd.h>
36#include <hsa/hsa_ven_amd_loader.h>
37
38#include <stdint.h>
39
40ROCPROFILER_EXTERN_C_INIT
41
42/**
43 * @defgroup CALLBACK_TRACING_SERVICE Synchronous Tracing Services
44 * @brief Receive immediate callbacks on the calling thread
45 *
46 * @{
47 */
48
49/**
50 * @brief ROCProfiler Enumeration for code object storage types (identical values to
51 * `hsa_ven_amd_loader_code_object_storage_type_t` enumeration)
52 */
53typedef enum
54{
55 ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_NONE = HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_NONE,
56 ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_FILE = HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_FILE,
58 HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_MEMORY,
61
62/**
63 * @brief ROCProfiler HSA API Callback Data.
64 */
71
72/**
73 * @brief ROCProfiler HIP runtime and compiler API Tracer Callback Data.
74 */
81
82/**
83 * @brief ROCProfiler OMPT Callback Data
84 */
85typedef struct
86{
87 uint64_t size; ///< size of this struct
90
91/**
92 * @brief ROCProfiler Marker Tracer Callback Data.
93 */
100
101/**
102 * @brief ROCProfiler RCCL API Callback Data.
103 */
110
111/**
112 * @brief ROCProfiler Code Object Load Tracer Callback Record.
113 */
114typedef struct
115{
116 uint64_t size; ///< size of this struct
117 uint64_t code_object_id; ///< unique code object identifier
118 rocprofiler_agent_id_t rocp_agent; ///< The agent on which this loaded code object is loaded
119 hsa_agent_t hsa_agent; ///< The agent on which this loaded code object is loaded
120 const char* uri; ///< The URI name from which the code object was loaded
121 uint64_t load_base; ///< The base memory address at which the code object is loaded. This is
122 ///< the base address of the allocation for the lowest addressed segment of
123 ///< the code object that is loaded. Note that any non-loaded segments
124 ///< before the first loaded segment are ignored.
125 uint64_t load_size; ///< The byte size of the loaded code objects contiguous memory allocation.
126 int64_t load_delta; ///< The signed byte address difference of the memory address at which the
127 ///< code object is loaded minus the virtual address specified in the code
128 ///< object that is loaded.
130 storage_type; ///< storage type of the code object reader used to load the loaded code
131 ///< object
132 union
133 {
134 struct
135 {
136 int storage_file; ///< file descriptor of the code object that was loaded. Access this
137 ///< field if @ref rocprofiler_code_object_storage_type_t is
138 ///< @ref ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_FILE
139 };
140 struct
141 {
142 uint64_t memory_base; ///< The memory address of the first byte of the code object that
143 ///< was loaded. Access this
144 ///< field if @ref rocprofiler_code_object_storage_type_t is
145 ///< @ref ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_MEMORY
146 uint64_t memory_size; ///< The memory size in bytes of the code object that was loaded.
147 ///< Access this field if @ref
148 ///< rocprofiler_code_object_storage_type_t is
149 ///< @ref ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_MEMORY
150 };
151 };
153
154/**
155 * @brief The NULL value of a code object id. Used when code object is unknown.
156 */
157#define ROCPROFILER_CODE_OBJECT_ID_NONE ROCPROFILER_UINT64_C(0)
158
159/**
160 * @brief ROCProfiler Code Object Kernel Symbol Tracer Callback Record.
161 *
162 */
163typedef struct
164{
165 uint64_t size; ///< size of this struct
166 uint64_t kernel_id; ///< unique symbol identifier value
167 uint64_t code_object_id; ///< parent unique code object identifier
168 const char* kernel_name; ///< name of the kernel
169 uint64_t kernel_object; ///< kernel object handle, used in the kernel dispatch packet
170 uint32_t kernarg_segment_size; ///< size of memory (in bytes) allocated for kernel arguments.
171 ///< Will be multiple of 16
172 uint32_t kernarg_segment_alignment; ///< Alignment (in bytes) of the buffer used to pass
173 ///< arguments to the kernel
174 uint32_t group_segment_size; ///< Size of static group segment memory required by the kernel
175 ///< (per work-group), in bytes. AKA: LDS size
176 uint32_t private_segment_size; ///< Size of static private, spill, and arg segment memory
177 ///< required by this kernel (per work-item), in bytes. AKA:
178 ///< scratch size
179 uint32_t sgpr_count; ///< Scalar general purpose register count
180 uint32_t arch_vgpr_count; ///< Architecture vector general purpose register count
181 uint32_t accum_vgpr_count; ///< Accum vector general purpose register count
182
184// rename struct
185
186typedef struct
187{
188 uint64_t size; ///< size of this struct
189 uint64_t host_function_id; ///< unique host function identifier value
190 uint64_t kernel_id; ///< unique symbol identifier value
191 uint64_t code_object_id; ///< parent unique code object identifier
192 rocprofiler_address_t host_function; ///< kernel host function pointer
193 rocprofiler_address_t modules; ///< reference address where modules will be loaded
194 const char* device_function;
195 uint32_t thread_limit; ///< thread limit
196 rocprofiler_dim3_t thread_ids; ///< thread ids address
197 rocprofiler_dim3_t block_ids; ///< block ids address
198 rocprofiler_dim3_t block_dims; ///< block dimensions address
199 rocprofiler_dim3_t grid_dims; ///< grid dimensions address
200 uint64_t workgroup_size; ///< workgroup size address
201
202 /// @var device_function
203 /// @brief device function name used to map the metadata during kernel launch
205
206/**
207 * @brief ROCProfiler Kernel Dispatch Callback Tracer Record.
208 *
209 */
217
218/**
219 * @brief ROCProfiler Memory Copy Callback Tracer Record.
220 *
221 * The timestamps in this record will only be non-zero in the ::ROCPROFILER_CALLBACK_PHASE_EXIT
222 * callback
223 */
224typedef struct
225{
226 uint64_t size; ///< size of this struct
227 rocprofiler_timestamp_t start_timestamp; ///< start time in nanoseconds
228 rocprofiler_timestamp_t end_timestamp; ///< end time in nanoseconds
229 rocprofiler_agent_id_t dst_agent_id; ///< destination agent of copy
230 rocprofiler_agent_id_t src_agent_id; ///< source agent of copy
231 uint64_t bytes; ///< bytes copied
233
234/**
235 * @brief ROCProfiler Memory Allocation Tracer Record.
236 */
237typedef struct
238{
239 uint64_t size; ///< size of this struct
240 rocprofiler_timestamp_t start_timestamp; ///< start time in nanoseconds
241 rocprofiler_timestamp_t end_timestamp; ///< end time in nanoseconds
242 rocprofiler_agent_id_t agent_id; ///< agent id for memory allocation
243 rocprofiler_address_t address; ///< starting address for memory allocation
244 uint64_t allocation_size; ///< size of memory allocation
246
247/**
248 * @brief ROCProfiler Scratch Memory Callback Data.
249 */
259
260/**
261 * @brief ROCProfiler Runtime Initialization Data.
262 */
264{
265 uint64_t size; ///< size of this struct
266 uint64_t version;
267 uint64_t instance; ///< Number of times this runtime had been loaded previously
268
269 /// @var version
270 /// @brief The version number of the library
271 ///
272 /// Version number is encoded as: (10000 * MAJOR) + (100 * MINOR) + PATCH
274
275/**
276 * @brief API Tracing callback function. This function is invoked twice per API function: once
277 * before the function is invoked and once after the function is invoked. The external correlation
278 * id value within the record is assigned the value at the top of the external correlation id stack.
279 * It is permissible to invoke @ref rocprofiler_push_external_correlation_id within the enter phase;
280 * when a new external correlation id is pushed during the enter phase, rocprofiler will use that
281 * external correlation id for any async events and provide the new external correlation id during
282 * the exit callback... In other words, pushing a new external correlation id within the enter
283 * callback will result in that external correlation id value in the exit callback (which may or may
284 * not be different from the external correlation id value in the enter callback). If a tool pushes
285 * new external correlation ids in the enter phase, it is recommended to pop the external
286 * correlation id in the exit callback.
287 *
288 * @param [in] record Callback record data
289 * @param [in,out] user_data This paramter can be used to retain information in between the enter
290 * and exit phases.
291 * @param [in] callback_data User data provided when configuring the callback tracing service
292 */
294 rocprofiler_user_data_t* user_data,
295 void* callback_data) ROCPROFILER_NONNULL(2);
296
297/**
298 * @brief Callback function for mapping @ref rocprofiler_callback_tracing_kind_t ids to
299 * string names. @see rocprofiler_iterate_callback_tracing_kind_names.
300 */
302 void* data);
303
304/**
305 * @brief Callback function for mapping the operations of a given @ref
306 * rocprofiler_callback_tracing_kind_t to string names. @see
307 * rocprofiler_iterate_callback_tracing_kind_operation_names.
308 */
312 void* data);
313
314/**
315 * @brief Callback function for iterating over the function arguments to a traced function.
316 * This function will be invoked for each argument.
317 * @see rocprofiler_iterate_callback_tracing_operation_args
318 *
319 * @param [in] kind domain
320 * @param [in] operation associated domain operation
321 * @param [in] arg_number the argument number, starting at zero
322 * @param [in] arg_value_addr the address of the argument stored by rocprofiler.
323 * @param [in] arg_indirection_count the total number of indirection levels for the argument, e.g.
324 * int == 0, int* == 1, int** == 2
325 * @param [in] arg_type the typeid name of the argument
326 * @param [in] arg_name the name of the argument in the prototype (or rocprofiler union)
327 * @param [in] arg_value_str conversion of the argument to a string, e.g. operator<< overload
328 * @param [in] arg_dereference_count the number of times the argument was dereferenced when it was
329 * converted to a string
330 * @param [in] data user data
331 */
335 uint32_t arg_number,
336 const void* const arg_value_addr,
337 int32_t arg_indirection_count,
338 const char* arg_type,
339 const char* arg_name,
340 const char* arg_value_str,
341 int32_t arg_dereference_count,
342 void* data);
343
344/**
345 * @brief Configure Callback Tracing Service. The callback tracing service provides two synchronous
346 * callbacks around an API function on the same thread as the application which is invoking the API
347 * function. This function can only be invoked once per @ref
348 * rocprofiler_callback_tracing_kind_t value, i.e. it can be invoked once for the HSA API,
349 * once for the HIP API, and so on but it will fail if it is invoked for the HSA API twice. Please
350 * note, the callback API does have the potentially non-trivial overhead of copying the function
351 * arguments into the record. If you are willing to let rocprofiler record the timestamps, do not
352 * require synchronous notifications of the API calls, and want to lowest possible overhead, use the
353 * @see BUFFER_TRACING_SERVICE.
354 *
355 * @param [in] context_id Context to associate the service with
356 * @param [in] kind The domain of the callback tracing service
357 * @param [in] operations Array of operations in the domain (i.e. enum values which identify
358 * specific API functions). If this is null, all API functions in the domain will be traced
359 * @param [in] operations_count If the operations array is non-null, set this to the size of the
360 * array.
361 * @param [in] callback The function to invoke before and after an API function
362 * @param [in] callback_args Data provided to every invocation of the callback function
363 * @return ::rocprofiler_status_t
364 * @retval ::ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED Invoked outside of the initialization
365 * function in @ref rocprofiler_tool_configure_result_t provided to rocprofiler via @ref
366 * rocprofiler_configure function
367 * @retval ::ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND The provided context is not valid/registered
368 * @retval ::ROCPROFILER_STATUS_ERROR_SERVICE_ALREADY_CONFIGURED if the same @ref
369 * rocprofiler_callback_tracing_kind_t value is provided more than once (per context) -- in
370 * other words, we do not support overriding or combining the operations in separate function calls.
371 *
372 */
376 const rocprofiler_tracing_operation_t* operations,
377 size_t operations_count,
379 void* callback_args) ROCPROFILER_API;
380
381/**
382 * @brief Query the name of the callback tracing kind. The name retrieved from this function is a
383 * string literal that is encoded in the read-only section of the binary (i.e. it is always
384 * "allocated" and never "deallocated").
385 *
386 * @param [in] kind Callback tracing domain
387 * @param [out] name If non-null and the name is a constant string that does not require dynamic
388 * allocation, this paramter will be set to the address of the string literal, otherwise it will
389 * be set to nullptr
390 * @param [out] name_len If non-null, this will be assigned the length of the name (regardless of
391 * the name is a constant string or requires dynamic allocation)
392 * @return ::rocprofiler_status_t
393 */
396 const char** name,
397 uint64_t* name_len) ROCPROFILER_API;
398
399/**
400 * @brief Query the name of the callback tracing kind. The name retrieved from this function is a
401 * string literal that is encoded in the read-only section of the binary (i.e. it is always
402 * "allocated" and never "deallocated").
403 *
404 * @param [in] kind Callback tracing domain
405 * @param [in] operation Enumeration id value which maps to a specific API function or event type
406 * @param [out] name If non-null and the name is a constant string that does not require dynamic
407 * allocation, this paramter will be set to the address of the string literal, otherwise it will
408 * be set to nullptr
409 * @param [out] name_len If non-null, this will be assigned the length of the name (regardless of
410 * the name is a constant string or requires dynamic allocation)
411 * @return ::rocprofiler_status_t
412 * @retval ::ROCPROFILER_STATUS_ERROR_KIND_NOT_FOUND Domain id is not valid
413 * @retval ::ROCPROFILER_STATUS_SUCCESS Valid domain provided, regardless if there is a constant
414 * string or not.
415 */
419 const char** name,
420 uint64_t* name_len) ROCPROFILER_API;
421
422/**
423 * @brief Iterate over all the mappings of the callback tracing kinds and get a callback for each
424 * kind.
425 *
426 * @param [in] callback Callback function invoked for each enumeration value in @ref
427 * rocprofiler_callback_tracing_kind_t with the exception of the `NONE` and `LAST` values.
428 * @param [in] data User data passed back into the callback
429 * @return ::rocprofiler_status_t
430 */
433 void* data) ROCPROFILER_API ROCPROFILER_NONNULL(1);
434
435/**
436 * @brief Iterates over all the mappings of the operations for a given @ref
437 * rocprofiler_callback_tracing_kind_t and invokes the callback with the kind id, operation
438 * id, and user-provided data.
439 *
440 * @param [in] kind which tracing callback kind operations to iterate over
441 * @param [in] callback Callback function invoked for each operation associated with @ref
442 * rocprofiler_callback_tracing_kind_t with the exception of the `NONE` and `LAST` values.
443 * @param [in] data User data passed back into the callback
444 * @return ::rocprofiler_status_t
445 * @retval ::ROCPROFILER_STATUS_ERROR_KIND_NOT_FOUND Invalid domain id
446 * @retval ::ROCPROFILER_STATUS_SUCCESS Valid domain
447 */
452 void* data) ROCPROFILER_API ROCPROFILER_NONNULL(2);
453
454/**
455 * @brief Iterates over all the arguments for the traced function (when available). This is
456 * particularly useful when tools want to annotate traces with the function arguments. See
457 * @example samples/api_callback_tracing/client.cpp for a usage example.
458 *
459 * It is recommended to use this function when the record phase is ::ROCPROFILER_CALLBACK_PHASE_EXIT
460 * or ::ROCPROFILER_CALLBACK_PHASE_NONE. When the phase is ::ROCPROFILER_CALLBACK_PHASE_ENTER, the
461 * function may have output parameters which have not set. In the case of an output parameter with
462 * one level of indirection, e.g. `int* output_len`, this is considered safe since the output
463 * parameter is either null or, in the worst case scenario, pointing to an uninitialized value which
464 * will result in garbage values to be stringified. However, if the output parameter has more than
465 * one level of indirection, e.g. `const char** output_name`, this can result in a segmentation
466 * fault because the dereferenced output parameter may be uninitialized and point to an invalid
467 address. E.g.:
468 *
469 * @code{.cpp}
470 * struct dim3
471 * {
472 * int x;
473 * int y;
474 * int z;
475 * };
476 *
477 * static dim3 default_dims = {.x = 1, .y = 1, .z = 1};
478 *
479 * void set_dim_x(int val, dim3* output_dims) { output_dims->x = val; }
480 *
481 * void get_default_dims(dim3** output_dims) { *output_dims = default_dims; }
482 *
483 * int main()
484 * {
485 * dim3 my_dims; // uninitialized value. x, y, and z may be set to random values
486 * dim3* current_dims; // uninitialized pointer. May be set to invalid address
487 *
488 * set_dim_x(3, &my_dims); // if rocprofiler-sdk wrapped this function and tried to stringify
489 * // in the enter phase, dereferencing my_dims is not problematic
490 * // since there is an actual dim3 allocation
491 *
492 * get_default_dims(&current_dims); // if rocprofiler-sdk wrapped this function,
493 * // and tried to stringify in the enter phase,
494 * // current_dims may point to an address outside
495 * // of the address space of this process and
496 * // cause a segfault
497 * }
498 * @endcode
499 *
500 *
501 * @param[in] record Record provided by service callback
502 * @param[in] callback The callback function which will be invoked for each argument
503 * @param[in] max_dereference_count In the callback enter phase, certain arguments may be output
504 * parameters which have not been set. When the output parameter has multiple levels of indirection,
505 * it may be invalid to dereference the output parameter more than once and doing so may result in a
506 * segmentation fault. Thus, it is recommended to set this parameter to a maximum value of 1 when
507 * the phase is ::ROCPROFILER_CALLBACK_PHASE_ENTER to ensure that output parameters which point to
508 * uninitialized pointers do not cause segmentation faults.
509 * @param[in] user_data Data to be passed to each invocation of the callback
510 */
515 int32_t max_dereference_count,
516 void* user_data) ROCPROFILER_API ROCPROFILER_NONNULL(2);
517
518/** @} */
519
520ROCPROFILER_EXTERN_C_FINI
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_status_t
Status codes.
Definition fwd.h:53
uint64_t rocprofiler_timestamp_t
ROCProfiler Timestamp.
Definition fwd.h:461
rocprofiler_callback_tracing_kind_t
Service Callback Tracing Kind.
Definition fwd.h:155
Agent Identifier.
Definition fwd.h:578
Context ID.
Definition fwd.h:539
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
Stores memory address for profiling.
Definition fwd.h:524
User-assignable data type.
Definition fwd.h:514
uint32_t group_segment_size
Size of static group segment memory required by the kernel (per work-group), in bytes....
const char * device_function
device function name used to map the metadata during kernel launch
uint64_t version
The version number of the library.
rocprofiler_agent_id_t rocp_agent
The agent on which this loaded code object is loaded.
rocprofiler_timestamp_t start_timestamp
start time in nanoseconds
uint64_t code_object_id
unique code object identifier
rocprofiler_address_t modules
reference address where modules will be loaded
uint32_t arch_vgpr_count
Architecture vector general purpose register count.
rocprofiler_timestamp_t end_timestamp
end time in nanoseconds
rocprofiler_kernel_dispatch_info_t dispatch_info
Dispatch info.
uint32_t kernarg_segment_alignment
Alignment (in bytes) of the buffer used to pass arguments to the kernel.
rocprofiler_timestamp_t start_timestamp
start time in nanoseconds
uint32_t accum_vgpr_count
Accum vector general purpose register count.
rocprofiler_agent_id_t agent_id
agent id for memory allocation
rocprofiler_code_object_storage_type_t storage_type
storage type of the code object reader used to load the loaded code object
uint64_t instance
Number of times this runtime had been loaded previously.
rocprofiler_timestamp_t start_timestamp
start time in nanoseconds
hsa_agent_t hsa_agent
The agent on which this loaded code object is loaded.
uint64_t load_base
The base memory address at which the code object is loaded. This is the base address of the allocatio...
uint32_t kernarg_segment_size
size of memory (in bytes) allocated for kernel arguments. Will be multiple of 16
uint64_t load_size
The byte size of the loaded code objects contiguous memory allocation.
uint32_t private_segment_size
Size of static private, spill, and arg segment memory required by this kernel (per work-item),...
rocprofiler_address_t address
starting address for memory allocation
const char * uri
The URI name from which the code object was loaded.
rocprofiler_timestamp_t end_timestamp
end time in nanoseconds
rocprofiler_timestamp_t end_timestamp
end time in nanoseconds
uint64_t kernel_object
kernel object handle, used in the kernel dispatch packet
rocprofiler_agent_id_t dst_agent_id
destination agent of copy
int64_t load_delta
The signed byte address difference of the memory address at which the code object is loaded minus the...
rocprofiler_agent_id_t src_agent_id
source agent of copy
rocprofiler_status_t rocprofiler_query_callback_tracing_kind_operation_name(rocprofiler_callback_tracing_kind_t kind, rocprofiler_tracing_operation_t operation, const char **name, uint64_t *name_len)
Query the name of the callback tracing kind. The name retrieved from this function is a string litera...
rocprofiler_status_t rocprofiler_iterate_callback_tracing_kind_operation_args(rocprofiler_callback_tracing_record_t record, rocprofiler_callback_tracing_operation_args_cb_t callback, int32_t max_dereference_count, void *user_data)
int(* rocprofiler_callback_tracing_kind_cb_t)(rocprofiler_callback_tracing_kind_t kind, void *data)
Callback function for mapping rocprofiler_callback_tracing_kind_t ids to string names.
int(* rocprofiler_callback_tracing_kind_operation_cb_t)(rocprofiler_callback_tracing_kind_t kind, rocprofiler_tracing_operation_t operation, void *data)
Callback function for mapping the operations of a given rocprofiler_callback_tracing_kind_t to string...
int(* rocprofiler_callback_tracing_operation_args_cb_t)(rocprofiler_callback_tracing_kind_t kind, rocprofiler_tracing_operation_t operation, uint32_t arg_number, const void *const arg_value_addr, int32_t arg_indirection_count, const char *arg_type, const char *arg_name, const char *arg_value_str, int32_t arg_dereference_count, void *data)
Callback function for iterating over the function arguments to a traced function. This function will ...
rocprofiler_status_t rocprofiler_configure_callback_tracing_service(rocprofiler_context_id_t context_id, rocprofiler_callback_tracing_kind_t kind, const rocprofiler_tracing_operation_t *operations, unsigned long operations_count, rocprofiler_callback_tracing_cb_t callback, void *callback_args)
Configure Callback Tracing Service. The callback tracing service provides two synchronous callbacks a...
rocprofiler_status_t rocprofiler_query_callback_tracing_kind_name(rocprofiler_callback_tracing_kind_t kind, const char **name, uint64_t *name_len)
Query the name of the callback tracing kind. The name retrieved from this function is a string litera...
rocprofiler_code_object_storage_type_t
ROCProfiler Enumeration for code object storage types (identical values to hsa_ven_amd_loader_code_ob...
rocprofiler_status_t rocprofiler_iterate_callback_tracing_kind_operations(rocprofiler_callback_tracing_kind_t kind, rocprofiler_callback_tracing_kind_operation_cb_t callback, void *data)
Iterates over all the mappings of the operations for a given rocprofiler_callback_tracing_kind_t and ...
rocprofiler_status_t rocprofiler_iterate_callback_tracing_kinds(rocprofiler_callback_tracing_kind_cb_t callback, void *data)
Iterate over all the mappings of the callback tracing kinds and get a callback for each kind.
void(* rocprofiler_callback_tracing_cb_t)(rocprofiler_callback_tracing_record_t record, rocprofiler_user_data_t *user_data, void *callback_data)
API Tracing callback function. This function is invoked twice per API function: once before the funct...
@ ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_NONE
@ ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_MEMORY
@ ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_FILE
@ ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_LAST
ROCProfiler Code Object Kernel Symbol Tracer Callback Record.
ROCProfiler Code Object Load Tracer Callback Record.
ROCProfiler HIP runtime and compiler API Tracer Callback Data.
ROCProfiler HSA API Callback Data.
ROCProfiler Kernel Dispatch Callback Tracer Record.
ROCProfiler Marker Tracer Callback Data.
ROCProfiler Memory Allocation Tracer Record.
ROCProfiler Memory Copy Callback Tracer Record.
ROCProfiler OMPT Callback Data.
ROCProfiler RCCL API Callback Data.
ROCProfiler Scratch Memory Callback Data.
rocprofiler_scratch_alloc_flag_t
Allocation flags for.