Skip to content

[SYCL][HIP] Fix crash from SYCL buffer destructors during shutdown#22679

Open
zjin-lcf wants to merge 1 commit into
intel:syclfrom
zjin-lcf:sycl-hip-fix-static-buffer-dtor
Open

[SYCL][HIP] Fix crash from SYCL buffer destructors during shutdown#22679
zjin-lcf wants to merge 1 commit into
intel:syclfrom
zjin-lcf:sycl-hip-fix-static-buffer-dtor

Conversation

@zjin-lcf

Copy link
Copy Markdown
Contributor

Summary

  • Fix a shutdown crash triggered by SYCL buffer destructors: make waits during shutdown optional and tolerate HIP errors (e.g. context already destroyed) during event release.
  • Re-enable Regression/static-buffer-dtor on HIP.

Test plan

  • Regression/static-buffer-dtor passes on gfx942 (repeated runs)

Relates to #22300

Static SYCL buffer destructors run after the runtime has begun shutting
down (their atexit handlers are registered before the runtime's). On HIP
this crashed twice: waiting on events whose stream/context was already
released segfaults inside libamdhip64, and releasing those events returns
hipErrorContextIsDestroyed which was thrown out of the event_impl
destructor.

- Scheduler::removeMemoryObject: extend the existing Windows-only
  "skip wait during shutdown" guard to all platforms. Host memory is not
  written back during shutdown anyway, so the wait serves no purpose and
  is unsafe against partially torn-down adapters.
- HIP adapter event release(): treat hipErrorContextIsDestroyed and
  hipErrorDeinitialized as a successful release instead of surfacing an
  exception.

Re-enables Regression/static-buffer-dtor.cpp on HIP.

Co-authored-by: Cursor <cursoragent@cursor.com>
@zjin-lcf
zjin-lcf requested review from a team as code owners July 19, 2026 21:50
// (e.g. the HIP runtime segfaults when synchronizing on an event whose
// stream/context has been released). During shutdown host memory is not
// written back either (see SYCLMemObjT::updateHostMemory), so waiting serves
// no purpose here and the memory WILL be reclaimed as the process exits.

@KseniyaTikhomirova KseniyaTikhomirova Jul 21, 2026

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.

AFAIK on Linux we were able to correctly shutdown SYCL RT and release its resources. This change disables proper resource release on Linux (on Windows due to thread's management it was disabled by default).
Could you please clarify what exact resources are released in unexpected order? And callstack where it happens.

@KseniyaTikhomirova KseniyaTikhomirova Jul 21, 2026

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.

@cperkinsintel could you please provide your opinion about this change? I found that this test was disabled for HIP ~5 years ago, I assume that there are some static vars in HIP runtime that have a conflict with SYCL RT internals in terms of order. It seems to be a HIP specific issue. I am pretty hesitant to disable proper shutdown cleanup on Linux because of that.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

The crash is a static-destruction-order problem between the SYCL RT and the HIP runtime (libamdhip64), not a case of us releasing SYCL RT resources in the wrong order. Static SYCL buffers register their atexit destructors before the runtime registers its shutdown handler, so (LIFO) the runtime's shutdown runs first and the static buffer's destructor runs afterwards, once the backend context/stream and the UR events bound to them have already been torn down.

When that late buffer destructor then goes through removeMemoryObject -> waitForRecordToFinish, it synchronizes on events whose stream/context no longer exists, which segfaults inside libamdhip64. Approximate callstack:

__run_exit_handlers / __cxa_finalize
  static sycl::buffer::~buffer
    detail::buffer_impl::~buffer_impl
      detail::SYCLMemObjT::~SYCLMemObjT
        Scheduler::removeMemoryObject          // allowWait == true on Linux (pre-patch)
          Scheduler::waitForRecordToFinish
            GraphProcessor::waitForEvent
              event_impl::wait -> urEventWait
                HIP adapter: hipEventSynchronize/hipStreamSynchronize
                  libamdhip64 (stream/context already destroyed) -> SIGSEGV

There is a second, related failure on the release path: hipEventDestroy returns hipErrorContextIsDestroyed, and UR_CHECK_ERROR throws out of ~event_impl, aborting the process. That is handled separately in the HIP adapter change.

On the concern about disabling proper cleanup on Linux: the change does not affect normal-lifetime buffers. allowWait is still true whenever GlobalHandler::isOkToDefer() is true (i.e. the runtime is alive) or the buffer has a user host pointer, so anything destroyed while the app is running still waits and writes back exactly as before. The wait is only skipped when the runtime is already shutting down and there is no user host memory - and in that state SYCLMemObjT::updateHostMemory does not write back anyway, so the wait has no functional effect; it only risks touching already-released backend resources. That is the same rationale as the existing Windows path, which is why I generalized it rather than adding a separate case.

You're right that the actual manifestation is HIP-specific (libamdhip64 static teardown). Since the shutdown-time wait is a no-op for correctness on every backend, extending the existing guard uniformly seemed lower-risk than a HIP-only branch, but I'm happy to gate it on the HIP backend specifically if you'd prefer to keep Linux behavior unchanged for other adapters.

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