clang/test/CodeGenHIP/spirv-amdgcn-dpp-const-fold.hip

   1 // RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fcuda-is-device -emit-llvm %s \
   2 // RUN:   -o - | FileCheck %s
   3
   4 constexpr static int OpCtrl()
   5 {
   6     return 15 + 1;
   7 }
   8
   9 constexpr static int RowMask()
  10 {
  11     return 3 + 1;
  12 }
  13
  14 constexpr static int BankMask()
  15 {
  16     return 2 + 1;
  17 }
  18
  19 constexpr 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 }