maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[CUDA] Report an error if code tries to mix incompatible CUDA attributes. 

Summary: Thanks to jhen for helping me figure this out.

Reviewers: tra, echristo

Subscribers: jhen

Differential Revision: http://reviews.llvm.org/D16129

llvm-svn: 257554
This commit is contained in:
Justin Lebar 2016-01-13 01:07:35 +00:00
parent c3340db77d
commit 3eaaf86397
4 changed files with 90 additions and 14 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

48
clang/lib/Sema/SemaDeclAttr.cpp
View File

@ -348,6 +348,25 @@ static void handleSimpleAttribute(Sema &S, Decl *D,
Attr.getAttributeSpellingListIndex()));
}
template <typename AttrType>
static void handleSimpleAttributeWithExclusions(Sema &S, Decl *D,
const AttributeList &Attr) {
handleSimpleAttribute<AttrType>(S, D, Attr);
}
/// \brief Applies the given attribute to the Decl so long as the Decl doesn't
/// already have one of the given incompatible attributes.
template <typename AttrType, typename IncompatibleAttrType,
typename... IncompatibleAttrTypes>
static void handleSimpleAttributeWithExclusions(Sema &S, Decl *D,
const AttributeList &Attr) {
if (checkAttrMutualExclusion<IncompatibleAttrType>(S, D, Attr.getRange(),
Attr.getName()))
return;
handleSimpleAttributeWithExclusions<AttrType, IncompatibleAttrTypes...>(S, D,
Attr);
}
/// \brief Check if the passed-in expression is of type int or bool.
static bool isIntOrBool(Expr *Exp) {
QualType QT = Exp->getType();
@ -3588,6 +3607,12 @@ static void handleOptimizeNoneAttr(Sema &S, Decl *D,
}
static void handleGlobalAttr(Sema &S, Decl *D, const AttributeList &Attr) {
if (checkAttrMutualExclusion<CUDADeviceAttr>(S, D, Attr.getRange(),
Attr.getName()) ||
checkAttrMutualExclusion<CUDAHostAttr>(S, D, Attr.getRange(),
Attr.getName())) {
return;
}
FunctionDecl *FD = cast<FunctionDecl>(D);
if (!FD->getReturnType()->isVoidType()) {
SourceRange RTRange = FD->getReturnTypeSourceRange();
@ -4558,14 +4583,6 @@ static void handleInterruptAttr(Sema &S, Decl *D, const AttributeList &Attr) {
handleARMInterruptAttr(S, D, Attr);
}
static void handleMips16Attribute(Sema &S, Decl *D, const AttributeList &Attr) {
if (checkAttrMutualExclusion<MipsInterruptAttr>(S, D, Attr.getRange(),
Attr.getName()))
return;
handleSimpleAttribute<Mips16Attr>(S, D, Attr);
}
static void handleAMDGPUNumVGPRAttr(Sema &S, Decl *D,
const AttributeList &Attr) {
uint32_t NumRegs;
@ -4955,7 +4972,8 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
handleDLLAttr(S, D, Attr);
break;
case AttributeList::AT_Mips16:
handleMips16Attribute(S, D, Attr);
handleSimpleAttributeWithExclusions<Mips16Attr, MipsInterruptAttr>(S, D,
Attr);
break;
case AttributeList::AT_NoMips16:
handleSimpleAttribute<NoMips16Attr>(S, D, Attr);
@ -5006,7 +5024,8 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
handleCommonAttr(S, D, Attr);
break;
case AttributeList::AT_CUDAConstant:
handleSimpleAttribute<CUDAConstantAttr>(S, D, Attr);
handleSimpleAttributeWithExclusions<CUDAConstantAttr, CUDASharedAttr>(S, D,
Attr);
break;
case AttributeList::AT_PassObjectSize:
handlePassObjectSizeAttr(S, D, Attr);
@ -5051,10 +5070,12 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
handleGlobalAttr(S, D, Attr);
break;
case AttributeList::AT_CUDADevice:
handleSimpleAttribute<CUDADeviceAttr>(S, D, Attr);
handleSimpleAttributeWithExclusions<CUDADeviceAttr, CUDAGlobalAttr>(S, D,
Attr);
break;
case AttributeList::AT_CUDAHost:
handleSimpleAttribute<CUDAHostAttr>(S, D, Attr);
handleSimpleAttributeWithExclusions<CUDAHostAttr, CUDAGlobalAttr>(S, D,
Attr);
break;
case AttributeList::AT_GNUInline:
handleGNUInlineAttr(S, D, Attr);
@ -5114,7 +5135,8 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
handleSimpleAttribute<NoThrowAttr>(S, D, Attr);
break;
case AttributeList::AT_CUDAShared:
handleSimpleAttribute<CUDASharedAttr>(S, D, Attr);
handleSimpleAttributeWithExclusions<CUDASharedAttr, CUDAConstantAttr>(S, D,
Attr);
break;
case AttributeList::AT_VecReturn:
handleVecReturnAttr(S, D, Attr);

4
clang/test/SemaCUDA/Inputs/cuda.h
View File

@ -2,6 +2,9 @@
#include <stddef.h>
// Make this file work with nvcc, for testing compatibility.
#ifndef __NVCC__
#define __constant__ __attribute__((constant))
#define __device__ __attribute__((device))
#define __global__ __attribute__((global))
@ -18,3 +21,4 @@ typedef struct cudaStream *cudaStream_t;
int cudaConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0,
cudaStream_t stream = 0);
#endif // !__NVCC__

3
clang/test/SemaCUDA/attributes.cu → clang/test/SemaCUDA/attributes-on-non-cuda.cu
View File

@ -1,4 +1,5 @@
// Tests handling of CUDA attributes.
// Tests that CUDA attributes are warnings when compiling C files, but not when
// compiling CUDA files.
//
// RUN: %clang_cc1 -fsyntax-only -verify %s
// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s

49
clang/test/SemaCUDA/bad-attributes.cu Normal file
View File

@ -0,0 +1,49 @@
// Tests handling of CUDA attributes that are bad either because they're
// applied to the wrong sort of thing, or because they're given in illegal
// combinations.
//
// You should be able to run this file through nvcc for compatibility testing.
//
// RUN: %clang_cc1 -fsyntax-only -verify %s
// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -verify %s
#include "Inputs/cuda.h"
// Try applying attributes to functions and variables. Some should generate
// warnings; others not.
__device__ int a1;
__device__ void a2();
__host__ int b1; // expected-warning {{attribute only applies to functions}}
__host__ void b2();
__constant__ int c1;
__constant__ void c2(); // expected-warning {{attribute only applies to variables}}
__shared__ int d1;
__shared__ void d2(); // expected-warning {{attribute only applies to variables}}
__global__ int e1; // expected-warning {{attribute only applies to functions}}
__global__ void e2();
// Try all pairs of attributes which can be present on a function or a
// variable. Check both orderings of the attributes, as that can matter in
// clang.
__device__ __host__ void z1();
__device__ __constant__ int z2;
__device__ __shared__ int z3;
__device__ __global__ void z4(); // expected-error {{attributes are not compatible}}
// expected-note@-1 {{conflicting attribute is here}}
__host__ __device__ void z5();
__host__ __global__ void z6(); // expected-error {{attributes are not compatible}}
// expected-note@-1 {{conflicting attribute is here}}
__constant__ __device__ int z7;
__constant__ __shared__ int z8; // expected-error {{attributes are not compatible}}
// expected-note@-1 {{conflicting attribute is here}}
__shared__ __device__ int z9;
__shared__ __constant__ int z10; // expected-error {{attributes are not compatible}}
// expected-note@-1 {{conflicting attribute is here}}
__global__ __device__ void z11(); // expected-error {{attributes are not compatible}}
// expected-note@-1 {{conflicting attribute is here}}
__global__ __host__ void z12(); // expected-error {{attributes are not compatible}}
// expected-note@-1 {{conflicting attribute is here}}