Skip to content

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

Draft
spatrang wants to merge 1 commit into
amd-stagingfrom
users/spatrang/gdb-rocm-omp-offload-tests
Draft

gdb/testsuite/gdb.rocm: add OpenMP GPU offload tests#134
spatrang wants to merge 1 commit into
amd-stagingfrom
users/spatrang/gdb-rocm-omp-offload-tests

Conversation

@spatrang
Copy link
Copy Markdown

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
Copy link
Copy Markdown
Contributor

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

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants