Index: include/llvm/IR/IntrinsicsAMDGPU.td
===================================================================
--- include/llvm/IR/IntrinsicsAMDGPU.td
+++ include/llvm/IR/IntrinsicsAMDGPU.td
@@ -389,6 +389,11 @@
   GCCBuiltin<"__builtin_amdgcn_lerp">,
   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 
+// llvm.amdgcn.cmp.ne
+def int_amdgcn_cmp_ne :
+  GCCBuiltin<"__builtin_amdgcn_cmp_ne">,
+  Intrinsic<[llvm_i64_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
+
 //===----------------------------------------------------------------------===//
 // CI+ Intrinsics
 //===----------------------------------------------------------------------===//
Index: lib/Target/AMDGPU/SIInstructions.td
===================================================================
--- lib/Target/AMDGPU/SIInstructions.td
+++ lib/Target/AMDGPU/SIInstructions.td
@@ -2358,6 +2358,14 @@
 >;
 
 //===----------------------------------------------------------------------===//
+// V_CMP_NE Intrinsic Pattern.
+//===----------------------------------------------------------------------===//
+def : Pat <
+  (int_amdgcn_cmp_ne i32:$src0, i32:$src1),
+  (V_CMP_NE_I32_e64 i32:$src0, i32:$src1)
+>;
+
+//===----------------------------------------------------------------------===//
 // SMRD Patterns
 //===----------------------------------------------------------------------===//
 
Index: test/CodeGen/AMDGPU/llvm.amdgcn.cmp.ne.ll
===================================================================
--- /dev/null
+++ test/CodeGen/AMDGPU/llvm.amdgcn.cmp.ne.ll
@@ -0,0 +1,14 @@
+; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
+; RUN: llc -march=amdgcn -mcpu=fiji -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
+
+declare i64 @llvm.amdgcn.cmp.ne(i32, i32) #0
+
+; GCN-LABEL: {{^}}v_cmp_ne:
+; GCN: v_cmp_ne_i32_e64
+define void @v_cmp_ne(i64 addrspace(1)* %out, i32 %src) nounwind {
+  %result = call i64 @llvm.amdgcn.cmp.ne(i32 %src, i32 100) #0
+  store i64 %result, i64 addrspace(1)* %out, align 4
+  ret void
+}
+
+attributes #0 = { nounwind readnone convergent }