Index: lib/CodeGen/MachineSink.cpp
===================================================================
--- lib/CodeGen/MachineSink.cpp
+++ lib/CodeGen/MachineSink.cpp
@@ -23,6 +23,7 @@
 #include "llvm/Analysis/AliasAnalysis.h"
 #include "llvm/CodeGen/MachineDominators.h"
 #include "llvm/CodeGen/MachineLoopInfo.h"
+#include "llvm/CodeGen/MachinePostDominators.h"
 #include "llvm/CodeGen/MachineRegisterInfo.h"
 #include "llvm/Support/CommandLine.h"
 #include "llvm/Support/Debug.h"
@@ -40,6 +41,16 @@
            cl::desc("Split critical edges during machine sinking"),
            cl::init(true), cl::Hidden);
 
+// FIXME: Using post dominator itself is not problematic but would trigger a
+// separate issue with TwoAddressInstructionPass and RegisterCoalescer
+// (PR20776). Therefore, we make this flag experimental and only enable it
+// using command line arguments.
+static cl::opt<bool> UsePostDominator(
+    "machine-sink-use-post-dominator",
+    cl::desc(
+        "Uses real post dominator to decide whether sinking is profitable"),
+    cl::init(false), cl::Hidden);
+
 STATISTIC(NumSunk,      "Number of machine instructions sunk");
 STATISTIC(NumSplit,     "Number of critical edges split");
 STATISTIC(NumCoalesces, "Number of copies coalesced");
@@ -48,8 +59,9 @@
   class MachineSinking : public MachineFunctionPass {
     const TargetInstrInfo *TII;
     const TargetRegisterInfo *TRI;
-    MachineRegisterInfo  *MRI;  // Machine register information
-    MachineDominatorTree *DT;   // Machine dominator tree
+    MachineRegisterInfo  *MRI;     // Machine register information
+    MachineDominatorTree *DT;      // Machine dominator tree
+    MachinePostDominatorTree *PDT; // Machine post dominator tree
     MachineLoopInfo *LI;
     AliasAnalysis *AA;
 
@@ -74,8 +86,12 @@
       MachineFunctionPass::getAnalysisUsage(AU);
       AU.addRequired<AliasAnalysis>();
       AU.addRequired<MachineDominatorTree>();
+      if (UsePostDominator)
+        AU.addRequired<MachinePostDominatorTree>();
       AU.addRequired<MachineLoopInfo>();
       AU.addPreserved<MachineDominatorTree>();
+      if (UsePostDominator)
+        AU.addPreserved<MachinePostDominatorTree>();
       AU.addPreserved<MachineLoopInfo>();
     }
 
@@ -236,6 +252,8 @@
   TRI = TM.getSubtargetImpl()->getRegisterInfo();
   MRI = &MF.getRegInfo();
   DT = &getAnalysis<MachineDominatorTree>();
+  if (UsePostDominator)
+    PDT = &getAnalysis<MachinePostDominatorTree>();
   LI = &getAnalysis<MachineLoopInfo>();
   AA = &getAnalysis<AliasAnalysis>();
 
@@ -454,9 +472,11 @@
 }
 
 /// isPostDominatedBy - Return true if A is post dominated by B.
-static bool isPostDominatedBy(MachineBasicBlock *A, MachineBasicBlock *B) {
+static bool isPostDominatedBy(MachinePostDominatorTree *PDT,
+                              MachineBasicBlock *A, MachineBasicBlock *B) {
+  if (UsePostDominator)
+    return PDT->dominates(B, A);
 
-  // FIXME - Use real post dominator.
   if (A->succ_size() != 2)
     return false;
   MachineBasicBlock::succ_iterator I = A->succ_begin();
@@ -481,8 +501,8 @@
     return false;
 
   // It is profitable if SuccToSinkTo does not post dominate current block.
-  if (!isPostDominatedBy(MBB, SuccToSinkTo))
-      return true;
+  if (!isPostDominatedBy(PDT, MBB, SuccToSinkTo))
+    return true;
 
   // Check if only use in post dominated block is PHI instruction.
   bool NonPHIUse = false;
Index: test/CodeGen/NVPTX/machine-sink.ll
===================================================================
--- /dev/null
+++ test/CodeGen/NVPTX/machine-sink.ll
@@ -0,0 +1,40 @@
+; RUN: llc < %s -march=nvptx -mcpu=sm_20 -machine-sink-use-post-dominator | FileCheck %s
+
+target datalayout = "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
+
+@scalar1 = internal addrspace(3) global float 0.000000e+00, align 4
+@scalar2 = internal addrspace(3) global float 0.000000e+00, align 4
+
+; We shouldn't sink mul.rn.f32 to BB %merge because BB %merge post-dominates
+; BB %entry. Over-sinking created more register pressure on this example. The
+; backend would sink the fmuls to BB %merge, but not the loads for being
+; conservative on sinking memory accesses. As a result, the loads and
+; the two fmuls would be separated to two basic blocks, causing two
+; cross-BB live ranges.
+define float @post_dominate(float %x, i1 %cond) {
+; CHECK-LABEL: post_dominate(
+entry:
+  %0 = load float* addrspacecast (float addrspace(3)* @scalar1 to float*), align 4
+  %1 = load float* addrspacecast (float addrspace(3)* @scalar2 to float*), align 4
+; CHECK: ld.shared.f32
+; CHECK: ld.shared.f32
+  %2 = fmul float %0, %0
+  %3 = fmul float %1, %2
+; CHECK-NOT: bra
+; CHECK: mul.rn.f32
+; CHECK: mul.rn.f32
+  br i1 %cond, label %then, label %merge
+
+then:
+  %z = fadd float %x, %x
+  br label %then2
+
+then2:
+  %z2 = fadd float %z, %z
+  br label %merge
+
+merge:
+  %y = phi float [ 0.0, %entry ], [ %z2, %then2 ]
+  %w = fadd float %y, %3
+  ret float %w
+}