1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
|
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx906 -x hip -fcuda-is-device -emit-llvm %s \
// RUN: -o - | FileCheck %s
constexpr static int OpCtrl()
{
return 15 + 1;
}
constexpr static int RowMask()
{
return 3 + 1;
}
constexpr static int BankMask()
{
return 2 + 1;
}
constexpr static bool BountCtrl()
{
return true & false;
}
// CHECK: call i32 @llvm.amdgcn.update.dpp.i32(i32 %0, i32 %1, i32 16, i32 0, i32 0, i1 false)
__attribute__((global)) void test_update_dpp_const_fold_imm_operand_2(int* out, int a, int b)
{
*out = __builtin_amdgcn_update_dpp(a, b, OpCtrl(), 0, 0, false);
}
// CHECK: call i32 @llvm.amdgcn.update.dpp.i32(i32 %0, i32 %1, i32 0, i32 4, i32 0, i1 false)
__attribute__((global)) void test_update_dpp_const_fold_imm_operand_3(int* out, int a, int b)
{
*out = __builtin_amdgcn_update_dpp(a, b, 0, RowMask(), 0, false);
}
// CHECK: call i32 @llvm.amdgcn.update.dpp.i32(i32 %0, i32 %1, i32 0, i32 0, i32 3, i1 false)
__attribute__((global)) void test_update_dpp_const_fold_imm_operand_4(int* out, int a, int b)
{
*out = __builtin_amdgcn_update_dpp(a, b, 0, 0, BankMask(), false);
}
// CHECK: call i32 @llvm.amdgcn.update.dpp.i32(i32 %0, i32 %1, i32 0, i32 0, i32 0, i1 false)
__attribute__((global)) void test_update_dpp_const_fold_imm_operand_5(int* out, int a, int b)
{
*out = __builtin_amdgcn_update_dpp(a, b, 0, 0, 0, BountCtrl());
}
|
