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

investigate if RAJA works without --extended-lambda flag #659

Closed
samuelpmishLLNL opened this issue Mar 28, 2022 · 3 comments
Closed

investigate if RAJA works without --extended-lambda flag #659

samuelpmishLLNL opened this issue Mar 28, 2022 · 3 comments
Labels
gpu GPU related

Comments

@samuelpmishLLNL
Copy link
Contributor

Last year, we ran into some really nasty bugs that came from using __host__ __device__ lambdas as q-functions in serac::Functional. Upon closer investigation, these bugs were caused by some of the caveats (specifically, #14) listed here:

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#extended-lambda-restrictions

The scary part about these caveats is that, in many cases: "no diagnostic will be generated". This means that code may compile without any warnings or errors, but violate some of the restrictions above, resulting in erroneous (and very hard to debug) binaries.

For this reason, I was hoping to avoid using the --extended-lambda flag, to explicitly prevent users from doing this.

It's unclear to me if RAJA needs this flag though. If it does, it might still be possible to restructure some of the files in serac to try and avoid encountering those restricted parts of extended lambdas.

@white238 white238 added the gpu GPU related label Mar 29, 2022
@white238
Copy link
Member

white238 commented Feb 6, 2023

Example created by @samuelpmishLLNL :

https://godbolt.org/z/7vabb56rx

"We take a user-provided function that gets inserted into an element loop. We intend to store the output, so we check the type of the output of that function beforehand.

Note how checking the return type of device functions from the host is inconsistently enforced by the compiler (lines 43-51). It's unclear if the compiler is silently failing with no diagnostic, or if it is actually doing the right thing."

@jamiebramwell
Copy link
Member

Here's a jist with some of the relevant issues and work arounds: https://gist.github.com/samuelpmish/5679a488c8724b215e96e107569b4fe9

@white238
Copy link
Member

Fixed in #1018

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

No branches or pull requests

3 participants