diff options
Diffstat (limited to 'gdb/testsuite/gdb.rocm')
| -rw-r--r-- | gdb/testsuite/gdb.rocm/precise-memory.cpp | 12 | ||||
| -rw-r--r-- | gdb/testsuite/gdb.rocm/precise-memory.exp | 38 |
2 files changed, 41 insertions, 9 deletions
diff --git a/gdb/testsuite/gdb.rocm/precise-memory.cpp b/gdb/testsuite/gdb.rocm/precise-memory.cpp index 769b58a..7a8c37e 100644 --- a/gdb/testsuite/gdb.rocm/precise-memory.cpp +++ b/gdb/testsuite/gdb.rocm/precise-memory.cpp @@ -31,7 +31,17 @@ __global__ void kernel () { - __builtin_amdgcn_s_sleep (1); + + /* Simple kernel which loads from address 0 to trigger a pagefault. + When precise memory is not enabled, it is expected that the memory fault + is reported after the s_nop instruction. With precise-memory, the + exception should be reported on the s_nop. */ + asm volatile ("s_mov_b64 [s10, s11], 0\n" + "s_load_dword s12, [s10, s11]\n" + "s_nop 0" + : + : + : "s10", "s11", "s12"); } int diff --git a/gdb/testsuite/gdb.rocm/precise-memory.exp b/gdb/testsuite/gdb.rocm/precise-memory.exp index f423a11..8c39f80 100644 --- a/gdb/testsuite/gdb.rocm/precise-memory.exp +++ b/gdb/testsuite/gdb.rocm/precise-memory.exp @@ -39,18 +39,40 @@ proc do_test { } { "AMDGPU precise memory access reporting is off \\(currently disabled\\)." \ "show precise-memory setting in CLI before" - if {[hip_devices_support_precise_memory]} { - gdb_test_no_output "set amdgpu precise-memory on" - set cli_effective_value "enabled" - } else { - gdb_test "set amdgpu precise-memory on" \ - "warning: AMDGPU precise memory access reporting could not be enabled." - set cli_effective_value "disabled" + # Assume precise-memory is available, unless GDB reports otherwise. + gdb_test_multiple "set amdgpu precise-memory on" "" { + -re -wrap "warning: AMDGPU precise memory access reporting could not be enabled\\." { + set cli_effective_value "disabled" + pass $gdb_test_name + } + -re -wrap "^" { + set cli_effective_value "enabled" + pass $gdb_test_name + } } gdb_test "show amdgpu precise-memory" \ - "AMDGPU precise memory access reporting is on \\(currently ${cli_effective_value}\\)." \ + "AMDGPU precise memory access reporting is on \\(currently ${cli_effective_value}\\)\\." \ "show precise-memory setting in CLI after" + + if { $cli_effective_value eq "disabled" } { + return + } + + # Get to the begining of the GPU kernel without precise memory enabled. + with_test_prefix "goto gpu code" { + gdb_test_no_output "set amdgpu precise-memory off" + gdb_breakpoint "kernel" allow-pending + gdb_test "continue" "Thread ${::decimal}.* hit Breakpoint .*" + gdb_test_no_output "set amdgpu precise-memory on" + } + + # If precise-memory is available, run until a SIGSEGV is reported. At + # that point, the PC should point to the s_nop instruction (the one + # following the one which caused the memory violation). + gdb_test "continue" "Thread ${::decimal}\[^\r\n\]* received signal SIGSEGV, Segmentation fault.*" + + gdb_test "x/i \$pc" "=> ${::hex} <_Z6kernelv\\+${::decimal}>:\[ \t\]+s_nop\[ \t\]+0" } } |