include/llvm/IR/IntrinsicsAMDGPU.td

   1 //===- IntrinsicsAMDGPU.td - Defines AMDGPU intrinsics -----*- tablegen -*-===//
   2 //
   3 //                     The LLVM Compiler Infrastructure
   4 //
   5 // This file is distributed under the University of Illinois Open Source
   6 // License. See LICENSE.TXT for details.
   7 //
   8 //===----------------------------------------------------------------------===//
   9 //
  10 // This file defines all of the R600-specific intrinsics.
  11 //
  12 //===----------------------------------------------------------------------===//
  13
  14 let TargetPrefix = "r600" in {
  15
  16 class R600ReadPreloadRegisterIntrinsic<string name>
  17   : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>,
  18     GCCBuiltin<name>;
  19
  20 multiclass R600ReadPreloadRegisterIntrinsic_xyz<string prefix> {
  21   def _x : R600ReadPreloadRegisterIntrinsic<!strconcat(prefix, "_x")>;
  22   def _y : R600ReadPreloadRegisterIntrinsic<!strconcat(prefix, "_y")>;
  23   def _z : R600ReadPreloadRegisterIntrinsic<!strconcat(prefix, "_z")>;
  24 }
  25
  26 defm int_r600_read_global_size : R600ReadPreloadRegisterIntrinsic_xyz <
  27                                        "__builtin_r600_read_global_size">;
  28 defm int_r600_read_local_size : R600ReadPreloadRegisterIntrinsic_xyz <
  29                                        "__builtin_r600_read_local_size">;
  30 defm int_r600_read_ngroups : R600ReadPreloadRegisterIntrinsic_xyz <
  31                                        "__builtin_r600_read_ngroups">;
  32 defm int_r600_read_tgid : R600ReadPreloadRegisterIntrinsic_xyz <
  33                                        "__builtin_r600_read_tgid">;
  34 defm int_r600_read_tidig : R600ReadPreloadRegisterIntrinsic_xyz <
  35                                        "__builtin_r600_read_tidig">;
  36
  37 def int_r600_rat_store_typed :
  38   // 1st parameter: Data
  39   // 2nd parameter: Index
  40   // 3rd parameter: Constant RAT ID
  41   Intrinsic<[], [llvm_v4i32_ty, llvm_v4i32_ty, llvm_i32_ty], []>,
  42   GCCBuiltin<"__builtin_r600_rat_store_typed">;
  43
  44 } // End TargetPrefix = "r600"
  45
  46 let TargetPrefix = "AMDGPU" in {
  47
  48 class AMDGPUReadPreloadRegisterIntrinsic<string name>
  49   : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>,
  50     GCCBuiltin<name>;
  51
  52 def int_AMDGPU_div_scale : GCCBuiltin<"__builtin_amdgpu_div_scale">,
  53   // 1st parameter: Numerator
  54   // 2nd parameter: Denominator
  55   // 3rd parameter: Constant to select select between first and
  56   //                second. (0 = first, 1 = second).
  57   Intrinsic<[llvm_anyfloat_ty, llvm_i1_ty],
  58             [LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty],
  59             [IntrNoMem]>;
  60
  61 def int_AMDGPU_div_fmas : GCCBuiltin<"__builtin_amdgpu_div_fmas">,
  62   Intrinsic<[llvm_anyfloat_ty],
  63             [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty],
  64             [IntrNoMem]>;
  65
  66 def int_AMDGPU_div_fixup : GCCBuiltin<"__builtin_amdgpu_div_fixup">,
  67   Intrinsic<[llvm_anyfloat_ty],
  68             [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>],
  69             [IntrNoMem]>;
  70
  71 def int_AMDGPU_trig_preop : GCCBuiltin<"__builtin_amdgpu_trig_preop">,
  72   Intrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty],
  73             [IntrNoMem]>;
  74
  75 def int_AMDGPU_rcp : GCCBuiltin<"__builtin_amdgpu_rcp">,
  76   Intrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]>;
  77
  78 def int_AMDGPU_rsq : GCCBuiltin<"__builtin_amdgpu_rsq">,
  79   Intrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]>;
  80
  81 def int_AMDGPU_rsq_clamped : GCCBuiltin<"__builtin_amdgpu_rsq_clamped">,
  82   Intrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]>;
  83
  84 def int_AMDGPU_ldexp : GCCBuiltin<"__builtin_amdgpu_ldexp">,
  85   Intrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty], [IntrNoMem]>;
  86
  87 def int_AMDGPU_class : GCCBuiltin<"__builtin_amdgpu_class">,
  88   Intrinsic<[llvm_i1_ty], [llvm_anyfloat_ty, llvm_i32_ty], [IntrNoMem]>;
  89
  90 def int_AMDGPU_read_workdim : AMDGPUReadPreloadRegisterIntrinsic <
  91                                        "__builtin_amdgpu_read_workdim">;
  92
  93 } // End TargetPrefix = "AMDGPU"
  94
  95 let TargetPrefix = "amdgcn" in {
  96
  97 // SI only
  98 def int_amdgcn_buffer_wbinvl1_sc :
  99   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_sc">,
 100   Intrinsic<[], [], []>;
 101
 102 // On CI+
 103 def int_amdgcn_buffer_wbinvl1_vol :
 104   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_vol">,
 105   Intrinsic<[], [], []>;
 106
 107 def int_amdgcn_buffer_wbinvl1 :
 108   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1">,
 109   Intrinsic<[], [], []>;
 110
 111 def int_amdgcn_s_dcache_inv :
 112   GCCBuiltin<"__builtin_amdgcn_s_dcache_inv">,
 113   Intrinsic<[], [], []>;
 114
 115 // CI+
 116 def int_amdgcn_s_dcache_inv_vol :
 117   GCCBuiltin<"__builtin_amdgcn_s_dcache_inv_vol">,
 118   Intrinsic<[], [], []>;
 119
 120 // VI
 121 def int_amdgcn_s_dcache_wb :
 122   GCCBuiltin<"__builtin_amdgcn_s_dcache_wb">,
 123   Intrinsic<[], [], []>;
 124
 125 // VI
 126 def int_amdgcn_s_dcache_wb_vol :
 127   GCCBuiltin<"__builtin_amdgcn_s_dcache_wb_vol">,
 128   Intrinsic<[], [], []>;
 129
 130 def int_amdgcn_dispatch_ptr :
 131   GCCBuiltin<"__builtin_amdgcn_dispatch_ptr">,
 132   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>;
 133
 134 // __builtin_amdgcn_interp_p1 <i>, <attr_chan>, <attr>, <m0>
 135 def int_amdgcn_interp_p1 :
 136   GCCBuiltin<"__builtin_amdgcn_interp_p1">,
 137   Intrinsic<[llvm_float_ty],
 138             [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
 139             [IntrNoMem]>;  // This intrinsic reads from lds, but the memory
 140                            // values are constant, so it behaves like IntrNoMem.
 141
 142 // __builtin_amdgcn_interp_p2 <p1>, <j>, <attr_chan>, <attr>, <m0>
 143 def int_amdgcn_interp_p2 :
 144   GCCBuiltin<"__builtin_amdgcn_interp_p2">,
 145   Intrinsic<[llvm_float_ty],
 146             [llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
 147             [IntrNoMem]>;  // See int_amdgcn_v_interp_p1 for why this is
 148                            // IntrNoMem.
 149
 150 def int_amdgcn_mbcnt_lo :
 151   GCCBuiltin<"__builtin_amdgcn_mbcnt_lo">,
 152   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 153
 154 def int_amdgcn_mbcnt_hi :
 155   GCCBuiltin<"__builtin_amdgcn_mbcnt_hi">,
 156   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 157 }