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

   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
   5 constexpr static int OpCtrl()
   6 {
   7     return 15 + 1;
   8 }
   9
  10 constexpr static int RowMask()
  11 {
  12     return 3 + 1;
  13 }
  14
  15 constexpr static int BankMask()
  16 {
  17     return 2 + 1;
  18 }
  19
  20 constexpr 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 }