r223490 - Workaround attribute ordering issue with kernel only attributes

[Matt Arsenault]

Author: arsenm
Date: Fri Dec  5 12:03:58 2014
New Revision: 223490

URL: http://llvm.org/viewvc/llvm-project?rev=223490&view=rev
Log:
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.

Modified:
    cfe/trunk/include/clang/Basic/Attr.td
    cfe/trunk/lib/Sema/SemaDeclAttr.cpp
    cfe/trunk/test/SemaCUDA/amdgpu-num-gpr-attr.cu
    cfe/trunk/test/SemaOpenCL/amdgpu-num-register-attrs.cl

Modified: cfe/trunk/include/clang/Basic/Attr.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/Attr.td?rev=223490&r1=223489&r2=223490&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/Attr.td (original)
+++ cfe/trunk/include/clang/Basic/Attr.td Fri Dec  5 12:03:58 2014
@@ -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">;
 }
 

Modified: cfe/trunk/lib/Sema/SemaDeclAttr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDeclAttr.cpp?rev=223490&r1=223489&r2=223490&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDeclAttr.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp Fri Dec  5 12:03:58 2014
@@ -4762,6 +4762,8 @@ void Sema::ProcessDeclAttributeList(Scop
   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(Scop
     } 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();
     }
   }
 }

Modified: cfe/trunk/test/SemaCUDA/amdgpu-num-gpr-attr.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/amdgpu-num-gpr-attr.cu?rev=223490&r1=223489&r2=223490&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/amdgpu-num-gpr-attr.cu (original)
+++ cfe/trunk/test/SemaCUDA/amdgpu-num-gpr-attr.cu Fri Dec  5 12:03:58 2014
@@ -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 at +2 {{'amdgpu_num_sgpr' attribute only applies to kernel functions}}
-// expected-error at +1 {{'amdgpu_num_vgpr' attribute only applies to kernel functions}}
+// fixme-expected-error at +3 {{'amdgpu_num_sgpr' attribute only applies to kernel functions}}
+// expected-error at +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() { }

Modified: cfe/trunk/test/SemaOpenCL/amdgpu-num-register-attrs.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/amdgpu-num-register-attrs.cl?rev=223490&r1=223489&r2=223490&view=diff
==============================================================================
--- cfe/trunk/test/SemaOpenCL/amdgpu-num-register-attrs.cl (original)
+++ cfe/trunk/test/SemaOpenCL/amdgpu-num-register-attrs.cl Fri Dec  5 12:03:58 2014
@@ -32,3 +32,9 @@ __attribute__((amdgpu_num_vgpr(429496729
 __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() {}