1//==- BuiltinsAMDGPU.def - AMDGPU Builtin function database ------*- C++ -*-==// 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 the AMDGPU-specific builtin function database. Users of 11// this file must define the BUILTIN macro to make use of this information. 12// 13//===----------------------------------------------------------------------===// 14 15// The format of this database matches clang/Basic/Builtins.def. 16 17#if defined(BUILTIN) && !defined(TARGET_BUILTIN) 18# define TARGET_BUILTIN(ID, TYPE, ATTRS, FEATURE) BUILTIN(ID, TYPE, ATTRS) 19#endif 20//===----------------------------------------------------------------------===// 21// SI+ only builtins. 22//===----------------------------------------------------------------------===// 23 24BUILTIN(__builtin_amdgcn_kernarg_segment_ptr, "Uc*2", "nc") 25BUILTIN(__builtin_amdgcn_implicitarg_ptr, "Uc*2", "nc") 26 27BUILTIN(__builtin_amdgcn_workgroup_id_x, "Ui", "nc") 28BUILTIN(__builtin_amdgcn_workgroup_id_y, "Ui", "nc") 29BUILTIN(__builtin_amdgcn_workgroup_id_z, "Ui", "nc") 30 31BUILTIN(__builtin_amdgcn_workitem_id_x, "Ui", "nc") 32BUILTIN(__builtin_amdgcn_workitem_id_y, "Ui", "nc") 33BUILTIN(__builtin_amdgcn_workitem_id_z, "Ui", "nc") 34 35//===----------------------------------------------------------------------===// 36// Instruction builtins. 37//===----------------------------------------------------------------------===// 38BUILTIN(__builtin_amdgcn_s_barrier, "v", "n") 39BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n") 40BUILTIN(__builtin_amdgcn_div_scalef, "fffbb*", "n") 41BUILTIN(__builtin_amdgcn_div_fmas, "ddddb", "nc") 42BUILTIN(__builtin_amdgcn_div_fmasf, "ffffb", "nc") 43BUILTIN(__builtin_amdgcn_div_fixup, "dddd", "nc") 44BUILTIN(__builtin_amdgcn_div_fixupf, "ffff", "nc") 45BUILTIN(__builtin_amdgcn_trig_preop, "ddi", "nc") 46BUILTIN(__builtin_amdgcn_trig_preopf, "ffi", "nc") 47BUILTIN(__builtin_amdgcn_rcp, "dd", "nc") 48BUILTIN(__builtin_amdgcn_rcpf, "ff", "nc") 49BUILTIN(__builtin_amdgcn_rsq, "dd", "nc") 50BUILTIN(__builtin_amdgcn_rsqf, "ff", "nc") 51BUILTIN(__builtin_amdgcn_rsq_clamp, "dd", "nc") 52BUILTIN(__builtin_amdgcn_rsq_clampf, "ff", "nc") 53BUILTIN(__builtin_amdgcn_sinf, "ff", "nc") 54BUILTIN(__builtin_amdgcn_cosf, "ff", "nc") 55BUILTIN(__builtin_amdgcn_log_clampf, "ff", "nc") 56BUILTIN(__builtin_amdgcn_ldexp, "ddi", "nc") 57BUILTIN(__builtin_amdgcn_ldexpf, "ffi", "nc") 58BUILTIN(__builtin_amdgcn_frexp_mant, "dd", "nc") 59BUILTIN(__builtin_amdgcn_frexp_mantf, "ff", "nc") 60BUILTIN(__builtin_amdgcn_frexp_exp, "id", "nc") 61BUILTIN(__builtin_amdgcn_frexp_expf, "if", "nc") 62BUILTIN(__builtin_amdgcn_fract, "dd", "nc") 63BUILTIN(__builtin_amdgcn_fractf, "ff", "nc") 64BUILTIN(__builtin_amdgcn_class, "bdi", "nc") 65BUILTIN(__builtin_amdgcn_classf, "bfi", "nc") 66BUILTIN(__builtin_amdgcn_cubeid, "ffff", "nc") 67BUILTIN(__builtin_amdgcn_cubesc, "ffff", "nc") 68BUILTIN(__builtin_amdgcn_cubetc, "ffff", "nc") 69BUILTIN(__builtin_amdgcn_cubema, "ffff", "nc") 70BUILTIN(__builtin_amdgcn_s_memtime, "LUi", "n") 71BUILTIN(__builtin_amdgcn_s_sleep, "vIi", "n") 72 73//===----------------------------------------------------------------------===// 74// VI+ only builtins. 75//===----------------------------------------------------------------------===// 76 77TARGET_BUILTIN(__builtin_amdgcn_s_memrealtime, "LUi", "n", "s-memrealtime") 78 79//===----------------------------------------------------------------------===// 80// Special builtins. 81//===----------------------------------------------------------------------===// 82BUILTIN(__builtin_amdgcn_read_exec, "LUi", "nc") 83 84//===----------------------------------------------------------------------===// 85// R600-NI only builtins. 86//===----------------------------------------------------------------------===// 87 88BUILTIN(__builtin_r600_implicitarg_ptr, "Uc*7", "nc") 89 90BUILTIN(__builtin_r600_read_tgid_x, "Ui", "nc") 91BUILTIN(__builtin_r600_read_tgid_y, "Ui", "nc") 92BUILTIN(__builtin_r600_read_tgid_z, "Ui", "nc") 93 94BUILTIN(__builtin_r600_read_tidig_x, "Ui", "nc") 95BUILTIN(__builtin_r600_read_tidig_y, "Ui", "nc") 96BUILTIN(__builtin_r600_read_tidig_z, "Ui", "nc") 97 98//===----------------------------------------------------------------------===// 99// Legacy names with amdgpu prefix 100//===----------------------------------------------------------------------===// 101 102BUILTIN(__builtin_amdgpu_rsq, "dd", "nc") 103BUILTIN(__builtin_amdgpu_rsqf, "ff", "nc") 104BUILTIN(__builtin_amdgpu_ldexp, "ddi", "nc") 105BUILTIN(__builtin_amdgpu_ldexpf, "ffi", "nc") 106 107#undef BUILTIN 108#undef TARGET_BUILTIN 109