This revision was automatically updated to reflect the committed changes.
Closed by commit rL291134: [CUDA] Add __declspec spellings for CUDA attributes.
(authored by jlebar).
Changed prior to commit:
https://reviews.llvm.org/D28321?vs=83128&id=83257#toc
Repository:
rL LLVM
https://reviews.llvm.org/D28321
Files:
cfe/trunk/include/clang/Basic/Attr.td
cfe/trunk/test/SemaCUDA/attr-declspec.cu
Index: cfe/trunk/test/SemaCUDA/attr-declspec.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/attr-declspec.cu
+++ cfe/trunk/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.
Index: cfe/trunk/include/clang/Basic/Attr.td
===================================================================
--- cfe/trunk/include/clang/Basic/Attr.td
+++ cfe/trunk/include/clang/Basic/Attr.td
@@ -601,49 +601,53 @@
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 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 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 NvWeak : IgnoredAttr {
+ // No Declspec spelling of this attribute; the CUDA headers use
+ // __attribute__((nv_weak)) unconditionally.
let Spellings = [GNU<"nv_weak">];
let LangOpts = [CUDA];
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
