[clang] [AMDGPU] Lower __builtin_amdgcn_read_exec_hi to use amdgcn_ballot (PR #69567)

[Matt Arsenault via cfe-commits]

================
@@ -7997,14 +7997,26 @@ enum SpecialRegisterAccessKind {
 
 static Value *EmitAMDGCNBallotForExec(CodeGenFunction &CGF, const CallExpr *E,
                                       llvm::Type *RegisterType,
-                                      llvm::Type *ValueType) {
+                                      llvm::Type *ValueType, bool isExecHi) {
   CodeGen::CGBuilderTy &Builder = CGF.Builder;
   CodeGen::CodeGenModule &CGM = CGF.CGM;
 
   llvm::Type *ResultType = CGF.ConvertType(E->getType());
-  Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_ballot, {ResultType});
-  llvm::Value *Call = Builder.CreateCall(F, {Builder.getInt1(true)});
-  return Call;
+  llvm::Value *Call;
+  Function *F;
+
+  if (isExecHi) {
+    F = CGM.getIntrinsic(Intrinsic::amdgcn_ballot, {RegisterType});
+    Call = Builder.CreateCall(F, {Builder.getInt1(true)});
+    Value *C1 = llvm::ConstantInt::get(ValueType, 32);
+    Value *Rt2 = Builder.CreateLShr(Call, C1);
+    Rt2 = Builder.CreateTruncOrBitCast(Rt2, CGF.Int32Ty);
+    return Rt2;
+  } else {
+    F = CGM.getIntrinsic(Intrinsic::amdgcn_ballot, {ResultType});
+    Call = Builder.CreateCall(F, {Builder.getInt1(true)});
----------------
arsenm wrote:

Can make this part common 

https://github.com/llvm/llvm-project/pull/69567