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