[AMDGPU] __builtin_amdgcn_read_exec_* should be implemented with llvm.amdgcn.ballot

[ranapratap55]

Summary

Clang provides builtin_amdgcn_read_exec, builtin_amdgcn_read_exec_lo and __builtin_amdgcn_read_exec_hi. These currently emit llvm.read_register with the corresponding registers.
This patch treats llvm.amdgcn.ballot(i1 true) as the canonical way to read a raw exec value.

[arsenm]

clang/lib/CodeGen/CGBuiltin.cpp

7786

You shouldn't be special casing EmitSpecialRegisterBuiltin, you should be specifically handling this in the __builtin_amdgcn_read_exec* handling

7788–7789

This should just always be getFalse

[ranapratap55]

Introduced EmitAMDGCNBallotForExec to create a ballot intrinsic for __builtin_amdgcn_read_exec*.

[arsenm]

clang/lib/CodeGen/CGBuiltin.cpp

17487–17488

Don't need the cast or setConvergent anymore

[ranapratap55]

clang/lib/CodeGen/CGBuiltin.cpp

7788–7789

ballot(i1 false) is resulting always to zero due to Line. So emitting ballot(i1 true) which results in exec AFAIK.

[arsenm]

clang/lib/CodeGen/CGBuiltin.cpp

7788–7789

Right this should be true

[arsenm]

As a follow up can you prepare an AMDGPUInstCombineIntrinsic patch to reduce ballot.i64 to ballot.i32 if wave32?

[ranapratap55]

clang/lib/CodeGen/CGBuiltin.cpp

17487–17488

removed cast & setConvergent.

clang/lib/CodeGen/CGBuiltin.cpp

static Value *EmitSpecialRegisterBuiltin(CodeGenFunction &CGF,
   llvm::Type *Types[] = { RegisterType };
 
   bool MixedTypes = RegisterType->isIntegerTy(64) && ValueType->isIntegerTy(32);
   assert(!(RegisterType->isIntegerTy(32) && ValueType->isIntegerTy(64))
             && "Can't fit 64-bit value in 32-bit register");
 
   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);
+    llvm::Value *Call = nullptr;
+    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});
+      Call = Builder.CreateCall(F, {Src});
+    }
 
     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;
   }
 
   llvm::Function *F = CGM.getIntrinsic(llvm::Intrinsic::write_register, Types);
   llvm::Value *ArgValue = CGF.EmitScalarExpr(E->getArg(1));
   if (MixedTypes) {
     // Extend 32 bit write value to 64 bit to pass to write.
     ArgValue = Builder.CreateZExt(ArgValue, RegisterType);
case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2f16: {
     llvm::Constant *ZeroI1 = llvm::ConstantInt::getIntegerValue(
         llvm::Type::getInt1Ty(getLLVMContext()), APInt(1, 0));
     llvm::Function *F = CGM.getIntrinsic(IID, {ArgTy});
     return Builder.CreateCall(F, {Addr, Val, ZeroI32, ZeroI32, ZeroI1});
   }
   case AMDGPU::BI__builtin_amdgcn_read_exec: {
     CallInst *CI = cast<CallInst>(
       EmitSpecialRegisterBuiltin(*this, E, Int64Ty, Int64Ty, NormalRead, "exec"));
     CI->setConvergent();
     return CI;
   }
   case AMDGPU::BI__builtin_amdgcn_read_exec_lo:
   case AMDGPU::BI__builtin_amdgcn_read_exec_hi: {
     StringRef RegName = BuiltinID == AMDGPU::BI__builtin_amdgcn_read_exec_lo ?
       "exec_lo" : "exec_hi";
     CallInst *CI = cast<CallInst>(
       EmitSpecialRegisterBuiltin(*this, E, Int32Ty, Int32Ty, NormalRead, RegName));
     CI->setConvergent();

clang/test/CodeGenOpenCL/builtins-amdgcn.cl

 
 // CHECK-LABEL: @test_cubema(
 // CHECK: call float @llvm.amdgcn.cubema(float %a, float %b, float %c)
 void test_cubema(global float* out, float a, float b, float c) {
   *out = __builtin_amdgcn_cubema(a, b, c);
 }
 
 // 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();
 }
 
 // CHECK-LABEL: @test_dispatch_ptr
 // CHECK: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
 void test_dispatch_ptr(__constant unsigned char ** out)
 {
void test_atomic_inc_dec(local uint *lptr, global uint *gptr, uint val) {
   res = __builtin_amdgcn_atomic_dec32(gptr, val, __ATOMIC_SEQ_CST, "");
 
   // CHECK: atomicrmw volatile udec_wrap ptr addrspace(1) %gptr, i32 %val seq_cst, align 4
   res = __builtin_amdgcn_atomic_dec32((volatile global uint*)gptr, val, __ATOMIC_SEQ_CST, "");
 }
 
 // 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"}