[llvm] r258559 - AMDGPU: Remove GCCBuiltin from intrinsics that need mangling

[Matt Arsenault via llvm-commits]

Author: arsenm
Date: Fri Jan 22 15:30:46 2016
New Revision: 258559

URL: http://llvm.org/viewvc/llvm-project?rev=258559&view=rev
Log:
AMDGPU: Remove GCCBuiltin from intrinsics that need mangling

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.

Modified:
    llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td

Modified: llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td?rev=258559&r1=258558&r2=258559&view=diff
==============================================================================
--- llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td (original)
+++ llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td Fri Jan 22 15:30:46 2016
@@ -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]>;
-
-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_fixup : GCCBuiltin<"__builtin_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_rcp : GCCBuiltin<"__builtin_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_clamped : GCCBuiltin<"__builtin_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_class : GCCBuiltin<"__builtin_amdgcn_class">,
-  Intrinsic<[llvm_i1_ty], [llvm_anyfloat_ty, llvm_i32_ty], [IntrNoMem]>;
+  [llvm_anyfloat_ty, llvm_i1_ty],
+  [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 : Intrinsic<[llvm_anyfloat_ty],
+  [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>],
+  [IntrNoMem]
+>;
+
+def int_amdgcn_trig_preop : Intrinsic<
+  [llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty], [IntrNoMem]
+>;
+
+def int_amdgcn_rcp : Intrinsic<
+  [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
+>;
+
+def int_amdgcn_rsq :  Intrinsic<
+  [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
+>;
+
+def int_amdgcn_rsq_clamped :  Intrinsic<
+  [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]>;
+
+def int_amdgcn_ldexp : Intrinsic<
+  [llvm_anyfloat_ty], [LLVMMatchType<0>, 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">;