replica
/
hanchenye-llvm-project
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

[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

llvm-svn: 291134
This commit is contained in:
Justin Lebar 2017-01-05 16:53:04 +00:00
parent 33387ea045
commit 0203f2c26e
2 changed files with 50 additions and 10 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

26
clang/include/clang/Basic/Attr.td
View File

@ -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];
}

34
clang/test/SemaCUDA/attr-declspec.cu Normal file
View File

@ -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.