Skip to content

[watch/rebar 1/3] Fix gdb.rocm/watch-gpu-global-from-host.exp on non-ReBAR systems#80

Open
palves wants to merge 1 commit intoamd-stagingfrom
spr/amd-staging/c99ccb02
Open

[watch/rebar 1/3] Fix gdb.rocm/watch-gpu-global-from-host.exp on non-ReBAR systems#80
palves wants to merge 1 commit intoamd-stagingfrom
spr/amd-staging/c99ccb02

Conversation

@palves
Copy link
Copy Markdown
Collaborator

@palves palves commented Apr 15, 2026

On a system with Resizable BAR (ReBAR) enabled, the BAR size is the
size of the whole VRAM, so the runtime is able to allocate rw mappings
on the CPU that map to /dev/dri/render nodes, and have device
variables mapped there. This means that on ReBAR systems, device
variables as directly host-accessible.

If ReBAR is off, though, the runtime can't do that, the BAR window is
only 256MB, and the runtime will only allocate CPU-visible rw mappings
for VRAM selectively, for managed memory, and not for device
variables.

So on a non-ReBAR system, accessing device variables from the CPU
side segfaults, and so gdb.rocm/watch-gpu-global-from-host.exp fails
with:

Thread 1 "watch-gpu-globa" received signal SIGSEGV, Segmentation fault.
[Switching to thread 1 (Thread 0x7ffff62d0f80 (LWP 590038))]
0x0000000000211c43 in main (argc=1, argv=0x7fffffffcf88) at /home/pedro/rocm/gdb/build/gdb/testsuite/../../../src/gdb/testsuite/gdb.rocm/watch-gpu-global-from-host.cpp:39
39 *devGlobal = 8;
(gdb) FAIL: gdb.rocm/watch-gpu-global-from-host.exp: hit watchpoint in main

I tried several approaches to try to determine whether the access
would work, before the access:

  • I tried several HIP APIs, but none work reliably.

  • I also tried making the testcase find the mapping that corresponds
    to the pointer, and check the mapping permissions, and it works, but
    then that's linux specific, and for Windows we'd need something
    else, or hardcode that the CPU can never access VRAM directly there,
    at least currently.

In the end, I concluded that just letting the access happen and seeing
the SIGSEGV is the best and simplest and the most portable, and
doesn't really have a downside.

So that's what this commit uses to detect the situation, and issue
UNSUPPORTED.

Tested on Linux ReBAR on and off, and on Windows.

Change-Id: I052f2858357ef34d0dbaaa6479d5131571cea112


Stack:

⚠️ Part of a stack created by spr. Do not merge manually using the UI - doing so may have unexpected results.

@palves palves requested a review from a team as a code owner April 15, 2026 13:20
@palves palves changed the title Fix gdb.rocm/watch-gpu-global-from-host.exp on non-ReBAR systems [1/3] Fix gdb.rocm/watch-gpu-global-from-host.exp on non-ReBAR systems Apr 15, 2026
@palves palves changed the title [1/3] Fix gdb.rocm/watch-gpu-global-from-host.exp on non-ReBAR systems [watch/rebar 1/3] Fix gdb.rocm/watch-gpu-global-from-host.exp on non-ReBAR systems Apr 15, 2026
Copy link
Copy Markdown
Collaborator

@lancesix lancesix left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM, thanks.

On a system with Resizable BAR (ReBAR) enabled, the BAR size is the
size of the whole VRAM, so the runtime is able to allocate rw mappings
on the CPU that map to /dev/dri/render nodes, and have __device__
variables mapped there.  This means that on ReBAR systems, __device__
variables as directly host-accessible.

If ReBAR is off, though, the runtime can't do that, the BAR window is
only 256MB, and the runtime will only allocate CPU-visible rw mappings
for VRAM selectively, for managed memory, and not for __device__
variables.

So on a non-ReBAR system, accessing __device__ variables from the CPU
side segfaults, and so gdb.rocm/watch-gpu-global-from-host.exp fails
with:

 Thread 1 "watch-gpu-globa" received signal SIGSEGV, Segmentation fault.
 [Switching to thread 1 (Thread 0x7ffff62d0f80 (LWP 590038))]
 0x0000000000211c43 in main (argc=1, argv=0x7fffffffcf88) at /home/pedro/rocm/gdb/build/gdb/testsuite/../../../src/gdb/testsuite/gdb.rocm/watch-gpu-global-from-host.cpp:39
 39        *devGlobal = 8;
 (gdb) FAIL: gdb.rocm/watch-gpu-global-from-host.exp: hit watchpoint in main

I tried several approaches to try to determine whether the access
would work, before the access:

- I tried several HIP APIs, but none work reliably.

- I also tried making the testcase find the mapping that corresponds
  to the pointer, and check the mapping permissions, and it works, but
  then that's linux specific, and for Windows we'd need something
  else, or hardcode that the CPU can never access VRAM directly there,
  at least currently.

In the end, I concluded that just letting the access happen and seeing
the SIGSEGV is the best and simplest and the most portable, and
doesn't really have a downside.

So that's what this commit uses to detect the situation, and issue
UNSUPPORTED.

Tested on Linux ReBAR on and off, and on Windows.

Change-Id: I052f2858357ef34d0dbaaa6479d5131571cea112
commit-id:c99ccb02
@palves palves force-pushed the spr/amd-staging/c99ccb02 branch from 7ddf723 to 241989a Compare April 16, 2026 18:22
@palves
Copy link
Copy Markdown
Collaborator Author

palves commented Apr 16, 2026

(fixed some typos in the commit log)

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants