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

Workaround attribute ordering issue with kernel only attributes 

Placing the attribute after the kernel keyword would incorrectly
reject the attribute, so use the smae workaround that other
kernel only attributes use.

Also add a FIXME because there are two different phrasings now
for the same error, althoug amdgpu_num_[sv]gpr uses a consistent one.

llvm-svn: 223490
This commit is contained in:
Matt Arsenault 2014-12-05 18:03:58 +00:00
parent 43cfcbca3f
commit b9e9dc5e89
4 changed files with 28 additions and 8 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -905,7 +905,11 @@ def AMDGPUNumVGPR : InheritableAttr {
let Spellings = [GNU<"amdgpu_num_vgpr">];
let Args = [UnsignedArgument<"NumVGPR">];
let Documentation = [AMDGPUNumVGPRDocs];
let Subjects = SubjectList<[OpenCLKernelFunction], ErrorDiag,
// FIXME: This should be for OpenCLKernelFunction, but is not to
// workaround needing to see kernel attribute before others to know if
// this should be rejected on non-kernels.
let Subjects = SubjectList<[Function], ErrorDiag,
"ExpectedKernelFunction">;
}
@ -913,7 +917,7 @@ def AMDGPUNumSGPR : InheritableAttr {
let Spellings = [GNU<"amdgpu_num_sgpr">];
let Args = [UnsignedArgument<"NumSGPR">];
let Documentation = [AMDGPUNumSGPRDocs];
let Subjects = SubjectList<[OpenCLKernelFunction], ErrorDiag,
let Subjects = SubjectList<[Function], ErrorDiag,
"ExpectedKernelFunction">;
}

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

@ -4762,6 +4762,8 @@ void Sema::ProcessDeclAttributeList(Scope *S, Decl *D,
if (!D->hasAttr<OpenCLKernelAttr>()) {
// These attributes cannot be applied to a non-kernel function.
if (Attr *A = D->getAttr<ReqdWorkGroupSizeAttr>()) {
// FIXME: This emits a different error message than
// diag::err_attribute_wrong_decl_type + ExpectedKernelFunction.
Diag(D->getLocation(), diag::err_opencl_kernel_attr) << A;
D->setInvalidDecl();
} else if (Attr *A = D->getAttr<WorkGroupSizeHintAttr>()) {
@ -4770,6 +4772,14 @@ void Sema::ProcessDeclAttributeList(Scope *S, Decl *D,
} else if (Attr *A = D->getAttr<VecTypeHintAttr>()) {
Diag(D->getLocation(), diag::err_opencl_kernel_attr) << A;
D->setInvalidDecl();
} else if (Attr *A = D->getAttr<AMDGPUNumVGPRAttr>()) {
Diag(D->getLocation(), diag::err_attribute_wrong_decl_type)
<< A << ExpectedKernelFunction;
D->setInvalidDecl();
} else if (Attr *A = D->getAttr<AMDGPUNumSGPRAttr>()) {
Diag(D->getLocation(), diag::err_attribute_wrong_decl_type)
<< A << ExpectedKernelFunction;
D->setInvalidDecl();
}
}
}

12
clang/test/SemaCUDA/amdgpu-num-gpr-attr.cu
View File

@ -2,13 +2,13 @@
#include "Inputs/cuda.h"
__attribute__((amdgpu_num_vgpr(64))) // expected-error {{'amdgpu_num_vgpr' attribute only applies to kernel functions}}
__global__ void test_num_vgpr() { }
__attribute__((amdgpu_num_vgpr(64)))
__global__ void test_num_vgpr() { } // expected-error {{'amdgpu_num_vgpr' attribute only applies to kernel functions}}
__attribute__((amdgpu_num_sgpr(32))) // expected-error {{'amdgpu_num_sgpr' attribute only applies to kernel functions}}
__global__ void test_num_sgpr() { }
__attribute__((amdgpu_num_sgpr(32)))
__global__ void test_num_sgpr() { } // expected-error {{'amdgpu_num_sgpr' attribute only applies to kernel functions}}
// expected-error@+2 {{'amdgpu_num_sgpr' attribute only applies to kernel functions}}
// expected-error@+1 {{'amdgpu_num_vgpr' attribute only applies to kernel functions}}
// fixme-expected-error@+3 {{'amdgpu_num_sgpr' attribute only applies to kernel functions}}
// expected-error@+2 {{'amdgpu_num_vgpr' attribute only applies to kernel functions}}
__attribute__((amdgpu_num_sgpr(32), amdgpu_num_vgpr(64)))
__global__ void test_num_vgpr_num_sgpr() { }

6
clang/test/SemaOpenCL/amdgpu-num-register-attrs.cl
View File

@ -32,3 +32,9 @@ __attribute__((amdgpu_num_vgpr(4294967296))) kernel void foo13() {} // expected-
__attribute__((amdgpu_num_sgpr(4294967296))) kernel void foo14() {} // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}
__attribute__((amdgpu_num_sgpr(4294967296), amdgpu_num_vgpr(4294967296))) kernel void foo15() {} // expected-error 2 {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}
// Make sure it is accepted with kernel keyword before the attribute.
kernel __attribute__((amdgpu_num_vgpr(40))) void foo16() {}
kernel __attribute__((amdgpu_num_sgpr(40))) void foo17() {}