gdb/testsuite/gdb.rocm: add OpenMP GPU offload tests

gdb/testsuite/gdb.rocm/omp-target-basic.c

@@ -0,0 +1,59 @@
+/* Copyright 2026 Free Software Foundation, Inc.
+   Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved.
+
+   This file is part of GDB.
+
+   This program is free software; you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3 of the License, or
+   (at your option) any later version.
+
+   This program is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with this program.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* Basic OpenMP "target" offload program shared by the breakpoint,
+   data-sharing (locals) and single-stepping checks in
+   omp-target-basic.exp.  It offloads a single '#pragma omp target'
+   region with a mapped input array (to:), a mapped output array
+   (tofrom:) and a firstprivate scalar, plus a few distinguishable
+   statements at the top of the region for single-stepping.  */
+
+#include <stdio.h>
+
+#define N 8
+
+int
+main (void)
+{
+  int in_arr[N];
+  int out_arr[N];
+
+  for (int i = 0; i < N; i++)
+    {
+      in_arr[i] = 100 + i;
+      out_arr[i] = 0;
+    }
+
+  int firstpriv_val = 42;
+
+  #pragma omp target						\
+	      map(to:in_arr[0:N]) map(tofrom:out_arr[0:N])	\
+	      firstprivate(firstpriv_val)
+  {
+    int priv_val = 7;			/* first-target-line */
+    int step_two = priv_val + 1;
+    int step_three = step_two + 1;
+
+    for (int i = 0; i < N; i++)
+      out_arr[i] = in_arr[i] + priv_val + firstpriv_val
+		   + step_two + step_three;
+  }
+
+  printf ("out[0]=%d out[%d]=%d\n", out_arr[0], N - 1, out_arr[N - 1]);
+  return 0;
+}

gdb/testsuite/gdb.rocm/omp-target-basic.exp

@@ -0,0 +1,81 @@
+# Copyright 2026 Free Software Foundation, Inc.
+# Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved.
+#
+# This file is part of GDB.
+#
+# This program is free software; you can redistribute it and/or modify
+# it under the terms of the GNU General Public License as published by
+# the Free Software Foundation; either version 3 of the License, or
+# (at your option) any later version.
+#
+# This program is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with this program.  If not, see <http://www.gnu.org/licenses/>.
+
+# Basic OpenMP-offload device-debugging checks, all sharing a single
+# '#pragma omp target' binary:
+#
+#   - "breakpoint": set a breakpoint inside the target region by
+#     file:line, hit it on the device, and observe an AMDGPU wave.
+#   - "locals": inspect variables in different data-sharing classes
+#     (map(to:), map(tofrom:), firstprivate).
+#   - "step": single-step through statements in the region, staying on
+#     the device.
+
+load_lib rocm.exp
+
+require allow_omp_offload_tests
+
+standard_testfile
+
+set first_line [gdb_get_line_number "first-target-line"]
+
+if {[gdb_compile_omp_offload [list $srcdir/$subdir/$srcfile] \
+	 [standard_output_file $testfile] executable {debug}] != ""} {
+    untested "failed to compile OpenMP-offload binary"
+    return -1
+}
+
+clean_restart $testfile
+
+with_rocm_gpu_lock {
+    if {![runto_main]} {
+	return
+    }
+
+    gdb_breakpoint "$srcfile:$first_line" allow-pending
+
+    with_test_prefix "breakpoint" {
+	rocm_omp_continue_to_breakpoint "continue to device breakpoint"
+	rocm_omp_check_on_wave "AMDGPU wave at OMP-target stop"
+    }
+
+    # Still stopped at the first statement in the region (out_arr is
+    # not written yet), so the data-sharing values are deterministic.
+    with_test_prefix "locals" {
+	gdb_test "print firstpriv_val" "= 42" "firstprivate value"
+	gdb_test "print in_arr\[0\]" "= 100" "mapped-to in_arr\[0\]"
+	gdb_test "print in_arr\[7\]" "= 107" "mapped-to in_arr\[7\]"
+	gdb_test "print out_arr\[0\]" "= 0" \
+	    "mapped-tofrom out_arr\[0\] before fill"
+    }
+
+    # Step forward through a few statements; each step must stay on a
+    # GPU wave.  Line numbers are not pinned because OpenMP-offload
+    # line tables vary across compiler versions.
+    with_test_prefix "step" {
+	foreach which {1 2 3} {
+	    with_test_prefix "next-$which" {
+		rocm_omp_next "next inside target region"
+		rocm_omp_check_on_wave "still on AMDGPU wave"
+	    }
+	}
+    }
+
+    delete_breakpoints
+    rocm_omp_continue_to_exit "continue to exit"
+}

gdb/testsuite/gdb.rocm/omp-target-cpp.cpp

@@ -0,0 +1,78 @@
+/* Copyright 2026 Free Software Foundation, Inc.
+   Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved.
+
+   This file is part of GDB.
+
+   This program is free software; you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3 of the License, or
+   (at your option) any later version.
+
+   This program is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with this program.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* C++-specific OpenMP-offload features.  Exercises a templated device
+   function and a callable functor used inside a '#pragma omp target'
+   region.  */
+
+#include <cstdio>
+
+#define N 16
+
+#pragma omp declare target
+template <typename T>
+static T
+add_t (T a, T b)
+{
+  return a + b;		/* tmpl-line */
+}
+
+struct multiplier
+{
+  int factor;
+
+  int
+  operator() (int x) const
+  {
+    return x * factor;	/* functor-line */
+  }
+};
+#pragma omp end declare target
+
+int
+main ()
+{
+  int a[N];
+  int b[N];
+  int sum[N];
+  int prod[N];
+
+  for (int i = 0; i < N; i++)
+    {
+      a[i] = i;
+      b[i] = 2 * i;
+      sum[i] = 0;
+      prod[i] = 0;
+    }
+
+  multiplier m { 3 };
+
+  #pragma omp target						\
+	      map(to:a[0:N], b[0:N], m)				\
+	      map(tofrom:sum[0:N], prod[0:N])
+  {
+    for (int i = 0; i < N; i++)
+      {
+	sum[i] = add_t<int> (a[i], b[i]);
+	prod[i] = m (a[i]);
+      }
+  }
+
+  std::printf ("sum[1]=%d prod[1]=%d\n", sum[1], prod[1]);
+  return 0;
+}

gdb/testsuite/gdb.rocm/omp-target-cpp.exp

@@ -0,0 +1,66 @@
+# Copyright 2026 Free Software Foundation, Inc.
+# Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved.
+#
+# This file is part of GDB.
+#
+# This program is free software; you can redistribute it and/or modify
+# it under the terms of the GNU General Public License as published by
+# the Free Software Foundation; either version 3 of the License, or
+# (at your option) any later version.
+#
+# This program is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with this program.  If not, see <http://www.gnu.org/licenses/>.
+
+# C++-specific OpenMP-offload tests.  Verify that GDB can break
+# inside a templated device function and a struct member function
+# used in a target region.
+
+load_lib rocm.exp
+
+require allow_omp_offload_tests
+
+standard_testfile .cpp
+
+set tmpl_line    [gdb_get_line_number "tmpl-line"]
+set functor_line [gdb_get_line_number "functor-line"]
+
+if {[gdb_compile_omp_offload_cpp [list $srcdir/$subdir/$srcfile] \
+	 [standard_output_file $testfile] executable {debug}] != ""} {
+    untested "failed to compile C++ OpenMP-offload binary"
+    return -1
+}
+
+clean_restart $testfile
+
+with_rocm_gpu_lock {
+    if {![runto_main]} {
+	return
+    }
+
+    with_test_prefix "templated add_t" {
+	gdb_breakpoint "$srcfile:$tmpl_line" allow-pending
+	rocm_omp_continue_to_breakpoint "continue to add_t<int>"
+	rocm_omp_check_on_wave "AMDGPU wave at add_t"
+	delete_breakpoints
+    }
+
+    with_test_prefix "multiplier::operator()" {
+	gdb_breakpoint "$srcfile:$functor_line" allow-pending
+	rocm_omp_continue_to_breakpoint "continue to multiplier::operator()"
+	rocm_omp_check_on_wave "AMDGPU wave at functor"
+
+	# 'this->factor' should be the value we set on the host (3),
+	# transferred via the implicit 'map(to:m)'.
+	gdb_test "print this->factor" \
+	    "= 3" \
+	    "factor visible via 'this'"
+	delete_breakpoints
+    }
+
+    rocm_omp_continue_to_exit "continue to exit"
+}

gdb/testsuite/gdb.rocm/omp-target-data.c

@@ -0,0 +1,66 @@
+/* Copyright 2026 Free Software Foundation, Inc.
+   Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved.
+
+   This file is part of GDB.
+
+   This program is free software; you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3 of the License, or
+   (at your option) any later version.
+
+   This program is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with this program.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* '#pragma omp target data' wraps two consecutive '#pragma omp target'
+   regions, sharing the device data buffer.  This is the natural way
+   OpenMP applications stage data to the GPU and exercise multiple
+   kernel launches against it.  */
+
+#include <stdio.h>
+
+#define N 32
+
+static void
+fill (int *p, int v)
+{
+  for (int i = 0; i < N; i++)
+    p[i] = v + i;
+}
+
+int
+main (void)
+{
+  int a[N];
+  int b[N];
+  int c[N];
+
+  fill (a, 0);
+  fill (b, 100);
+  for (int i = 0; i < N; i++)
+    c[i] = 0;
+
+  #pragma omp target data map(to:a[0:N], b[0:N]) map(tofrom:c[0:N])
+  {
+    /* First kernel: c = a + b.  */
+    #pragma omp target
+    {
+      for (int i = 0; i < N; i++)
+	c[i] = a[i] + b[i];	/* data-k1 */
+    }
+
+    /* Second kernel: c *= 2 (data buffer reused, no host transfer).  */
+    #pragma omp target
+    {
+      for (int i = 0; i < N; i++)
+	c[i] = c[i] * 2;	/* data-k2 */
+    }
+  }
+
+  printf ("c[0]=%d c[%d]=%d\n", c[0], N - 1, c[N - 1]);
+  return 0;
+}