AMDGPU: Add builtin to read exec mask

[• tstellarAMD]

LGTM.

[arsenm]

r273965

include/clang/Basic/BuiltinsAMDGPU.def

 //===----------------------------------------------------------------------===//
 
 // The format of this database matches clang/Basic/Builtins.def.
 
 #if defined(BUILTIN) && !defined(TARGET_BUILTIN)
 #   define TARGET_BUILTIN(ID, TYPE, ATTRS, FEATURE) BUILTIN(ID, TYPE, ATTRS)
 #endif
 
+//===----------------------------------------------------------------------===//
+// Instruction builtins.
+//===----------------------------------------------------------------------===//
 BUILTIN(__builtin_amdgcn_s_barrier, "v", "n")
 BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n")
 BUILTIN(__builtin_amdgcn_div_scalef, "fffbb*", "n")
 BUILTIN(__builtin_amdgcn_div_fmas, "ddddb", "nc")
 BUILTIN(__builtin_amdgcn_div_fmasf, "ffffb", "nc")
 BUILTIN(__builtin_amdgcn_div_fixup, "dddd", "nc")
 BUILTIN(__builtin_amdgcn_div_fixupf, "ffff", "nc")
 BUILTIN(__builtin_amdgcn_trig_preop, "ddi", "nc")

 
 //===----------------------------------------------------------------------===//
 // VI+ only builtins.
 //===----------------------------------------------------------------------===//
 
 TARGET_BUILTIN(__builtin_amdgcn_s_memrealtime, "LUi", "n", "s-memrealtime")
 
 //===----------------------------------------------------------------------===//
+// Special builtins.
+//===----------------------------------------------------------------------===//
+BUILTIN(__builtin_amdgcn_read_exec, "LUi", "nc")
+
+//===----------------------------------------------------------------------===//
 // Legacy names with amdgpu prefix
 //===----------------------------------------------------------------------===//
 
 BUILTIN(__builtin_amdgpu_rsq, "dd", "nc")
 BUILTIN(__builtin_amdgpu_rsqf, "ff", "nc")
 BUILTIN(__builtin_amdgpu_ldexp, "ddi", "nc")
 BUILTIN(__builtin_amdgpu_ldexpf, "ffi", "nc")
 
 #undef BUILTIN
 #undef TARGET_BUILTIN

lib/CodeGen/CGBuiltin.cpp

 }
 
 // Generates the IR for the read/write special register builtin,
 // ValueType is the type of the value that is to be written or read,
 // RegisterType is the type of the register being written to or read from.
 static Value *EmitSpecialRegisterBuiltin(CodeGenFunction &CGF,
                                          const CallExpr *E,
                                          llvm::Type *RegisterType,
-                                         llvm::Type *ValueType, bool IsRead) {
+                                         llvm::Type *ValueType,
+                                         bool IsRead,
+                                         StringRef SysReg = "") {
   // write and register intrinsics only support 32 and 64 bit operations.
   assert((RegisterType->isIntegerTy(32) || RegisterType->isIntegerTy(64))
           && "Unsupported size for register.");
 
   CodeGen::CGBuilderTy &Builder = CGF.Builder;
   CodeGen::CodeGenModule &CGM = CGF.CGM;
   LLVMContext &Context = CGM.getLLVMContext();
 
+  if (SysReg.empty()) {
     const Expr *SysRegStrExpr = E->getArg(0)->IgnoreParenCasts();
-  StringRef SysReg = cast<StringLiteral>(SysRegStrExpr)->getString();
+    SysReg = cast<StringLiteral>(SysRegStrExpr)->getString();
+  }
 
   llvm::Metadata *Ops[] = { llvm::MDString::get(Context, SysReg) };
   llvm::MDNode *RegName = llvm::MDNode::get(Context, Ops);
   llvm::Value *Metadata = llvm::MetadataAsValue::get(Context, RegName);
 
   llvm::Type *Types[] = { RegisterType };
 
   bool MixedTypes = RegisterType->isIntegerTy(64) && ValueType->isIntegerTy(32);
Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
   }
   case AMDGPU::BI__builtin_amdgcn_fract:
   case AMDGPU::BI__builtin_amdgcn_fractf:
     return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_fract);
   case AMDGPU::BI__builtin_amdgcn_class:
   case AMDGPU::BI__builtin_amdgcn_classf:
     return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_class);
 
-    // Legacy amdgpu prefix
+  case AMDGPU::BI__builtin_amdgcn_read_exec: {
+    CallInst *CI = cast<CallInst>(
+      EmitSpecialRegisterBuiltin(*this, E, Int64Ty, Int64Ty, true, "exec"));
+    CI->setConvergent();
+    return CI;
+  }
   case AMDGPU::BI__builtin_amdgpu_rsq:
   case AMDGPU::BI__builtin_amdgpu_rsqf: {
     if (getTarget().getTriple().getArch() == Triple::amdgcn)
       return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_rsq);
     return emitUnaryBuiltin(*this, E, Intrinsic::r600_rsq);
   }
   case AMDGPU::BI__builtin_amdgpu_ldexp:
   case AMDGPU::BI__builtin_amdgpu_ldexpf: {

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]+]]
+void test_read_exec(global ulong* out) {
+  *out = __builtin_amdgcn_read_exec();
+}
+
+// CHECK: declare i64 @llvm.read_register.i64(metadata) #[[NOUNWIND_READONLY:[0-9]+]]
+
 // Legacy intrinsics with AMDGPU prefix
 
 // CHECK-LABEL: @test_legacy_rsq_f32
 // CHECK: call float @llvm.amdgcn.rsq.f32
 void test_legacy_rsq_f32(global float* out, float a)
 {
   *out = __builtin_amdgpu_rsqf(a);
 }

 }
 
 // CHECK-LABEL: @test_legacy_ldexp_f64
 // CHECK: call double @llvm.amdgcn.ldexp.f64
 void test_legacy_ldexp_f64(global double* out, double a, int b)
 {
   *out = __builtin_amdgpu_ldexp(a, b);
 }
+
+// CHECK-DAG: attributes #[[NOUNWIND_READONLY:[0-9]+]] = { nounwind readonly }
+// CHECK-DAG: attributes #[[READ_EXEC_ATTRS]] = { convergent }
+// CHECK: ![[EXEC]] = !{!"exec"}