[PATCH] D117487: AMDGPU: Remove llvm.amdgcn.alignbit and handle bitcode upgrade to fshr

[Matt Arsenault via Phabricator via llvm-commits]

arsenm created this revision.
arsenm added a reviewer: AMDGPU.
Herald added subscribers: foad, dexonsmith, kerbowa, hiraditya, t-tye, tpr, dstuttard, yaxunl, nhaehnle, jvesely, kzhuravl.
arsenm requested review of this revision.
Herald added a subscriber: wdng.
Herald added a project: LLVM.

https://reviews.llvm.org/D117487

Files:
  llvm/include/llvm/IR/IntrinsicsAMDGPU.td
  llvm/lib/IR/AutoUpgrade.cpp
  llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
  llvm/lib/Target/AMDGPU/SIISelLowering.cpp
  llvm/test/Bitcode/amdgcn-alignbit.ll
  llvm/test/CodeGen/AMDGPU/llvm.amdgcn.alignb.ll
  llvm/test/CodeGen/AMDGPU/llvm.amdgcn.alignbyte.ll


Index: llvm/test/CodeGen/AMDGPU/llvm.amdgcn.alignbyte.ll
===================================================================
--- llvm/test/CodeGen/AMDGPU/llvm.amdgcn.alignbyte.ll
+++ llvm/test/CodeGen/AMDGPU/llvm.amdgcn.alignbyte.ll
@@ -1,16 +1,7 @@
 ; 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 {
Index: llvm/test/Bitcode/amdgcn-alignbit.ll
===================================================================
--- /dev/null
+++ llvm/test/Bitcode/amdgcn-alignbit.ll
@@ -0,0 +1,12 @@
+; 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
Index: llvm/lib/Target/AMDGPU/SIISelLowering.cpp
===================================================================
--- llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -6933,9 +6933,6 @@
                                 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));
Index: llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
===================================================================
--- llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
+++ llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
@@ -4083,7 +4083,6 @@
     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:
Index: llvm/lib/IR/AutoUpgrade.cpp
===================================================================
--- llvm/lib/IR/AutoUpgrade.cpp
+++ llvm/lib/IR/AutoUpgrade.cpp
@@ -727,6 +727,13 @@
         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;
   }
 
Index: llvm/include/llvm/IR/IntrinsicsAMDGPU.td
===================================================================
--- llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -1514,12 +1514,6 @@
   [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]


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D117487.400536.patch
Type: text/x-patch
Size: 4053 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20220117/7bfc282d/attachment-0001.bin>