Skip to content

test: re-enable hipFFT and rocFFT callback tests - #500

Open
mapatel-amd wants to merge 3 commits into
amd-stagingfrom
users/mapatel-amd/reenable-fft-callback-tests
Open

test: re-enable hipFFT and rocFFT callback tests#500
mapatel-amd wants to merge 3 commits into
amd-stagingfrom
users/mapatel-amd/reenable-fft-callback-tests

Conversation

@mapatel-amd

@mapatel-amd mapatel-amd commented Aug 18, 2026

Copy link
Copy Markdown
Contributor

Summary

Re-enables the hipfft_callback and rocfft_callback tests by removing their
entries from .github/build_tools/skip_manifest.py.

Why they were skipped

rocFFT's default store callback was a static device function, so the compiler
appended a .static.<hash> suffix to its mangled name. HIP CLR could not resolve
that symbol through hipModuleGetGlobal, and both examples aborted at runtime:

:0:.../clr/hipamd/src/hip_global.cpp:209 : Cannot create GlobalVar Obj for symbol:
_ZL31store_cb_default_complex_double.static.4d22c59b5a98cebc

Tracked as
ROCm/rocm-systems#7263.

Why it is safe to re-enable

The issue was fixed by [ROCm/rocm-libraries@004a100615a](https://github.com/ROCm/r
ocm-libraries/commit/004a100615a) — "rocFFT: give default FFT callbacks external
linkage to fix hipModuleGlobal (#8477)"
, merged 2026-07-10; #7263 was closed
the same day.

The default callbacks are now declared with external linkage rather than static:

// projects/rocfft/library/src/device/kernels/callback.h:108
inline __device__ auto store_cb_default_complex_double =
store_cb_default<rocfft_complex<double>>;

No .static.<hash> suffix, so hipModuleGetGlobal resolves the symbol. The
commit is an ancestor of develop and predates the current nightly package by
roughly five weeks, so the fix is present in what CI tests against.

Changes

Removed both entries from SKIP_MANIFEST (15 lines, one file). The examples
themselves are unchanged.

Verification

Generator output for --channel nightly --target gfx1100 --distro ubuntu-24.04 --install-method tarball-multi-arch:

skip_tests.txt SKIP_FROM_TEST
before hipfft_callback, rocfft_callback `Libraries/hipFFT/callback
Libraries/rocFFT/callback`
after (empty) (empty)

All other manifest entries (rocDecode and the rest) are unaffected.

Empty-exclude-list edge case: these were the only entries carrying a ctest
key, so skip_tests.txt is now empty while the workflow still passes
--exclude-from-file unconditionally. Both halves check out:

  • generate_skip_tests.py opens the file outside the if skip_tests: guard, so
    it is still created, just zero bytes.
  • ctest --exclude-from-file <empty file> runs normally and exits 0 (verified on
    ctest 3.28.3).

Not verified on hardware

Every ROCm SDK available locally predates the fix (newest 7.15.0a20260702; the
fix landed 2026-07-10), so the examples could not be run directly. CI on the next
nightly is the real check.

If either test still aborts with the GlobalVar Obj error, the right response is
to reopen
ROCm/rocm-systems#7263 with
the new log rather than silently restoring the skip — the manifest entry would
otherwise hide a regression in a fix that is supposed to have shipped.

These two tests were skipped because rocFFT's default store callback was a
static device function, and HIP CLR could not resolve the resulting
`.static.<hash>` mangled symbol through hipModuleGetGlobal:

    Cannot create GlobalVar Obj for symbol:
    _ZL31store_cb_default_complex_double.static.4d22c59b5a98cebc

ROCm/rocm-systems#7263 was fixed by ROCm/rocm-libraries@004a100615a
("rocFFT: give default FFT callbacks external linkage to fix
hipModuleGlobal", 2026-07-10) and closed the same day. The default
callbacks are now declared `inline __device__` rather than static, so the
symbol has external linkage and resolves. The fix predates the current
nightly by several weeks.

Drop both manifest entries so the tests run again under ctest and
`make test`.

Co-Authored-By: Claude <noreply@anthropic.com>
@mapatel-amd
mapatel-amd requested a review from a team as a code owner August 18, 2026 13:51
The build workflows filter on paths and listed .github/workflows/** but not
.github/build_tools/**, so changes to the CI tooling did not start a build.
That tooling decides what gets built and tested: skip_manifest.py and
generate_skip_tests.py produce the ctest exclude list and the make skip
variables, and configure_ci.py builds the job matrix.

A pull request that only edits the skip manifest therefore merges with no
build or test signal at all. This one is an example: re-enabling the hipFFT
and rocFFT callback tests touches only skip_manifest.py, so it ran the
linter and nothing else, and had to be dispatched by hand to get coverage.

Add .github/build_tools/** to the push and pull_request path filters in
both workflows so changes to the skip machinery are exercised by the same
matrix they control.

Co-Authored-By: Claude <noreply@anthropic.com>

@regan-amd regan-amd left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

ROCm/rocm-libraries@004a100615a is expected to resolve the issue that motivated this test to be disabled in the first place, indeed.
However, the function-pointer callback feature that these samples illustrate are expected to fail on gfx1250 (known issue). Ideally, the re-enabled tests should be skipped/disabled/not tested on gfx1250, if possible.

@zichguan-amd

Copy link
Copy Markdown
Collaborator

We don't run gfx1250 in our CI workflows so if it works for other archs then it should be fine to reenable them.

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.

3 participants