AMDGPU: Remove llvm.amdgcn.alignbit and handle bitcode upgrade to fshr

[foad]

LGTM, thanks!

[arsenm]

82de129ab8f723ba94d0026b54d76b11b2a9e4f9

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

def int_amdgcn_writelane :
   Intrinsic<[llvm_i32_ty], [
     llvm_i32_ty,    // uniform value to write: returned by the selected lane
     llvm_i32_ty,    // uniform lane select
     llvm_i32_ty     // returned by all lanes other than the selected one
   ],
   [IntrNoMem, IntrConvergent, IntrWillReturn]
 >;
 
-// FIXME: Deprecated. This is equivalent to llvm.fshr
-def int_amdgcn_alignbit : Intrinsic<[llvm_i32_ty],
-  [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
-  [IntrNoMem, IntrSpeculatable, IntrWillReturn]
->;
-
 def int_amdgcn_alignbyte : GCCBuiltin<"__builtin_amdgcn_alignbyte">,
   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
   [IntrNoMem, IntrSpeculatable, IntrWillReturn]
 >;
 
 def int_amdgcn_mul_i24 : Intrinsic<[llvm_i32_ty],
   [llvm_i32_ty, llvm_i32_ty],
   [IntrNoMem, IntrSpeculatable, IntrWillReturn]

llvm/lib/IR/AutoUpgrade.cpp

if (Name == "arm.mve.mull.int.predicated.v2i64.v4i32.v4i1" ||
         Name == "arm.cde.vcx1q.predicated.v2i64.v4i1" ||
         Name == "arm.cde.vcx1qa.predicated.v2i64.v4i1" ||
         Name == "arm.cde.vcx2q.predicated.v2i64.v4i1" ||
         Name == "arm.cde.vcx2qa.predicated.v2i64.v4i1" ||
         Name == "arm.cde.vcx3q.predicated.v2i64.v4i1" ||
         Name == "arm.cde.vcx3qa.predicated.v2i64.v4i1")
       return true;
 
+    if (Name == "amdgcn.alignbit") {
+      // Target specific intrinsic became redundant
+      NewFn = Intrinsic::getDeclaration(F->getParent(), Intrinsic::fshr,
+                                        {F->getReturnType()});
+      return true;
+    }
+
     break;
   }
 
   case 'c': {
     if (Name.startswith("ctlz.") && F->arg_size() == 1) {
       rename(F);
       NewFn = Intrinsic::getDeclaration(F->getParent(), Intrinsic::ctlz,
                                         F->arg_begin()->getType());

llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp

case AMDGPU::G_INTRINSIC: {
     case Intrinsic::amdgcn_sad_u8:
     case Intrinsic::amdgcn_msad_u8:
     case Intrinsic::amdgcn_sad_hi_u8:
     case Intrinsic::amdgcn_sad_u16:
     case Intrinsic::amdgcn_qsad_pk_u16_u8:
     case Intrinsic::amdgcn_mqsad_pk_u16_u8:
     case Intrinsic::amdgcn_mqsad_u32_u8:
     case Intrinsic::amdgcn_cvt_pk_u8_f32:
-    case Intrinsic::amdgcn_alignbit:
     case Intrinsic::amdgcn_alignbyte:
     case Intrinsic::amdgcn_perm:
     case Intrinsic::amdgcn_fdot2:
     case Intrinsic::amdgcn_sdot2:
     case Intrinsic::amdgcn_udot2:
     case Intrinsic::amdgcn_sdot4:
     case Intrinsic::amdgcn_udot4:
     case Intrinsic::amdgcn_sdot8:

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

case Intrinsic::amdgcn_is_private: {
     SDValue Aperture = getSegmentAperture(AS, SL, DAG);
     SDValue SrcVec = DAG.getNode(ISD::BITCAST, DL, MVT::v2i32,
                                  Op.getOperand(1));
 
     SDValue SrcHi = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, SL, MVT::i32, SrcVec,
                                 DAG.getConstant(1, SL, MVT::i32));
     return DAG.getSetCC(SL, MVT::i1, SrcHi, Aperture, ISD::SETEQ);
   }
-  case Intrinsic::amdgcn_alignbit:
-    return DAG.getNode(ISD::FSHR, DL, VT,
-                       Op.getOperand(1), Op.getOperand(2), Op.getOperand(3));
   case Intrinsic::amdgcn_perm:
     return DAG.getNode(AMDGPUISD::PERM, DL, MVT::i32, Op.getOperand(1),
                        Op.getOperand(2), Op.getOperand(3));
   case Intrinsic::amdgcn_reloc_constant: {
     Module *M = const_cast<Module *>(MF.getFunction().getParent());
     const MDNode *Metadata = cast<MDNodeSDNode>(Op.getOperand(1))->getMD();
     auto SymbolName = cast<MDString>(Metadata->getOperand(0))->getString();
     auto RelocSymbol = cast<GlobalVariable>(

llvm/test/Bitcode/amdgcn-alignbit.ll

+; RUN: llvm-as < %s | llvm-dis | FileCheck %s
+
+define i32 @user(i32 %a, i32 %b, i32 %c) {
+  ; CHECK: %call = call i32 @llvm.fshr.i32(i32 %a, i32 %b, i32 %c)
+  ; CHECK-NOT: amdgcn.alignbit
+  %call = call i32 @llvm.amdgcn.alignbit(i32 %a, i32 %b, i32 %c)
+  ret i32 %call
+}
+
+declare i32 @llvm.amdgcn.alignbit(i32, i32, i32)
+; CHECK: declare i32 @llvm.fshr.i32(i32, i32, i32) #0
+; CHECK-NOT: amdgcn.alignbit

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.alignbyte.ll

 ; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
 
-declare i32 @llvm.amdgcn.alignbit(i32, i32, i32) #0
 declare i32 @llvm.amdgcn.alignbyte(i32, i32, i32) #0
 
-; GCN-LABEL: {{^}}v_alignbit_b32:
-; GCN: v_alignbit_b32 {{[vs][0-9]+}}, {{[vs][0-9]+}}, {{[vs][0-9]+}}
-define amdgpu_kernel void @v_alignbit_b32(i32 addrspace(1)* %out, i32 %src1, i32 %src2, i32 %src3) #1 {
-  %val = call i32 @llvm.amdgcn.alignbit(i32 %src1, i32 %src2, i32 %src3) #0
-  store i32 %val, i32 addrspace(1)* %out
-  ret void
-}
-
 ; GCN-LABEL: {{^}}v_alignbyte_b32:
 ; GCN: v_alignbyte_b32 {{[vs][0-9]+}}, {{[vs][0-9]+}}, {{[vs][0-9]+}}
 define amdgpu_kernel void @v_alignbyte_b32(i32 addrspace(1)* %out, i32 %src1, i32 %src2, i32 %src3) #1 {
   %val = call i32 @llvm.amdgcn.alignbyte(i32 %src1, i32 %src2, i32 %src3) #0
   store i32 %val, i32 addrspace(1)* %out
   ret void
 }
 
 attributes #0 = { nounwind readnone }
 attributes #1 = { nounwind }