This is an archive of the discontinued LLVM Phabricator instance.

[CUDA] Add __declspec spellings for CUDA attributes.
ClosedPublic

Authored by jlebar on Jan 4 2017, 1:58 PM.

Event Timeline

jlebar updated this revision to Diff 83128.Jan 4 2017, 1:58 PM
jlebar retitled this revision from to [CUDA] Add __declspec spellings for CUDA attributes..
jlebar updated this object.
jlebar added a reviewer: tra.
jlebar added subscribers: rnk, cfe-commits.
aaron.ballman added a subscriber: aaron.ballman.
aaron.ballman added inline comments.
clang/include/clang/Basic/Attr.td
604

For my own edification, do you have a link to some documentation of what CUDA attributes are spelled with __declspec?

610

Now would be a good time to add the documentation for these attributes, otherwise there's a lot less chance users will know what ways they can spell the attributes (or what the attribute do).

jlebar added inline comments.Jan 4 2017, 2:09 PM
clang/include/clang/Basic/Attr.td
604

I do not believe such documentation exists. It is an implementation detail in the CUDA headers -- users never write __attribute__((device)) or __declspec(__device__). Instead, they write __device__.

610

See above, these are an implementation detail.

jlebar updated this object.Jan 4 2017, 2:10 PM
aaron.ballman added inline comments.Jan 4 2017, 2:20 PM
clang/include/clang/Basic/Attr.td
604

Then why are we introducing __declspec(__device__) rather than a keyword attribute __device__?

My biggest concern is: I would like to verify that these actually should be supported as a __declspec attribute. From my simple testing in MSVC, it does not appear to support __declspec(__device__), but perhaps I am doing it wrong (I'm mostly unfamiliar with CUDA). If this isn't something MSVC supports, then it is the first attribute we're supporting with a __declspec spelling that is not actually a Microsoft attribute, which is something worth discussing.

610

We can still document that we support the attributes under their macro names, or do users not typically think of these macros as being attributes? I am mostly concerned about discoverability of the attributes -- how is a user to know what Clang does or does not support?

jlebar added inline comments.Jan 4 2017, 2:26 PM
clang/include/clang/Basic/Attr.td
604

Then why are we introducing declspec(device) rather than a keyword attribute device__?

__device__ is a macro defined in the CUDA headers, which must include and we are not able to modify.

610

how is a user to know what Clang does or does not support?

These macros are fundamental to CUDA support. The statement "you can compile CUDA code with clang" will immediately imply to every CUDA developer in existence the statement "you can use __device__ in your code".

aaron.ballman added inline comments.Jan 4 2017, 2:43 PM
clang/include/clang/Basic/Attr.td
604

Okay, it makes sense as to why we can't have a keyword spelling, but it also doesn't answer why we need the __declspec spelling for it. Are you saying that there are CUDA headers out there where this attribute is spelled with __attribute__ and others with __declspec? If so, then this change makes a bit more sense to me.

610

Yet we add new CUDA attributes periodically, and CUDA comes out with new versions and new features.

Looking at the CUDA docs, I see __managed__, but I don't see such an attribute in Clang. How is a user to know whether we do/don't support such a construct?

aaron.ballman accepted this revision.Jan 4 2017, 3:01 PM
aaron.ballman edited edge metadata.

After getting some realtime clarifications in IRC, I now understand better why this is needed. This patch LGTM! The documentation points I raised are still valid, but are by no means required for this patch to go in.

This revision is now accepted and ready to land.Jan 4 2017, 3:01 PM
jlebar added a comment.Jan 4 2017, 3:04 PM

Thank you for the review!

clang/include/clang/Basic/Attr.td
604

(Phab bug; I can't delete this comment. Please ignore.)

610

This is a fair point, I agree on this basis that we should add documentation here.

This revision was automatically updated to reflect the committed changes.