Index: clang/lib/CodeGen/CGBuiltin.cpp
===================================================================
--- clang/lib/CodeGen/CGBuiltin.cpp
+++ clang/lib/CodeGen/CGBuiltin.cpp
@@ -7782,21 +7782,15 @@
 
   if (AccessKind != Write) {
     assert(AccessKind == NormalRead || AccessKind == VolatileRead);
-    llvm::Function *F = CGM.getIntrinsic(
-        AccessKind == VolatileRead ? llvm::Intrinsic::read_volatile_register
-                                   : llvm::Intrinsic::read_register,
-        Types);
-    llvm::Value *Call = Builder.CreateCall(F, Metadata);
-
-    if (MixedTypes)
-      // Read into 64 bit register and then truncate result to 32 bit.
-      return Builder.CreateTrunc(Call, ValueType);
-
-    if (ValueType->isPointerTy())
-      // Have i32/i64 result (Call) but want to return a VoidPtrTy (i8*).
-      return Builder.CreateIntToPtr(Call, ValueType);
 
-    return Call;
+    if (SysReg == "exec" || SysReg == "exec_lo" || SysReg == "exec_hi") {
+      llvm::Type *ResultType = CGF.ConvertType(E->getType());
+      llvm::Value *Src = (SysReg == "exec_hi") ? Builder.getInt1(false)
+                                               : Builder.getInt1(true);
+      Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_ballot, {ResultType});
+      llvm::Value *Call = Builder.CreateCall(F, {Src});
+      return Call;
+    }
   }
 
   llvm::Function *F = CGM.getIntrinsic(llvm::Intrinsic::write_register, Types);
Index: clang/test/CodeGenOpenCL/builtins-amdgcn.cl
===================================================================
--- clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -510,21 +510,21 @@
 }
 
 // CHECK-LABEL: @test_read_exec(
-// CHECK: call i64 @llvm.read_register.i64(metadata ![[$EXEC:[0-9]+]]) #[[$READ_EXEC_ATTRS:[0-9]+]]
+// CHECK: call i64 @llvm.amdgcn.ballot.i64(i1 true) #[[$READ_EXEC_ATTRS:[0-9]+]]
 void test_read_exec(global ulong* out) {
   *out = __builtin_amdgcn_read_exec();
 }
 
-// CHECK: declare i64 @llvm.read_register.i64(metadata) #[[$NOUNWIND_READONLY:[0-9]+]]
+// CHECK: declare i64 @llvm.amdgcn.ballot.i64(i1) #[[$NOUNWIND_READONLY:[0-9]+]]
 
 // CHECK-LABEL: @test_read_exec_lo(
-// CHECK: call i32 @llvm.read_register.i32(metadata ![[$EXEC_LO:[0-9]+]]) #[[$READ_EXEC_ATTRS]]
+// CHECK: call i32 @llvm.amdgcn.ballot.i32(i1 true) #[[$READ_EXEC_ATTRS]]
 void test_read_exec_lo(global uint* out) {
   *out = __builtin_amdgcn_read_exec_lo();
 }
 
 // CHECK-LABEL: @test_read_exec_hi(
-// CHECK: call i32 @llvm.read_register.i32(metadata ![[$EXEC_HI:[0-9]+]]) #[[$READ_EXEC_ATTRS]]
+// CHECK: store i32 0, ptr addrspace(1) %out, align 4
 void test_read_exec_hi(global uint* out) {
   *out = __builtin_amdgcn_read_exec_hi();
 }
@@ -830,8 +830,5 @@
 
 // CHECK-DAG: [[$WI_RANGE]] = !{i32 0, i32 1024}
 // CHECK-DAG: [[$WS_RANGE]] = !{i16 1, i16 1025}
-// CHECK-DAG: attributes #[[$NOUNWIND_READONLY]] = { mustprogress nocallback nofree nosync nounwind willreturn memory(read) }
+// CHECK-DAG: attributes #[[$NOUNWIND_READONLY]] = { convergent mustprogress nocallback nofree nounwind willreturn memory(none) }
 // CHECK-DAG: attributes #[[$READ_EXEC_ATTRS]] = { convergent }
-// CHECK-DAG: ![[$EXEC]] = !{!"exec"}
-// CHECK-DAG: ![[$EXEC_LO]] = !{!"exec_lo"}
-// CHECK-DAG: ![[$EXEC_HI]] = !{!"exec_hi"}