1 // REQUIRES
: amdgpu-registered-target
2 // RUN
: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck %s
3 // RUN
: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1011 -S -emit-llvm -o - %s | FileCheck %s
4 // RUN
: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1012 -S -emit-llvm -o - %s | FileCheck %s
6 typedef unsigned int uint
;
7 typedef unsigned long ulong
;
9 // CHECK-LABEL
: @test_permlane16
(
10 // CHECK
: call i32
@llvm.amdgcn.permlane16
(i32 %a
, i32 %b
, i32 %c
, i32 %d
, i1 false
, i1 false
)
11 void test_permlane16
(global uint
* out
, uint a
, uint b
, uint c
, uint d
) {
12 *out
= __builtin_amdgcn_permlane16
(a, b
, c
, d
, 0, 0);
15 // CHECK-LABEL
: @test_permlanex16
(
16 // CHECK
: call i32
@llvm.amdgcn.permlanex16
(i32 %a
, i32 %b
, i32 %c
, i32 %d
, i1 false
, i1 false
)
17 void test_permlanex16
(global uint
* out
, uint a
, uint b
, uint c
, uint d
) {
18 *out
= __builtin_amdgcn_permlanex16
(a, b
, c
, d
, 0, 0);
21 // CHECK-LABEL
: @test_mov_dpp8
(
22 // CHECK
: call i32
@llvm.amdgcn.mov.dpp8.i32
(i32 %a
, i32
1)
23 void test_mov_dpp8
(global uint
* out
, uint a
) {
24 *out
= __builtin_amdgcn_mov_dpp8
(a, 1);
27 // CHECK-LABEL
: @test_s_memtime
28 // CHECK
: call i64
@llvm.amdgcn.s.memtime
()
29 void test_s_memtime
(global ulong
* out
)
31 *out
= __builtin_amdgcn_s_memtime
();
34 // CHECK-LABEL
: @test_groupstaticsize
35 // CHECK
: call i32
@llvm.amdgcn.groupstaticsize
()
36 void test_groupstaticsize
(global uint
* out
)
38 *out
= __builtin_amdgcn_groupstaticsize
();