Failed to capture trace of HIP application

[jeromew]

os - windows 11 gpu - Radeon 6700 XT driver - up to date with Adrenalin 24.12.1 - driver version 32.0.12033.1030 / 27/11/2024

when I try to profile a HIP application with RGP, the application is detected, the profiling starts and then at the end I get a "Failed to capture trace" in the logs.

typical log is

(10:25:08.832) INFO [RGP Trace Source - PID: 13420] Client connected [222 HIP] (10:25:08.926) INFO [RGP Trace Source - PID: 13420] Client reached init state [222 HIP] (10:25:09.035) INFO [RGP Trace Source - PID: 13420] Successfully enabled tracing [222 HIP] (10:25:09.035) INFO [RGP Trace Source - PID: 13420] Initialized new client [222 HIP] (10:25:09.097) INFO [DDToolConn] Successfully initialized driver (connection id: 222). (10:25:09.224) INFO [RGP Trace Source - PID: 13420] Successfully queried SPM counters [222 HIP] (10:25:09.224) INFO [RGP Trace Source - PID: 13420] Successfully updated SPM counters [222 HIP] (10:25:09.463) INFO [RGP Trace Source - PID: 13420] Successfully began trace [222 HIP] (10:25:09.675) INFO [RGP Trace Source - PID: 13420] Client disconnected [222 HIP] (10:25:09.714) ERROR [RGP Trace Source - PID: 13420] Failed to capture trace [222 HIP] (10:25:09.714) INFO [RGP Trace Source - PID: 13420] Finished disconnecting client [222 HIP]

I tried with a minimalist example from the ROCm examples like

#include <hip/hip_runtime.h>

#include <iostream>

#define HIP_CHECK(expression)                  \
{                                              \
    const hipError_t status = expression;      \
    if(status != hipSuccess){                  \
        std::cerr << "HIP error "              \
                  << status << ": "            \
                  << hipGetErrorString(status) \
                  << " at " << __FILE__ << ":" \
                  << __LINE__ << std::endl;    \
    }                                          \
}


__device__ unsigned int get_thread_idx()
{
    return threadIdx.x;
}

__host__ void print_hello_host()
{
    std::cout << "Hello world from host!" << std::endl;
}

__device__ __host__ void print_hello()
{
    printf("Hello world from device or host!\n");
}

__global__ void helloworld_kernel()
{
    unsigned int thread_idx = get_thread_idx();
    unsigned int block_idx = blockIdx.x;

    print_hello();

    printf("Hello world from device kernel block %u thread %u!\n", block_idx, thread_idx);
}

int main()
{
    print_hello_host();

    print_hello();

    helloworld_kernel<<<dim3(2), // 3D grid specifying number of blocks to launch: (2, 1, 1)
                        dim3(2), // 3D grid specifying number of threads to launch: (2, 1, 1)
                        0, // number of bytes of additional shared memory to allocate
                        hipStreamDefault // stream where the kernel should execute: default stream
                        >>>();

    HIP_CHECK(hipDeviceSynchronize());
}

compiled it with hipcc.bat --offload-arch=gfx1031 .\example.hip

and it does not work either I get the same "Failed to capture trace" message.

are there special compile flags to use for enabling HIP application profiling ? Is there something I could to make it work or is this card ?

I tested after compiling with both HIP SDK 6.1.2 and HIP SDK 6.2.4.

Do I need a specific driver ?

I saw the the HIP SDK installer mentions (but did not install) a driver called "PRO 24.20.36 launched on 19.11.2024" I installed it and tried again => same problem "Failed to capture trace"