RPCS3/llvm-mirror
1
0
Fork 0
mirror of https://github.com/RPCS3/llvm-mirror.git synced 2024-11-25 12:12:47 +01:00
Code Issues Projects Releases Wiki Activity

AMDGPU: Remove GCCBuiltin from intrinsics that need mangling

Browse Source 
If the intrinsic is overloaded and works on multiple types,
it cannot resolve to a single corresponding builtin and requires
handling in clang. This just causes crashes now.

llvm-svn: 258559
This commit is contained in:
Matt Arsenault 2016-01-22 21:30:46 +00:00
parent 3913a77bb9
commit 919da0fa79
1 changed files with 30 additions and 25 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

55
include/llvm/IR/IntrinsicsAMDGPU.td
View File

@ -71,43 +71,48 @@ let TargetPrefix = "amdgcn" in {
def int_amdgcn_s_barrier : GCCBuiltin<"__builtin_amdgcn_s_barrier">,
Intrinsic<[], [], [IntrConvergent]>;
def int_amdgcn_div_scale : GCCBuiltin<"__builtin_amdgcn_div_scale">,
def int_amdgcn_div_scale : Intrinsic<
// 1st parameter: Numerator
// 2nd parameter: Denominator
// 3rd parameter: Constant to select select between first and
// second. (0 = first, 1 = second).
Intrinsic<[llvm_anyfloat_ty, llvm_i1_ty],
[LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty],
[IntrNoMem]>;
[llvm_anyfloat_ty, llvm_i1_ty],
[LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty],
[IntrNoMem]
>;
def int_amdgcn_div_fmas : GCCBuiltin<"__builtin_amdgcn_div_fmas">,
Intrinsic<[llvm_anyfloat_ty],
[LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty],
[IntrNoMem]>;
def int_amdgcn_div_fmas : Intrinsic<[llvm_anyfloat_ty],
[LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty],
[IntrNoMem]
>;
def int_amdgcn_div_fixup : GCCBuiltin<"__builtin_amdgcn_div_fixup">,
Intrinsic<[llvm_anyfloat_ty],
[LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>],
[IntrNoMem]>;
def int_amdgcn_div_fixup : Intrinsic<[llvm_anyfloat_ty],
[LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>],
[IntrNoMem]
>;
def int_amdgcn_trig_preop : GCCBuiltin<"__builtin_amdgcn_trig_preop">,
Intrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty],
[IntrNoMem]>;
def int_amdgcn_trig_preop : Intrinsic<
[llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty], [IntrNoMem]
>;
def int_amdgcn_rcp : GCCBuiltin<"__builtin_amdgcn_rcp">,
Intrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]>;
def int_amdgcn_rcp : Intrinsic<
[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
>;
def int_amdgcn_rsq : GCCBuiltin<"__builtin_amdgcn_rsq">,
Intrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]>;
def int_amdgcn_rsq : Intrinsic<
[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
>;
def int_amdgcn_rsq_clamped : GCCBuiltin<"__builtin_amdgcn_rsq_clamped">,
Intrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]>;
def int_amdgcn_rsq_clamped : Intrinsic<
[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]>;
def int_amdgcn_ldexp : GCCBuiltin<"__builtin_amdgcn_ldexp">,
Intrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty], [IntrNoMem]>;
def int_amdgcn_ldexp : Intrinsic<
[llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty], [IntrNoMem]
>;
def int_amdgcn_class : GCCBuiltin<"__builtin_amdgcn_class">,
Intrinsic<[llvm_i1_ty], [llvm_anyfloat_ty, llvm_i32_ty], [IntrNoMem]>;
def int_amdgcn_class : Intrinsic<
[llvm_i1_ty], [llvm_anyfloat_ty, llvm_i32_ty], [IntrNoMem]
>;
def int_amdgcn_read_workdim : AMDGPUReadPreloadRegisterIntrinsic <
"__builtin_amdgcn_read_workdim">;