gdb/testsuite: Test the effect of amdgpu-precise memory

The gdb.rocm/precise-memory.exp test currently checks that the "amdgpu
precise-memory" setting can be set.  It does not test that this setting
has any meaningful effect.

This patch extends this test to ensure that precise-memory has the
expected behaviour.

Change-Id: I58f72a51a566f04fc89114b94ee656c2e7ac35bb
Approved-by: Pedro Alves <pedro@palves.net>

diff --git a/gdb/testsuite/gdb.rocm/precise-memory.cpp b/gdb/testsuite/gdb.rocm/precise-memory.cpp
index 769b58a72b5c2fd5e8fbbc974de14f67e360182b..7a8c37eeb57288f716b89ed4bb5417a9650de874 100644 (file)
--- 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 209bae641790cbd4a82545a0dca87674a2483a27..8c39f808706840a5259c5d46d114acdfa1aea109 100644 (file)
--- a/gdb/testsuite/gdb.rocm/precise-memory.exp
+++ b/gdb/testsuite/gdb.rocm/precise-memory.exp
@@ -54,6 +54,25 @@ proc do_test { } {
        gdb_test "show amdgpu precise-memory" \
            "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"
     }
 }