2 files changed, 30 insertions, 1 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 209bae6..8c39f80 100644
--- 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"
     }
 }