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