revert instrinsic change

Author: broxigarchen

clang/lib/CodeGen/CGBuiltin.cpp

@@ -20418,14 +20418,6 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
     Function *F = CGM.getIntrinsic(Intrinsic::fshr, Src0->getType());
     return Builder.CreateCall(F, { Src0, Src1, Src2 });
   }
-  case AMDGPU::BI__builtin_amdgcn_alignbyte: {
-    llvm::Value *Src0 = EmitScalarExpr(E->getArg(0));
-    llvm::Value *Src1 = EmitScalarExpr(E->getArg(1));
-    llvm::Value *Src2 = EmitScalarExpr(E->getArg(2));
-    llvm::Function *F =
-        CGM.getIntrinsic(Intrinsic::amdgcn_alignbyte, Src2->getType());
-    return Builder.CreateCall(F, {Src0, Src1, Src2});
-  }
   case AMDGPU::BI__builtin_amdgcn_fence: {
     ProcessOrderScopeAMDGCN(EmitScalarExpr(E->getArg(0)),
                             EmitScalarExpr(E->getArg(1)), AO, SSID);

clang/test/CodeGenOpenCL/builtins-amdgcn.cl

@@ -734,7 +734,7 @@ kernel void test_alignbit(global uint* out, uint src0, uint src1, uint src2) {
 }
 
 // CHECK-LABEL: @test_alignbyte(
-// CHECK: tail call{{.*}} i32 @llvm.amdgcn.alignbyte.i32(i32 %src0, i32 %src1, i32 %src2)
+// CHECK: tail call{{.*}} i32 @llvm.amdgcn.alignbyte(i32 %src0, i32 %src1, i32 %src2)
 kernel void test_alignbyte(global uint* out, uint src0, uint src1, uint src2) {
   *out = __builtin_amdgcn_alignbyte(src0, src1, src2);
 }

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

@@ -2352,8 +2352,8 @@ def int_amdgcn_writelane :
   [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree]
 >;
 
-def int_amdgcn_alignbyte : DefaultAttrsIntrinsic<[llvm_i32_ty],
-  [llvm_i32_ty, llvm_i32_ty, llvm_anyint_ty],
+def int_amdgcn_alignbyte : ClangBuiltin<"__builtin_amdgcn_alignbyte">,
+  DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
   [IntrNoMem, IntrSpeculatable]
 >;
 

llvm/unittests/IR/IntrinsicsTest.cpp

@@ -87,6 +87,7 @@ TEST(IntrinsicNameLookup, ClangBuiltinLookup) {
       {"__builtin_adjust_trampoline", "", adjust_trampoline},
       {"__builtin_trap", "", trap},
       {"__builtin_arm_chkfeat", "aarch64", aarch64_chkfeat},
+      {"__builtin_amdgcn_alignbyte", "amdgcn", amdgcn_alignbyte},
       {"__builtin_amdgcn_workgroup_id_z", "amdgcn", amdgcn_workgroup_id_z},
       {"__builtin_arm_cdp", "arm", arm_cdp},
       {"__builtin_bpf_preserve_type_info", "bpf", bpf_preserve_type_info},