[AMDGPU] Extend readlane, writelane and readfirstlane intrinsic lowering for generic types

clang/test/CodeGenOpenCL/builtins-amdgcn.cl

@@ -308,14 +308,14 @@ void test_ds_bpermute(global int* out, int a, int b)
 }
 
 // CHECK-LABEL: @test_readfirstlane
-// CHECK: call i32 @llvm.amdgcn.readfirstlane.i32(i32 %a)
+// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.readfirstlane.i32(i32 %a)
 void test_readfirstlane(global int* out, int a)
 {
   *out = __builtin_amdgcn_readfirstlane(a);
 }
 
 // CHECK-LABEL: @test_readlane
-// CHECK: call i32 @llvm.amdgcn.readlane.i32(i32 %a, i32 %b)
+// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.readlane.i32(i32 %a, i32 %b)
 void test_readlane(global int* out, int a, int b)
 {
   *out = __builtin_amdgcn_readlane(a, b);

llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp

@@ -904,9 +904,6 @@ bool AMDGPUTargetLowering::isSDNodeAlwaysUniform(const SDNode *N) const {
     return false;
   case AMDGPUISD::SETCC: // ballot-style instruction
     return true;
-  case AMDGPUISD::READFIRSTLANE:
-  case AMDGPUISD::READLANE:
-    return true;
   }
   return false;
 }
@@ -5511,9 +5508,6 @@ const char* AMDGPUTargetLowering::getTargetNodeName(unsigned Opcode) const {
   NODE_NAME_CASE(LDS)
   NODE_NAME_CASE(FPTRUNC_ROUND_UPWARD)
   NODE_NAME_CASE(FPTRUNC_ROUND_DOWNWARD)
-  NODE_NAME_CASE(READLANE)
-  NODE_NAME_CASE(READFIRSTLANE)
-  NODE_NAME_CASE(WRITELANE)
   NODE_NAME_CASE(DUMMY_CHAIN)
   case AMDGPUISD::FIRST_MEM_OPCODE_NUMBER: break;
   NODE_NAME_CASE(LOAD_D16_HI)

llvm/lib/Target/AMDGPU/AMDGPUISelLowering.h

@@ -558,10 +558,6 @@ enum NodeType : unsigned {
   FPTRUNC_ROUND_UPWARD,
   FPTRUNC_ROUND_DOWNWARD,
 
-  READLANE,
-  READFIRSTLANE,
-  WRITELANE,
-
   DUMMY_CHAIN,
   FIRST_MEM_OPCODE_NUMBER = ISD::FIRST_TARGET_MEMORY_OPCODE,
   LOAD_D16_HI,

llvm/lib/Target/AMDGPU/AMDGPUInstrInfo.td

@@ -342,22 +342,6 @@ def AMDGPUfdot2_impl : SDNode<"AMDGPUISD::FDOT2",
 
 def AMDGPUperm_impl : SDNode<"AMDGPUISD::PERM", AMDGPUDTIntTernaryOp, []>;
 
-def AMDGPUReadfirstlaneOp : SDTypeProfile<1, 1, [
-  SDTCisSameAs<0, 1>
-]>;
-
-def AMDGPUReadlaneOp : SDTypeProfile<1, 2, [
-  SDTCisSameAs<0, 1>, SDTCisInt<2>
-]>;
-
-def AMDGPUDWritelaneOp : SDTypeProfile<1, 3, [
-  SDTCisSameAs<0, 1>, SDTCisInt<2>, SDTCisSameAs<0, 3>
-]>;
-
-def AMDGPUreadlane_impl : SDNode<"AMDGPUISD::READLANE", AMDGPUReadlaneOp, [SDNPOptInGlue]>;
-def AMDGPUreadfirstlane_impl : SDNode<"AMDGPUISD::READFIRSTLANE", AMDGPUReadfirstlaneOp, [SDNPOptInGlue]>;
-def AMDGPUwritelane_impl : SDNode<"AMDGPUISD::WRITELANE", AMDGPUDWritelaneOp, [SDNPOptInGlue]>;
-
 // SI+ export
 def AMDGPUExportOp : SDTypeProfile<0, 8, [
   SDTCisInt<0>,       // i8 tgt
@@ -523,15 +507,4 @@ def AMDGPUperm : PatFrags<(ops node:$src0, node:$src1, node:$src2),
   [(int_amdgcn_perm node:$src0, node:$src1, node:$src2),
    (AMDGPUperm_impl node:$src0, node:$src1, node:$src2)]>;
 
-def AMDGPUreadlane : PatFrags<(ops node:$src0, node:$src1),
-  [(int_amdgcn_readlane node:$src0, node:$src1),
-   (AMDGPUreadlane_impl node:$src0, node:$src1)]>;
-
-def AMDGPUreadfirstlane : PatFrags<(ops node:$src),
-  [(int_amdgcn_readfirstlane node:$src),
-   (AMDGPUreadfirstlane_impl node:$src)]>;
-
-def AMDGPUwritelane : PatFrags<(ops node:$src0, node:$src1, node:$src2),
-  [(int_amdgcn_writelane node:$src0, node:$src1, node:$src2),
-   (AMDGPUwritelane_impl node:$src0, node:$src1, node:$src2)]>;
 

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

@@ -6088,23 +6088,44 @@ static SDValue lowerLaneOp(const SITargetLowering &TLI, SDNode *N,
                            SelectionDAG &DAG) {
   EVT VT = N->getValueType(0);
   unsigned ValSize = VT.getSizeInBits();
-  unsigned IntrinsicID = N->getConstantOperandVal(0);
-  SDValue Src0 = N->getOperand(1);
+  unsigned IID = N->getConstantOperandVal(0);
   SDLoc SL(N);
   MVT IntVT = MVT::getIntegerVT(ValSize);
 
-  auto createLaneOp = [&DAG, &SL](SDValue Src0, SDValue Src1, SDValue Src2,
-                                  MVT VT) -> SDValue {
-    return (Src2 ? DAG.getNode(AMDGPUISD::WRITELANE, SL, VT, {Src0, Src1, Src2})
-            : Src1 ? DAG.getNode(AMDGPUISD::READLANE, SL, VT, {Src0, Src1})
-                   : DAG.getNode(AMDGPUISD::READFIRSTLANE, SL, VT, {Src0}));
+  auto createLaneOp = [&DAG, &SL, N, IID](SDValue Src0, SDValue Src1,
+                                          SDValue Src2, MVT ValT) -> SDValue {
+    SmallVector<SDValue, 8> Operands;
+    Operands.push_back(DAG.getTargetConstant(IID, SL, MVT::i32));
+    switch (IID) {
+    case Intrinsic::amdgcn_readfirstlane:
+      Operands.push_back(Src0);
+      break;
+    case Intrinsic::amdgcn_readlane:
+      Operands.push_back(Src0);
+      Operands.push_back(Src1);
+      break;
+    case Intrinsic::amdgcn_writelane:
+      Operands.push_back(Src0);
+      Operands.push_back(Src1);
+      Operands.push_back(Src2);
+      break;
+    }
+
+    if (SDNode *GL = N->getGluedNode()) {
+      assert(GL->getOpcode() == ISD::CONVERGENCECTRL_GLUE);
+      GL = GL->getOperand(0).getNode();
+      Operands.push_back(DAG.getNode(ISD::CONVERGENCECTRL_GLUE, SL, MVT::Glue,
+                                     SDValue(GL, 0)));
+    }
+
+    return DAG.getNode(ISD::INTRINSIC_WO_CHAIN, SL, ValT, Operands);
   };
 
+  SDValue Src0 = N->getOperand(1);
   SDValue Src1, Src2;
-  if (IntrinsicID == Intrinsic::amdgcn_readlane ||
-      IntrinsicID == Intrinsic::amdgcn_writelane) {
+  if (IID == Intrinsic::amdgcn_readlane || IID == Intrinsic::amdgcn_writelane) {
     Src1 = N->getOperand(2);
-    if (IntrinsicID == Intrinsic::amdgcn_writelane)
+    if (IID == Intrinsic::amdgcn_writelane)
       Src2 = N->getOperand(3);
   }
 
@@ -6129,13 +6150,55 @@ static SDValue lowerLaneOp(const SITargetLowering &TLI, SDNode *N,
   if (ValSize % 32 != 0)
     return SDValue();
 
+  auto unrollLaneOp = [&DAG, &SL](SDNode *N) -> SDValue {
+    EVT VT = N->getValueType(0);
+    unsigned NE = VT.getVectorNumElements();
+    EVT EltVT = VT.getVectorElementType();
+    SmallVector<SDValue, 8> Scalars;
+    unsigned NumOperands = N->getNumOperands();
+    SmallVector<SDValue, 4> Operands(NumOperands);
+    SDNode *GL = N->getGluedNode();
+
+    if (GL) {
+      // only handle convegrencectrl_glue
+      assert(GL->getOpcode() == ISD::CONVERGENCECTRL_GLUE);
+    }
+
+    for (unsigned i = 0; i != NE; ++i) {
+      for (unsigned j = 0, e = GL ? NumOperands - 1 : NumOperands; j != e;
+           ++j) {
+        SDValue Operand = N->getOperand(j);
+        EVT OperandVT = Operand.getValueType();
+        if (OperandVT.isVector()) {
+          // A vector operand; extract a single element.
+          EVT OperandEltVT = OperandVT.getVectorElementType();
+          Operands[j] = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, SL, OperandEltVT,
+                                    Operand, DAG.getVectorIdxConstant(i, SL));
+        } else {
+          // A scalar operand; just use it as is.
+          Operands[j] = Operand;
+        }
+      }
+
+      if (GL)
+        Operands[NumOperands - 1] =
+            DAG.getNode(ISD::CONVERGENCECTRL_GLUE, SL, MVT::Glue,
+                        SDValue(GL->getOperand(0).getNode(), 0));
+
+      Scalars.push_back(DAG.getNode(N->getOpcode(), SL, EltVT, Operands));
+    }
+
+    EVT VecVT = EVT::getVectorVT(*DAG.getContext(), EltVT, NE);
+    return DAG.getBuildVector(VecVT, SL, Scalars);
+  };
+
   if (VT.isVector()) {
     switch (MVT::SimpleValueType EltTy =
                 VT.getVectorElementType().getSimpleVT().SimpleTy) {
     case MVT::i32:
     case MVT::f32: {
       SDValue LaneOp = createLaneOp(Src0, Src1, Src2, VT.getSimpleVT());
-      return DAG.UnrollVectorOp(LaneOp.getNode());
+      return unrollLaneOp(LaneOp.getNode());
     }
     case MVT::i16:
     case MVT::f16:
@@ -6170,7 +6233,7 @@ static SDValue lowerLaneOp(const SITargetLowering &TLI, SDNode *N,
     Src2 = DAG.getBitcast(VecVT, Src2);
 
   SDValue LaneOp = createLaneOp(Src0, Src1, Src2, VecVT);
-  SDValue UnrolledLaneOp = DAG.UnrollVectorOp(LaneOp.getNode());
+  SDValue UnrolledLaneOp = unrollLaneOp(LaneOp.getNode());
   return DAG.getBitcast(VT, UnrolledLaneOp);
 }
 
@@ -15931,8 +15994,6 @@ bool SITargetLowering::isSDNodeSourceOfDivergence(const SDNode *N,
   case AMDGPUISD::BUFFER_ATOMIC_FMAX:
     // Target-specific read-modify-write atomics are sources of divergence.
     return true;
-  case AMDGPUISD::WRITELANE:
-    return true;
   default:
     if (auto *A = dyn_cast<AtomicSDNode>(N)) {
       // Generic read-modify-write atomics are sources of divergence.

llvm/lib/Target/AMDGPU/SIInstructions.td

@@ -3383,7 +3383,7 @@ def : GCNPat<
 // FIXME: Should also do this for readlane, but tablegen crashes on
 // the ignored src1.
 def : GCNPat<
-  (i32 (AMDGPUreadfirstlane (i32 imm:$src))),
+  (i32 (int_amdgcn_readfirstlane (i32 imm:$src))),
   (S_MOV_B32 SReg_32:$src)
 >;
 

llvm/lib/Target/AMDGPU/VOP1Instructions.td

@@ -255,7 +255,7 @@ def V_READFIRSTLANE_B32 : VOP1_Pseudo <"v_readfirstlane_b32", VOP_READFIRSTLANE,
 }
 
 foreach vt = Reg32Types.types in {
-  def : GCNPat<(vt (AMDGPUreadfirstlane (vt VRegOrLdsSrc_32:$src0))),
+  def : GCNPat<(vt (int_amdgcn_readfirstlane (vt VRegOrLdsSrc_32:$src0))),
         (V_READFIRSTLANE_B32 (vt VRegOrLdsSrc_32:$src0))
   >;
 }

llvm/lib/Target/AMDGPU/VOP2Instructions.td

@@ -779,19 +779,21 @@ defm V_SUBREV_U32 : VOP2Inst <"v_subrev_u32", VOP_I32_I32_I32_ARITH, null_frag,
 } // End isCommutable = 1
 
 // These are special and do not read the exec mask.
-let isConvergent = 1, Uses = []<Register> in {
-def V_READLANE_B32 : VOP2_Pseudo<"v_readlane_b32", VOP_READLANE,[]>;
+let isConvergent = 1, Uses = []<Register>, IsInvalidSingleUseConsumer = 1 in {
+def V_READLANE_B32 : VOP2_Pseudo<"v_readlane_b32", VOP_READLANE, []>;
 let IsNeverUniform = 1, Constraints = "$vdst = $vdst_in", DisableEncoding="$vdst_in" in {
-def V_WRITELANE_B32 : VOP2_Pseudo<"v_writelane_b32", VOP_WRITELANE, []>;
+def V_WRITELANE_B32 : VOP2_Pseudo<"v_writelane_b32", VOP_WRITELANE, []> {
+    let IsInvalidSingleUseProducer = 1;
+  }
 } // End IsNeverUniform, $vdst = $vdst_in, DisableEncoding $vdst_in
 } // End isConvergent = 1
 
 foreach vt = Reg32Types.types in {
-  def : GCNPat<(vt (AMDGPUreadlane vt:$src0, i32:$src1)),
+  def : GCNPat<(vt (int_amdgcn_readlane vt:$src0, i32:$src1)),
         (V_READLANE_B32 VRegOrLdsSrc_32:$src0, SCSrc_b32:$src1)
   >;
 
-  def : GCNPat<(vt (AMDGPUwritelane vt:$src0, i32:$src1, vt:$src2)),
+  def : GCNPat<(vt (int_amdgcn_writelane vt:$src0, i32:$src1, vt:$src2)),
         (V_WRITELANE_B32 SCSrc_b32:$src0, SCSrc_b32:$src1, VGPR_32:$src2)
   >;
 }

llvm/test/CodeGen/AMDGPU/convergence-laneops.ll

@@ -0,0 +1,65 @@
+; RUN: llc -stop-after=amdgpu-isel -mtriple=amdgcn-- -mcpu=gfx1100 -verify-machineinstrs -o - %s | FileCheck --check-prefixes=CHECK,ISEL %s
+
+; CHECK-LABEL: name:            basic_readfirstlane_i64
+;       CHECK:    [[TOKEN:%[0-9]+]]{{[^ ]*}} = CONVERGENCECTRL_ANCHOR
+;        ISEL:    CONVERGENCECTRL_GLUE [[TOKEN]]
+;        ISEL:    {{.*}} = V_READFIRSTLANE_B32 {{.*}}, implicit [[TOKEN]]
+;        ISEL:    CONVERGENCECTRL_GLUE [[TOKEN]]
+;        ISEL:    {{.*}} = V_READFIRSTLANE_B32 {{.*}}, implicit [[TOKEN]]
+define i64 @basic_readfirstlane_i64(i64 %src, i1 %cond) #0 {
+entry:
+  %t = call token @llvm.experimental.convergence.anchor()
+  %x = add i64 %src, 1
+  br i1 %cond, label %then, label %else
+
+then:
+  %r = call i64 @llvm.amdgcn.readfirstlane.i64(i64 %x) [ "convergencectrl"(token %t) ]
+  br label %else
+
+else:
+  %p = phi i64 [%r, %then], [%x, %entry]
+  ret i64 %p
+}
+
+; CHECK-LABEL: name:            basic_readlane_i64
+;       CHECK:    [[TOKEN:%[0-9]+]]{{[^ ]*}} = CONVERGENCECTRL_ANCHOR
+;        ISEL:    CONVERGENCECTRL_GLUE [[TOKEN]]
+;        ISEL:    {{.*}} = V_READLANE_B32 {{.*}}, implicit [[TOKEN]]
+;        ISEL:    CONVERGENCECTRL_GLUE [[TOKEN]]
+;        ISEL:    {{.*}} = V_READLANE_B32 {{.*}}, implicit [[TOKEN]]
+define i64 @basic_readlane_i64(i64 %src, i32 %lane, i1 %cond) #0 {
+entry:
+  %t = call token @llvm.experimental.convergence.anchor()
+  %x = add i64 %src, 1
+  br i1 %cond, label %then, label %else
+
+then:
+  %r = call i64 @llvm.amdgcn.readlane.i64(i64 %x, i32 %lane) [ "convergencectrl"(token %t) ]
+  br label %else
+
+else:
+  %p = phi i64 [%r, %then], [%x, %entry]
+  ret i64 %p
+}
+
+; CHECK-LABEL: name:            basic_writelane_i64
+;       CHECK:    [[TOKEN:%[0-9]+]]{{[^ ]*}} = CONVERGENCECTRL_ANCHOR
+;        ISEL:    CONVERGENCECTRL_GLUE [[TOKEN]]
+;        ISEL:    {{.*}} = V_WRITELANE_B32 {{.*}}, implicit [[TOKEN]]
+;        ISEL:    CONVERGENCECTRL_GLUE [[TOKEN]]
+;        ISEL:    {{.*}} = V_WRITELANE_B32 {{.*}}, implicit [[TOKEN]]
+define i64 @basic_writelane_i64(i64 %src, i1 %cond, i32 %lane, ptr addrspace(1) %out) #0 {
+entry:
+  %old = load i64, ptr addrspace(1) %out
+  %t = call token @llvm.experimental.convergence.anchor()
+  %x = add i64 %src, 1
+  br i1 %cond, label %then, label %else
+
+then:
+  %r = call i64 @llvm.amdgcn.writelane.i64(i64 %x, i32 %lane, i64 %old) [ "convergencectrl"(token %t) ]
+  br label %else
+
+else:
+  %p = phi i64 [%r, %then], [%x, %entry]
+  ret i64 %p
+}