48 lines · plain
1// REQUIRES: amdgpu-registered-target2// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx906 -x hip -fcuda-is-device -emit-llvm %s \3// RUN: -o - | FileCheck %s4 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