From: Matt Arsenault <Matthew.Arsenault@amd.com>
Date: Fri, 5 Dec 2014 18:03:58 +0000 (+0000)
Subject: Workaround attribute ordering issue with kernel only attributes
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=dcbb8fe6d53e6737e198954850aaa33cf19bdd1d;p=clang

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.

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@223490 91177308-0d34-0410-b5e6-96231b3b80d8
---

diff --git a/include/clang/Basic/Attr.td b/include/clang/Basic/Attr.td
index 63bb7dd168..8437461114 100644
--- a/include/clang/Basic/Attr.td
+++ b/include/clang/Basic/Attr.td
@@ -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">;
 }
 
diff --git a/lib/Sema/SemaDeclAttr.cpp b/lib/Sema/SemaDeclAttr.cpp
index af7cf5713d..ffb8b77a7d 100644
--- a/lib/Sema/SemaDeclAttr.cpp
+++ b/lib/Sema/SemaDeclAttr.cpp
@@ -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();
     }
   }
 }
diff --git a/test/SemaCUDA/amdgpu-num-gpr-attr.cu b/test/SemaCUDA/amdgpu-num-gpr-attr.cu
index f6e28454b7..83acbc5007 100644
--- a/test/SemaCUDA/amdgpu-num-gpr-attr.cu
+++ b/test/SemaCUDA/amdgpu-num-gpr-attr.cu
@@ -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() { }
diff --git a/test/SemaOpenCL/amdgpu-num-register-attrs.cl b/test/SemaOpenCL/amdgpu-num-register-attrs.cl
index 0967700d10..ec8f8c072a 100644
--- a/test/SemaOpenCL/amdgpu-num-register-attrs.cl
+++ b/test/SemaOpenCL/amdgpu-num-register-attrs.cl
@@ -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() {}