contrib/llvm/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 class AMDGPUReadPreloadRegisterIntrinsic
  15   : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>;
  16
  17 class AMDGPUReadPreloadRegisterIntrinsicNamed<string name>
  18   : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>, GCCBuiltin<name>;
  19
  20 let TargetPrefix = "r600" in {
  21
  22 multiclass AMDGPUReadPreloadRegisterIntrinsic_xyz {
  23   def _x : AMDGPUReadPreloadRegisterIntrinsic;
  24   def _y : AMDGPUReadPreloadRegisterIntrinsic;
  25   def _z : AMDGPUReadPreloadRegisterIntrinsic;
  26 }
  27
  28 multiclass AMDGPUReadPreloadRegisterIntrinsic_xyz_named<string prefix> {
  29   def _x : AMDGPUReadPreloadRegisterIntrinsicNamed<!strconcat(prefix, "_x")>;
  30   def _y : AMDGPUReadPreloadRegisterIntrinsicNamed<!strconcat(prefix, "_y")>;
  31   def _z : AMDGPUReadPreloadRegisterIntrinsicNamed<!strconcat(prefix, "_z")>;
  32 }
  33
  34 defm int_r600_read_global_size : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
  35                                  <"__builtin_r600_read_global_size">;
  36 defm int_r600_read_ngroups : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
  37                              <"__builtin_r600_read_ngroups">;
  38 defm int_r600_read_tgid : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
  39                           <"__builtin_r600_read_tgid">;
  40
  41 defm int_r600_read_local_size : AMDGPUReadPreloadRegisterIntrinsic_xyz;
  42 defm int_r600_read_tidig : AMDGPUReadPreloadRegisterIntrinsic_xyz;
  43
  44 def int_r600_read_workdim : AMDGPUReadPreloadRegisterIntrinsic;
  45
  46 def int_r600_group_barrier : GCCBuiltin<"__builtin_r600_group_barrier">,
  47   Intrinsic<[], [], [IntrConvergent]>;
  48
  49 // AS 7 is PARAM_I_ADDRESS, used for kernel arguments
  50 def int_r600_implicitarg_ptr :
  51   GCCBuiltin<"__builtin_r600_implicitarg_ptr">,
  52   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 7>], [], [IntrNoMem]>;
  53
  54 def int_r600_rat_store_typed :
  55   // 1st parameter: Data
  56   // 2nd parameter: Index
  57   // 3rd parameter: Constant RAT ID
  58   Intrinsic<[], [llvm_v4i32_ty, llvm_v4i32_ty, llvm_i32_ty], []>,
  59   GCCBuiltin<"__builtin_r600_rat_store_typed">;
  60
  61 def int_r600_recipsqrt_ieee :  Intrinsic<
  62   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
  63 >;
  64
  65 def int_r600_recipsqrt_clamped : Intrinsic<
  66   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
  67 >;
  68
  69 } // End TargetPrefix = "r600"
  70
  71 let TargetPrefix = "amdgcn" in {
  72
  73 defm int_amdgcn_workitem_id : AMDGPUReadPreloadRegisterIntrinsic_xyz;
  74 defm int_amdgcn_workgroup_id : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
  75                                <"__builtin_amdgcn_workgroup_id">;
  76
  77 def int_amdgcn_s_barrier : GCCBuiltin<"__builtin_amdgcn_s_barrier">,
  78   Intrinsic<[], [], [IntrConvergent]>;
  79
  80 def int_amdgcn_s_waitcnt : Intrinsic<[], [llvm_i32_ty], []>;
  81
  82 def int_amdgcn_div_scale : Intrinsic<
  83   // 1st parameter: Numerator
  84   // 2nd parameter: Denominator
  85   // 3rd parameter: Constant to select select between first and
  86   //                second. (0 = first, 1 = second).
  87   [llvm_anyfloat_ty, llvm_i1_ty],
  88   [LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty],
  89   [IntrNoMem]
  90 >;
  91
  92 def int_amdgcn_div_fmas : Intrinsic<[llvm_anyfloat_ty],
  93   [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty],
  94   [IntrNoMem]
  95 >;
  96
  97 def int_amdgcn_div_fixup : Intrinsic<[llvm_anyfloat_ty],
  98   [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>],
  99   [IntrNoMem]
 100 >;
 101
 102 def int_amdgcn_trig_preop : Intrinsic<
 103   [llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty], [IntrNoMem]
 104 >;
 105
 106 def int_amdgcn_sin : Intrinsic<
 107   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 108 >;
 109
 110 def int_amdgcn_cos : Intrinsic<
 111   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 112 >;
 113
 114 def int_amdgcn_log_clamp : Intrinsic<
 115   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 116 >;
 117
 118 def int_amdgcn_rcp : Intrinsic<
 119   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 120 >;
 121
 122 def int_amdgcn_rsq :  Intrinsic<
 123   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 124 >;
 125
 126 def int_amdgcn_rsq_legacy :  GCCBuiltin<"__builtin_amdgcn_rsq_legacy">,
 127   Intrinsic<
 128   [llvm_float_ty], [llvm_float_ty], [IntrNoMem]
 129 >;
 130
 131 def int_amdgcn_rsq_clamp : Intrinsic<
 132   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]>;
 133
 134 def int_amdgcn_ldexp : Intrinsic<
 135   [llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty], [IntrNoMem]
 136 >;
 137
 138 def int_amdgcn_frexp_mant : Intrinsic<
 139   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 140 >;
 141
 142 def int_amdgcn_frexp_exp : Intrinsic<
 143   [llvm_i32_ty], [llvm_anyfloat_ty], [IntrNoMem]
 144 >;
 145
 146 // v_fract is buggy on SI/CI. It mishandles infinities, may return 1.0
 147 // and always uses rtz, so is not suitable for implementing the OpenCL
 148 // fract function. It should be ok on VI.
 149 def int_amdgcn_fract : Intrinsic<
 150   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 151 >;
 152
 153 def int_amdgcn_class : Intrinsic<
 154   [llvm_i1_ty], [llvm_anyfloat_ty, llvm_i32_ty], [IntrNoMem]
 155 >;
 156
 157 def int_amdgcn_cubeid : GCCBuiltin<"__builtin_amdgcn_cubeid">,
 158   Intrinsic<[llvm_float_ty],
 159     [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem]
 160 >;
 161
 162 def int_amdgcn_cubema : GCCBuiltin<"__builtin_amdgcn_cubema">,
 163   Intrinsic<[llvm_float_ty],
 164   [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem]
 165 >;
 166
 167 def int_amdgcn_cubesc : GCCBuiltin<"__builtin_amdgcn_cubesc">,
 168   Intrinsic<[llvm_float_ty],
 169     [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem]
 170 >;
 171
 172 def int_amdgcn_cubetc : GCCBuiltin<"__builtin_amdgcn_cubetc">,
 173   Intrinsic<[llvm_float_ty],
 174     [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem]
 175 >;
 176
 177 // TODO: Do we want an ordering for these?
 178 def int_amdgcn_atomic_inc : Intrinsic<[llvm_anyint_ty],
 179   [llvm_anyptr_ty, LLVMMatchType<0>],
 180   [IntrArgMemOnly, NoCapture<0>]
 181 >;
 182
 183 def int_amdgcn_atomic_dec : Intrinsic<[llvm_anyint_ty],
 184   [llvm_anyptr_ty, LLVMMatchType<0>],
 185   [IntrArgMemOnly, NoCapture<0>]
 186 >;
 187
 188 class AMDGPUImageLoad : Intrinsic <
 189   [llvm_v4f32_ty],    // vdata(VGPR)
 190   [llvm_anyint_ty,    // vaddr(VGPR)
 191    llvm_v8i32_ty,     // rsrc(SGPR)
 192    llvm_i32_ty,       // dmask(imm)
 193    llvm_i1_ty,        // r128(imm)
 194    llvm_i1_ty,        // da(imm)
 195    llvm_i1_ty,        // glc(imm)
 196    llvm_i1_ty],       // slc(imm)
 197   [IntrReadMem]>;
 198
 199 def int_amdgcn_image_load : AMDGPUImageLoad;
 200 def int_amdgcn_image_load_mip : AMDGPUImageLoad;
 201
 202 class AMDGPUImageStore : Intrinsic <
 203   [],
 204   [llvm_v4f32_ty,     // vdata(VGPR)
 205    llvm_anyint_ty,    // vaddr(VGPR)
 206    llvm_v8i32_ty,     // rsrc(SGPR)
 207    llvm_i32_ty,       // dmask(imm)
 208    llvm_i1_ty,        // r128(imm)
 209    llvm_i1_ty,        // da(imm)
 210    llvm_i1_ty,        // glc(imm)
 211    llvm_i1_ty],       // slc(imm)
 212   []>;
 213
 214 def int_amdgcn_image_store : AMDGPUImageStore;
 215 def int_amdgcn_image_store_mip : AMDGPUImageStore;
 216
 217 class AMDGPUImageAtomic : Intrinsic <
 218   [llvm_i32_ty],
 219   [llvm_i32_ty,       // vdata(VGPR)
 220    llvm_anyint_ty,    // vaddr(VGPR)
 221    llvm_v8i32_ty,     // rsrc(SGPR)
 222    llvm_i1_ty,        // r128(imm)
 223    llvm_i1_ty,        // da(imm)
 224    llvm_i1_ty],       // slc(imm)
 225   []>;
 226
 227 def int_amdgcn_image_atomic_swap : AMDGPUImageAtomic;
 228 def int_amdgcn_image_atomic_add : AMDGPUImageAtomic;
 229 def int_amdgcn_image_atomic_sub : AMDGPUImageAtomic;
 230 def int_amdgcn_image_atomic_smin : AMDGPUImageAtomic;
 231 def int_amdgcn_image_atomic_umin : AMDGPUImageAtomic;
 232 def int_amdgcn_image_atomic_smax : AMDGPUImageAtomic;
 233 def int_amdgcn_image_atomic_umax : AMDGPUImageAtomic;
 234 def int_amdgcn_image_atomic_and : AMDGPUImageAtomic;
 235 def int_amdgcn_image_atomic_or : AMDGPUImageAtomic;
 236 def int_amdgcn_image_atomic_xor : AMDGPUImageAtomic;
 237 def int_amdgcn_image_atomic_inc : AMDGPUImageAtomic;
 238 def int_amdgcn_image_atomic_dec : AMDGPUImageAtomic;
 239 def int_amdgcn_image_atomic_cmpswap : Intrinsic <
 240   [llvm_i32_ty],
 241   [llvm_i32_ty,       // src(VGPR)
 242    llvm_i32_ty,       // cmp(VGPR)
 243    llvm_anyint_ty,    // vaddr(VGPR)
 244    llvm_v8i32_ty,     // rsrc(SGPR)
 245    llvm_i1_ty,        // r128(imm)
 246    llvm_i1_ty,        // da(imm)
 247    llvm_i1_ty],       // slc(imm)
 248   []>;
 249
 250 class AMDGPUBufferLoad : Intrinsic <
 251   [llvm_anyfloat_ty],
 252   [llvm_v4i32_ty,     // rsrc(SGPR)
 253    llvm_i32_ty,       // vindex(VGPR)
 254    llvm_i32_ty,       // offset(SGPR/VGPR/imm)
 255    llvm_i1_ty,        // glc(imm)
 256    llvm_i1_ty],       // slc(imm)
 257   [IntrReadMem]>;
 258 def int_amdgcn_buffer_load_format : AMDGPUBufferLoad;
 259 def int_amdgcn_buffer_load : AMDGPUBufferLoad;
 260
 261 class AMDGPUBufferStore : Intrinsic <
 262   [],
 263   [llvm_anyfloat_ty,  // vdata(VGPR) -- can currently only select f32, v2f32, v4f32
 264    llvm_v4i32_ty,     // rsrc(SGPR)
 265    llvm_i32_ty,       // vindex(VGPR)
 266    llvm_i32_ty,       // offset(SGPR/VGPR/imm)
 267    llvm_i1_ty,        // glc(imm)
 268    llvm_i1_ty],       // slc(imm)
 269   [IntrWriteMem]>;
 270 def int_amdgcn_buffer_store_format : AMDGPUBufferStore;
 271 def int_amdgcn_buffer_store : AMDGPUBufferStore;
 272
 273 class AMDGPUBufferAtomic : Intrinsic <
 274   [llvm_i32_ty],
 275   [llvm_i32_ty,       // vdata(VGPR)
 276    llvm_v4i32_ty,     // rsrc(SGPR)
 277    llvm_i32_ty,       // vindex(VGPR)
 278    llvm_i32_ty,       // offset(SGPR/VGPR/imm)
 279    llvm_i1_ty],       // slc(imm)
 280   []>;
 281 def int_amdgcn_buffer_atomic_swap : AMDGPUBufferAtomic;
 282 def int_amdgcn_buffer_atomic_add : AMDGPUBufferAtomic;
 283 def int_amdgcn_buffer_atomic_sub : AMDGPUBufferAtomic;
 284 def int_amdgcn_buffer_atomic_smin : AMDGPUBufferAtomic;
 285 def int_amdgcn_buffer_atomic_umin : AMDGPUBufferAtomic;
 286 def int_amdgcn_buffer_atomic_smax : AMDGPUBufferAtomic;
 287 def int_amdgcn_buffer_atomic_umax : AMDGPUBufferAtomic;
 288 def int_amdgcn_buffer_atomic_and : AMDGPUBufferAtomic;
 289 def int_amdgcn_buffer_atomic_or : AMDGPUBufferAtomic;
 290 def int_amdgcn_buffer_atomic_xor : AMDGPUBufferAtomic;
 291 def int_amdgcn_buffer_atomic_cmpswap : Intrinsic<
 292   [llvm_i32_ty],
 293   [llvm_i32_ty,       // src(VGPR)
 294    llvm_i32_ty,       // cmp(VGPR)
 295    llvm_v4i32_ty,     // rsrc(SGPR)
 296    llvm_i32_ty,       // vindex(VGPR)
 297    llvm_i32_ty,       // offset(SGPR/VGPR/imm)
 298    llvm_i1_ty],       // slc(imm)
 299   []>;
 300
 301 def int_amdgcn_read_workdim : AMDGPUReadPreloadRegisterIntrinsic;
 302
 303
 304 def int_amdgcn_buffer_wbinvl1_sc :
 305   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_sc">,
 306   Intrinsic<[], [], []>;
 307
 308 def int_amdgcn_buffer_wbinvl1 :
 309   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1">,
 310   Intrinsic<[], [], []>;
 311
 312 def int_amdgcn_s_dcache_inv :
 313   GCCBuiltin<"__builtin_amdgcn_s_dcache_inv">,
 314   Intrinsic<[], [], []>;
 315
 316 def int_amdgcn_s_memtime :
 317   GCCBuiltin<"__builtin_amdgcn_s_memtime">,
 318   Intrinsic<[llvm_i64_ty], [], []>;
 319
 320 def int_amdgcn_s_sleep :
 321   GCCBuiltin<"__builtin_amdgcn_s_sleep">,
 322   Intrinsic<[], [llvm_i32_ty], []> {
 323 }
 324
 325 def int_amdgcn_s_getreg :
 326   GCCBuiltin<"__builtin_amdgcn_s_getreg">,
 327   Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrReadMem]>;
 328
 329 def int_amdgcn_groupstaticsize :
 330   GCCBuiltin<"__builtin_amdgcn_groupstaticsize">,
 331   Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>;
 332
 333 def int_amdgcn_dispatch_ptr :
 334   GCCBuiltin<"__builtin_amdgcn_dispatch_ptr">,
 335   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>;
 336
 337 def int_amdgcn_queue_ptr :
 338   GCCBuiltin<"__builtin_amdgcn_queue_ptr">,
 339   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>;
 340
 341 def int_amdgcn_kernarg_segment_ptr :
 342   GCCBuiltin<"__builtin_amdgcn_kernarg_segment_ptr">,
 343   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>;
 344
 345 def int_amdgcn_implicitarg_ptr :
 346   GCCBuiltin<"__builtin_amdgcn_implicitarg_ptr">,
 347   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>;
 348
 349 // __builtin_amdgcn_interp_p1 <i>, <attr_chan>, <attr>, <m0>
 350 def int_amdgcn_interp_p1 :
 351   GCCBuiltin<"__builtin_amdgcn_interp_p1">,
 352   Intrinsic<[llvm_float_ty],
 353             [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
 354             [IntrNoMem]>;  // This intrinsic reads from lds, but the memory
 355                            // values are constant, so it behaves like IntrNoMem.
 356
 357 // __builtin_amdgcn_interp_p2 <p1>, <j>, <attr_chan>, <attr>, <m0>
 358 def int_amdgcn_interp_p2 :
 359   GCCBuiltin<"__builtin_amdgcn_interp_p2">,
 360   Intrinsic<[llvm_float_ty],
 361             [llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
 362             [IntrNoMem]>;  // See int_amdgcn_v_interp_p1 for why this is
 363                            // IntrNoMem.
 364
 365 // Pixel shaders only: whether the current pixel is live (i.e. not a helper
 366 // invocation for derivative computation).
 367 def int_amdgcn_ps_live : Intrinsic <
 368   [llvm_i1_ty],
 369   [],
 370   [IntrNoMem]>;
 371
 372 def int_amdgcn_mbcnt_lo :
 373   GCCBuiltin<"__builtin_amdgcn_mbcnt_lo">,
 374   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 375
 376 def int_amdgcn_mbcnt_hi :
 377   GCCBuiltin<"__builtin_amdgcn_mbcnt_hi">,
 378   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 379
 380 // llvm.amdgcn.ds.swizzle src offset
 381 def int_amdgcn_ds_swizzle :
 382   GCCBuiltin<"__builtin_amdgcn_ds_swizzle">,
 383   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
 384
 385 // llvm.amdgcn.lerp
 386 def int_amdgcn_lerp :
 387   GCCBuiltin<"__builtin_amdgcn_lerp">,
 388   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 389
 390 //===----------------------------------------------------------------------===//
 391 // CI+ Intrinsics
 392 //===----------------------------------------------------------------------===//
 393
 394 def int_amdgcn_s_dcache_inv_vol :
 395   GCCBuiltin<"__builtin_amdgcn_s_dcache_inv_vol">,
 396   Intrinsic<[], [], []>;
 397
 398 def int_amdgcn_buffer_wbinvl1_vol :
 399   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_vol">,
 400   Intrinsic<[], [], []>;
 401
 402 //===----------------------------------------------------------------------===//
 403 // VI Intrinsics
 404 //===----------------------------------------------------------------------===//
 405
 406 // llvm.amdgcn.mov.dpp.i32 <src> <dpp_ctrl> <row_mask> <bank_mask> <bound_ctrl>
 407 def int_amdgcn_mov_dpp :
 408   Intrinsic<[llvm_anyint_ty],
 409             [LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
 410              llvm_i1_ty], [IntrNoMem, IntrConvergent]>;
 411
 412 def int_amdgcn_s_dcache_wb :
 413   GCCBuiltin<"__builtin_amdgcn_s_dcache_wb">,
 414   Intrinsic<[], [], []>;
 415
 416 def int_amdgcn_s_dcache_wb_vol :
 417   GCCBuiltin<"__builtin_amdgcn_s_dcache_wb_vol">,
 418   Intrinsic<[], [], []>;
 419
 420 def int_amdgcn_s_memrealtime :
 421   GCCBuiltin<"__builtin_amdgcn_s_memrealtime">,
 422   Intrinsic<[llvm_i64_ty], [], []>;
 423
 424 // llvm.amdgcn.ds.permute <index> <src>
 425 def int_amdgcn_ds_permute :
 426   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
 427
 428 // llvm.amdgcn.ds.bpermute <index> <src>
 429 def int_amdgcn_ds_bpermute :
 430   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
 431
 432 }