diff options
| -rw-r--r-- | clang/include/clang/Basic/Attr.td | 8 | ||||
| -rw-r--r-- | clang/lib/Sema/SemaDeclAttr.cpp | 10 | ||||
| -rw-r--r-- | clang/test/SemaCUDA/amdgpu-num-gpr-attr.cu | 12 | ||||
| -rw-r--r-- | clang/test/SemaOpenCL/amdgpu-num-register-attrs.cl | 6 |
4 files changed, 28 insertions, 8 deletions
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 63bb7dd1687..84374611146 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/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/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index af7cf5713de..ffb8b77a7d0 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/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/clang/test/SemaCUDA/amdgpu-num-gpr-attr.cu b/clang/test/SemaCUDA/amdgpu-num-gpr-attr.cu index f6e28454b74..83acbc5007e 100644 --- a/clang/test/SemaCUDA/amdgpu-num-gpr-attr.cu +++ b/clang/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/clang/test/SemaOpenCL/amdgpu-num-register-attrs.cl b/clang/test/SemaOpenCL/amdgpu-num-register-attrs.cl index 0967700d109..ec8f8c072a2 100644 --- a/clang/test/SemaOpenCL/amdgpu-num-register-attrs.cl +++ b/clang/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() {} |
