Skip unsupported tests in case of optimized code

motokultivator · spatrang · commit 42d183ed9bd1 · 2026-05-28T01:50:03.000-05:00
For tests whose semantics cannot be reconciled with optimised
code, gate them so they UNSUPPORTED-out cleanly instead of
failing.

  - "require no_optimized_code" at the top of: aspace-user-
    input, aspace-watchpoint, convenience_variables, device-
    barrier, finish, fork-exec-non-gpu-to-gpu, lane-execution,
    mi-aspace, mi-lanes, nonstop-displaced, precise-memory-
    warning-watchpoint, register-watchpoint, runtime-core,
    shared-memory.

  - "require no_lto" at the top of: hip-builtin-variables,
    hip-lang-detect, mi-aspace, watchpoint-at-end-of-shader.

  - watchpoint-basic.exp: proc-level skip of
    test_non_write_watchpoint_before_runtime_load's rwatch/
    awatch sub-cases under optimisation.  The HIP runtime is
    initialised before main() reaches "Break before runtime
    load", so amd-dbgapi is already attached and rejects the
    host-side read/access watchpoint at insert time.  Only
    this sub-test depends on the "before runtime load"
    premise; the other procs still work under optimisation,
    so use a proc-level check_for_any_flags skip rather than
    a file-level "require no_optimized_code".

Co-Authored-By: Andrey Kasaurov &lt;andrey.kasaurov@amd.com&gt;
Co-Authored-By: Luis Machado &lt;luis.machado@amd.com&gt;

Change-Id: I750d64e842413429e281bb8f7d85c4f46319a627
Signed-off-by: Sarang Patrange &lt;spatrang@amd.com&gt;
diff --git a/gdb/testsuite/gdb.rocm/alu-exceptions.cpp b/gdb/testsuite/gdb.rocm/alu-exceptions.cpp
@@ -21,12 +21,20 @@
 #include <limits>
 #include <cstdlib>
 
-__device__ void
+/* The debugger sets a breakpoint on this function and uses an inferior
+   call to enable ALU exceptions before each kernel launches.  We need:
+     - noinline, so the breakpoint reliably fires at each call site, and
+     - an observable side effect the optimiser cannot prove away,
+       otherwise IPO removes every call (the function would appear pure)
+       and then dead-strips the now-uncalled function.
+   An empty volatile asm with a memory clobber is the cheapest way to
+   give the function an effect the compiler must preserve, and codegens
+   to nothing on amdgcn.  */
+
+__device__ __attribute__ ((noinline)) void
 enable_alu_exceptions ()
 {
-  /* By default, ALU exceptions are not enabled.  Break here and use the
-     debugger to enable exceptions.  */
-  volatile int i = 0;
+  asm volatile ("" ::: "memory");
 }
 
 __global__ void
diff --git a/gdb/testsuite/gdb.rocm/aspace-user-input.exp b/gdb/testsuite/gdb.rocm/aspace-user-input.exp
@@ -22,6 +22,7 @@
 load_lib rocm.exp
 
 require allow_hipcc_tests
+require no_optimized_code
 
 standard_testfile bit-extract.cpp
 
diff --git a/gdb/testsuite/gdb.rocm/aspace-watchpoint.exp b/gdb/testsuite/gdb.rocm/aspace-watchpoint.exp
@@ -25,6 +25,7 @@
 load_lib rocm.exp
 
 require allow_hipcc_tests
+require no_optimized_code
 
 standard_testfile .cpp
 
diff --git a/gdb/testsuite/gdb.rocm/convenience_variables.exp b/gdb/testsuite/gdb.rocm/convenience_variables.exp
@@ -22,6 +22,7 @@
 load_lib rocm.exp
 
 require allow_hipcc_tests
+require no_optimized_code
 
 standard_testfile .cpp
 
diff --git a/gdb/testsuite/gdb.rocm/corefile.cpp b/gdb/testsuite/gdb.rocm/corefile.cpp
@@ -33,8 +33,8 @@
 	}                                                                    \
     } while (0)
 
-__device__ int
-bar () __attribute__((optnone))
+__device__ int __attribute__((optnone))
+bar ()
 {
   return threadIdx.x;
 }
diff --git a/gdb/testsuite/gdb.rocm/deep-stack.exp b/gdb/testsuite/gdb.rocm/deep-stack.exp
@@ -47,17 +47,17 @@ with_rocm_gpu_lock {
     gdb_test "continue" "Thread $decimal \"hip_deep\" hit Breakpoint $decimal.*"
     gdb_test "backtrace" \
 	[multi_line \
-	    "#0 .* base_case .*" \
-	    "#1 .* deep<0u> .*" \
-	    "#2 .* deep<1u> .*" \
-	    "#3 .* deep<2u> .*" \
-	    "#4 .* deep<3u> .*" \
-	    "#5 .* deep<4u> .*" \
-	    "#6 .* deep<5u> .*" \
-	    "#7 .* deep<6u> .*" \
-	    "#8 .* deep<7u> .*" \
-	    "#9 .* deep<8u> .*" \
-	    "#10 .*deep<9u> .*" \
-	    "#11 .*deep<10u> .*" \
-	    "#12 .*hip_deep .*"]
+	    "#0 +((0x.* in )?base_case) .*" \
+	    "#1 +((0x.* in )?deep<0u>) .*" \
+	    "#2 +((0x.* in )?deep<1u>) .*" \
+	    "#3 +((0x.* in )?deep<2u>) .*" \
+	    "#4 +((0x.* in )?deep<3u>) .*" \
+	    "#5 +((0x.* in )?deep<4u>) .*" \
+	    "#6 +((0x.* in )?deep<5u>) .*" \
+	    "#7 +((0x.* in )?deep<6u>) .*" \
+	    "#8 +((0x.* in )?deep<7u>) .*" \
+	    "#9 +((0x.* in )?deep<8u>) .*" \
+	    "#10 +((0x.* in )?deep<9u>) .*" \
+	    "#11 +((0x.* in )?deep<10u>) .*" \
+	    "#12 +((0x.* in )?hip_deep) .*"]
 }
diff --git a/gdb/testsuite/gdb.rocm/deref-scoped-pointer.cpp b/gdb/testsuite/gdb.rocm/deref-scoped-pointer.cpp
@@ -32,16 +32,12 @@
       }									\
   }
 
-/* Non-static + noinline so a line-breakpoint at the "done ()" call
-   site below has a stable instruction address under -O1/-O2/-O3.
-   Without these, the empty body is inlined into its caller, so
-   there is no instruction at the call-site line for GDB to insert
-   the breakpoint at.  */
+/* Non-static so a line breakpoint at the "done ()" call site below
+   has a stable address at -O0.  */
 
-__device__ void __attribute__((noinline))
+__device__ void
 done ()
 {
-  asm volatile ("" : : : "memory");
 }
 
 /* This is initialized by the kernel function below.  */
@@ -50,16 +46,10 @@ __device__ size_t *global_ptr = nullptr;
 /* This is initialized by GDB.  */
 __device__ size_t *global_ptr2 = nullptr;
 
-/* kernel() is marked __attribute__((optnone)) and "local_var" is
-   declared "volatile" so the variable remains addressable from the
-   "done" call site through the end of the function under
-   -O1/-O2/-O3.  The test takes the address of local_var via GDB
-   expressions, which requires it to have an lvalue/storage.  */
-
-__global__ void __attribute__((optnone))
+__global__ void
 kernel ()
 {
-  volatile size_t local_var = blockIdx.x * blockDim.x + threadIdx.x;
+  size_t local_var = blockIdx.x * blockDim.x + threadIdx.x;
 
   /* This will be zero for the first HIP thread, or lane.  */
   if (local_var == 0)
diff --git a/gdb/testsuite/gdb.rocm/deref-scoped-pointer.exp b/gdb/testsuite/gdb.rocm/deref-scoped-pointer.exp
@@ -28,6 +28,15 @@ load_lib rocm.exp
 
 require allow_hipcc_tests
 
+# At -O1+ the HIP compiler drops the DWARF name->location mapping for
+# kernel locals (e.g. "local_var") at arbitrary line breakpoints, even
+# when their address is escaped and the storage is kept live, so GDB
+# reports them as <optimized out>.  This test relies on resolving
+# "local_var" symbolically at the "set breakpoint here" location, so
+# restrict it to non-optimised builds until the toolchain improves
+# debug info for HIP kernels under optimisation.
+require no_optimized_code
+
 standard_testfile .cpp
 
 if {[build_executable "failed to prepare" $testfile $srcfile {debug hip}]} {
diff --git a/gdb/testsuite/gdb.rocm/device-barrier.exp b/gdb/testsuite/gdb.rocm/device-barrier.exp
@@ -20,6 +20,7 @@
 load_lib rocm.exp
 
 require allow_hipcc_tests
+require no_optimized_code
 
 standard_testfile .cpp
 
diff --git a/gdb/testsuite/gdb.rocm/finish.exp b/gdb/testsuite/gdb.rocm/finish.exp
@@ -23,6 +23,7 @@
 load_lib rocm.exp
 
 require allow_hipcc_tests
+require no_optimized_code
 
 standard_testfile .cpp
 
diff --git a/gdb/testsuite/gdb.rocm/fork-exec-gpu-to-non-gpu-execee.cpp b/gdb/testsuite/gdb.rocm/fork-exec-gpu-to-non-gpu-execee.cpp
@@ -15,10 +15,13 @@
    You should have received a copy of the GNU General Public License
    along with this program.  If not, see <http://www.gnu.org/licenses/>.  */
 
-static void
+/* noinline so the breakpoint reliably fires; the asm gives the
+   function a side effect the optimiser cannot prove away.  */
+
+static void __attribute__ ((noinline))
 break_here_execee (void)
 {
-  volatile int i = 0;
+  asm volatile ("" ::: "memory");
 }
 
 int
diff --git a/gdb/testsuite/gdb.rocm/fork-exec-gpu-to-non-gpu-execer.cpp b/gdb/testsuite/gdb.rocm/fork-exec-gpu-to-non-gpu-execer.cpp
@@ -26,10 +26,13 @@ __global__ static void
 kernel1 ()
 {}
 
-__device__ static void
+/* noinline so the breakpoint reliably fires; the asm gives the
+   function a side effect the optimiser cannot prove away.  */
+
+__device__ static void __attribute__ ((noinline))
 break_here_execer ()
 {
-  volatile int j = 0;
+  asm volatile ("" ::: "memory");
 }
 
 __global__ static void
diff --git a/gdb/testsuite/gdb.rocm/fork-exec-non-gpu-to-gpu.exp b/gdb/testsuite/gdb.rocm/fork-exec-non-gpu-to-gpu.exp
@@ -21,6 +21,7 @@ load_lib rocm.exp
 
 require allow_hipcc_tests
 require allow_fork_tests
+require no_optimized_code
 
 standard_testfile -execer.cpp -execee.cpp
 
diff --git a/gdb/testsuite/gdb.rocm/hip-builtin-completions.exp b/gdb/testsuite/gdb.rocm/hip-builtin-completions.exp
@@ -27,6 +27,7 @@
 load_lib rocm.exp
 
 require allow_hipcc_tests
+require no_lto
 
 standard_testfile hip-builtin-variables.cpp
 
diff --git a/gdb/testsuite/gdb.rocm/hip-builtin-variables.cpp b/gdb/testsuite/gdb.rocm/hip-builtin-variables.cpp
@@ -52,23 +52,23 @@ inner_func2 ()
 __global__ void __attribute__((optnone))
 kern ()
 {
-  volatile int thread_idx_x = threadIdx.x;
-  volatile int thread_idx_y = threadIdx.y;
-  volatile int thread_idx_z = threadIdx.z;
+  int thread_idx_x = threadIdx.x;
+  int thread_idx_y = threadIdx.y;
+  int thread_idx_z = threadIdx.z;
 
-  volatile int group_idx_x = blockIdx.x;
-  volatile int group_idx_y = blockIdx.y;
-  volatile int group_idx_z = blockIdx.z;
+  int group_idx_x = blockIdx.x;
+  int group_idx_y = blockIdx.y;
+  int group_idx_z = blockIdx.z;
 
-  volatile int group_size_x = blockDim.x;
-  volatile int group_size_y = blockDim.y;
-  volatile int group_size_z = blockDim.z;
+  int group_size_x = blockDim.x;
+  int group_size_y = blockDim.y;
+  int group_size_z = blockDim.z;
 
-  volatile int grid_size_x = gridDim.x;
-  volatile int grid_size_y = gridDim.y;
-  volatile int grid_size_z = gridDim.z;
+  int grid_size_x = gridDim.x;
+  int grid_size_y = gridDim.y;
+  int grid_size_z = gridDim.z;
 
-  volatile int wave_size = warpSize;
+  int wave_size = warpSize;
 
   /* Break here.  */
   inner_func1 ();
diff --git a/gdb/testsuite/gdb.rocm/hip-builtin-variables.exp b/gdb/testsuite/gdb.rocm/hip-builtin-variables.exp
@@ -21,6 +21,7 @@
 load_lib rocm.exp
 
 require allow_hipcc_tests
+require no_lto
 
 standard_testfile .cpp
 
diff --git a/gdb/testsuite/gdb.rocm/hip-lang-detect.exp b/gdb/testsuite/gdb.rocm/hip-lang-detect.exp
@@ -22,6 +22,7 @@
 load_lib rocm.exp
 
 require allow_hipcc_tests
+require no_lto
 
 standard_testfile hip-builtin-variables.cpp
 
diff --git a/gdb/testsuite/gdb.rocm/instruction-stepping-commands.cpp b/gdb/testsuite/gdb.rocm/instruction-stepping-commands.cpp
@@ -35,19 +35,24 @@
     }								\
   while (0)
 
-/* Make sure the function isn't inlined, for the nexti test.  */
-__device__ static __attribute__ ((noinline))
+/* "noinline" so the nexti test can step into a real function call;
+   "optnone" so the body has stable instructions for stepi/nexti and
+   the locally-initialized value is not folded into the return.  */
+__device__ static __attribute__ ((noinline, optnone))
 int
 return_zero ()
 {
-  volatile int ret = 0;
+  int ret = 0;
   return ret;
 }
 
-__global__ void
+/* "optnone" so each statement below maps to its own instruction
+   sequence and the "var += 1" lines are not folded together --
+   stepi/nexti rely on stepping through them one at a time.  */
+__global__ void __attribute__ ((optnone))
 kernel ()
 {
-  volatile int var = return_zero ();
+  int var = return_zero ();
   var += 1; /* next line */
   var += 1;
 }
diff --git a/gdb/testsuite/gdb.rocm/instruction-stepping-commands.exp b/gdb/testsuite/gdb.rocm/instruction-stepping-commands.exp
@@ -44,7 +44,7 @@ proc test_step_until {cmd until_regexp} {
     set insn_count_max 100
 
     gdb_test_multiple $cmd "" {
-	-re -wrap "$::hex\\s+$::decimal\\s+volatile int var = return_zero \\(\\);" {
+	-re -wrap "$::hex\\s+$::decimal\\s+int var = return_zero \\(\\);" {
 	    # Note: we expect to see a leading PC address as we're
 	    # stopping at instructions that map to the middle of a
 	    # source line.
diff --git a/gdb/testsuite/gdb.rocm/lane-execution.exp b/gdb/testsuite/gdb.rocm/lane-execution.exp
@@ -35,6 +35,7 @@
 load_lib rocm.exp
 
 require allow_hipcc_tests
+require no_optimized_code
 
 standard_testfile .cpp
 
diff --git a/gdb/testsuite/gdb.rocm/line-breakpoint-in-kernel.cpp b/gdb/testsuite/gdb.rocm/line-breakpoint-in-kernel.cpp
@@ -31,10 +31,10 @@
       }                                                                      \
   }
 
-__global__ void
+__global__ void __attribute__((optnone))
 kernel ()
 {
-  volatile int x = 0;
+  int x = 0;
   x++; /* break here */
 }
 
diff --git a/gdb/testsuite/gdb.rocm/mi-aspace.exp b/gdb/testsuite/gdb.rocm/mi-aspace.exp
@@ -21,6 +21,8 @@ load_lib mi-support.exp
 set MIFLAGS "-i=mi"
 
 require allow_hipcc_tests
+require no_optimized_code
+require no_lto
 
 standard_testfile .cpp
 
diff --git a/gdb/testsuite/gdb.rocm/mi-lanes.exp b/gdb/testsuite/gdb.rocm/mi-lanes.exp
@@ -21,6 +21,7 @@ load_lib mi-support.exp
 set MIFLAGS "-i=mi"
 
 require allow_hipcc_tests
+require no_optimized_code
 
 standard_testfile lane-execution.cpp
 
diff --git a/gdb/testsuite/gdb.rocm/multi-inferior-fork.cpp b/gdb/testsuite/gdb.rocm/multi-inferior-fork.cpp
@@ -36,10 +36,11 @@ __global__ void
 kernel ()
 {}
 
-static void
+static void __attribute__((optnone))
 child_after_fork ()
 {
-  volatile int j = 0;
+  int j = 0;
+  (void) j;
 }
 
 int
diff --git a/gdb/testsuite/gdb.rocm/multi-inferior-gpu.cpp b/gdb/testsuite/gdb.rocm/multi-inferior-gpu.cpp
@@ -32,7 +32,7 @@ kern ()
 
 /* Spawn one child process per detected GPU.  */
 
-static int
+static int __attribute__((optnone))
 parent (int argc, char **argv)
 {
   /* Identify how many GPUs we have, and spawn one child for each.  */
@@ -73,7 +73,7 @@ parent (int argc, char **argv)
     }
 
   /* Last break here.  */
-  volatile int ret = 0;
+  int ret = 0;
   return ret;
 }
 
diff --git a/gdb/testsuite/gdb.rocm/nonstop-displaced.exp b/gdb/testsuite/gdb.rocm/nonstop-displaced.exp
@@ -21,6 +21,7 @@
 load_lib rocm.exp
 
 require allow_hipcc_tests
+require no_optimized_code
 
 standard_testfile .cpp
 
diff --git a/gdb/testsuite/gdb.rocm/precise-memory-exec.c b/gdb/testsuite/gdb.rocm/precise-memory-exec.c
@@ -19,10 +19,13 @@
 #include <stdlib.h>
 #include <stdio.h>
 
-static void
+/* noinline so the breakpoint reliably fires; the asm gives the
+   function a side effect the optimiser cannot prove away.  */
+
+static void __attribute__ ((noinline))
 second (void)
 {
-  volatile int i = 0;
+  asm volatile ("" ::: "memory");
 }
 
 int
diff --git a/gdb/testsuite/gdb.rocm/precise-memory-fork.c b/gdb/testsuite/gdb.rocm/precise-memory-fork.c
diff --git a/gdb/testsuite/gdb.rocm/precise-memory-warning-watchpoint.exp b/gdb/testsuite/gdb.rocm/precise-memory-warning-watchpoint.exp
diff --git a/gdb/testsuite/gdb.rocm/register-watchpoint.exp b/gdb/testsuite/gdb.rocm/register-watchpoint.exp
diff --git a/gdb/testsuite/gdb.rocm/runtime-core.exp b/gdb/testsuite/gdb.rocm/runtime-core.exp
diff --git a/gdb/testsuite/gdb.rocm/shared-memory.exp b/gdb/testsuite/gdb.rocm/shared-memory.exp
diff --git a/gdb/testsuite/gdb.rocm/static-global.cpp b/gdb/testsuite/gdb.rocm/static-global.cpp
diff --git a/gdb/testsuite/gdb.rocm/watchpoint-at-end-of-shader.exp b/gdb/testsuite/gdb.rocm/watchpoint-at-end-of-shader.exp
diff --git a/gdb/testsuite/gdb.rocm/watchpoint-basic.exp b/gdb/testsuite/gdb.rocm/watchpoint-basic.exp
diff --git a/gdb/testsuite/lib/rocm.exp b/gdb/testsuite/lib/rocm.exp

Original file line number	Diff line number	Diff line change
`@@ -33,8 +33,8 @@`
`33`	`33`	`} \`
`34`	`34`	`} while (0)`
`35`	`35`
`36`		`-__device__ int`
`37`		`-bar () __attribute__((optnone))`
	`36`	`+__device__ int __attribute__((optnone))`
	`37`	`+bar ()`
`38`	`38`	`{`
`39`	`39`	`return threadIdx.x;`
`40`	`40`	`}`