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
14let TargetPrefix = "r600" in {
15
16class R600ReadPreloadRegisterIntrinsic<string name>
17  : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>,
18    GCCBuiltin<name>;
19
20multiclass 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
26defm int_r600_read_global_size : R600ReadPreloadRegisterIntrinsic_xyz <
27                                       "__builtin_r600_read_global_size">;
28defm int_r600_read_local_size : R600ReadPreloadRegisterIntrinsic_xyz <
29                                       "__builtin_r600_read_local_size">;
30defm int_r600_read_ngroups : R600ReadPreloadRegisterIntrinsic_xyz <
31                                       "__builtin_r600_read_ngroups">;
32defm int_r600_read_tgid : R600ReadPreloadRegisterIntrinsic_xyz <
33                                       "__builtin_r600_read_tgid">;
34defm int_r600_read_tidig : R600ReadPreloadRegisterIntrinsic_xyz <
35                                       "__builtin_r600_read_tidig">;
36
37def 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
46let TargetPrefix = "AMDGPU" in {
47
48class AMDGPUReadPreloadRegisterIntrinsic<string name>
49  : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>,
50    GCCBuiltin<name>;
51
52def 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
61def 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
66def int_AMDGPU_div_fixup : GCCBuiltin<"__builtin_amdgpu_div_fixup">,
67  Intrinsic<[llvm_anyfloat_ty],
68            [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>],
69            [IntrNoMem]>;
70
71def int_AMDGPU_trig_preop : GCCBuiltin<"__builtin_amdgpu_trig_preop">,
72  Intrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty],
73            [IntrNoMem]>;
74
75def int_AMDGPU_rcp : GCCBuiltin<"__builtin_amdgpu_rcp">,
76  Intrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]>;
77
78def int_AMDGPU_rsq : GCCBuiltin<"__builtin_amdgpu_rsq">,
79  Intrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]>;
80
81def int_AMDGPU_rsq_clamped : GCCBuiltin<"__builtin_amdgpu_rsq_clamped">,
82  Intrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]>;
83
84def int_AMDGPU_ldexp : GCCBuiltin<"__builtin_amdgpu_ldexp">,
85  Intrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty], [IntrNoMem]>;
86
87def int_AMDGPU_class : GCCBuiltin<"__builtin_amdgpu_class">,
88  Intrinsic<[llvm_i1_ty], [llvm_anyfloat_ty, llvm_i32_ty], [IntrNoMem]>;
89
90def int_AMDGPU_read_workdim : AMDGPUReadPreloadRegisterIntrinsic <
91                                       "__builtin_amdgpu_read_workdim">;
92
93} // End TargetPrefix = "AMDGPU"
94
95let TargetPrefix = "amdgcn" in {
96
97// SI only
98def int_amdgcn_buffer_wbinvl1_sc :
99  GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_sc">,
100  Intrinsic<[], [], []>;
101
102// On CI+
103def int_amdgcn_buffer_wbinvl1_vol :
104  GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_vol">,
105  Intrinsic<[], [], []>;
106
107def int_amdgcn_buffer_wbinvl1 :
108  GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1">,
109  Intrinsic<[], [], []>;
110
111def int_amdgcn_s_dcache_inv :
112  GCCBuiltin<"__builtin_amdgcn_s_dcache_inv">,
113  Intrinsic<[], [], []>;
114
115// CI+
116def int_amdgcn_s_dcache_inv_vol :
117  GCCBuiltin<"__builtin_amdgcn_s_dcache_inv_vol">,
118  Intrinsic<[], [], []>;
119
120// VI
121def int_amdgcn_s_dcache_wb :
122  GCCBuiltin<"__builtin_amdgcn_s_dcache_wb">,
123  Intrinsic<[], [], []>;
124
125// VI
126def int_amdgcn_s_dcache_wb_vol :
127  GCCBuiltin<"__builtin_amdgcn_s_dcache_wb_vol">,
128  Intrinsic<[], [], []>;
129
130def 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>
135def 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>
143def 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
150def int_amdgcn_mbcnt_lo :
151  GCCBuiltin<"__builtin_amdgcn_mbcnt_lo">,
152  Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
153
154def int_amdgcn_mbcnt_hi :
155  GCCBuiltin<"__builtin_amdgcn_mbcnt_hi">,
156  Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
157}
158