[Issue] `rocprof-sys-run` roctxRangePush("name") / Pop Broken

[ax3l]

When tracing an application in ROCm 6.4.1 that uses RTX range annotations of the form:

roctxRangePush("name");

// user code goes here
// ... nesting ranges with unique names possible, too
// compute compute compute

roctxRangePop();

then rocprof-sys-run -- ./exe does not record the RTX ranges correctly. Instead of creating a range until pop, the push command drops the range immediately after the roctxRangePush("name"); returns.

Running instead with rocprofv3 --hip-trace --marker-trace --output-format=pftrace -- ./exe shows proper ranges.

Example traces using the same binary:

rocprofv3 example

rocprof-sys-run

note the missing stack

If I zoom in, individual RTX annotations show as very-short blocks that end immediately again.

[emedwedeff]

Here's a simple example code of this:

#include <hip/hip_runtime.h>
#include <rocprofiler-sdk-roctx/roctx.h>

__global__ void simple_kernel()
{
    unsigned int block_idx = blockIdx.x;
}

int main()
{
    roctxMark("My Test");
    int rangeId = roctxRangeStart("simple_kernel  range");
    roctxRangePush("launch_simple_kernel");
    simple_kernel<<<dim3(2), dim3(2), 0, hipStreamDefault>>>();
    roctxRangePop();
    roctxRangeStop(rangeId);
    return 0;
}

you can build it simply using:

hipcc simple.hip -lroctx64

or

/opt/rocm-6.4.1/lib/llvm/bin/clang++ --offload-arch=gfx942 --offload-arch=gfx942 --offload-arch=gfx942 --offload-arch=gfx942 -O3 --driver-mode=g++ -O3 --hip-link -x hip simple.hip -lroctx64

and then run:

rocprof-sys-run -- ./a.out

and

rocprofv3 --hip-trace --marker-trace --output-format=pftrace -- ./a.out

Please reach out me to directly if you have any questions.