Index: include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- include/clang/Basic/BuiltinsAMDGPU.def
+++ include/clang/Basic/BuiltinsAMDGPU.def
@@ -70,6 +70,10 @@
 BUILTIN(__builtin_amdgcn_cubema, "ffff", "nc")
 BUILTIN(__builtin_amdgcn_s_memtime, "LUi", "n")
 BUILTIN(__builtin_amdgcn_s_sleep, "vIi", "n")
+BUILTIN(__builtin_amdgcn_icmp, "Wiiii", "nc")
+BUILTIN(__builtin_amdgcn_icmpw, "WiWiWii", "nc")
+BUILTIN(__builtin_amdgcn_fcmp, "Widdi", "nc")
+BUILTIN(__builtin_amdgcn_fcmpf, "Wiffi", "nc")
 
 //===----------------------------------------------------------------------===//
 // VI+ only builtins.
Index: lib/CodeGen/CGBuiltin.cpp
===================================================================
--- lib/CodeGen/CGBuiltin.cpp
+++ lib/CodeGen/CGBuiltin.cpp
@@ -7688,6 +7688,12 @@
     return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_fract);
   case AMDGPU::BI__builtin_amdgcn_lerp:
     return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_lerp);
+  case AMDGPU::BI__builtin_amdgcn_icmp:
+  case AMDGPU::BI__builtin_amdgcn_icmpw:
+    return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_icmp);
+  case AMDGPU::BI__builtin_amdgcn_fcmp:
+  case AMDGPU::BI__builtin_amdgcn_fcmpf:
+    return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_fcmp);
   case AMDGPU::BI__builtin_amdgcn_class:
   case AMDGPU::BI__builtin_amdgcn_classf:
     return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_class);
Index: test/CodeGenOpenCL/builtins-amdgcn.cl
===================================================================
--- test/CodeGenOpenCL/builtins-amdgcn.cl
+++ test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -4,6 +4,7 @@
 #pragma OPENCL EXTENSION cl_khr_fp64 : enable
 
 typedef unsigned long ulong;
+typedef long int lint;
 
 // CHECK-LABEL: @test_div_scale_f64
 // CHECK: call { double, i1 } @llvm.amdgcn.div.scale.f64(double %a, double %b, i1 true)
@@ -199,6 +200,34 @@
   *out = __builtin_amdgcn_lerp(a, b, c);
 }
 
+// CHECK-LABEL: @test_icmp_i32
+// CHECK: call i64 @llvm.amdgcn.icmp.i32
+void test_icmp_i32(global lint* out, int a, int b, int c)
+{
+  *out = __builtin_amdgcn_icmp(a, b, c);
+}
+
+// CHECK-LABEL: @test_icmp_i64
+// CHECK: call i64 @llvm.amdgcn.icmp.i64
+void test_icmp_i64(global lint* out, lint a, lint b, int c)
+{
+  *out = __builtin_amdgcn_icmpw(a, b, c);
+}
+
+// CHECK-LABEL: @test_fcmp_f32
+// CHECK: call i64 @llvm.amdgcn.fcmp.f32
+void test_fcmp_f32(global lint* out, float a, float b, int c)
+{
+  *out = __builtin_amdgcn_fcmpf(a, b, c);
+}
+
+// CHECK-LABEL: @test_fcmp_f64
+// CHECK: call i64 @llvm.amdgcn.fcmp.f64
+void test_fcmp_f64(global lint* out, double a, double b, int c)
+{
+  *out = __builtin_amdgcn_fcmp(a, b, c);
+}
+
 // CHECK-LABEL: @test_class_f32
 // CHECK: call i1 @llvm.amdgcn.class.f32
 void test_class_f32(global float* out, float a, int b)