| // 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 %1, i32 %2, 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 %1, i32 %2, 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 %1, i32 %2, 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 %1, i32 %2, 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()); |
| } |