diff options
| author | Yaxun Liu <Yaxun.Liu@amd.com> | 2018-06-12 23:58:59 +0000 |
|---|---|---|
| committer | Yaxun Liu <Yaxun.Liu@amd.com> | 2018-06-12 23:58:59 +0000 |
| commit | aa24601f9890c5cbc9b33caf8a1d058776ece0af (patch) | |
| tree | ce4fc30864508d99c2340e1c537e1f8d931503db /clang/test/SemaOpenCL | |
| parent | e6fdb6a28c5dbaabfd15e9b991eeb1b160870cf2 (diff) | |
| download | bcm5719-llvm-aa24601f9890c5cbc9b33caf8a1d058776ece0af.tar.gz bcm5719-llvm-aa24601f9890c5cbc9b33caf8a1d058776ece0af.zip | |
[CUDA][HIP] Allow CUDA __global__ functions to have amdgpu kernel attributes
There are HIP applications e.g. Tensorflow 1.3 using amdgpu kernel attributes, however
currently they are only allowed on OpenCL kernel functions.
This patch will allow amdgpu kernel attributes to be applied to CUDA/HIP __global__
functions.
Differential Revision: https://reviews.llvm.org/D47958
llvm-svn: 334561
Diffstat (limited to 'clang/test/SemaOpenCL')
| -rw-r--r-- | clang/test/SemaOpenCL/invalid-kernel-attrs.cl | 8 |
1 files changed, 4 insertions, 4 deletions
diff --git a/clang/test/SemaOpenCL/invalid-kernel-attrs.cl b/clang/test/SemaOpenCL/invalid-kernel-attrs.cl index 1f1359cceea..a19a989f41b 100644 --- a/clang/test/SemaOpenCL/invalid-kernel-attrs.cl +++ b/clang/test/SemaOpenCL/invalid-kernel-attrs.cl @@ -14,11 +14,11 @@ kernel __attribute__((work_group_size_hint(8,16,32,4))) void kernel6() {} //expe kernel __attribute__((work_group_size_hint(1,2,3))) __attribute__((work_group_size_hint(3,2,1))) void kernel7() {} //expected-warning{{attribute 'work_group_size_hint' is already applied with different parameters}} -__attribute__((reqd_work_group_size(8,16,32))) void kernel8(){} // expected-error {{attribute 'reqd_work_group_size' can only be applied to a kernel}} +__attribute__((reqd_work_group_size(8,16,32))) void kernel8(){} // expected-error {{attribute 'reqd_work_group_size' can only be applied to an OpenCL kernel}} -__attribute__((work_group_size_hint(8,16,32))) void kernel9(){} // expected-error {{attribute 'work_group_size_hint' can only be applied to a kernel}} +__attribute__((work_group_size_hint(8,16,32))) void kernel9(){} // expected-error {{attribute 'work_group_size_hint' can only be applied to an OpenCL kernel}} -__attribute__((vec_type_hint(char))) void kernel10(){} // expected-error {{attribute 'vec_type_hint' can only be applied to a kernel}} +__attribute__((vec_type_hint(char))) void kernel10(){} // expected-error {{attribute 'vec_type_hint' can only be applied to an OpenCL kernel}} constant int foo1 __attribute__((reqd_work_group_size(8,16,32))) = 0; // expected-error {{'reqd_work_group_size' attribute only applies to functions}} @@ -34,6 +34,6 @@ kernel __attribute__((reqd_work_group_size(1,2,0))) void kernel11(){} // expecte kernel __attribute__((reqd_work_group_size(1,0,2))) void kernel12(){} // expected-error {{'reqd_work_group_size' attribute must be greater than 0}} kernel __attribute__((reqd_work_group_size(0,1,2))) void kernel13(){} // expected-error {{'reqd_work_group_size' attribute must be greater than 0}} -__attribute__((intel_reqd_sub_group_size(8))) void kernel14(){} // expected-error {{attribute 'intel_reqd_sub_group_size' can only be applied to a kernel}} +__attribute__((intel_reqd_sub_group_size(8))) void kernel14(){} // expected-error {{attribute 'intel_reqd_sub_group_size' can only be applied to an OpenCL kernel}} kernel __attribute__((intel_reqd_sub_group_size(0))) void kernel15(){} // expected-error {{'intel_reqd_sub_group_size' attribute must be greater than 0}} kernel __attribute__((intel_reqd_sub_group_size(8))) __attribute__((intel_reqd_sub_group_size(16))) void kernel16() {} //expected-warning{{attribute 'intel_reqd_sub_group_size' is already applied with different parameters}} |

