Index: include/llvm/IR/IntrinsicsAMDGPU.td
===================================================================
--- include/llvm/IR/IntrinsicsAMDGPU.td
+++ include/llvm/IR/IntrinsicsAMDGPU.td
@@ -83,3 +83,21 @@
                                        "__builtin_amdgpu_read_workdim">;
 
 } // End TargetPrefix = "AMDGPU"
+
+let TargetPrefix = "amdgcn" in {
+
+// SI only
+def int_amdgcn_buffer_wbinvl1_sc :
+  GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_sc">,
+  Intrinsic<[], [], []>;
+
+// On CI+
+def int_amdgcn_buffer_wbinvl1_vol :
+  GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_vol">,
+  Intrinsic<[], [], []>;
+
+def int_amdgcn_buffer_wbinvl1 :
+  GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1">,
+  Intrinsic<[], [], []>;
+
+}
Index: lib/Target/AMDGPU/CIInstructions.td
===================================================================
--- lib/Target/AMDGPU/CIInstructions.td
+++ lib/Target/AMDGPU/CIInstructions.td
@@ -43,6 +43,14 @@
 >;
 
 //===----------------------------------------------------------------------===//
+// MUBUF Instructions
+//===----------------------------------------------------------------------===//
+
+defm BUFFER_WBINVL1_VOL : MUBUF_Invalidate <mubuf<0x70, 0x3f>,
+  "buffer_wbinvl1_vol", int_amdgcn_buffer_wbinvl1_vol
+>;
+
+//===----------------------------------------------------------------------===//
 // Flat Instructions
 //===----------------------------------------------------------------------===//
 
Index: lib/Target/AMDGPU/SIInstrInfo.td
===================================================================
--- lib/Target/AMDGPU/SIInstrInfo.td
+++ lib/Target/AMDGPU/SIInstrInfo.td
@@ -2440,6 +2440,23 @@
   } // End mayLoad = 0, mayStore = 1
 }
 
+// For cache invalidation instructions.
+multiclass MUBUF_Invalidate <mubuf op, string opName, SDPatternOperator node> {
+  let hasSideEffects = 1, mayStore = 1, AsmMatchConverter = "" in {
+    def "" : MUBUF_Pseudo <opName, (outs), (ins), [(node)]>;
+
+    // Set everything to 0.
+    let offset = 0, offen = 0, idxen = 0, glc = 0, vaddr = 0,
+        vdata = 0, srsrc = 0, slc = 0, tfe = 0, soffset = 0 in {
+      let addr64 = 0 in {
+        def _si : MUBUF_Real_si <op, opName, (outs), (ins), opName>;
+      }
+
+      def _vi : MUBUF_Real_vi <op, opName, (outs), (ins), opName>;
+    }
+  } // End hasSideEffects = 1, mayStore = 1, AsmMatchConverter = ""
+}
+
 class FLAT_Load_Helper <bits<7> op, string asm, RegisterClass regClass> :
       FLAT <op, (outs regClass:$vdst),
                 (ins VReg_64:$addr, glc_flat:$glc, slc_flat:$slc, tfe_flat:$tfe),
Index: lib/Target/AMDGPU/SIInstructions.td
===================================================================
--- lib/Target/AMDGPU/SIInstructions.td
+++ lib/Target/AMDGPU/SIInstructions.td
@@ -30,7 +30,9 @@
                       ">= AMDGPUSubtarget::SOUTHERN_ISLANDS">,
             AssemblerPredicate<"FeatureGCN">;
 def isSI : Predicate<"Subtarget->getGeneration() "
-                      "== AMDGPUSubtarget::SOUTHERN_ISLANDS">;
+                      "== AMDGPUSubtarget::SOUTHERN_ISLANDS">,
+           AssemblerPredicate<"FeatureSouthernIslands">;
+
 
 def has16BankLDS : Predicate<"Subtarget->getLDSBankCount() == 16">;
 def has32BankLDS : Predicate<"Subtarget->getLDSBankCount() == 32">;
@@ -1036,9 +1038,12 @@
 //def BUFFER_ATOMIC_FCMPSWAP_X2 : MUBUF_X2 <mubuf<0x5e>, "buffer_atomic_fcmpswap_x2", []>; // isn't on VI
 //def BUFFER_ATOMIC_FMIN_X2 : MUBUF_X2 <mubuf<0x5f>, "buffer_atomic_fmin_x2", []>; // isn't on VI
 //def BUFFER_ATOMIC_FMAX_X2 : MUBUF_X2 <mubuf<0x60>, "buffer_atomic_fmax_x2", []>; // isn't on VI
-//def BUFFER_WBINVL1_SC : MUBUF_WBINVL1 <mubuf<0x70>, "buffer_wbinvl1_sc", []>; // isn't on CI & VI
-//def BUFFER_WBINVL1_VOL : MUBUF_WBINVL1 <mubuf<0x70, 0x3f>, "buffer_wbinvl1_vol", []>; // isn't on SI
-//def BUFFER_WBINVL1 : MUBUF_WBINVL1 <mubuf<0x71, 0x3e>, "buffer_wbinvl1", []>;
+
+let SubtargetPredicate = isSI in {
+defm BUFFER_WBINVL1_SC : MUBUF_Invalidate <mubuf<0x70>, "buffer_wbinvl1_sc", int_amdgcn_buffer_wbinvl1_sc>; // isn't on CI & VI
+}
+
+defm BUFFER_WBINVL1 : MUBUF_Invalidate <mubuf<0x71, 0x3e>, "buffer_wbinvl1", int_amdgcn_buffer_wbinvl1>;
 
 //===----------------------------------------------------------------------===//
 // MTBUF Instructions
Index: test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.ll
===================================================================
--- /dev/null
+++ test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.ll
@@ -0,0 +1,16 @@
+; RUN: llc -march=amdgcn -mcpu=tahiti -show-mc-encoding < %s | FileCheck -check-prefix=GCN -check-prefix=SI %s
+; RUN: llc -march=amdgcn -mcpu=fiji -show-mc-encoding < %s | FileCheck -check-prefix=GCN -check-prefix=VI %s
+
+declare void @llvm.amdgcn.buffer.wbinvl1() #0
+
+; GCN-LABEL: {{^}}test_buffer_wbinvl1:
+; GCN-NEXT: ; BB#0:
+; SI-NEXT: buffer_wbinvl1 ; encoding: [0x00,0x00,0xc4,0xe1,0x00,0x00,0x00,0x00]
+; VI-NEXT: buffer_wbinvl1 ; encoding: [0x00,0x00,0xf8,0xe0,0x00,0x00,0x00,0x00]
+; GCN-NEXT: s_endpgm
+define void @test_buffer_wbinvl1() #0 {
+  call void @llvm.amdgcn.buffer.wbinvl1()
+  ret void
+}
+
+attributes #0 = { nounwind }
Index: test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.sc.ll
===================================================================
--- /dev/null
+++ test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.sc.ll
@@ -0,0 +1,14 @@
+; RUN: llc -march=amdgcn -mcpu=tahiti -show-mc-encoding < %s | FileCheck -check-prefix=SI %s
+
+declare void @llvm.amdgcn.buffer.wbinvl1.sc() #0
+
+; SI-LABEL: {{^}}test_buffer_wbinvl1_sc:
+; SI-NEXT: ; BB#0:
+; SI-NEXT: buffer_wbinvl1_sc ; encoding: [0x00,0x00,0xc0,0xe1,0x00,0x00,0x00,0x00]
+; SI-NEXT: s_endpgm
+define void @test_buffer_wbinvl1_sc() #0 {
+  call void @llvm.amdgcn.buffer.wbinvl1.sc()
+  ret void
+}
+
+attributes #0 = { nounwind }
Index: test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.vol.ll
===================================================================
--- /dev/null
+++ test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.vol.ll
@@ -0,0 +1,16 @@
+; RUN: llc -march=amdgcn -mcpu=bonaire -show-mc-encoding < %s | FileCheck -check-prefix=GCN -check-prefix=CI %s
+; RUN: llc -march=amdgcn -mcpu=tonga -show-mc-encoding < %s | FileCheck -check-prefix=GCN -check-prefix=VI %s
+
+declare void @llvm.amdgcn.buffer.wbinvl1.vol() #0
+
+; GCN-LABEL: {{^}}test_buffer_wbinvl1_vol:
+; GCN-NEXT: ; BB#0:
+; CI-NEXT: buffer_wbinvl1_vol ; encoding: [0x00,0x00,0xc0,0xe1,0x00,0x00,0x00,0x00]
+; VI-NEXT: buffer_wbinvl1_vol ; encoding: [0x00,0x00,0xfc,0xe0,0x00,0x00,0x00,0x00]
+; GCN-NEXT: s_endpgm
+define void @test_buffer_wbinvl1_vol() #0 {
+  call void @llvm.amdgcn.buffer.wbinvl1.vol()
+  ret void
+}
+
+attributes #0 = { nounwind }
Index: test/MC/AMDGPU/mubuf-ci.s
===================================================================
--- /dev/null
+++ test/MC/AMDGPU/mubuf-ci.s
@@ -0,0 +1,8 @@
+// RUN: llvm-mc -arch=amdgcn -mcpu=bonaire -show-encoding %s | FileCheck -check-prefix=CI %s
+// XUN: llvm-mc -arch=amdgcn -mcpu=tonga -show-encoding %s | FileCheck -check-prefix=VI %s
+// RUN: not llvm-mc -arch=amdgcn -mcpu=tahiti %s 2>&1 | FileCheck -check-prefix=NOSI %s
+
+buffer_wbinvl1_vol
+// CI: buffer_wbinvl1_vol ; encoding: [0x00,0x00,0xc0,0xe1,0x00,0x00,0x00,0x00]
+// XVI: buffer_wbinvl1_vol ; encoding: [0x00,0x00,0xfc,0xe0,0x00,0x00,0x00,0x00]
+// NOSI: error: instruction not supported on this GPU
Index: test/MC/AMDGPU/mubuf-si.s
===================================================================
--- /dev/null
+++ test/MC/AMDGPU/mubuf-si.s
@@ -0,0 +1,13 @@
+// RUN: llvm-mc -arch=amdgcn -mcpu=SI -show-encoding %s | FileCheck %s
+// RUN: not llvm-mc -arch=amdgcn -mcpu=bonaire %s 2>&1 | FileCheck -check-prefix=NOCI %s
+// RUN: not llvm-mc -arch=amdgcn -mcpu=tonga %s 2>&1 | FileCheck -check-prefix=NOVI %s
+
+//===----------------------------------------------------------------------===//
+// Cache invalidation
+//===----------------------------------------------------------------------===//
+
+buffer_wbinvl1_sc
+// CHECK: buffer_wbinvl1_sc ; encoding: [0x00,0x00,0xc0,0xe1,0x00,0x00,0x00,0x00]
+// NOCI: error: instruction not supported on this GPU
+// NOVI: error: instruction not supported on this GPU
+
Index: test/MC/AMDGPU/mubuf.s
===================================================================
--- test/MC/AMDGPU/mubuf.s
+++ test/MC/AMDGPU/mubuf.s
@@ -349,4 +349,11 @@
 buffer_store_dwordx4 v[1:4], s[4:7], s1
 // CHECK: buffer_store_dwordx4 v[1:4], s[4:7], s1 ; encoding: [0x00,0x00,0x78,0xe0,0x00,0x01,0x01,0x01]
 
+//===----------------------------------------------------------------------===//
+// Cache invalidation
+//===----------------------------------------------------------------------===//
+
+buffer_wbinvl1
+// CHECK: buffer_wbinvl1   ; encoding: [0x00,0x00,0xc4,0xe1,0x00,0x00,0x00,0x00]
+
 // TODO: Atomics