[CUDA] Add __declspec spellings for CUDA attributes.

[jlebar]

Summary

CUDA attributes are spelled __declspec(__foo__) on Windows.

[aaron.ballman]

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]

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.

[aaron.ballman]

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]

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]

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]

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.

[jlebar]

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.

clang/include/clang/Basic/Attr.td

 
 def Constructor : InheritableAttr {
   let Spellings = [GCC<"constructor">];
   let Args = [DefaultIntArgument<"Priority", 65535>];
   let Subjects = SubjectList<[Function]>;
   let Documentation = [Undocumented];
 }
 
+// CUDA attributes are spelled __attribute__((attr)) or __declspec(__attr__).
+
 def CUDAConstant : InheritableAttr {
-  let Spellings = [GNU<"constant">];
+  let Spellings = [GNU<"constant">, Declspec<"__constant__">];
   let Subjects = SubjectList<[Var]>;
   let LangOpts = [CUDA];
   let Documentation = [Undocumented];
 }
 
 def CUDACudartBuiltin : IgnoredAttr {
-  let Spellings = [GNU<"cudart_builtin">];
+  let Spellings = [GNU<"cudart_builtin">, Declspec<"__cudart_builtin__">];
   let LangOpts = [CUDA];
 }
 
 def CUDADevice : InheritableAttr {
-  let Spellings = [GNU<"device">];
+  let Spellings = [GNU<"device">, Declspec<"__device__">];
   let Subjects = SubjectList<[Function, Var]>;
   let LangOpts = [CUDA];
   let Documentation = [Undocumented];
 }
 
 def CUDADeviceBuiltin : IgnoredAttr {
-  let Spellings = [GNU<"device_builtin">];
+  let Spellings = [GNU<"device_builtin">, Declspec<"__device_builtin__">];
   let LangOpts = [CUDA];
 }
 
 def CUDADeviceBuiltinSurfaceType : IgnoredAttr {
-  let Spellings = [GNU<"device_builtin_surface_type">];
+  let Spellings = [GNU<"device_builtin_surface_type">,
+                   Declspec<"__device_builtin_surface_type__">];
   let LangOpts = [CUDA];
 }
 
 def CUDADeviceBuiltinTextureType : IgnoredAttr {
-  let Spellings = [GNU<"device_builtin_texture_type">];
+  let Spellings = [GNU<"device_builtin_texture_type">,
+                   Declspec<"__device_builtin_texture_type__">];
   let LangOpts = [CUDA];
 }
 
 def CUDAGlobal : InheritableAttr {
-  let Spellings = [GNU<"global">];
+  let Spellings = [GNU<"global">, Declspec<"__global__">];
   let Subjects = SubjectList<[Function]>;
   let LangOpts = [CUDA];
   let Documentation = [Undocumented];
 }
 
 def CUDAHost : InheritableAttr {
-  let Spellings = [GNU<"host">];
+  let Spellings = [GNU<"host">, Declspec<"__host__">];
   let Subjects = SubjectList<[Function]>;
   let LangOpts = [CUDA];
   let Documentation = [Undocumented];
 }
 
 def CUDAInvalidTarget : InheritableAttr {
   let Spellings = [];
   let Subjects = SubjectList<[Function]>;
   let LangOpts = [CUDA];
   let Documentation = [Undocumented];
 }
 
 def CUDALaunchBounds : InheritableAttr {
-  let Spellings = [GNU<"launch_bounds">];
+  let Spellings = [GNU<"launch_bounds">, Declspec<"__launch_bounds__">];
   let Args = [ExprArgument<"MaxThreads">, ExprArgument<"MinBlocks", 1>];
   let LangOpts = [CUDA];
   let Subjects = SubjectList<[ObjCMethod, FunctionLike], WarnDiag,
                              "ExpectedFunctionOrMethod">;
   // An AST node is created for this attribute, but is not used by other parts
   // of the compiler. However, this node needs to exist in the AST because
   // non-LLVM backends may be relying on the attribute's presence.
   let Documentation = [Undocumented];
 }
 
 def CUDAShared : InheritableAttr {
-  let Spellings = [GNU<"shared">];
+  let Spellings = [GNU<"shared">, Declspec<"__shared__">];
   let Subjects = SubjectList<[Var]>;
   let LangOpts = [CUDA];
   let Documentation = [Undocumented];
 }
 
 def C11NoReturn : InheritableAttr {
   let Spellings = [Keyword<"_Noreturn">];
   let Subjects = SubjectList<[Function], ErrorDiag>;

 }
 
 def NoThrow : InheritableAttr {
   let Spellings = [GCC<"nothrow">, Declspec<"nothrow">];
   let Documentation = [Undocumented];
 }
 
 def NvWeak : IgnoredAttr {
+  // No Declspec spelling of this attribute; the CUDA headers use
+  // __attribute__((nv_weak)) unconditionally.
   let Spellings = [GNU<"nv_weak">];
   let LangOpts = [CUDA];
 }
 
 def ObjCBridge : InheritableAttr {
   let Spellings = [GNU<"objc_bridge">];
   let Subjects = SubjectList<[Record, TypedefName], ErrorDiag,
                              "ExpectedStructOrUnionOrTypedef">;

clang/test/SemaCUDA/attr-declspec.cu

+// Test the __declspec spellings of CUDA attributes.
+//
+// RUN: %clang_cc1 -fsyntax-only -fms-extensions -verify %s
+// RUN: %clang_cc1 -fsyntax-only -fms-extensions -fcuda-is-device -verify %s
+// Now pretend that we're compiling a C file. There should be warnings.
+// RUN: %clang_cc1 -DEXPECT_WARNINGS -fms-extensions -fsyntax-only -verify -x c %s
+
+#if defined(EXPECT_WARNINGS)
+// expected-warning@+12 {{'__device__' attribute ignored}}
+// expected-warning@+12 {{'__global__' attribute ignored}}
+// expected-warning@+12 {{'__constant__' attribute ignored}}
+// expected-warning@+12 {{'__shared__' attribute ignored}}
+// expected-warning@+12 {{'__host__' attribute ignored}}
+//
+// (Currently we don't for the other attributes. They are implemented with
+// IgnoredAttr, which is ignored irrespective of any LangOpts.)
+#else
+// expected-no-diagnostics
+#endif
+
+__declspec(__device__) void f_device();
+__declspec(__global__) void f_global();
+__declspec(__constant__) int* g_constant;
+__declspec(__shared__) float *g_shared;
+__declspec(__host__) void f_host();
+__declspec(__device_builtin__) void f_device_builtin();
+typedef __declspec(__device_builtin__) const void *t_device_builtin;
+enum __declspec(__device_builtin__) e_device_builtin {E};
+__declspec(__device_builtin__) int v_device_builtin;
+__declspec(__cudart_builtin__) void f_cudart_builtin();
+__declspec(__device_builtin_surface_type__) unsigned long long surface_var;
+__declspec(__device_builtin_texture_type__) unsigned long long texture_var;
+
+// Note that there's no __declspec spelling of nv_weak.