Update some gdb.rocm tests to support optimized code

Adapt several gdb.rocm tests so they pass under -O1/-O2/-O3:

  - Apply __attribute__((optnone)) to device functions whose
    locals or call structure the tests rely on, and to host
    main() where its store-then-exit pattern would otherwise
    be tail-optimised.

  - Make breakpoint-anchor helper functions ("done", "foo",
    "bar") non-static and noinline, with an asm volatile body,
    so the symbol survives the optimiser and the breakpoint
    can be inserted.

  - Mark local variables volatile where the test reads them
    via convenience variables or takes their address.

  - Restore "__global__ void" qualifier on device-attach
    kernel, and mark its spin-loop guard volatile so the loop
    survives optimisation.

  - Relax regex/line-number expectations under optimisation
    in deep-stack, hip-builtin-variables, snapshot-objfile-on-
    load, and step-schedlock-spurious-waves.

  - Brace expr-substitutions in deep-stack and until-tests
    (tclint hygiene).

Co-Authored-By: Luis Machado <luis.machado@amd.com>

Change-Id: I750d64e842413429e281bb8f7d85c4f46319a627
Signed-off-by: Sarang Patrange <spatrang@amd.com>

Author: kasaurov

gdb/testsuite/gdb.rocm/alu-exceptions.cpp

@@ -26,62 +26,63 @@ enable_alu_exceptions ()
 {
   /* By default, ALU exceptions are not enabled.  Break here and use the
      debugger to enable exceptions.  */
+  volatile int i = 0;
 }
 
 __global__ void
 raise_invalid ()
 {
   enable_alu_exceptions ();
   volatile float a = -1;
-  float b = sqrt (a); /* Break here for invalid.  */
+  volatile float b = sqrt (a); /* Break here for invalid.  */
 }
 
 __global__ void
 raise_denorm ()
 {
   enable_alu_exceptions ();
   volatile float inp = std::numeric_limits<float>::denorm_min ();
-  float b = inp * inp; /* Break here for denorm.  */
+  volatile float b = inp * inp; /* Break here for denorm.  */
 }
 
 __global__ void
 raise_float_div0 ()
 {
   enable_alu_exceptions ();
   volatile float a = 1, b = 0;
-  float c = a / b; /* Break here for float_div0.  */
+  volatile float c = a / b; /* Break here for float_div0.  */
 }
 
 __global__ void
 raise_overflow ()
 {
   enable_alu_exceptions ();
   volatile float max = std::numeric_limits<float>::max ();
-  float b = max * 2.0f; /* Break here for overflow.  */
+  volatile float b = max * 2.0f; /* Break here for overflow.  */
 }
 
 __global__ void
 raise_underflow ()
 {
   enable_alu_exceptions ();
   volatile float min = std::numeric_limits<float>::min ();
-  float c = min * min; /* Break here for underflow.  */
+  volatile float c = min * min; /* Break here for underflow.  */
 }
 
 __global__ void
 raise_inexact ()
 {
   enable_alu_exceptions ();
   volatile float f = 1.1f;
-  float r = f * f; /* Break here for inexact.  */
+  volatile float r = f * f; /* Break here for inexact.  */
 }
 
 __global__ void
 raise_int_div0 ()
 {
   enable_alu_exceptions ();
   volatile int a = 1, b = 0;
-  int c = a / b; /* Break here for int_div0.  */
+  volatile int c = a / b; /* Break here for int_div0.  */
 }
 
 int

gdb/testsuite/gdb.rocm/corefile.cpp

@@ -34,12 +34,12 @@
     } while (0)
 
 __device__ int
-bar ()
+bar () __attribute__((optnone))
 {
   return threadIdx.x;
 }
 
-__device__ int
+__device__ int __attribute__((optnone))
 baz ()
 {
   return 0;

gdb/testsuite/gdb.rocm/deep-stack.exp

@@ -38,6 +38,10 @@ with_rocm_gpu_lock {
     }
 
     set line [gdb_get_line_number "break here"]
+    if {[check_for_any_flags {-O1 -O2 -O3}]} {
+	# return is optimized out, stop one statement before.
+	set line [expr {$line - 1}]
+    }
     gdb_breakpoint "$line" allow-pending
 
     gdb_test "continue" "Thread $decimal \"hip_deep\" hit Breakpoint $decimal.*"
@@ -53,7 +57,7 @@ with_rocm_gpu_lock {
 	    "#7 .* deep<6u> .*" \
 	    "#8 .* deep<7u> .*" \
 	    "#9 .* deep<8u> .*" \
-	    "#10 .* deep<9u> .*" \
-	    "#11 .* deep<10u> .*" \
-	    "#12 .* hip_deep .*"]
+	    "#10 .*deep<9u> .*" \
+	    "#11 .*deep<10u> .*" \
+	    "#12 .*hip_deep .*"]
 }

gdb/testsuite/gdb.rocm/deref-scoped-pointer.cpp

@@ -32,9 +32,16 @@
       }									\
   }
 
-__device__ static void
+/* 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.  */
+
+__device__ void __attribute__((noinline))
 done ()
 {
+  asm volatile ("" : : : "memory");
 }
 
 /* This is initialized by the kernel function below.  */
@@ -43,10 +50,16 @@ __device__ size_t *global_ptr = nullptr;
 /* This is initialized by GDB.  */
 __device__ size_t *global_ptr2 = nullptr;
 
-__global__ void
+/* 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))
 kernel ()
 {
-  size_t local_var = blockIdx.x * blockDim.x + threadIdx.x;
+  volatile 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)

gdb/testsuite/gdb.rocm/device-attach.cpp

@@ -37,7 +37,7 @@ bit_extract_kernel (uint32_t *C_d, const uint32_t *A_d, size_t N)
 {
   size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
   size_t stride = blockDim.x * gridDim.x;
-  int loop_forever = 1;
+  volatile int loop_forever = 1;
   int a = 9;
   while (loop_forever)
     a = 8;

gdb/testsuite/gdb.rocm/finish.cpp

@@ -104,14 +104,14 @@ struct WithStaticFields
 };
 __device__ int WithStaticFields::c = 12;
 
-__device__ static Empty
+__device__ __attribute__((noinline)) Empty
 returnEmpty ()
 {
   return {};
 }
 
 template<int N>
-__device__ static Custom<char, N>
+__device__ __attribute__((noinline)) Custom<char, N>
 returnSized ()
 {
   Custom<char, N> ret;
@@ -122,7 +122,7 @@ returnSized ()
 }
 
 template<int N>
-__device__ static Custom<Union, N>
+__device__ __attribute__((noinline)) Custom<Union, N>
 returnCustomWithUnion ()
 {
   Custom<Union, N> ret;
@@ -131,15 +131,15 @@ returnCustomWithUnion ()
   return ret;
 }
 
-__device__ static FlexibleArrayMember
+__device__ __attribute__((noinline)) FlexibleArrayMember
 returnFlexibleArrayMember ()
 {
   return {
       .non_flexible = static_cast<int> (threadIdx.x)
   };
 }
 
-__device__ static NotPackedWithBitField
+__device__ __attribute__((noinline)) NotPackedWithBitField
 returnNotPackedWithBitField ()
 {
   const int val = static_cast<int> (threadIdx.x);
@@ -157,7 +157,7 @@ returnNotPackedWithBitField ()
   };
 }
 
-__device__ static PackedWithBitField
+__device__ __attribute__((noinline)) PackedWithBitField
 returnPackedWithBitField ()
 {
   const int val = static_cast<int> (threadIdx.x);
@@ -168,7 +168,7 @@ returnPackedWithBitField ()
   };
 }
 
-__device__ static const Custom<char, 8> *
+__device__ __attribute__((noinline)) const Custom<char, 8> *
 returnPtr ()
 {
   static Custom<char, 8> vv {
@@ -177,13 +177,13 @@ returnPtr ()
   return &vv;
 }
 
-__device__ static int *
+__device__ __attribute__((noinline)) int *
 returnPtr2 (int *a)
 {
   return a;
 }
 
-__device__ static const Custom<char, 8> &
+__device__ __attribute__((noinline)) const Custom<char, 8> &
 returnRef ()
 {
   static Custom<char, 8> vv {
@@ -192,13 +192,13 @@ returnRef ()
   return vv;
 };
 
-__device__ static int &
+__device__ __attribute__((noinline)) int &
 returnRef2 (int &a)
 {
   return a;
 }
 
-__device__ static WithStaticFields
+__device__ __attribute__((noinline)) WithStaticFields
 returnWithStatic ()
 {
   OnlyStatic::something = 12;
@@ -210,7 +210,7 @@ returnWithStatic ()
   };
 }
 
-__device__ static OnlyStatic
+__device__ __attribute__((noinline)) OnlyStatic
 returnEmptyWithStatic ()
 {
   OnlyStatic::something = 56;
@@ -230,7 +230,7 @@ returnMixedSmallSruct ()
 
 __device__ int someGlobal = 42;
 
-__global__ void
+__global__ void __attribute__((optnone))
 kernel ()
 {
   returnEmpty ();

gdb/testsuite/gdb.rocm/fork-exec-gpu-to-non-gpu-execee.cpp

@@ -17,7 +17,9 @@
 
 static void
 break_here_execee (void)
-{}
+{
+  volatile int i = 0;
+}
 
 int
 main (void)

gdb/testsuite/gdb.rocm/fork-exec-gpu-to-non-gpu-execer.cpp

@@ -29,6 +29,7 @@ kernel1 ()
 __device__ static void
 break_here_execer ()
 {
+  volatile int j = 0;
 }
 
 __global__ static void

gdb/testsuite/gdb.rocm/hip-builtin-variables.cpp

@@ -34,41 +34,41 @@
 /* A function to be called by "kern ()", without any references to
    threadIdx, blockIdx, blockDim, gridDim, or warpSize variables.  */
 
-__device__ void
+__device__ void __attribute__((optnone))
 inner_func1 ()
 {
   return;
 }
 
 /* A function to be called by "kern ()", with its own "blockIdx".  */
 
-__device__ void
+__device__ void __attribute__((optnone))
 inner_func2 ()
 {
   int blockIdx = 0x42424242;
   return;
 }
 
-__global__ void
+__global__ void __attribute__((optnone))
 kern ()
 {
-  const int thread_idx_x = threadIdx.x;
-  const int thread_idx_y = threadIdx.y;
-  const int thread_idx_z = threadIdx.z;
+  volatile int thread_idx_x = threadIdx.x;
+  volatile int thread_idx_y = threadIdx.y;
+  volatile int thread_idx_z = threadIdx.z;
 
-  const int group_idx_x = blockIdx.x;
-  const int group_idx_y = blockIdx.y;
-  const int group_idx_z = blockIdx.z;
+  volatile int group_idx_x = blockIdx.x;
+  volatile int group_idx_y = blockIdx.y;
+  volatile int group_idx_z = blockIdx.z;
 
-  const int group_size_x = blockDim.x;
-  const int group_size_y = blockDim.y;
-  const int group_size_z = blockDim.z;
+  volatile int group_size_x = blockDim.x;
+  volatile int group_size_y = blockDim.y;
+  volatile int group_size_z = blockDim.z;
 
-  const int grid_size_x = gridDim.x;
-  const int grid_size_y = gridDim.y;
-  const int grid_size_z = gridDim.z;
+  volatile int grid_size_x = gridDim.x;
+  volatile int grid_size_y = gridDim.y;
+  volatile int grid_size_z = gridDim.z;
 
-  const int wave_size = warpSize;
+  volatile int wave_size = warpSize;
 
   /* Break here.  */
   inner_func1 ();

gdb/testsuite/gdb.rocm/hip-builtin-variables.exp

@@ -226,7 +226,16 @@ with_rocm_gpu_lock {
 	# Get into inner_func2() with its own local blockIdx variable.
 	gdb_test "step" "inner_func2 \\(\\) at.*" \
 	    "step into inner_func2() with its own blockIdx"
-	gdb_test "next" "$::decimal\[ \t\]+return;" \
+	# Accept landing on either the "return;" statement or the
+	# closing brace.  At -O0 Clang emits a separate line entry
+	# for the explicit "return;", so "next" lands there.  Under
+	# __attribute__((optnone)) (which we use to keep the local
+	# blockIdx observable at -O1/-O2/-O3) Clang merges the line
+	# entry of the "return;" with the function's closing brace,
+	# so "next" lands on the brace instead.  Either is fine for
+	# this test - the point is just to step past the local
+	# initialisation before reading blockIdx.
+	gdb_test "next" "$::decimal\[ \t\]+(return;|\})" \
 	    "let the initialization happen"
 	gdb_test "p/x blockIdx" "0x42424242" \
 	    "local blockIdx shadowing the built-in"