Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

use constexpr in device function #1580

Closed
s9105947 opened this issue Jan 24, 2022 · 3 comments
Closed

use constexpr in device function #1580

s9105947 opened this issue Jan 24, 2022 · 3 comments

Comments

@s9105947
Copy link
Contributor

Dear Maintainers,

when I refer to a C++ constexpr function from inside a (inline-) device function, should I mark the constexpr as a device function as well?

In a short discussion @j-stephan noted that this (probably) depends on the used compiler.

Apart from the question: Where/How should this be noted in the documentation?
I'd suggest a Notes/Misc/Q&A section, but I haven't found any.

@fwyzard
Copy link
Contributor

fwyzard commented Jan 24, 2022

nvcc --expt-relaxed-constexpr will consider any constexpr function as __host__ __device__

I'm not sure what clang does, nor if HIP has a similar flag.

@j-stephan
Copy link
Member

As @fwyzard said CUDA allows this with an extra compiler flag. According to ROCm/HIP#374 this is explicitly not supported for HIP.

So right now constexpr functions should also be marked as device functions.

@s9105947
Copy link
Contributor Author

Copy-paste-able text for documentation (I don't know where to put it):

Using constexpr from device functions
-------------------------------------

When calling a ``constexpr`` function from inside a device function, also mark the called function as a device function, e.g. by prepending ``ALPAKA_FN_ACC``.

Note that some compilers do that by default, but not all.
For details please refer to `#1580 <https://github.com/alpaka-group/alpaka/issues/1580>`_ .

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

No branches or pull requests

4 participants