ROCprofiler-SDK

https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/
Admin
ROCprofiler-SDK is a tooling infrastructure for profiling GPU compute applications on ROCm that...
Tokens:528,941
Snippets:5,426
Trust Score:10
Update:5 months ago
Show doc for...
Tokens:
Raw
### Tool Initialization and Context Setup

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/_doxygen/rocprofiler-sdk/html/samples_2api_buffered_tracing_2client_8cpp-example.html

Initializes the ROCProfiler client, configures tracing services, creates a buffer, and starts the profiling context. Includes error checking for context validity.

```cpp
int
tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data)
{
assert(tool_data != nullptr);
auto* call_stack_v = static_cast<call_stack_t*>(tool_data);
call_stack_v->emplace_back(source_location{__FUNCTION__, __FILE__, __LINE__, ""});
client_name_info = common::get_buffer_tracing_names();
for(const auto& itr : client_name_info)
{
auto name_idx = std::stringstream{};
name_idx << " [" << std::setw(3) << itr.value << "]";
call_stack_v->emplace_back(
source_location{"rocprofiler_buffer_tracing_kind_names " + name_idx.str(),
__FILE__,
__LINE__,
std::string{itr.name}});
for(auto [didx, ditr] : itr.items())
{
auto operation_idx = std::stringstream{};
operation_idx << " [" << std::setw(3) << didx << "]";
call_stack_v->emplace_back(source_location{
"rocprofiler_buffer_tracing_kind_operation_names" + operation_idx.str(),
__FILE__,
__LINE__,
std::string{"-" + std::string{*ditr}}});
}
}
client_fini_func = fini_func;
ROCPROFILER_CALL(rocprofiler_create_context(&client_ctx), "context creation");
auto code_object_ops = std::vector<rocprofiler_tracing_operation_t>{
ROCPROFILER_CODE_OBJECT_DEVICE_KERNEL_SYMBOL_REGISTER};
ROCPROFILER_CALL(
rocprofiler_configure_callback_tracing_service(client_ctx,
ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT,
code_object_ops.data(),
code_object_ops.size(),
tool_code_object_callback,
nullptr),
"code object tracing service configure");
constexpr auto buffer_size_bytes = 4096;
constexpr auto buffer_watermark_bytes = buffer_size_bytes - (buffer_size_bytes / 8);
ROCPROFILER_CALL(rocprofiler_create_buffer(client_ctx,
buffer_size_bytes,
buffer_watermark_bytes,
ROCPROFILER_BUFFER_POLICY_LOSSLESS,
tool_tracing_callback,
tool_data,
&client_buffer),
"buffer creation");
for(auto itr :
{ROCPROFILER_BUFFER_TRACING_HSA_CORE_API, ROCPROFILER_BUFFER_TRACING_HSA_AMD_EXT_API})
{
ROCPROFILER_CALL(
rocprofiler_configure_buffer_tracing_service(
client_ctx, itr, nullptr, 0, client_buffer),
"buffer tracing service configure");
}
ROCPROFILER_CALL(
rocprofiler_configure_buffer_tracing_service(
client_ctx, ROCPROFILER_BUFFER_TRACING_HIP_RUNTIME_API, nullptr, 0, client_buffer),
"buffer tracing service configure");
ROCPROFILER_CALL(
rocprofiler_configure_buffer_tracing_service(
client_ctx, ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH, nullptr, 0, client_buffer),
"buffer tracing service for kernel dispatch configure");
ROCPROFILER_CALL(
rocprofiler_configure_buffer_tracing_service(
client_ctx, ROCPROFILER_BUFFER_TRACING_MEMORY_COPY, nullptr, 0, client_buffer),
"buffer tracing service for memory copy configure");
ROCPROFILER_CALL(
rocprofiler_configure_buffer_tracing_service(
client_ctx, ROCPROFILER_BUFFER_TRACING_SCRATCH_MEMORY, nullptr, 0, client_buffer),
"buffer tracing service for scratch memory configure");
auto client_thread = rocprofiler_callback_thread_t{};
ROCPROFILER_CALL(rocprofiler_create_callback_thread(&client_thread),
"creating callback thread");
ROCPROFILER_CALL(rocprofiler_assign_callback_thread(client_buffer, client_thread),
"assignment of thread for buffer");
int valid_ctx = 0;
ROCPROFILER_CALL(rocprofiler_context_is_valid(client_ctx, &valid_ctx),
"context validity check");
if(valid_ctx == 0)
{
// notify rocprofiler that initialization failed
// and all the contexts, buffers, etc. created
// should be ignored
return -1;
}
ROCPROFILER_CALL(rocprofiler_start_context(client_ctx), "rocprofiler context start");
return 0;
}
```

--------------------------------

### Start Target Application

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/how-to/using-rocprofv3-process-attachment.html

Start your target application in the background before attaching the profiler.

```bash
./myapp -n 1 &
```

--------------------------------

### Tool Initialization Function Example

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/api-reference/rocprofiler-sdk_api/modules/tool_registration.html

Example of a tool's initialization function. It receives a finalize function pointer and tool-specific data. It creates and starts a rocProfiler context, returning -1 on failure to notify rocProfiler to ignore created resources.

```c++
#include <rocprofiler-sdk/registration.h>

static rocprofiler_client_id_t       my_client_id;
static rocprofiler_client_finalize_t my_fini_func;
static int                           my_tool_data = 1234;

static int my_init_func(rocprofiler_client_finalize_t fini_func,
                        void* tool_data)
{
     my_fini_func = fini_func;

     assert(*static_cast<int*>(tool_data) == 1234 && "tool_data is wrong");

     rocprofiler_context_id_t ctx{0};
     rocprofiler_create_context(&ctx);

     if(int valid_ctx = 0;
        rocprofiler_context_is_valid(ctx, &valid_ctx) != ROCPROFILER_STATUS_SUCCESS ||
        valid_ctx != 0)
     {
         // notify rocprofiler that initialization failed
         // and all the contexts, buffers, etc. created
         // should be ignored
         return -1;
     }

     if(rocprofiler_start_context(ctx) != ROCPROFILER_STATUS_SUCCESS)
     {
         // notify rocprofiler that initialization failed
         // and all the contexts, buffers, etc. created
         // should be ignored
         return -1;
     }

     // no errors
     return 0;
}

static int my_fini_func(void* tool_data)
{
     assert(*static_cast<int*>(tool_data) == 1234 && "tool_data is wrong");
}

rocprofiler_tool_configure_result_t*

```

--------------------------------

### Install ROCprofiler-SDK

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/install/installation.html

Install ROCprofiler-SDK from the build directory using CMake.

```bash
cmake --build rocprofiler-sdk-build --target install
```

--------------------------------

### Trace Decoder Setup

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/api-reference/thread_trace.html

Provides functions to create and initialize a Trace Decoder for processing raw thread trace data. Requires the ROCprof Trace Decoder package to be installed separately.

```APIDOC
## Trace Decoder Setup

### Description
This section covers the setup and usage of the ROCprof Trace Decoder, which is necessary for processing raw thread trace data. The decoder library must be installed separately.

### Create Trace Decoder

#### Method
`rocprofiler_thread_trace_decoder_create`

#### Parameters
- `decoder` (rocprofiler_thread_trace_decoder_id_t*) - Pointer to store the created decoder ID.
- `decoder_lib_path` (const char*) - Path to the ROCprof Trace Decoder library.

#### Request Example
```c++
rocprofiler_thread_trace_decoder_id_t decoder{};

// Create the Trace Decoder with the path to the decoder library
ROCPROFILER_CALL(
    rocprofiler_thread_trace_decoder_create(&decoder, "/opt/rocm/lib"),
    "thread trace decoder creation");
```

### Code Object Load Information

#### Method
`rocprofiler_thread_trace_decoder_codeobj_load`

#### Parameters
- `decoder` (rocprofiler_thread_trace_decoder_id_t) - The decoder ID.
- `code_object_id` (uint64_t) - The ID of the code object.
- `load_delta` (uint64_t) - The load delta.
- `load_size` (uint64_t) - The load size.
- `data` (const void*) - Pointer to the code object data.
- `datasize` (size_t) - Size of the code object data.

#### Request Example
```c++
// Adds code object load information, reported by the code object tracing service
ROCPROFILER_CALL(rocprofiler_thread_trace_decoder_codeobj_load(decoder,
                                                               code_object_id,
                                                               load_delta,
                                                               load_size,
                                                               data,
                                                               datasize),
                "code object load");
```
```

--------------------------------

### Install Debian/Ubuntu Dependencies

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/install/installation.html

Install necessary development libraries for ROCprofiler-SDK on Debian/Ubuntu systems.

```bash
sudo apt install -y libdw-dev libsqlite3-dev
```

--------------------------------

### rocprofv3 Version Output Example

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/how-to/using-rocprofv3.html

This is an example output of the `rocprofv3 -v` command, showing detailed build and system information.

```bash
$ rocprofv3 -v
             version: 1.0.0
        git_revision: a1b2c3d4e5f6789012345678901234567890abcd
        library_arch: x86_64-linux-gnu
         system_name: Linux
    system_processor: x86_64
      system_version: 6.8.0-57-generic
         compiler_id: GNU
    compiler_version: 11.4.0
        rocm_version: 6.2.0

```

--------------------------------

### Install SUSE Dependencies and Set Compilers

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/install/installation.html

Install build tools and development libraries on SUSE systems, and set C++ and C compiler environment variables.

```bash
sudo zypper install gcc12 gcc12-c++ cmake make python3-devel elfutils sqlite3-devel libelf-devel libdw-devel
export CXX=/usr/bin/g++-12
export CC=/usr/bin/gcc-12
```

--------------------------------

### Check rocprofv3 Version and Help

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/quick_guide.html

Verify the installed rocprofv3 version and access its help documentation. Useful for confirming installation and understanding available options.

```bash
rocprofv3 --version
rocprofv3 --help

```

--------------------------------

### Install ROCprofiler-SDK via APT

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/install/installation.html

Install a prebuilt version of ROCprofiler-SDK on Ubuntu using the APT package manager, assuming ROCm 6.2 or later is installed.

```bash
$ sudo apt install rocprofiler-sdk
```

--------------------------------

### Example Usage: api_registration_callback and rocprofiler_configure

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/_doxygen/rocprofiler-sdk/html/intercept__table_8h_source.html

Demonstrates how to use `rocprofiler_at_intercept_table_registration` within the `rocprofiler_configure` function to intercept HSA API calls and install wrappers.

```APIDOC
## Example: Intercepting HSA API Calls

This example shows a typical implementation of `rocprofiler_configure` and an `api_registration_callback` to intercept HSA runtime functions.

### `rocprofiler_configure` Function

```cpp
extern "C" rocprofiler_tool_configure_result_t*
rocprofiler_configure(uint32_t version,
 const char* runtime_version,
 uint32_t priority,
 rocprofiler_client_id_t* id)
{
 // set the client name
 id->name = "ExampleTool";

 // specify that we only want to intercept the HSA library
 rocprofiler_at_intercept_table_registration(api_registration_callback,
 ROCPROFILER_HSA_TABLE, nullptr);

 return nullptr;
}
```

### `api_registration_callback` Function

This function is called by rocprofiler when the HSA intercept table is registered. It iterates through the HSA API functions and installs wrappers using the `GENERATE_WRAPPER` macro.

```cpp
void
api_registration_callback(rocprofiler_intercept_table_t type,
 uint64_t lib_version,
 uint64_t lib_instance,
 void** tables,
 uint64_t num_tables,
 void* user_data)
{
 if(type != ROCPROFILER_HSA_TABLE)
 throw std::runtime_error{"unexpected library type: " +
 std::to_string(static_cast<int>(type))};
 if(lib_instance != 0) throw std::runtime_error{"multiple instances of HSA runtime library"};
 if(num_tables != 1) throw std::runtime_error{"expected only one table of type HsaApiTable"};

 auto* hsa_api_table = static_cast<HsaApiTable*>(tables[0]);
 // Example of wrapping a few HSA functions:
 GENERATE_WRAPPER(hsa_api_table->core_, hsa_agent_get_info);
 GENERATE_WRAPPER(hsa_api_table->core_, hsa_queue_create);
 GENERATE_WRAPPER(hsa_api_table->core_, hsa_signal_wait_scacquire);
 // ... other GENERATE_WRAPPER calls for various HSA functions ...
}
```

### `GENERATE_WRAPPER` Macro and `generate_wrapper` Function

The `GENERATE_WRAPPER` macro and the `generate_wrapper` template function are internal helpers used to create and install wrappers around original library functions. The wrapper typically logs the function call and then invokes the original function.
```

--------------------------------

### Get PC Sampling Configurations

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/how-to/using-rocprofv3-avail.html

Use 'info --pc-sampling' to get a list of supported PC sampling configurations for each agent. The '-d' option is not applicable here.

```bash
rocprofv3-avail info --pc-sampling

```

--------------------------------

### Example Output File Names

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/how-to/using-rocprofv3-with-mpi.html

These are example file names generated by rocprofv3 with the specified output directory format for MPI processes.

```text
ubuntu-latest.3000020.1/ubuntu-latest/3000020_agent_info.csv
ubuntu-latest.3000019.0/ubuntu-latest/3000019_agent_info.csv
ubuntu-latest.3000020.1/ubuntu-latest/3000020_hip_api_trace.csv
ubuntu-latest.3000019.0/ubuntu-latest/3000019_hip_api_trace.csv

```

--------------------------------

### Install RHEL/Fedora Dependencies

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/install/installation.html

Install build tools, development libraries, and Python packages required for ROCprofiler-SDK on Red Hat-based systems.

```bash
sudo dnf install elfutils elfutils-devel sqlite-devel clang-tools-extra gcc gcc-c++ cmake make openssl-devel
python3 -m pip install --upgrade pip
python3 -m pip install scikit-build
```

--------------------------------

### Tool Registration Function Example

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/_doxygen/rocprofiler-sdk/html/registration_8h_source.html

Example implementation of the `rocprofiler_configure` function, which tools define to enable rocProfiler support. It demonstrates setting up initialization and finalization callbacks, creating contexts, and handling initialization errors.

```cpp
#include <rocprofiler-sdk/registration.h>

static rocprofiler_client_id_t my_client_id;
static rocprofiler_client_finalize_t my_fini_func;
static int my_tool_data = 1234;

static int my_init_func(rocprofiler_client_finalize_t fini_func,
void* tool_data)
{
my_fini_func = fini_func;

assert(*static_cast<int*>(tool_data) == 1234 && "tool_data is wrong");

rocprofiler_context_id_t ctx{0};
rocprofiler_create_context(&ctx);

if(int valid_ctx = 0;
rocprofiler_context_is_valid(ctx, &valid_ctx) != ROCPROFILER_STATUS_SUCCESS ||
valid_ctx != 0)
{
// notify rocprofiler that initialization failed
// and all the contexts, buffers, etc. created
// should be ignored
return -1;
}

if(rocprofiler_start_context(ctx) != ROCPROFILER_STATUS_SUCCESS)
{
// notify rocprofiler that initialization failed
// and all the contexts, buffers, etc. created
// should be ignored
return -1;
}

// no errors
return 0;
}

static int my_fini_func(void* tool_data)
{
assert(*static_cast<int*>(tool_data) == 1234 && "tool_data is wrong");
}
```

--------------------------------

### Build and Install ROCTx Module

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/how-to/using-rocprofiler-sdk-roctx.html

These commands show how to build and install the `roctx` Python module using CMake. Adjust `ROCPROFILER_PYTHON_VERSIONS` and `CMAKE_INSTALL_PREFIX` as needed for your environment.

```bash
cmake -B build-sdk -DCMAKE_INSTALL_PREFIX=/opt/rocm -DROCPROFILER_PYTHON_VERSIONS="3.10" -DCMAKE_PREFIX_PATH=/opt/rocm

```

```bash
cmake -B build-sdk -DCMAKE_INSTALL_PREFIX=/opt/rocm -DROCPROFILER_PYTHON_VERSIONS="3.8;3.9;3.10;3.11;3.12" -DCMAKE_PREFIX_PATH=/opt/rocm

```

--------------------------------

### Get Agent and Counter Information

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/how-to/using-rocprofv3-avail.html

Use the 'info' command to get agent information and a list of all counters supported on each agent.

```bash
rocprofv3-avail info

```

--------------------------------

### PC Sampling Stochastic Output Example

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/how-to/using-pc-sampling.html

This example shows the additional fields generated by the stochastic PC sampling method compared to host trap sampling. These fields help in understanding instruction issuance and stall reasons.

```text
Sample_Timestamp | Exec_Mask | Dispatch_Id | Instruction | Instruction_Comment | Correlation_Id | Wave_Issued_Instruction | Instruction_Type | Stall_Reason | Wave_Count  
---|---|---|---|---|---|---|---|---|---
390705261841337 | 18446744073709551615 | 24 | s_waitcnt vmcnt(0) | /home/vlaindic/git/rocprofiler-sdk-internal/tests/bin/vector-operations/vector-ops.cpp:133 | 24 | 0 | ROCPROFILER_PC_SAMPLING_INSTRUCTION_TYPE_NO_INST | ROCPROFILER_PC_SAMPLING_INSTRUCTION_NOT_ISSUED_REASON_WAITCNT | 4  
390705261924637 | 18446744073709551615 | 29 | v_max_i32_e32 v1, v2, v0 | /home/vlaindic/git/rocprofiler-sdk-internal/tests/bin/vector-operations/vector-ops.cpp:77 | 29 | 1 | ROCPROFILER_PC_SAMPLING_INSTRUCTION_TYPE_VALU | ROCPROFILER_PC_SAMPLING_INSTRUCTION_NOT_ISSUED_REASON_OTHER_WAIT | 6  
390705694732429 | 18446744073709551615 | 53 | v_mad_u64_u32 v[0:1], s[2:3], v0, s2, v[2:3] | /home/vlaindic/git/rocprofiler-sdk-internal/tests/bin/vector-operations/vector-ops.cpp:80 | 53 | 1 | ROCPROFILER_PC_SAMPLING_INSTRUCTION_TYPE_VALU | ROCPROFILER_PC_SAMPLING_INSTRUCTION_NOT_ISSUED_REASON_OTHER_WAIT | 6  
390705694744189 | 18446744073709551615 | 54 | v_lshl_add_u64 v[0:1], s[4:5], 0, v[0:1] | /home/vlaindic/git/rocprofiler-sdk-internal/tests/bin/vector-operations/vector-ops.cpp:99 | 54 | 1 | ROCPROFILER_PC_SAMPLING_INSTRUCTION_TYPE_VALU | ROCPROFILER_PC_SAMPLING_INSTRUCTION_NOT_ISSUED_REASON_OTHER_WAIT | 4  
390705694769549 | 18446744073709551615 | 56 | s_waitcnt vmcnt(0) | /home/vlaindic/git/rocprofiler-sdk-internal/tests/bin/vector-operations/vector-ops.cpp:133 | 56 | 0 | ROCPROFILER_PC_SAMPLING_INSTRUCTION_TYPE_NO_INST | ROCPROFILER_PC_SAMPLING_INSTRUCTION_NOT_ISSUED_REASON_WAITCNT | 6  
390705694772089 | 18446744073709551615 | 56 | s_waitcnt lgkmcnt(0) | /usr/include/hip/amd_detail/amd_hip_runtime.h:275 | 56 | 0 | ROCPROFILER_PC_SAMPLING_INSTRUCTION_TYPE_NO_INST | ROCPROFILER_PC_SAMPLING_INSTRUCTION_NOT_ISSUED_REASON_WAITCNT | 5  
390705694810449 | 18446744073709551615 | 58 | v_cmp_gt_i32_e32 vcc, s2, v1 | /home/vlaindic/git/rocprofiler-sdk-internal/tests/bin/vector-operations/vector-ops.cpp:94 | 58 | 1 | ROCPROFILER_PC_SAMPLING_INSTRUCTION_TYPE_VALU | ROCPROFILER_PC_SAMPLING_INSTRUCTION_NOT_ISSUED_REASON_OTHER_WAIT | 3  

```

--------------------------------

### Example PC Sampling Configuration

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/api-reference/pc_sampling.html

An example structure for a PC sampling configuration, specifying the method, unit, and interval ranges. This is typically populated by `rocprofiler_query_pc_sampling_agent_configurations`.

```c++
rocprofiler_pc_sampling_configuration_t {
    .method = ROCPROFILER_PC_SAMPLING_METHOD_HOST_TRAP,
    .unit = ROCPROFILER_PC_SAMPLING_UNIT_TIME,
    .min_interval = 1,
    .max_interval = 10000
};


```

--------------------------------

### Buffer Tracing Record Example

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/api-reference/buffered_services.html

An example demonstrating how to process buffer tracing records within the callback function, casting the payload based on the record's category and kind.

```APIDOC
## Buffer Tracing Record Processing Example

### Description
Example of how to handle different types of buffer tracing records within a callback function.

### Code Example
```c
void
buffer_callback_func(
    rocprofiler_context_id_t      context,
    rocprofiler_buffer_id_t       buffer_id,
    rocprofiler_record_header_t** headers,
    size_t                        num_headers,
    void*                         user_data,
    uint64_t                      drop_count)
{
    for(size_t i = 0; i < num_headers; ++i)
    {
        auto* header = headers[i];

        if(header->category == ROCPROFILER_BUFFER_CATEGORY_TRACING &&
           header->kind == ROCPROFILER_BUFFER_TRACING_HIP_RUNTIME_API)
        {
            auto* record =
                static_cast<rocprofiler_buffer_tracing_hip_api_record_t*>(header->payload);

            // Process HIP API record
        }
        else if(header->category == ROCPROFILER_BUFFER_CATEGORY_TRACING &&
                header->kind == ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH)
        {
            auto* record =
                static_cast<rocprofiler_buffer_tracing_kernel_dispatch_record_t*>(header->payload);

            // Process Kernel Dispatch record
        }
        else
        {
            // Handle unhandled record types
            throw std::runtime_error("unhandled record header category + kind");
        }
    }
}
```
```

--------------------------------

### Example PC Sampling JSON Output

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/how-to/using-pc-sampling.html

This is a trimmed-down example of the JSON output generated by host-trap PC sampling. It shows multiple records with hardware and PC information.

```json
{
  "pc_sample_host_trap": [
    {
      "record": {
        "hw_id": {
          "chiplet": 0,
          "wave_id": 0,
          "simd_id": 2,
          "pipe_id": 0,
          "cu_or_wgp_id": 1,
          "shader_array_id": 0,
          "shader_engine_id": 2,
          "workgroup_id": 0,
          "vm_id": 3,
          "queue_id": 2,
          "microengine_id": 1
        },
        "pc": {
          "code_object_id": 1,
          "code_object_offset": 20228
        },
        "exec_mask": 18446744073709551615,
        "timestamp": 51040126667689,
        "dispatch_id": 1,
        "corr_id": {
          "internal": 1,
          "external": 0
        },
        "wrkgrp_id": {
          "x": 182,
          "y": 0,
          "z": 0
        },
        "wave_in_grp": 1
      },
      "inst_index": 0
    },
    {
      "record": {
        "hw_id": {
          "chiplet": 0,
          "wave_id": 0,
          "simd_id": 2,
          "pipe_id": 0,
          "cu_or_wgp_id": 0,
          "shader_array_id": 0,
          "shader_engine_id": 2,
          "workgroup_id": 0,
          "vm_id": 3,
          "queue_id": 2,
          "microengine_id": 1
        },
        "pc": {
          "code_object_id": 1,
          "code_object_offset": 20236
        },
        "exec_mask": 18446744073709551615,
        "timestamp": 51040126667689,
        "dispatch_id": 1,
        "corr_id": {
          "internal": 1,
          "external": 0
        },
        "wrkgrp_id": {
          "x": 158,
          "y": 0,
          "z": 0
        },
        "wave_in_grp": 2
      },
      "inst_index": 1
    }
  ]
}
```

--------------------------------

### tool_init() setup for thread trace

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/api-reference/thread_trace.html

This snippet demonstrates the necessary steps to initialize the thread trace service using `tool_init()`. It includes context creation, configuring callback tracing for code objects, querying available GPU agents, and optionally setting thread trace parameters.

```APIDOC
## tool_init() setup for thread trace

### Description
This section describes how to use the ROCprofiler-SDK thread trace API to configure and use the thread trace service.

### Steps:

1.  **Configure callback tracing for code objects:**
    ```c++
    rocprofiler_context_id_t ctx{0};
    ROCPROFILER_CALL(rocprofiler_create_context(&ctx), "context creation failed");

    ROCPROFILER_CALL(
        rocprofiler_configure_callback_tracing_service(ctx,
                                                       ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT,
                                                       nullptr,
                                                       0,
                                                       tool_codeobj_tracing_callback,
                                                       nullptr),
        "code object tracing service configure");
    ```

2.  **Query available GPU agents:**
    ```c++
    std::vector<rocprofiler_agent_id_t> agents{};

    ROCPROFILER_CALL(
        rocprofiler_query_available_agents(
            ROCPROFILER_AGENT_INFO_VERSION_0,
            [](rocprofiler_agent_version_t, const void** _agents, size_t _num_agents, void* _data) {
                auto* agent_v = static_cast<std::vector<rocprofiler_agent_id_t>*>(_data);
                for(size_t i = 0; i < _num_agents; ++i)
                {
                    auto* agent = static_cast<const rocprofiler_agent_v0_t*>(_agents[i]);
                    if(agent->type == ROCPROFILER_AGENT_TYPE_GPU)
                        agent_v->emplace_back(agent->id);
                }
                return ROCPROFILER_STATUS_SUCCESS;
            },
            sizeof(rocprofiler_agent_v0_t),
            &agents),
        "Failed to iterate agents");
    ```

3.  **Optionally, specify configuration parameters:**
    ```c++
    std::vector<rocprofiler_thread_trace_parameter_t> params{};

    params.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_SHADER_ENGINE_MASK, 0xF});

    params.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_TARGET_CU, 0});

    params.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_SIMD_SELECT, 0xF});

    params.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_BUFFER_SIZE, 1u<<30}); // 1 GB
    ```

### Configuration Parameters:

*   **ROCPROFILER_THREAD_TRACE_PARAMETER_SHADER_ENGINE_MASK**: Configures the Shader Engine (SE) mask. Recommended to trace only one SE at a time.
*   **ROCPROFILER_THREAD_TRACE_PARAMETER_TARGET_CU**: Configures the target Compute Unit (CU) or WGP for instruction tracing.
*   **ROCPROFILER_THREAD_TRACE_PARAMETER_SIMD_SELECT**: Configures SIMD selection for tracing.
*   **ROCPROFILER_THREAD_TRACE_PARAMETER_BUFFER_SIZE**: Configures the buffer size for trace data.
```

--------------------------------

### Initialize PC Sampling Context and Buffer

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/api-reference/pc_sampling.html

Demonstrates the initialization steps for the PC sampling service, including creating a context and a buffer with specified policies and a callback function.

```APIDOC
## tool_init() setup

Here are the steps to set up `tool_init()`:

```cpp
rocprofiler_context_id_t ctx{0};
rocprofiler_buffer_id_t buff;
ROCPROFILER_CALL(rocprofiler_create_context(&ctx), "context creation failed");
ROCPROFILER_CALL(rocprofiler_create_buffer(ctx,
                                            8192,
                                            2048,
                                            ROCPROFILER_BUFFER_POLICY_LOSSLESS,
                                            pc_sampling_callback, // Callback to process PC samples
                                            user_data,
                                            &buff),
                    "buffer creation failed");

```

For more details on buffer creation, see ROCprofiler-SDK buffered services.
```

--------------------------------

### Example: Configure Buffered Tracing with Dedicated Thread

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/api-reference/buffered_services.html

Demonstrates the workflow for creating a context, a buffer, a dedicated callback thread, and assigning the buffer to that thread.

```cpp
{
    // create a context
    auto context_id = rocprofiler_context_id_t{0};
    rocprofiler_create_context(&context_id);

    // create a buffer associated with the context
    auto buffer_id  = rocprofiler_buffer_id_t{};
    rocprofiler_create_buffer(context_id, ..., &buffer_id);

    // specify that a new callback thread should be created and provide
    // and assign the identifier for it to the "thr_id" variable
    auto thr_id = rocprofiler_callback_thread_t{};
    rocprofiler_create_callback_thread(&thr_id);

    // assign the buffer callback to be delivered on this thread
    rocprofiler_assign_callback_thread(buffer_id, thr_id);
}

```

--------------------------------

### Example PC Sampling Configuration Output

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/how-to/using-pc-sampling.html

This output demonstrates a typical configuration for PC sampling on a supported GPU, showing the method, unit, and interval parameters.

```text
GPU:0
NAME:gfx90a
configs:
   Method              :host_trap
   Unit                :time
   Min_Interval        :1
   Max_Interval        :18446744073709551615
   Flags               :none

```

--------------------------------

### Install ROCprofiler-SDK via DNF

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/install/installation.html

Install a prebuilt version of ROCprofiler-SDK on Red Hat Enterprise Linux using the DNF package manager, assuming ROCm 6.2 or later is installed.

```bash
$ sudo dnf install rocprofiler-sdk
```

--------------------------------

### Example: Configuring Buffer Tracing for HIP API and Kernel Dispatches

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/api-reference/buffered_services.html

This example demonstrates how to configure buffer tracing for both HIP runtime API functions and kernel dispatches using the same buffer. Ensure the buffer is created before configuring tracing services.

```c++
{
    auto ctx = rocprofiler_context_id_t{};
    // ... creation of context, etc. ...

    // buffer parameters
    constexpr auto KB          = 1024;  // 1024 bytes
    constexpr auto buffer_size = 16 * KB;
    constexpr auto watermark   = 15 * KB;
    constexpr auto policy      = ROCPROFILER_BUFFER_POLICY_LOSSLESS;

    // buffer handle
    auto buffer_id = rocprofiler_buffer_id_t{};

    // create a buffer associated with the context
    rocprofiler_create_buffer(
        context_id, buffer_size, watermark, policy, callback_func, nullptr, &buffer_id);

    // configure HIP runtime API function records to be placed in buffer
    rocprofiler_configure_buffer_tracing_service(
        ctx, ROCPROFILER_BUFFER_TRACING_HIP_RUNTIME_API, nullptr, 0, buffer_id);

    // configure kernel dispatch records to be placed in buffer
    // (more than one service can use the same buffer)
    rocprofiler_configure_buffer_tracing_service(
        ctx, ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH, nullptr, 0, buffer_id);

    // ... etc. ...
}

```

--------------------------------

### Variables starting with 'i'

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/_doxygen/rocprofiler-sdk/html/functions_vars_i.html

This section details the variables that start with the letter 'i' and the structures they belong to.

```APIDOC
## Variables starting with 'i'

### id
- **Description**: Identifier for various entities.
- **Structures**: `rocprofiler_counter_info_v0_t`, `rocprofiler_counter_info_v1_t`, `rocprofiler_counter_record_dimension_info_t`, `rocprofiler_counter_record_t`

### index
- **Description**: Index of a dimension.
- **Structures**: `rocprofiler_counter_dimension_info_t`

### initialize
- **Description**: Flag or function pointer to indicate initialization status or action.
- **Structures**: `rocprofiler_tool_configure_result_t`

### inst_type
- **Description**: Type of instruction.
- **Structures**: `rocprofiler_pc_sampling_record_stochastic_v0_t`

### instance
- **Description**: Represents an instance of a runtime initialization record or data.
- **Structures**: `rocprofiler_buffer_tracing_runtime_initialization_record_t`, `rocprofiler_callback_tracing_runtime_initialization_data_t`

### instance_id
- **Description**: Unique identifier for a dimension instance.
- **Structures**: `rocprofiler_counter_record_dimension_instance_info_t`

### instructions_array
- **Description**: Array containing instruction data.
- **Structures**: `rocprofiler_thread_trace_decoder_wave_t`

### instructions_size
- **Description**: Size of the instruction data.
- **Structures**: `rocprofiler_thread_trace_decoder_wave_t`

### internal
- **Description**: Internal correlation ID.
- **Structures**: `rocprofiler_async_correlation_id_t`, `rocprofiler_correlation_id_t`

### internal_correlation_id
- **Description**: Specific internal correlation ID for record retirement.
- **Structures**: `rocprofiler_buffer_tracing_correlation_id_retirement_record_t`

### is_constant
- **Description**: Flag indicating if the counter is constant.
- **Structures**: `rocprofiler_counter_info_v0_t`, `rocprofiler_counter_info_v1_t`

### is_derived
- **Description**: Flag indicating if the counter is derived.
- **Structures**: `rocprofiler_counter_info_v0_t`, `rocprofiler_counter_info_v1_t`
```

--------------------------------

### Variables starting with 'f'

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/_doxygen/rocprofiler-sdk/html/functions_vars_f.html

This section details variables that start with the letter 'f'.

```APIDOC
## Variable: finalize

### Description
Represents the finalization status or action for a tool configuration.

### Type
`rocprofiler_tool_configure_result_t`
```

```APIDOC
## Variable: flags

### Description
Contains various flag settings related to agent I/O links, PC sampling configurations, and thread trace decoder shader data.

### Types
- `rocprofiler_agent_io_link_t`
- `rocprofiler_pc_sampling_configuration_t`
- `rocprofiler_pc_sampling_record_stochastic_v0_t`
- `rocprofiler_thread_trace_decoder_shaderdata_t`
```

--------------------------------

### Locate ROCprofiler-SDK Samples

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/how-to/samples.html

This command shows the default installation path for ROCprofiler-SDK sample programs.

```bash
/opt/rocm/share/rocprofiler-sdk/samples
```

--------------------------------

### Install Test Requirements

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/install/installation.html

Install pandas and pytest, which are required for running some of the ROCprofiler-SDK tests.

```bash
/usr/local/bin/python -m pip install -r requirements.txt
```

--------------------------------

### Run rocprofv3 for Application Tracing

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/how-to/using-rocprofv3.html

Use this command to initiate application tracing with rocprofv3. Replace `<tracing_option>` with specific tracing flags and `<application_path>` with the path to your application. The `--output-format csv` option is optional; omitting it defaults to the `rocpd` format.

```bash
rocprofv3 <tracing_option> -- <application_path>

```

--------------------------------

### Variables starting with 't'

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/_doxygen/rocprofiler-sdk/html/functions_vars_t.html

This section details variables that start with the letter 't' and the structures they belong to.

```APIDOC
## Variables starting with 't'

### thread_id

- **Description**: Identifier for a thread.
- **Structures**: `rocprofiler_buffer_tracing_hip_api_ext_record_t`, `rocprofiler_buffer_tracing_hip_api_record_t`, `rocprofiler_buffer_tracing_hsa_api_record_t`, `rocprofiler_buffer_tracing_kernel_dispatch_record_t`, `rocprofiler_buffer_tracing_marker_api_record_t`, `rocprofiler_buffer_tracing_memory_allocation_record_t`, `rocprofiler_buffer_tracing_memory_copy_record_t`, `rocprofiler_buffer_tracing_ompt_record_t`, `rocprofiler_buffer_tracing_rccl_api_record_t`, `rocprofiler_buffer_tracing_rocdecode_api_ext_record_t`, `rocprofiler_buffer_tracing_rocdecode_api_record_t`, `rocprofiler_buffer_tracing_rocjpeg_api_record_t`, `rocprofiler_buffer_tracing_runtime_initialization_record_t`, `rocprofiler_buffer_tracing_scratch_memory_record_t`

### time

- **Description**: Represents a point in time or a duration.
- **Structures**: `rocprofiler_thread_trace_decoder_inst_t`, `rocprofiler_thread_trace_decoder_occupancy_t`, `rocprofiler_thread_trace_decoder_perfevent_t`

### timeline_array

- **Description**: An array related to timeline data.
- **Structures**: `rocprofiler_thread_trace_decoder_wave_t`

### timeline_size

- **Description**: The size of the timeline data.
- **Structures**: `rocprofiler_thread_trace_decoder_wave_t`

### timestamp

- **Description**: A timestamp indicating an event's occurrence.
- **Structures**: `rocprofiler_buffer_tracing_correlation_id_retirement_record_t`, `rocprofiler_buffer_tracing_kfd_event_dropped_events_record_t`, `rocprofiler_buffer_tracing_kfd_event_page_fault_record_t`, `rocprofiler_buffer_tracing_kfd_event_page_migrate_record_t`, `rocprofiler_buffer_tracing_kfd_event_queue_record_t`, `rocprofiler_buffer_tracing_kfd_event_unmap_from_gpu_record_t`, `rocprofiler_buffer_tracing_runtime_initialization_record_t`, `rocprofiler_pc_sampling_record_host_trap_v0_t`, `rocprofiler_pc_sampling_record_stochastic_v0_t`

### tool_attach

- **Description**: Information related to tool attachment.
- **Structures**: `rocprofiler_tool_configure_attach_result_t`

### tool_data

- **Description**: Data associated with a tool.
- **Structures**: `rocprofiler_tool_configure_attach_result_t`, `rocprofiler_tool_configure_result_t`

### tool_detach

- **Description**: Information related to tool detachment.
- **Structures**: `rocprofiler_tool_configure_attach_result_t`

### type

- **Description**: The type or category of an item.
- **Structures**: `rocprofiler_agent_io_link_t`, `rocprofiler_thread_trace_decoder_wave_state_t`
```

--------------------------------

### Rocprofiler Start Context Function

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/_doxygen/rocprofiler-sdk/html/api_buffered_tracing_2client_8cpp-example.html

Starts the rocprofiler context, enabling tracing and profiling operations.

```cpp
void
start()
{
ROCPROFILER_CALL(rocprofiler_start_context(client_ctx), "context start");
}
```

--------------------------------

### List Basic Agent Info

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/how-to/using-rocprofv3-avail.html

Use 'list --agent' to list basic info for all agents. Use '-d' to list info only for the specified device.

```bash
rocprofv3-avail list --agent

```

--------------------------------

### Variables starting with 'c'

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/_doxygen/rocprofiler-sdk/html/functions_vars_c.html

This section details variables that start with the letter 'c', along with the structures they belong to.

```APIDOC
## Variables starting with 'c'

### cache_line_size
- **Type**: Undefined
- **Description**: Belongs to `rocprofiler_agent_cache_t`.

### cache_lines_per_tag
- **Type**: Undefined
- **Description**: Belongs to `rocprofiler_agent_cache_t`.

### category
- **Type**: Undefined
- **Description**: Belongs to `rocprofiler_thread_trace_decoder_inst_t`.

### chiplet
- **Type**: Undefined
- **Description**: Belongs to `rocprofiler_pc_sampling_hw_id_v0_t`.

### code_object_id
- **Type**: Undefined
- **Description**: Belongs to `rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t`, `rocprofiler_callback_tracing_code_object_load_data_t`, `rocprofiler_pc_t`, `rocprofiler_thread_trace_decoder_pc_t`.

### code_object_offset
- **Type**: Undefined
- **Description**: Belongs to `rocprofiler_pc_t`.

### contexts
- **Type**: Undefined
- **Description**: Belongs to `rocprofiler_thread_trace_decoder_wave_t`.

### correlation_id
- **Type**: Undefined
- **Description**: Belongs to `rocprofiler_buffer_tracing_hip_api_ext_record_t`, `rocprofiler_buffer_tracing_hip_api_record_t`, `rocprofiler_buffer_tracing_hsa_api_record_t`, `rocprofiler_buffer_tracing_kernel_dispatch_record_t`, `rocprofiler_buffer_tracing_marker_api_record_t`, `rocprofiler_buffer_tracing_memory_allocation_record_t`, `rocprofiler_buffer_tracing_memory_copy_record_t`, `rocprofiler_buffer_tracing_ompt_record_t`, `rocprofiler_buffer_tracing_rccl_api_record_t`, `rocprofiler_buffer_tracing_rocdecode_api_ext_record_t`, `rocprofiler_buffer_tracing_rocdecode_api_record_t`, `rocprofiler_buffer_tracing_rocjpeg_api_record_t`, `rocprofiler_buffer_tracing_runtime_initialization_record_t`, `rocprofiler_buffer_tracing_scratch_memory_record_t`, `rocprofiler_dispatch_counting_service_data_t`, `rocprofiler_dispatch_counting_service_record_t`, `rocprofiler_pc_sampling_record_host_trap_v0_t`, `rocprofiler_pc_sampling_record_stochastic_v0_t`.

### counter_id
- **Type**: Undefined
- **Description**: Belongs to `rocprofiler_counter_record_dimension_instance_info_t`.

### counter_value
- **Type**: Undefined
- **Description**: Belongs to `rocprofiler_counter_record_t`.

### counters
- **Type**: Undefined
- **Description**: Belongs to `rocprofiler_spm_record_t`.

### cu
- **Type**: Undefined
- **Description**: Belongs to `rocprofiler_thread_trace_decoder_occupancy_t`, `rocprofiler_thread_trace_decoder_shaderdata_t`, `rocprofiler_thread_trace_decoder_wave_t`.

### CU
- **Type**: Undefined
- **Description**: Belongs to `rocprofiler_thread_trace_decoder_perfevent_t`.

### cu_or_wgp_id
- **Type**: Undefined
- **Description**: Belongs to `rocprofiler_pc_sampling_hw_id_v0_t`.
```

--------------------------------

### Data Fields starting with 'n'

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/_doxygen/rocprofiler-sdk/html/functions_n.html

This section details the data fields that start with the letter 'n' in the ROCprofiler-SDK.

```APIDOC
## Data Fields starting with 'n'

### name
- **Type**: rocprofiler_client_id_t, rocprofiler_counter_info_v0_t, rocprofiler_counter_info_v1_t
- **Description**: Represents the name of a client ID or counter information.

### node_from
- **Type**: rocprofiler_agent_io_link_t
- **Description**: Represents the source node in an agent I/O link.

### node_to
- **Type**: rocprofiler_agent_io_link_t
- **Description**: Represents the destination node in an agent I/O link.

### num_records
- **Type**: rocprofiler_dispatch_counting_service_record_t
- **Description**: Indicates the number of records in a dispatch counting service.
```

--------------------------------

### roctxRangePushA

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/_doxygen/roctx/html/roctx_8h_source.html

Starts a new nested range, which is local to the current CPU thread. Returns the level at which the range was started.

```APIDOC
## roctxRangePushA

### Description
Starts a new nested range. Nested ranges are stacked and local to the current CPU thread. The function returns the 0-based level of the nested range.

### Parameters
#### Path Parameters
None

#### Query Parameters
None

#### Request Body
None

### Parameters
- **message** (const char*) - The message associated with this range. This parameter is required and must be null-terminated.

### Returns
- **int** - The 0-based level at which this nested range is started.
```

--------------------------------

### Data Fields starting with 'w'

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/_doxygen/rocprofiler-sdk/html/functions_w.html

This section details the data fields that start with the letter 'w' within the ROCprofiler-SDK.

```APIDOC
## Data Fields starting with 'w'

This section lists data fields that begin with the letter 'w'.

### Fields:

*   **wave_count** : `rocprofiler_pc_sampling_record_stochastic_v0_t`
    *   Description: Number of waves.

*   **wave_id** : `rocprofiler_pc_sampling_hw_id_v0_t`, `rocprofiler_thread_trace_decoder_occupancy_t`, `rocprofiler_thread_trace_decoder_shaderdata_t`, `rocprofiler_thread_trace_decoder_wave_t`
    *   Description: Identifier for a wave.

*   **wave_in_group** : `rocprofiler_pc_sampling_record_host_trap_v0_t`, `rocprofiler_pc_sampling_record_stochastic_v0_t`
    *   Description: Indicates if a wave is part of a group.

*   **wave_issued** : `rocprofiler_pc_sampling_record_stochastic_v0_t`
    *   Description: Indicates if a wave has been issued.

*   **weight** : `rocprofiler_agent_io_link_t`
    *   Description: Weight associated with an I/O link.

*   **width** : `rocprofiler_agent_mem_bank_t`
    *   Description: Width of a memory bank.

*   **workgroup_id** : `rocprofiler_pc_sampling_hw_id_v0_t`, `rocprofiler_pc_sampling_record_host_trap_v0_t`, `rocprofiler_pc_sampling_record_stochastic_v0_t`
    *   Description: Identifier for a workgroup.

*   **workgroup_size** : `rocprofiler_kernel_dispatch_info_t`
    *   Description: Size of a workgroup.
```

--------------------------------

### Data Fields starting with 'l'

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/_doxygen/rocprofiler-sdk/html/functions_l.html

This section details the data fields that start with the letter 'l' within the ROCprofiler-SDK.

```APIDOC
## Data Fields - l

### latency
- **Type**: `rocprofiler_agent_cache_t`
- **Description**: Represents the latency associated with the agent cache.

### level
- **Type**: `rocprofiler_agent_cache_t`
- **Description**: Represents the level of the agent cache.

### load_base
- **Type**: `rocprofiler_callback_tracing_code_object_load_data_t`
- **Description**: Represents the base address for code object loading.

### load_cnt
- **Type**: `rocprofiler_pc_sampling_memory_counters_t`
- **Description**: Represents the count of memory loads.

### load_delta
- **Type**: `rocprofiler_callback_tracing_code_object_load_data_t`
- **Description**: Represents the change in code object loading.

### load_size
- **Type**: `rocprofiler_callback_tracing_code_object_load_data_t`
- **Description**: Represents the size of the code object being loaded.
```

--------------------------------

### hipLaunchKernel_spt

Source: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/_doxygen/rocprofiler-sdk/html/hip_2api__args_8h_source.html

Structure for launching a kernel.

```APIDOC
## hipLaunchKernel_spt

### Description
Structure used to define parameters for launching a kernel.

### Parameters
- **function_address** (const void*) - The address of the kernel function.
- **numBlocks** (rocprofiler_dim3_t) - The grid dimensions (number of thread blocks).
- **dimBlocks** (rocprofiler_dim3_t) - The block dimensions (threads per block).
- **args** (void**) - Array of kernel arguments.
- **sharedMemBytes** (size_t) - The amount of shared memory to allocate for the kernel.
- **stream** (hipStream_t) - The HIP stream to launch the kernel on.
```
Tokens:
Raw