[CUDA] Add __declspec spellings for CUDA attributes.

Summary: CUDA attributes are spelled __declspec(__foo__) on Windows.

Reviewers: tra

Subscribers: cfe-commits, rnk

Differential Revision: https://reviews.llvm.org/D28321

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@291134 91177308-0d34-0410-b5e6-96231b3b80d8

include/clang/Basic/Attr.td

@@ -601,49 +601,53 @@ def Constructor : InheritableAttr {
   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];
@@ -657,7 +661,7 @@ def CUDAInvalidTarget : InheritableAttr {
 }
 
 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,
@@ -669,7 +673,7 @@ def CUDALaunchBounds : InheritableAttr {
 }
 
 def CUDAShared : InheritableAttr {
-  let Spellings = [GNU<"shared">];
+  let Spellings = [GNU<"shared">, Declspec<"__shared__">];
   let Subjects = SubjectList<[Var]>;
   let LangOpts = [CUDA];
   let Documentation = [Undocumented];
@@ -1195,6 +1199,8 @@ def NoThrow : InheritableAttr {
 }
 
 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];
 }

test/SemaCUDA/attr-declspec.cu

@@ -0,0 +1,34 @@
+// 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.