[PATCH] Workaround attribute ordering issue with kernel only attributes

[Matt Arsenault]

Hi aaron.ballman,

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.

http://reviews.llvm.org/D6542

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

Index: include/clang/Basic/Attr.td
===================================================================
--- include/clang/Basic/Attr.td
+++ include/clang/Basic/Attr.td
@@ -905,15 +905,19 @@
   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">;
 }
 
 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">;
 }
 
Index: lib/Sema/SemaDeclAttr.cpp
===================================================================
--- lib/Sema/SemaDeclAttr.cpp
+++ lib/Sema/SemaDeclAttr.cpp
@@ -4762,14 +4762,24 @@
   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>()) {
       Diag(D->getLocation(), diag::err_opencl_kernel_attr) << A;
       D->setInvalidDecl();
     } 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();
     }
   }
 }
Index: test/SemaCUDA/amdgpu-num-gpr-attr.cu
===================================================================
--- test/SemaCUDA/amdgpu-num-gpr-attr.cu
+++ 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() { }
Index: test/SemaOpenCL/amdgpu-num-register-attrs.cl
===================================================================
--- test/SemaOpenCL/amdgpu-num-register-attrs.cl
+++ test/SemaOpenCL/amdgpu-num-register-attrs.cl
@@ -32,3 +32,9 @@
 __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() {}

_______________________________________________
cfe-commits mailing list
cfe-commits@cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits