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_group_barrier : GCCBuiltin<"__builtin_r600_group_barrier">,
  45   Intrinsic<[], [], [IntrConvergent]>;
  46
  47 // AS 7 is PARAM_I_ADDRESS, used for kernel arguments
  48 def int_r600_implicitarg_ptr :
  49   GCCBuiltin<"__builtin_r600_implicitarg_ptr">,
  50   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 7>], [], [IntrNoMem]>;
  51
  52 def int_r600_rat_store_typed :
  53   // 1st parameter: Data
  54   // 2nd parameter: Index
  55   // 3rd parameter: Constant RAT ID
  56   Intrinsic<[], [llvm_v4i32_ty, llvm_v4i32_ty, llvm_i32_ty], []>,
  57   GCCBuiltin<"__builtin_r600_rat_store_typed">;
  58
  59 def int_r600_recipsqrt_ieee :  Intrinsic<
  60   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
  61 >;
  62
  63 def int_r600_recipsqrt_clamped : Intrinsic<
  64   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
  65 >;
  66
  67 } // End TargetPrefix = "r600"
  68
  69 let TargetPrefix = "amdgcn" in {
  70
  71 //===----------------------------------------------------------------------===//
  72 // ABI Special Intrinsics
  73 //===----------------------------------------------------------------------===//
  74
  75 defm int_amdgcn_workitem_id : AMDGPUReadPreloadRegisterIntrinsic_xyz;
  76 defm int_amdgcn_workgroup_id : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
  77                                <"__builtin_amdgcn_workgroup_id">;
  78
  79 def int_amdgcn_dispatch_ptr :
  80   GCCBuiltin<"__builtin_amdgcn_dispatch_ptr">,
  81   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>;
  82
  83 def int_amdgcn_queue_ptr :
  84   GCCBuiltin<"__builtin_amdgcn_queue_ptr">,
  85   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>;
  86
  87 def int_amdgcn_kernarg_segment_ptr :
  88   GCCBuiltin<"__builtin_amdgcn_kernarg_segment_ptr">,
  89   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>;
  90
  91 def int_amdgcn_implicitarg_ptr :
  92   GCCBuiltin<"__builtin_amdgcn_implicitarg_ptr">,
  93   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>;
  94
  95 def int_amdgcn_groupstaticsize :
  96   GCCBuiltin<"__builtin_amdgcn_groupstaticsize">,
  97   Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>;
  98
  99 def int_amdgcn_dispatch_id :
 100   GCCBuiltin<"__builtin_amdgcn_dispatch_id">,
 101   Intrinsic<[llvm_i64_ty], [], [IntrNoMem]>;
 102
 103 //===----------------------------------------------------------------------===//
 104 // Instruction Intrinsics
 105 //===----------------------------------------------------------------------===//
 106
 107 def int_amdgcn_s_barrier : GCCBuiltin<"__builtin_amdgcn_s_barrier">,
 108   Intrinsic<[], [], [IntrConvergent]>;
 109
 110 def int_amdgcn_wave_barrier : GCCBuiltin<"__builtin_amdgcn_wave_barrier">,
 111   Intrinsic<[], [], [IntrConvergent]>;
 112
 113 def int_amdgcn_s_waitcnt : Intrinsic<[], [llvm_i32_ty], []>;
 114
 115 def int_amdgcn_div_scale : Intrinsic<
 116   // 1st parameter: Numerator
 117   // 2nd parameter: Denominator
 118   // 3rd parameter: Constant to select select between first and
 119   //                second. (0 = first, 1 = second).
 120   [llvm_anyfloat_ty, llvm_i1_ty],
 121   [LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty],
 122   [IntrNoMem]
 123 >;
 124
 125 def int_amdgcn_div_fmas : Intrinsic<[llvm_anyfloat_ty],
 126   [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty],
 127   [IntrNoMem]
 128 >;
 129
 130 def int_amdgcn_div_fixup : Intrinsic<[llvm_anyfloat_ty],
 131   [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>],
 132   [IntrNoMem]
 133 >;
 134
 135 def int_amdgcn_trig_preop : Intrinsic<
 136   [llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty], [IntrNoMem]
 137 >;
 138
 139 def int_amdgcn_sin : Intrinsic<
 140   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 141 >;
 142
 143 def int_amdgcn_cos : Intrinsic<
 144   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 145 >;
 146
 147 def int_amdgcn_log_clamp : Intrinsic<
 148   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 149 >;
 150
 151 def int_amdgcn_fmul_legacy : GCCBuiltin<"__builtin_amdgcn_fmul_legacy">,
 152   Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem]
 153 >;
 154
 155 def int_amdgcn_rcp : Intrinsic<
 156   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 157 >;
 158
 159 def int_amdgcn_rcp_legacy : GCCBuiltin<"__builtin_amdgcn_rcp_legacy">,
 160   Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]
 161 >;
 162
 163 def int_amdgcn_rsq :  Intrinsic<
 164   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 165 >;
 166
 167 def int_amdgcn_rsq_legacy :  GCCBuiltin<"__builtin_amdgcn_rsq_legacy">,
 168   Intrinsic<
 169   [llvm_float_ty], [llvm_float_ty], [IntrNoMem]
 170 >;
 171
 172 def int_amdgcn_rsq_clamp : Intrinsic<
 173   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]>;
 174
 175 def int_amdgcn_ldexp : Intrinsic<
 176   [llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty], [IntrNoMem]
 177 >;
 178
 179 def int_amdgcn_frexp_mant : Intrinsic<
 180   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 181 >;
 182
 183 def int_amdgcn_frexp_exp : Intrinsic<
 184   [llvm_anyint_ty], [llvm_anyfloat_ty], [IntrNoMem]
 185 >;
 186
 187 // v_fract is buggy on SI/CI. It mishandles infinities, may return 1.0
 188 // and always uses rtz, so is not suitable for implementing the OpenCL
 189 // fract function. It should be ok on VI.
 190 def int_amdgcn_fract : Intrinsic<
 191   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 192 >;
 193
 194 def int_amdgcn_class : Intrinsic<
 195   [llvm_i1_ty], [llvm_anyfloat_ty, llvm_i32_ty], [IntrNoMem]
 196 >;
 197
 198 def int_amdgcn_cubeid : GCCBuiltin<"__builtin_amdgcn_cubeid">,
 199   Intrinsic<[llvm_float_ty],
 200     [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem]
 201 >;
 202
 203 def int_amdgcn_cubema : GCCBuiltin<"__builtin_amdgcn_cubema">,
 204   Intrinsic<[llvm_float_ty],
 205   [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem]
 206 >;
 207
 208 def int_amdgcn_cubesc : GCCBuiltin<"__builtin_amdgcn_cubesc">,
 209   Intrinsic<[llvm_float_ty],
 210     [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem]
 211 >;
 212
 213 def int_amdgcn_cubetc : GCCBuiltin<"__builtin_amdgcn_cubetc">,
 214   Intrinsic<[llvm_float_ty],
 215     [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem]
 216 >;
 217
 218 // v_ffbh_i32, as opposed to v_ffbh_u32. For v_ffbh_u32, llvm.ctlz
 219 // should be used.
 220 def int_amdgcn_sffbh :
 221   Intrinsic<[llvm_anyint_ty], [LLVMMatchType<0>], [IntrNoMem]>;
 222
 223 // TODO: Do we want an ordering for these?
 224 def int_amdgcn_atomic_inc : Intrinsic<[llvm_anyint_ty],
 225   [llvm_anyptr_ty, LLVMMatchType<0>],
 226   [IntrArgMemOnly, NoCapture<0>]
 227 >;
 228
 229 def int_amdgcn_atomic_dec : Intrinsic<[llvm_anyint_ty],
 230   [llvm_anyptr_ty, LLVMMatchType<0>],
 231   [IntrArgMemOnly, NoCapture<0>]
 232 >;
 233
 234 class AMDGPUImageLoad : Intrinsic <
 235   [llvm_anyfloat_ty], // vdata(VGPR)
 236   [llvm_anyint_ty,    // vaddr(VGPR)
 237    llvm_anyint_ty,    // rsrc(SGPR)
 238    llvm_i32_ty,       // dmask(imm)
 239    llvm_i1_ty,        // glc(imm)
 240    llvm_i1_ty,        // slc(imm)
 241    llvm_i1_ty,        // lwe(imm)
 242    llvm_i1_ty],       // da(imm)
 243   [IntrReadMem]>;
 244
 245 def int_amdgcn_image_load : AMDGPUImageLoad;
 246 def int_amdgcn_image_load_mip : AMDGPUImageLoad;
 247 def int_amdgcn_image_getresinfo : AMDGPUImageLoad;
 248
 249 class AMDGPUImageStore : Intrinsic <
 250   [],
 251   [llvm_anyfloat_ty,  // vdata(VGPR)
 252    llvm_anyint_ty,    // vaddr(VGPR)
 253    llvm_anyint_ty,    // rsrc(SGPR)
 254    llvm_i32_ty,       // dmask(imm)
 255    llvm_i1_ty,        // glc(imm)
 256    llvm_i1_ty,        // slc(imm)
 257    llvm_i1_ty,        // lwe(imm)
 258    llvm_i1_ty],       // da(imm)
 259   []>;
 260
 261 def int_amdgcn_image_store : AMDGPUImageStore;
 262 def int_amdgcn_image_store_mip : AMDGPUImageStore;
 263
 264 class AMDGPUImageSample : Intrinsic <
 265     [llvm_anyfloat_ty], // vdata(VGPR)
 266     [llvm_anyfloat_ty,  // vaddr(VGPR)
 267      llvm_anyint_ty,    // rsrc(SGPR)
 268      llvm_v4i32_ty,     // sampler(SGPR)
 269      llvm_i32_ty,       // dmask(imm)
 270      llvm_i1_ty,        // unorm(imm)
 271      llvm_i1_ty,        // glc(imm)
 272      llvm_i1_ty,        // slc(imm)
 273      llvm_i1_ty,        // lwe(imm)
 274      llvm_i1_ty],       // da(imm)
 275      [IntrReadMem]>;
 276
 277 // Basic sample
 278 def int_amdgcn_image_sample : AMDGPUImageSample;
 279 def int_amdgcn_image_sample_cl : AMDGPUImageSample;
 280 def int_amdgcn_image_sample_d : AMDGPUImageSample;
 281 def int_amdgcn_image_sample_d_cl : AMDGPUImageSample;
 282 def int_amdgcn_image_sample_l : AMDGPUImageSample;
 283 def int_amdgcn_image_sample_b : AMDGPUImageSample;
 284 def int_amdgcn_image_sample_b_cl : AMDGPUImageSample;
 285 def int_amdgcn_image_sample_lz : AMDGPUImageSample;
 286 def int_amdgcn_image_sample_cd : AMDGPUImageSample;
 287 def int_amdgcn_image_sample_cd_cl : AMDGPUImageSample;
 288
 289 // Sample with comparison
 290 def int_amdgcn_image_sample_c : AMDGPUImageSample;
 291 def int_amdgcn_image_sample_c_cl : AMDGPUImageSample;
 292 def int_amdgcn_image_sample_c_d : AMDGPUImageSample;
 293 def int_amdgcn_image_sample_c_d_cl : AMDGPUImageSample;
 294 def int_amdgcn_image_sample_c_l : AMDGPUImageSample;
 295 def int_amdgcn_image_sample_c_b : AMDGPUImageSample;
 296 def int_amdgcn_image_sample_c_b_cl : AMDGPUImageSample;
 297 def int_amdgcn_image_sample_c_lz : AMDGPUImageSample;
 298 def int_amdgcn_image_sample_c_cd : AMDGPUImageSample;
 299 def int_amdgcn_image_sample_c_cd_cl : AMDGPUImageSample;
 300
 301 // Sample with offsets
 302 def int_amdgcn_image_sample_o : AMDGPUImageSample;
 303 def int_amdgcn_image_sample_cl_o : AMDGPUImageSample;
 304 def int_amdgcn_image_sample_d_o : AMDGPUImageSample;
 305 def int_amdgcn_image_sample_d_cl_o : AMDGPUImageSample;
 306 def int_amdgcn_image_sample_l_o : AMDGPUImageSample;
 307 def int_amdgcn_image_sample_b_o : AMDGPUImageSample;
 308 def int_amdgcn_image_sample_b_cl_o : AMDGPUImageSample;
 309 def int_amdgcn_image_sample_lz_o : AMDGPUImageSample;
 310 def int_amdgcn_image_sample_cd_o : AMDGPUImageSample;
 311 def int_amdgcn_image_sample_cd_cl_o : AMDGPUImageSample;
 312
 313 // Sample with comparison and offsets
 314 def int_amdgcn_image_sample_c_o : AMDGPUImageSample;
 315 def int_amdgcn_image_sample_c_cl_o : AMDGPUImageSample;
 316 def int_amdgcn_image_sample_c_d_o : AMDGPUImageSample;
 317 def int_amdgcn_image_sample_c_d_cl_o : AMDGPUImageSample;
 318 def int_amdgcn_image_sample_c_l_o : AMDGPUImageSample;
 319 def int_amdgcn_image_sample_c_b_o : AMDGPUImageSample;
 320 def int_amdgcn_image_sample_c_b_cl_o : AMDGPUImageSample;
 321 def int_amdgcn_image_sample_c_lz_o : AMDGPUImageSample;
 322 def int_amdgcn_image_sample_c_cd_o : AMDGPUImageSample;
 323 def int_amdgcn_image_sample_c_cd_cl_o : AMDGPUImageSample;
 324
 325 // Basic gather4
 326 def int_amdgcn_image_gather4 : AMDGPUImageSample;
 327 def int_amdgcn_image_gather4_cl : AMDGPUImageSample;
 328 def int_amdgcn_image_gather4_l : AMDGPUImageSample;
 329 def int_amdgcn_image_gather4_b : AMDGPUImageSample;
 330 def int_amdgcn_image_gather4_b_cl : AMDGPUImageSample;
 331 def int_amdgcn_image_gather4_lz : AMDGPUImageSample;
 332
 333 // Gather4 with comparison
 334 def int_amdgcn_image_gather4_c : AMDGPUImageSample;
 335 def int_amdgcn_image_gather4_c_cl : AMDGPUImageSample;
 336 def int_amdgcn_image_gather4_c_l : AMDGPUImageSample;
 337 def int_amdgcn_image_gather4_c_b : AMDGPUImageSample;
 338 def int_amdgcn_image_gather4_c_b_cl : AMDGPUImageSample;
 339 def int_amdgcn_image_gather4_c_lz : AMDGPUImageSample;
 340
 341 // Gather4 with offsets
 342 def int_amdgcn_image_gather4_o : AMDGPUImageSample;
 343 def int_amdgcn_image_gather4_cl_o : AMDGPUImageSample;
 344 def int_amdgcn_image_gather4_l_o : AMDGPUImageSample;
 345 def int_amdgcn_image_gather4_b_o : AMDGPUImageSample;
 346 def int_amdgcn_image_gather4_b_cl_o : AMDGPUImageSample;
 347 def int_amdgcn_image_gather4_lz_o : AMDGPUImageSample;
 348
 349 // Gather4 with comparison and offsets
 350 def int_amdgcn_image_gather4_c_o : AMDGPUImageSample;
 351 def int_amdgcn_image_gather4_c_cl_o : AMDGPUImageSample;
 352 def int_amdgcn_image_gather4_c_l_o : AMDGPUImageSample;
 353 def int_amdgcn_image_gather4_c_b_o : AMDGPUImageSample;
 354 def int_amdgcn_image_gather4_c_b_cl_o : AMDGPUImageSample;
 355 def int_amdgcn_image_gather4_c_lz_o : AMDGPUImageSample;
 356
 357 def int_amdgcn_image_getlod : AMDGPUImageSample;
 358
 359 class AMDGPUImageAtomic : Intrinsic <
 360   [llvm_i32_ty],
 361   [llvm_i32_ty,       // vdata(VGPR)
 362    llvm_anyint_ty,    // vaddr(VGPR)
 363    llvm_v8i32_ty,     // rsrc(SGPR)
 364    llvm_i1_ty,        // r128(imm)
 365    llvm_i1_ty,        // da(imm)
 366    llvm_i1_ty],       // slc(imm)
 367   []>;
 368
 369 def int_amdgcn_image_atomic_swap : AMDGPUImageAtomic;
 370 def int_amdgcn_image_atomic_add : AMDGPUImageAtomic;
 371 def int_amdgcn_image_atomic_sub : AMDGPUImageAtomic;
 372 def int_amdgcn_image_atomic_smin : AMDGPUImageAtomic;
 373 def int_amdgcn_image_atomic_umin : AMDGPUImageAtomic;
 374 def int_amdgcn_image_atomic_smax : AMDGPUImageAtomic;
 375 def int_amdgcn_image_atomic_umax : AMDGPUImageAtomic;
 376 def int_amdgcn_image_atomic_and : AMDGPUImageAtomic;
 377 def int_amdgcn_image_atomic_or : AMDGPUImageAtomic;
 378 def int_amdgcn_image_atomic_xor : AMDGPUImageAtomic;
 379 def int_amdgcn_image_atomic_inc : AMDGPUImageAtomic;
 380 def int_amdgcn_image_atomic_dec : AMDGPUImageAtomic;
 381 def int_amdgcn_image_atomic_cmpswap : Intrinsic <
 382   [llvm_i32_ty],
 383   [llvm_i32_ty,       // src(VGPR)
 384    llvm_i32_ty,       // cmp(VGPR)
 385    llvm_anyint_ty,    // vaddr(VGPR)
 386    llvm_v8i32_ty,     // rsrc(SGPR)
 387    llvm_i1_ty,        // r128(imm)
 388    llvm_i1_ty,        // da(imm)
 389    llvm_i1_ty],       // slc(imm)
 390   []>;
 391
 392 class AMDGPUBufferLoad : Intrinsic <
 393   [llvm_anyfloat_ty],
 394   [llvm_v4i32_ty,     // rsrc(SGPR)
 395    llvm_i32_ty,       // vindex(VGPR)
 396    llvm_i32_ty,       // offset(SGPR/VGPR/imm)
 397    llvm_i1_ty,        // glc(imm)
 398    llvm_i1_ty],       // slc(imm)
 399   [IntrReadMem]>;
 400 def int_amdgcn_buffer_load_format : AMDGPUBufferLoad;
 401 def int_amdgcn_buffer_load : AMDGPUBufferLoad;
 402
 403 class AMDGPUBufferStore : Intrinsic <
 404   [],
 405   [llvm_anyfloat_ty,  // vdata(VGPR) -- can currently only select f32, v2f32, v4f32
 406    llvm_v4i32_ty,     // rsrc(SGPR)
 407    llvm_i32_ty,       // vindex(VGPR)
 408    llvm_i32_ty,       // offset(SGPR/VGPR/imm)
 409    llvm_i1_ty,        // glc(imm)
 410    llvm_i1_ty],       // slc(imm)
 411   [IntrWriteMem]>;
 412 def int_amdgcn_buffer_store_format : AMDGPUBufferStore;
 413 def int_amdgcn_buffer_store : AMDGPUBufferStore;
 414
 415 class AMDGPUBufferAtomic : Intrinsic <
 416   [llvm_i32_ty],
 417   [llvm_i32_ty,       // vdata(VGPR)
 418    llvm_v4i32_ty,     // rsrc(SGPR)
 419    llvm_i32_ty,       // vindex(VGPR)
 420    llvm_i32_ty,       // offset(SGPR/VGPR/imm)
 421    llvm_i1_ty],       // slc(imm)
 422   []>;
 423 def int_amdgcn_buffer_atomic_swap : AMDGPUBufferAtomic;
 424 def int_amdgcn_buffer_atomic_add : AMDGPUBufferAtomic;
 425 def int_amdgcn_buffer_atomic_sub : AMDGPUBufferAtomic;
 426 def int_amdgcn_buffer_atomic_smin : AMDGPUBufferAtomic;
 427 def int_amdgcn_buffer_atomic_umin : AMDGPUBufferAtomic;
 428 def int_amdgcn_buffer_atomic_smax : AMDGPUBufferAtomic;
 429 def int_amdgcn_buffer_atomic_umax : AMDGPUBufferAtomic;
 430 def int_amdgcn_buffer_atomic_and : AMDGPUBufferAtomic;
 431 def int_amdgcn_buffer_atomic_or : AMDGPUBufferAtomic;
 432 def int_amdgcn_buffer_atomic_xor : AMDGPUBufferAtomic;
 433 def int_amdgcn_buffer_atomic_cmpswap : Intrinsic<
 434   [llvm_i32_ty],
 435   [llvm_i32_ty,       // src(VGPR)
 436    llvm_i32_ty,       // cmp(VGPR)
 437    llvm_v4i32_ty,     // rsrc(SGPR)
 438    llvm_i32_ty,       // vindex(VGPR)
 439    llvm_i32_ty,       // offset(SGPR/VGPR/imm)
 440    llvm_i1_ty],       // slc(imm)
 441   []>;
 442
 443 def int_amdgcn_buffer_wbinvl1_sc :
 444   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_sc">,
 445   Intrinsic<[], [], []>;
 446
 447 def int_amdgcn_buffer_wbinvl1 :
 448   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1">,
 449   Intrinsic<[], [], []>;
 450
 451 def int_amdgcn_s_dcache_inv :
 452   GCCBuiltin<"__builtin_amdgcn_s_dcache_inv">,
 453   Intrinsic<[], [], []>;
 454
 455 def int_amdgcn_s_memtime :
 456   GCCBuiltin<"__builtin_amdgcn_s_memtime">,
 457   Intrinsic<[llvm_i64_ty], [], []>;
 458
 459 def int_amdgcn_s_sleep :
 460   GCCBuiltin<"__builtin_amdgcn_s_sleep">,
 461   Intrinsic<[], [llvm_i32_ty], []> {
 462 }
 463
 464 def int_amdgcn_s_incperflevel :
 465   GCCBuiltin<"__builtin_amdgcn_s_incperflevel">,
 466   Intrinsic<[], [llvm_i32_ty], []> {
 467 }
 468
 469 def int_amdgcn_s_decperflevel :
 470   GCCBuiltin<"__builtin_amdgcn_s_decperflevel">,
 471   Intrinsic<[], [llvm_i32_ty], []> {
 472 }
 473
 474 def int_amdgcn_s_getreg :
 475   GCCBuiltin<"__builtin_amdgcn_s_getreg">,
 476   Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrReadMem]>;
 477
 478 // __builtin_amdgcn_interp_mov <param>, <attr_chan>, <attr>, <m0>
 479 // param values: 0 = P10, 1 = P20, 2 = P0
 480 def int_amdgcn_interp_mov :
 481   GCCBuiltin<"__builtin_amdgcn_interp_mov">,
 482   Intrinsic<[llvm_float_ty],
 483             [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
 484             [IntrNoMem]>;
 485
 486 // __builtin_amdgcn_interp_p1 <i>, <attr_chan>, <attr>, <m0>
 487 def int_amdgcn_interp_p1 :
 488   GCCBuiltin<"__builtin_amdgcn_interp_p1">,
 489   Intrinsic<[llvm_float_ty],
 490             [llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
 491             [IntrNoMem]>;  // This intrinsic reads from lds, but the memory
 492                            // values are constant, so it behaves like IntrNoMem.
 493
 494 // __builtin_amdgcn_interp_p2 <p1>, <j>, <attr_chan>, <attr>, <m0>
 495 def int_amdgcn_interp_p2 :
 496   GCCBuiltin<"__builtin_amdgcn_interp_p2">,
 497   Intrinsic<[llvm_float_ty],
 498             [llvm_float_ty, llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
 499             [IntrNoMem]>;  // See int_amdgcn_v_interp_p1 for why this is
 500                            // IntrNoMem.
 501
 502 // Pixel shaders only: whether the current pixel is live (i.e. not a helper
 503 // invocation for derivative computation).
 504 def int_amdgcn_ps_live : Intrinsic <
 505   [llvm_i1_ty],
 506   [],
 507   [IntrNoMem]>;
 508
 509 def int_amdgcn_mbcnt_lo :
 510   GCCBuiltin<"__builtin_amdgcn_mbcnt_lo">,
 511   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 512
 513 def int_amdgcn_mbcnt_hi :
 514   GCCBuiltin<"__builtin_amdgcn_mbcnt_hi">,
 515   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 516
 517 // llvm.amdgcn.ds.swizzle src offset
 518 def int_amdgcn_ds_swizzle :
 519   GCCBuiltin<"__builtin_amdgcn_ds_swizzle">,
 520   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
 521
 522 // llvm.amdgcn.lerp
 523 def int_amdgcn_lerp :
 524   GCCBuiltin<"__builtin_amdgcn_lerp">,
 525   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 526
 527 def int_amdgcn_sad_u8 :
 528   GCCBuiltin<"__builtin_amdgcn_sad_u8">,
 529   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 530
 531 def int_amdgcn_msad_u8 :
 532   GCCBuiltin<"__builtin_amdgcn_msad_u8">,
 533   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 534
 535 def int_amdgcn_sad_hi_u8 :
 536   GCCBuiltin<"__builtin_amdgcn_sad_hi_u8">,
 537   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 538
 539 def int_amdgcn_sad_u16 :
 540   GCCBuiltin<"__builtin_amdgcn_sad_u16">,
 541   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 542
 543 def int_amdgcn_qsad_pk_u16_u8 :
 544   GCCBuiltin<"__builtin_amdgcn_qsad_pk_u16_u8">,
 545   Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty], [IntrNoMem]>;
 546
 547 def int_amdgcn_mqsad_pk_u16_u8 :
 548   GCCBuiltin<"__builtin_amdgcn_mqsad_pk_u16_u8">,
 549   Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty], [IntrNoMem]>;
 550
 551 def int_amdgcn_mqsad_u32_u8 :
 552   GCCBuiltin<"__builtin_amdgcn_mqsad_u32_u8">,
 553   Intrinsic<[llvm_v4i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_v4i32_ty], [IntrNoMem]>;
 554
 555 def int_amdgcn_cvt_pk_u8_f32 :
 556   GCCBuiltin<"__builtin_amdgcn_cvt_pk_u8_f32">,
 557   Intrinsic<[llvm_i32_ty], [llvm_float_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 558
 559 def int_amdgcn_icmp :
 560   Intrinsic<[llvm_i64_ty], [llvm_anyint_ty, LLVMMatchType<0>, llvm_i32_ty],
 561             [IntrNoMem, IntrConvergent]>;
 562
 563 def int_amdgcn_fcmp :
 564   Intrinsic<[llvm_i64_ty], [llvm_anyfloat_ty, LLVMMatchType<0>, llvm_i32_ty],
 565             [IntrNoMem, IntrConvergent]>;
 566
 567 def int_amdgcn_readfirstlane :
 568   GCCBuiltin<"__builtin_amdgcn_readfirstlane">,
 569   Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
 570
 571 def int_amdgcn_readlane :
 572   GCCBuiltin<"__builtin_amdgcn_readlane">,
 573   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
 574
 575 //===----------------------------------------------------------------------===//
 576 // CI+ Intrinsics
 577 //===----------------------------------------------------------------------===//
 578
 579 def int_amdgcn_s_dcache_inv_vol :
 580   GCCBuiltin<"__builtin_amdgcn_s_dcache_inv_vol">,
 581   Intrinsic<[], [], []>;
 582
 583 def int_amdgcn_buffer_wbinvl1_vol :
 584   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_vol">,
 585   Intrinsic<[], [], []>;
 586
 587 //===----------------------------------------------------------------------===//
 588 // VI Intrinsics
 589 //===----------------------------------------------------------------------===//
 590
 591 // llvm.amdgcn.mov.dpp.i32 <src> <dpp_ctrl> <row_mask> <bank_mask> <bound_ctrl>
 592 def int_amdgcn_mov_dpp :
 593   Intrinsic<[llvm_anyint_ty],
 594             [LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
 595              llvm_i1_ty], [IntrNoMem, IntrConvergent]>;
 596
 597 def int_amdgcn_s_dcache_wb :
 598   GCCBuiltin<"__builtin_amdgcn_s_dcache_wb">,
 599   Intrinsic<[], [], []>;
 600
 601 def int_amdgcn_s_dcache_wb_vol :
 602   GCCBuiltin<"__builtin_amdgcn_s_dcache_wb_vol">,
 603   Intrinsic<[], [], []>;
 604
 605 def int_amdgcn_s_memrealtime :
 606   GCCBuiltin<"__builtin_amdgcn_s_memrealtime">,
 607   Intrinsic<[llvm_i64_ty], [], []>;
 608
 609 // llvm.amdgcn.ds.permute <index> <src>
 610 def int_amdgcn_ds_permute :
 611   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
 612
 613 // llvm.amdgcn.ds.bpermute <index> <src>
 614 def int_amdgcn_ds_bpermute :
 615   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
 616
 617 }