Page MenuHomePhabricator

[HIP] Support <functional> in device code
Needs ReviewPublic

Authored by yaxunl on May 14 2021, 9:10 AM.

Details

Reviewers
tra
Summary

This patch adds wrapper headers for <functional>
and a few others which is required to support
<functional>.

The basic idea is to make template functions
defined in these headers host device by pragmas.

Since this only works for libc++. The code is conditioned
for libc++ only. For libstdc++ it is NFC.

A test is added to llvm-test-suite for testing this:
https://reviews.llvm.org/D102508

Diff Detail

Event Timeline

yaxunl created this revision.May 14 2021, 9:10 AM
yaxunl requested review of this revision.May 14 2021, 9:10 AM
yaxunl edited the summary of this revision. (Show Details)May 14 2021, 9:22 AM
tra added a subscriber: rsmith.May 14 2021, 9:54 AM

In effect this patch applies __host__ __device__ to a subset of the standard library headers and whatever headers *they* happen to include. While it may happen to work, I'm not at all confident that it does not create interesting issues.

Considering that the patch only works with libc++ anyways, perhaps it's time to make (parts) of libc++ itself usable from CUDA/HIP, instead of hacking around it in the wrappers?

@rsmith Richard, who would be the right person to discuss the standard library changes we may need?

In effect this patch applies __host__ __device__ to a subset of the standard library headers and whatever headers *they* happen to include. While it may happen to work, I'm not at all confident that it does not create interesting issues.

Considering that the patch only works with libc++ anyways, perhaps it's time to make (parts) of libc++ itself usable from CUDA/HIP, instead of hacking around it in the wrappers?

@rsmith Richard, who would be the right person to discuss the standard library changes we may need?

ping.

If we are allowed to make changes to libc++ we may have cleaner implementation for supporting libc++ in HIP device functions.

Currently by default libc++ functions are host functions except constexpr functions. Except constexpr functions, we can't call libc++ host functions in HIP device functions. Our goal is to make libc++ functions __host__ __device__ functions so that they can be called in HIP device functions. We may not be able to support all libc++ functions, e.g. file I/O, threads, but at least we should be able to support some of them, e.g. type_traits, functional, containers. We do this by supporting the underlying functions e.g. malloc/free on device.

The change will be NFC for other languages.

rsmith added a subscriber: ldionne.Tue, Jun 1, 1:35 PM
rsmith added a comment.Tue, Jun 1, 1:40 PM

@ldionne How should we go about establishing whether libc++ would be prepared to officially support CUDA? Right now, Clang's CUDA support is patching in attributes onto libc++ functions from the outside, which doesn't seem like a sustainable model.