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

[spatrang]

Summary

Add OpenMP GPU offload coverage to the gdb.rocm/ testsuite.
The existing gdb.rocm/ tests cover HIP-style kernels (__global__
functions launched via hipLaunchKernelGGL), but there is no coverage
for OpenMP target offload — an increasingly common way of programming
AMD GPUs from C, C++ and Fortran. Without these tests, regressions in
GDB's handling of:

• the __omp_offloading_<hash>_<func>_l<line> outliner-frame symbols,
• OpenMP-runtime-loaded device ELFs (loaded via libomptarget plugins,
  not libamdhip64),
• DWARF emitted for OpenMP map-clause variables,
  can land silently while the HIP path keeps working.

Motivation

gdb.threads/ already has CPU-side OpenMP tests (omp-par-scope.exp,
omp-task.exp), but nothing in the testsuite exercises the target
construct on an actual GPU. This PR adds those tests and the dejagnu
framework hooks needed to drive an OpenMP-offload-capable Clang/Flang
toolchain.

What changed

Framework hooks in gdb/testsuite/lib/rocm.exp (+186 lines)

Reusable procs that future OpenMP-offload tests can also rely on:

• `rocm_find_llvm_tool` — locate a tool under `$ROCM_PATH/llvm/bin`,
  then `$ROCM_PATH/lib/llvm/bin`, then `PATH`.
• `find_amdclang`, `find_amdclangpp`, `find_amdflang` — cached
  wrappers for the three offload-capable compiler drivers.
• `allow_omp_offload_tests`, `allow_omp_offload_fortran_tests` —
  `require`-compatible predicates: amd-dbgapi support, toolchain
  availability, ≥ 1 AMD GPU.
• `rocm_omp_offload_flags ` — build the
  `-fopenmp -fopenmp-targets=amdgcn-amd-amdhsa --offload-arch=…`
  flag set; omits the Clang-only `-Wno-unused-command-line-argument`
  for Fortran (the Flang driver rejects that diagnostic option).
• `gdb_compile_omp_offload`, `gdb_compile_omp_offload_cpp`,
  `gdb_compile_omp_offload_fortran` — `gdb_compile` wrappers that
  swap in the right offload-capable driver and add the offload flags.
• `with_rocm_omp_gpu_lock` — alias for `with_rocm_gpu_lock` so
  OpenMP tests serialize with HIP tests on the GPU.

New tests under `gdb/testsuite/gdb.rocm/` (16 files)

• `omp-target-break` — basic `#pragma omp target` breakpoint by
  `file:line`, AMDGPU wave check.
• `omp-target-step` — `next` stepping inside the offloaded region
  keeps the program on a wave.
• `omp-target-locals` — inspect `map(to:)`, `map(tofrom:)`,
  `firstprivate` and a device-local automatic variable.
• `omp-target-teams` — `target teams distribute parallel for`
  produces multiple AMDGPU waves at one breakpoint; `thread N`
  switches to a wave; `print i` works inside it.
• `omp-target-data` — `target data` enclosing two consecutive
  `target` regions; the device buffer is reused between kernels
  (verified via `c[0] == 100` in the second kernel).
• `omp-target-multi-kernel` — two `declare target` device functions
  launched from separate target regions; named breakpoints land on the
  right kernel.
• `omp-target-cpp` — templated device function (`add_t`) and
  struct functor (`multiplier::operator()`) compiled with the C++
  driver; `print this->factor` works.
• `omp-target-fortran` — `!$omp target` compiled with the Fortran
  driver; breakpoint inside the loop hits an AMDGPU wave.
  Each `.exp` follows the established `gdb.rocm/` patterns:
  `load_lib rocm.exp`, `require allow_omp_offload_tests`,
  `with_rocm_omp_gpu_lock { … }`, and uses the `gdb_test_multiple` +
  `exp_continue` idiom (same as `gdb.rocm/shared-memory.exp`) so the
  patterns survive multi-line GDB output.

Testing

```
make -C testsuite check TESTS="gdb.rocm/omp-target-*.exp"
…

of expected passes 47

of unexpected failures 0

```
Per-test PASS / FAIL counts:

• `omp-target-break.exp` — 3 PASS, 0 FAIL
• `omp-target-step.exp` — 9 PASS, 0 FAIL
• `omp-target-locals.exp` — 7 PASS, 0 FAIL
• `omp-target-teams.exp` — 6 PASS, 0 FAIL
• `omp-target-data.exp` — 8 PASS, 0 FAIL
• `omp-target-multi-kernel.exp` — 5 PASS, 0 FAIL
• `omp-target-cpp.exp` — 6 PASS, 0 FAIL
• `omp-target-fortran.exp` — 3 PASS, 0 FAIL
• Total — 47 PASS, 0 FAIL
  On hosts without an offload-capable Clang/Flang toolchain or without an
  AMD GPU, the `require allow_omp_offload_tests` /
  `allow_omp_offload_fortran_tests` gates report the tests as
  `UNSUPPORTED` rather than `FAIL`, so this PR is safe to land in CI
  immediately.

Backwards compatibility

• `lib/rocm.exp` is only extended; no existing proc is renamed and
  no signature changes. Existing HIP tests are unaffected.
• New procs are uniquely named (e.g. `find_amdclang` vs the existing
  `find_hipcc`) so there is no collision.

[motokultivator]

With flang-23 I get:
gdb compile failed, flang-23: error: unknown argument: '-fno-stack-protector'
for omp-target-fortran.exp.
It comes from this part of gdb.exp:

    if { !$getting_compiler_info
         && [test_compiler_info {gcc-*-*}]
         && !([test_compiler_info {gcc-[0-3]-*}]
              || [test_compiler_info {gcc-4-0-*}])
         && [lsearch -exact $options rust] == -1
         && [lsearch -exact $options hip] == -1} {
        # Put it at the front to not override any user-provided value.
        lappend new_options "early_flags=-fno-stack-protector"
    }

flang --version
AMD flang version 23.0.0git (https://github.com/ROCm/llvm-project.git 43215c73116c407735c85a180d174f718798c328+PATCHED:2506c552d8428e2cc1778bef048b20f818e06bb3)
Target: x86_64-unknown-linux-gnu
Thread model: posix

[spatrang]

With flang-23 I get:
gdb compile failed, flang-23: error: unknown argument: '-fno-stack-protector'
for omp-target-fortran.exp.
It comes from this part of gdb.exp:

    if { !$getting_compiler_info
         && [test_compiler_info {gcc-*-*}]
         && !([test_compiler_info {gcc-[0-3]-*}]
              || [test_compiler_info {gcc-4-0-*}])
         && [lsearch -exact $options rust] == -1
         && [lsearch -exact $options hip] == -1} {
        # Put it at the front to not override any user-provided value.
        lappend new_options "early_flags=-fno-stack-protector"
    }

flang --version
AMD flang version 23.0.0git (https://github.com/ROCm/llvm-project.git 43215c73116c407735c85a180d174f718798c328+PATCHED:2506c552d8428e2cc1778bef048b20f818e06bb3)
Target: x86_64-unknown-linux-gnu
Thread model: posix

Applied fix in 1d0d3cd this commit

[lumachad]

I think we're on the right track, but it seems this patch does things that should be done by #62 or some other rework patch before this one.

[Copilot]

Pull request overview

Adds OpenMP GPU-offload coverage to the gdb.rocm testsuite by extending the ROCm DejaGnu helpers to build/run OpenMP target-offload programs (C/C++/Fortran) and introducing new OpenMP-offload-focused testcases.

Changes:

• Extend gdb/testsuite/lib/rocm.exp with OpenMP offload toolchain discovery, build flag construction, compile wrappers, GPU-lock aliasing, and reusable “continue/step/check wave” helpers.
• Adjust gdb_compile in gdb/testsuite/lib/gdb.exp to avoid adding GCC-specific flags for non-GNU Fortran drivers (notably affecting f90 with flang/amdflang/etc.).
• Add new gdb.rocm/omp-target-* tests covering basic target break/locals/step, teams+waves, multi-kernel named breakpoints, target data reuse, C++ templates/functors, and a Fortran target loop.

Reviewed changes

Copilot reviewed 14 out of 14 changed files in this pull request and generated 4 comments.

Show a summary per file

File	Description
gdb/testsuite/lib/rocm.exp	Adds OpenMP-offload tool/flag/compile helpers and common stop/step assertions for OpenMP target regions.
gdb/testsuite/lib/gdb.exp	Refines -fno-stack-protector injection logic to avoid breaking non-gfortran f90 compilers.
gdb/testsuite/gdb.rocm/omp-target-basic.exp	New basic OpenMP target offload debugging checks (breakpoint/locals/step).
gdb/testsuite/gdb.rocm/omp-target-basic.c	Test program for basic target breakpoint/locals/step coverage.
gdb/testsuite/gdb.rocm/omp-target-teams.exp	New OpenMP teams/distribute/parallel-for wave/thread switching test.
gdb/testsuite/gdb.rocm/omp-target-teams.c	Test program intended to create multiple waves at a target breakpoint.
gdb/testsuite/gdb.rocm/omp-target-data.exp	New target data reuse test with two consecutive target regions.
gdb/testsuite/gdb.rocm/omp-target-data.c	Test program exercising device buffer reuse across kernels under target data.
gdb/testsuite/gdb.rocm/omp-target-multi-kernel.exp	New test for distinct declare target device functions hit by named breakpoints.
gdb/testsuite/gdb.rocm/omp-target-multi-kernel.c	Test program with square_dev and cube_dev launched from separate target regions.
gdb/testsuite/gdb.rocm/omp-target-cpp.exp	New C++-specific offload test for template and functor debugging.
gdb/testsuite/gdb.rocm/omp-target-cpp.cpp	C++ test program using a templated device function and struct functor in a target region.
gdb/testsuite/gdb.rocm/omp-target-fortran.exp	New Fortran OpenMP-offload breakpoint/wave test.
gdb/testsuite/gdb.rocm/omp-target-fortran.f90	Fortran offload test program for !$omp target loop.

---

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

[spatrang]

Thanks for the review @lumachad. Following your note about not blocking #62: I'll keep the OpenMP-offload series here in #134 and base it on top of #62's branch, so #62 stays independent and #134 is reviewed as close-to-final code. The OMP-specific parts (offload flags, test helpers, the allow_* gate) stay here and build on #62's reworked compiler discovery/search-order rather than being pushed into #62. Replied inline on the discovery comments with specifics.

Already addressed in the meantime:

• Dropped the with_rocm_omp_gpu_lock alias; the tests now use with_rocm_gpu_lock directly.
• Reworded the omp-target-teams "multiple waves" comments to match the "at least one wave" assertion.

[aktemur]

I looked at the omp-target-teams test and rocm.exp only. I think we should first resolve infrastructure questions.

[spatrang]

@aktemur Agreed on settling the infrastructure. Per @lumachad's note about keeping #62 unblocked, I'll base #134 on top of #62's series and keep the OpenMP-offload infra + tests here so it's reviewed close-to-final, rather than holding it or pushing pieces into #62. The shared bits (common omp/rocm build option, allow_* reuse, search order) build on #62's foundation. The test-level cleanups (braces, standard_testfile, simplified teams block) are ready to apply.

[spatrang]

Changelog — 6ce4ae405be (gdb/testsuite: filter benign OpenMP-offload linker warnings)

Pushed a follow-up commit addressing the review feedback on the OpenMP-offload tests.

Test infrastructure (gdb/testsuite/lib/rocm.exp)

• Added rocm_omp_prune_benign_warnings to drop the known, harmless ROCm 7.14 ld.lld __keep_alive "local memory global used by non-kernel function" warning. The match is anchored to a single line so genuine compiler/linker errors are never swallowed.
• Added rocm_omp_compile, which compiles quietly, prunes the benign warning, and still logs any remaining output via verbose -log so real failures stay diagnosable.
• Routed gdb_compile_omp_offload{,_cpp,_fortran} through rocm_omp_compile.
• Dropped the with_rocm_omp_gpu_lock alias; tests call with_rocm_gpu_lock directly.
• Fixed rocm → ROCm spelling in comments.

Tests (gdb.rocm/omp-target-*)

• Switched the .exp files to with_rocm_gpu_lock and standard_testfile (no redundant .c).
• omp-target-teams.c: added braces to the for-loop bodies for stable breakpoint line numbers.
• omp-target-teams.exp: removed the explicit wave thread-switch (the breakpoint hit already selects the wave context), dropped the redundant >= 1 guard, reworded the comments/assertion to "at least one AMDGPU wave", and kept a plain gdb_test for print i >= 0 (valid for any lane, including lane 0).

Why the warning filter
The binary builds and runs correctly on ROCm 7.14; only the stderr warning tripped DejaGnu into marking omp-target-teams.exp UNTESTED. With the filter the test now runs.

Verification

• ROCm 7.14.0: all 6 omp-target-* tests pass (39 PASS, 0 FAIL, 0 UNTESTED).
• ROCm 7.13.0: all 6 pass (39 PASS), no regression.

Note: the compiler-discovery and omp/rocm build-option items remain tracked against the #62-based rework, as discussed in the resolved threads.

[lumachad]

Set to @aktemur as he left the earlier comments.

[spatrang]

Rebased on amd-staging + review fixes (a5e145b6d9e)

Rebased onto latest amd-staging (the hipcc→amdclang++ switch), squashed into one commit, and applied review feedback.

Rebase

• Clean rebase; renamed allow_hipcc_tests → allow_hip_tests reference.
• OpenMP tests use their own amdclang discovery, so they're unaffected by the HIP board change.

Review fixes

• allow_omp_offload_tests now reuses allow_hip_tests for the shared GPU/dbgapi gate instead of duplicating it.
• Dropped the dead f77 clause in gdb.exp.
• Prune regex now also drops trailing >>> lld note lines (still anchored to __keep_alive).
• Reworded rocm_omp_compile log to "compile output".
• Removed unused sq-line/cb-line markers.
• Documented why the custom rocm_omp_continue_* helpers exist.

Note on \r\n anchoring: skipped — gdb_test here already appends \r\n$gdb_prompt, so adding it broke every value test. The implicit anchor already covers the check-read1 concern.

Testing: ROCm 7.14 — all 6 tests pass (40 passes, 0 failures).

[aktemur]

Nit: '+' should be indented two spaces inside wrt out_arr. A preferred alternative is to parenthesize the arithmetic expression, like this:

    for (int i = 0; i < N; i++)
      out_arr[i] = (in_arr[i] + priv_val + firstpriv_val
		    + step_two + step_three);

[aktemur]

I don't think using custom procs for these GDB commands will scale well. Can we please try to use the usual gdb_breakpoint, gdb_test, etc. procs? This applies to procs like rocm_omp_check_on_wave, rocm_omp_next, rocm_omp_continue_to_exit`, etc.

[aktemur]

We don't need this prefix; there is next-$which below already.

[aktemur]

My first expectation would be that we use the usual build_executable/prepare_for_testing by passing it the proper option (e.g. omp rocm), like we do for HIP tests. Then the testing framework finds the right compiler and sets the right flags, again similar to how HIP testing is done.

[aktemur]

By passing the proper omp and rocm options to gdb_compile, I think all can be handled in gdb_compile. So, this custom proc and the related ones below can be removed.

[aktemur]

Ideally, the custom procs should all go away and we should be able to use the usual gdb_... procs in .exp files.

[spatrang]

@aktemur Thank you for the recent review.

On moving the OpenMP-offload tests to build_executable/prepare_for_testing with an omp rocm-style build option (mirroring how HIP works): agreed, I'll do that. One design question before I implement it.

Since amdclang isn't the only OpenMP compiler and the driver differs by language (C -> amdclang, C++ -> amdclang++, Fortran -> amdflang), would you prefer omp to set the common flags (-fopenmp -fopenmp-targets=amdgcn-amd-amdhsa, --offload-arch=...) and rocm to select the AMD compiler, with the language coming from the usual c++ / f90 options? I'll follow whatever option split you think reads best.

The remaining comments are clear and I'll address them in the next revision (fold the compile wrappers into gdb_compile, drop the now-redundant find_* helpers since #62 is merged, switch to the stock gdb_breakpoint/gdb_test/gdb_continue_to_breakpoint/gdb_continue_to_termination procs, remove the redundant with_test_prefix, and fix the omp-target-basic.c indentation).

[aktemur]

Since amdclang isn't the only OpenMP compiler and the driver differs by language (C -> amdclang, C++ -> amdclang++, Fortran -> amdflang), would you prefer omp to set the common flags (-fopenmp -fopenmp-targets=amdgcn-amd-amdhsa, --offload-arch=...) and rocm to select the AMD compiler, with the language coming from the usual c++ / f90 options? I'll follow whatever option split you think reads best.

Hi @spatrang, as the first approach, I'd suggest having two arguments: (1) language, (2) omp-rocm. Please take this as a mere suggestion and not as a "must". If you see a better split, we can certainly evaluate. I think we should avoid over-engineering and prefer simplicity. Splitting omp-rocm into two arguments, omp and rocm, may complicate the flow unnecessarily, because then we need to think about having "rocm" but no "omp" and vice-versa, too.