diff options
| author | Yaxun Liu <Yaxun.Liu@amd.com> | 2017-03-10 01:30:46 +0000 |
|---|---|---|
| committer | Yaxun Liu <Yaxun.Liu@amd.com> | 2017-03-10 01:30:46 +0000 |
| commit | 7679a84815fd5a48f3e3d4e4c2a28882ec4d6574 (patch) | |
| tree | 5e56573896d3c14e685b0d5fb941f2b4617663ab /test/SemaOpenCL | |
| parent | 4b236427388eb74bd6e29130da2a4c9e6fd7424a (diff) | |
[AMDGPU] Add builtin functions readlane ds_permute mov_dpp
Differential Revision: https://reviews.llvm.org/D30551
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@297436 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'test/SemaOpenCL')
| -rw-r--r-- | test/SemaOpenCL/builtins-amdgcn-error.cl | 20 |
1 files changed, 15 insertions, 5 deletions
diff --git a/test/SemaOpenCL/builtins-amdgcn-error.cl b/test/SemaOpenCL/builtins-amdgcn-error.cl index 6cb2f38822..2639bf2775 100644 --- a/test/SemaOpenCL/builtins-amdgcn-error.cl +++ b/test/SemaOpenCL/builtins-amdgcn-error.cl @@ -1,16 +1,17 @@ // REQUIRES: amdgpu-registered-target // RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s -// FIXME: We only get one error if the functions are the other order in the -// file. - #pragma OPENCL EXTENSION cl_khr_fp64 : enable typedef unsigned long ulong; typedef unsigned int uint; -ulong test_s_memrealtime() +// To get all errors for feature checking we need to put them in one function +// since Clang will stop codegen for the next function if it finds error during +// codegen of the previous function. +void test_target_builtin(global int* out, int a) { - return __builtin_amdgcn_s_memrealtime(); // expected-error {{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}} + __builtin_amdgcn_s_memrealtime(); // expected-error {{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}} + *out = __builtin_amdgcn_mov_dpp(a, 0, 0, 0, false); // expected-error {{'__builtin_amdgcn_mov_dpp' needs target feature dpp}} } void test_s_sleep(int x) @@ -92,3 +93,12 @@ void test_s_getreg(global int* out, int a) { *out = __builtin_amdgcn_s_getreg(a); // expected-error {{argument to '__builtin_amdgcn_s_getreg' must be a constant integer}} } + +void test_mov_dpp2(global int* out, int a, int b, int c, int d, bool e) +{ + *out = __builtin_amdgcn_mov_dpp(a, b, 0, 0, false); // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}} + *out = __builtin_amdgcn_mov_dpp(a, 0, c, 0, false); // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}} + *out = __builtin_amdgcn_mov_dpp(a, 0, 0, d, false); // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}} + *out = __builtin_amdgcn_mov_dpp(a, 0, 0, 0, e); // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}} +} + |