[funcattrs] Infer writeonly argument attribute

[reames]

Summary

This change extends the current logic for inferring readonly and readnone argument attributes to also infer writeonly.

This change is deliberately minimal; there's a couple of areas for follow up.

1. I left out all call handling and thus any benefit from the SCC walk. When examining the test changes, I realized the existing code is imprecise, and am going to fix that in it's own revision before adding in the writeonly handling. (Mostly because updating the tests is hard when I, the human, can't figure out whether the result is correct.)
2. I left out handling for storing a value (as opposed to storing to a pointer). This should benefit readonly/readnone as well, and applies to a bunch of other instructions. Seemed worth having as a separate review.

[aeubanks]

some clang tests are failing

clang/test/CodeGen/SystemZ/systemz-inline-asm.c

double test_f64(double f, double g) {
   return f;
 // CHECK-LABEL: define{{.*}} double @test_f64(double %f, double %g)
 // CHECK: call double asm "adbr $0, $2", "=f,0,f"(double %f, double %g)
 }
 
 long double test_f128(long double f, long double g) {
   asm("axbr %0, %2" : "=f" (f) : "0" (f), "f" (g));
   return f;
-// CHECK: define{{.*}} void @test_f128(fp128* noalias nocapture sret(fp128) align 8 [[DEST:%.*]], fp128* nocapture readonly %0, fp128* nocapture readonly %1)
+// CHECK: define{{.*}} void @test_f128(fp128* noalias nocapture writeonly sret(fp128) align 8 [[DEST:%.*]], fp128* nocapture readonly %0, fp128* nocapture readonly %1)
 // CHECK: %f = load fp128, fp128* %0
 // CHECK: %g = load fp128, fp128* %1
 // CHECK: [[RESULT:%.*]] = tail call fp128 asm "axbr $0, $2", "=f,0,f"(fp128 %f, fp128 %g)
 // CHECK: store fp128 [[RESULT]], fp128* [[DEST]]
 }
 
 // Test that there are no tied physreg uses. TwoAddress pass cannot deal with them.
 int test_physregs(void) {

clang/test/CodeGen/aarch64-sve-acle-__ARM_FEATURE_SVE_VECTOR_OPERATORS.c

 // CHECK128-NEXT:  entry:
 // CHECK128-NEXT:    [[TMP0:%.*]] = call <vscale x 16 x i1> @llvm.aarch64.sve.ptrue.nxv16i1(i32 31)
 // CHECK128-NEXT:    [[CASTSCALABLESVE:%.*]] = call <vscale x 16 x i8> @llvm.experimental.vector.insert.nxv16i8.v16i8(<vscale x 16 x i8> undef, <16 x i8> [[X:%.*]], i64 0)
 // CHECK128-NEXT:    [[TMP1:%.*]] = call <vscale x 16 x i8> @llvm.aarch64.sve.asrd.nxv16i8(<vscale x 16 x i1> [[TMP0]], <vscale x 16 x i8> [[CASTSCALABLESVE]], i32 1)
 // CHECK128-NEXT:    [[CASTFIXEDSVE:%.*]] = call <16 x i8> @llvm.experimental.vector.extract.v16i8.nxv16i8(<vscale x 16 x i8> [[TMP1]], i64 0)
 // CHECK128-NEXT:    ret <16 x i8> [[CASTFIXEDSVE]]
 
 // CHECK-LABEL: define{{.*}} void @f2(
-// CHECK-SAME:   <[[#div(VBITS,8)]] x i8>* noalias nocapture sret(<[[#div(VBITS,8)]] x i8>) align 16 %agg.result, <[[#div(VBITS,8)]] x i8>* nocapture readonly %0)
+// CHECK-SAME:   <[[#div(VBITS,8)]] x i8>* noalias nocapture writeonly sret(<[[#div(VBITS,8)]] x i8>) align 16 %agg.result, <[[#div(VBITS,8)]] x i8>* nocapture readonly %0)
 // CHECK-NEXT: entry:
 // CHECK-NEXT:   [[X:%.*]] = load <[[#div(VBITS,8)]] x i8>, <[[#div(VBITS,8)]] x i8>* [[TMP0:%.*]], align 16, [[TBAA6:!tbaa !.*]]
 // CHECK-NEXT:   [[TMP1:%.*]] = call <vscale x 16 x i1> @llvm.aarch64.sve.ptrue.nxv16i1(i32 31)
 // CHECK-NEXT:   [[CASTSCALABLESVE:%.*]] = call <vscale x 16 x i8> @llvm.experimental.vector.insert.nxv16i8.v[[#div(VBITS,8)]]i8(<vscale x 16 x i8> undef, <[[#div(VBITS,8)]] x i8> [[X]], i64 0)
 // CHECK-NEXT:   [[TMP2:%.*]] = call <vscale x 16 x i8> @llvm.aarch64.sve.asrd.nxv16i8(<vscale x 16 x i1> [[TMP1]], <vscale x 16 x i8> [[CASTSCALABLESVE]], i32 1)
 // CHECK-NEXT:   [[CASTFIXEDSVE:%.*]] = call <[[#div(VBITS,8)]] x i8> @llvm.experimental.vector.extract.v[[#div(VBITS,8)]]i8.nxv16i8(<vscale x 16 x i8> [[TMP2]], i64 0)
 // CHECK-NEXT:   store <[[#div(VBITS,8)]] x i8> [[CASTFIXEDSVE]], <[[#div(VBITS,8)]] x i8>* [[AGG_RESULT:%.*]], align 16, [[TBAA6]]
 // CHECK-NEXT:   ret void

clang/test/CodeGen/arm-vfp16-arguments2.cpp

struct S4 : B1 {
   B2 M[1];
 };
 
 // S5 does not contain any FP16 vectors
 struct S5 : B1 {
   B1 M[1];
 };
 
-// CHECK-SOFT: define{{.*}} void @_Z2f12S1(%struct.S1* noalias nocapture sret(%struct.S1) align 8 %agg.result, [2 x i64] %s1.coerce)
+// CHECK-SOFT: define{{.*}} void @_Z2f12S1(%struct.S1* noalias nocapture writeonly sret(%struct.S1) align 8 %agg.result, [2 x i64] %s1.coerce)
 // CHECK-HARD: define{{.*}} arm_aapcs_vfpcc [2 x <2 x i32>] @_Z2f12S1([2 x <2 x i32>] returned %s1.coerce)
 // CHECK-FULL: define{{.*}} arm_aapcs_vfpcc %struct.S1 @_Z2f12S1(%struct.S1 returned %s1.coerce)
 struct S1 f1(struct S1 s1) { return s1; }
 
-// CHECK-SOFT: define{{.*}} void @_Z2f22S2(%struct.S2* noalias nocapture sret(%struct.S2) align 8 %agg.result, [4 x i32] %s2.coerce)
+// CHECK-SOFT: define{{.*}} void @_Z2f22S2(%struct.S2* noalias nocapture writeonly sret(%struct.S2) align 8 %agg.result, [4 x i32] %s2.coerce)
 // CHECK-HARD: define{{.*}} arm_aapcs_vfpcc [2 x <2 x i32>] @_Z2f22S2([2 x <2 x i32>] returned %s2.coerce)
 // CHECK-FULL: define{{.*}} arm_aapcs_vfpcc %struct.S2 @_Z2f22S2(%struct.S2 returned %s2.coerce)
 struct S2 f2(struct S2 s2) { return s2; }
 
-// CHECK-SOFT: define{{.*}} void @_Z2f32S3(%struct.S3* noalias nocapture sret(%struct.S3) align 8 %agg.result, [2 x i64] %s3.coerce)
+// CHECK-SOFT: define{{.*}} void @_Z2f32S3(%struct.S3* noalias nocapture writeonly sret(%struct.S3) align 8 %agg.result, [2 x i64] %s3.coerce)
 // CHECK-HARD: define{{.*}} arm_aapcs_vfpcc [2 x <2 x i32>] @_Z2f32S3([2 x <2 x i32>] returned %s3.coerce)
 // CHECK-FULL: define{{.*}} arm_aapcs_vfpcc %struct.S3 @_Z2f32S3(%struct.S3 returned %s3.coerce)
 struct S3 f3(struct S3 s3) { return s3; }
 
-// CHECK-SOFT: define{{.*}} void @_Z2f42S4(%struct.S4* noalias nocapture sret(%struct.S4) align 8 %agg.result, [2 x i64] %s4.coerce)
+// CHECK-SOFT: define{{.*}} void @_Z2f42S4(%struct.S4* noalias nocapture writeonly sret(%struct.S4) align 8 %agg.result, [2 x i64] %s4.coerce)
 // CHECK-HARD: define{{.*}} arm_aapcs_vfpcc [2 x <2 x i32>] @_Z2f42S4([2 x <2 x i32>] returned %s4.coerce)
 // CHECK-FULL: define{{.*}} arm_aapcs_vfpcc %struct.S4 @_Z2f42S4(%struct.S4 returned %s4.coerce)
 struct S4 f4(struct S4 s4) { return s4; }
 
-// CHECK-SOFT: define{{.*}} void @_Z2f52S5(%struct.S5* noalias nocapture sret(%struct.S5) align 8 %agg.result, [2 x i64] %s5.coerce)
+// CHECK-SOFT: define{{.*}} void @_Z2f52S5(%struct.S5* noalias nocapture writeonly sret(%struct.S5) align 8 %agg.result, [2 x i64] %s5.coerce)
 // CHECK-HARD: define{{.*}} arm_aapcs_vfpcc %struct.S5 @_Z2f52S5(%struct.S5 returned %s5.coerce)
 // CHECK-FULL: define{{.*}} arm_aapcs_vfpcc %struct.S5 @_Z2f52S5(%struct.S5 returned %s5.coerce)
 struct S5 f5(struct S5 s5) { return s5; }

clang/test/CodeGen/mips-vector-return.c

 // RUN: %clang -target mipsel-unknown-linux -O3 -S -o - -emit-llvm %s | FileCheck %s -check-prefix=O32
 // RUN: %clang -target mips64el-unknown-linux -O3 -S -mabi=n64 -o - -emit-llvm %s | FileCheck %s -check-prefix=N64
 
 // vectors larger than 16-bytes are returned via the hidden pointer argument. 
 // N64/N32 returns vectors whose size is equal to or smaller than 16-bytes in
 // integer registers. 
 typedef float  v4sf __attribute__ ((__vector_size__ (16)));
 typedef double v4df __attribute__ ((__vector_size__ (32)));
 typedef int v4i32 __attribute__ ((__vector_size__ (16)));
 
-// O32-LABEL: define dso_local void @test_v4sf(<4 x float>* noalias nocapture sret
+// O32-LABEL: define dso_local void @test_v4sf(<4 x float>* noalias nocapture writeonly sret
 // N64: define inreg { i64, i64 } @test_v4sf
 v4sf test_v4sf(float a) {
   return (v4sf){0.0f, a, 0.0f, 0.0f};
 }
 
-// O32-LABEL: define dso_local void @test_v4df(<4 x double>* noalias nocapture sret
-// N64-LABEL: define void @test_v4df(<4 x double>* noalias nocapture sret
+// O32-LABEL: define dso_local void @test_v4df(<4 x double>* noalias nocapture writeonly sret
+// N64-LABEL: define void @test_v4df(<4 x double>* noalias nocapture writeonly sret
 v4df test_v4df(double a) {
   return (v4df){0.0, a, 0.0, 0.0};
 }
 
 // O32 returns integer vectors whose size is equal to or smaller than 16-bytes
 // in integer registers.
 //
 // O32: define dso_local inreg { i32, i32, i32, i32 } @test_v4i32
 // N64: define inreg { i64, i64 } @test_v4i32
 v4i32 test_v4i32(int a) {
   return (v4i32){0, a, 0, 0};
 }
 

clang/test/CodeGen/mips64-nontrivial-return.cpp

 // RUN: %clang -target mips64el-unknown-linux -O3 -S -mabi=n64 -o - -emit-llvm %s | FileCheck %s
 
 class B {
 public:
   virtual ~B() {}
 };
 
 class D : public B {
 };
 
 extern D gd0;
 
-// CHECK: _Z4foo1v(%class.D* noalias nocapture sret
+// CHECK: _Z4foo1v(%class.D* noalias nocapture writeonly sret
 
 D foo1(void) {
   return gd0;
 }

clang/test/CodeGen/ms-mixed-ptr-sizes.c

 // RUN: %clang_cc1 -triple x86_64-windows-msvc -fms-extensions -emit-llvm -O2 < %s | FileCheck %s --check-prefixes=X64,ALL
 // RUN: %clang_cc1 -triple i386-pc-win32 -fms-extensions -emit-llvm -O2 < %s | FileCheck %s --check-prefixes=X86,ALL
 
 struct Foo {
   int * __ptr32 p32;
   int * __ptr64 p64;
 };
 void use_foo(struct Foo *f);
 void test_sign_ext(struct Foo *f, int * __ptr32 __sptr i) {
-// X64-LABEL: define dso_local void @test_sign_ext({{.*}}i32 addrspace(270)* %i)
-// X86-LABEL: define dso_local void @test_sign_ext(%struct.Foo* %f, i32* %i)
+// X64-LABEL: define dso_local void @test_sign_ext({{.*}}i32 addrspace(270)* writeonly %i)
+// X86-LABEL: define dso_local void @test_sign_ext(%struct.Foo* %f, i32* writeonly %i)
 // X64: %{{.+}} = addrspacecast i32 addrspace(270)* %i to i32*
 // X86: %{{.+}} = addrspacecast i32* %i to i32 addrspace(272)*
   f->p64 = i;
   use_foo(f);
 }
 void test_zero_ext(struct Foo *f, int * __ptr32 __uptr i) {
-// X64-LABEL: define dso_local void @test_zero_ext({{.*}}i32 addrspace(271)* %i)
-// X86-LABEL: define dso_local void @test_zero_ext({{.*}}i32 addrspace(271)* %i)
+// X64-LABEL: define dso_local void @test_zero_ext({{.*}}i32 addrspace(271)* writeonly %i)
+// X86-LABEL: define dso_local void @test_zero_ext({{.*}}i32 addrspace(271)* writeonly %i)
 // X64: %{{.+}} = addrspacecast i32 addrspace(271)* %i to i32*
 // X86: %{{.+}} = addrspacecast i32 addrspace(271)* %i to i32 addrspace(272)*
   f->p64 = i;
   use_foo(f);
 }
 void test_trunc(struct Foo *f, int * __ptr64 i) {
-// X64-LABEL: define dso_local void @test_trunc(%struct.Foo* %f, i32* %i)
-// X86-LABEL: define dso_local void @test_trunc({{.*}}i32 addrspace(272)* %i)
+// X64-LABEL: define dso_local void @test_trunc(%struct.Foo* %f, i32* writeonly %i)
+// X86-LABEL: define dso_local void @test_trunc({{.*}}i32 addrspace(272)* writeonly %i)
 // X64: %{{.+}} = addrspacecast i32* %i to i32 addrspace(270)*
 // X86: %{{.+}} = addrspacecast i32 addrspace(272)* %i to i32*
   f->p32 = i;
   use_foo(f);
 }
 void test_noop(struct Foo *f, int * __ptr32 i) {
-// X64-LABEL: define dso_local void @test_noop({{.*}}i32 addrspace(270)* %i)
-// X86-LABEL: define dso_local void @test_noop({{.*}}i32* %i)
+// X64-LABEL: define dso_local void @test_noop({{.*}}i32 addrspace(270)* writeonly %i)
+// X86-LABEL: define dso_local void @test_noop({{.*}}i32* writeonly %i)
 // X64-NOT: addrspacecast
 // X86-NOT: addrspacecast
   f->p32 = i;
   use_foo(f);
 }
 
 void test_other(struct Foo *f, __attribute__((address_space(10))) int *i) {
-// X64-LABEL: define dso_local void @test_other({{.*}}i32 addrspace(10)* %i)
-// X86-LABEL: define dso_local void @test_other({{.*}}i32 addrspace(10)* %i)
+// X64-LABEL: define dso_local void @test_other({{.*}}i32 addrspace(10)* writeonly %i)
+// X86-LABEL: define dso_local void @test_other({{.*}}i32 addrspace(10)* writeonly %i)
 // X64: %{{.+}} = addrspacecast i32 addrspace(10)* %i to i32 addrspace(270)*
 // X86: %{{.+}} = addrspacecast i32 addrspace(10)* %i to i32*
   f->p32 = (int * __ptr32)i;
   use_foo(f);
 }
 
 int test_compare1(int *__ptr32 __uptr i, int *__ptr64 j) {
   // ALL-LABEL: define dso_local i32 @test_compare1

clang/test/CodeGenOpenCL/amdgpu-abi-struct-coerce.cl

 
 // CHECK: define{{.*}} %struct.different_size_type_pair @func_different_size_type_pair_ret()
 different_size_type_pair func_different_size_type_pair_ret()
 {
   different_size_type_pair s = { 0 };
   return s;
 }
 
-// CHECK: define{{.*}} void @func_flexible_array_ret(%struct.flexible_array addrspace(5)* noalias nocapture sret(%struct.flexible_array) align 4 %agg.result)
+// CHECK: define{{.*}} void @func_flexible_array_ret(%struct.flexible_array addrspace(5)* noalias nocapture writeonly sret(%struct.flexible_array) align 4 %agg.result)
 flexible_array func_flexible_array_ret()
 {
   flexible_array s = { 0 };
   return s;
 }
 
 // CHECK: define{{.*}} void @func_reg_state_lo(<4 x i32> %arg0, <4 x i32> %arg1, <4 x i32> %arg2, i32 %arg3, i32 %s.coerce0, float %s.coerce1, i32 %s.coerce2)
 void func_reg_state_lo(int4 arg0, int4 arg1, int4 arg2, int arg3, struct_arg_t s) { }

clang/test/CodeGenOpenCL/amdgpu-call-kernel.cl

 // REQUIRES: amdgpu-registered-target
 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s
-// CHECK: define{{.*}} amdgpu_kernel void @test_call_kernel(i32 addrspace(1)* nocapture %out)
+// CHECK: define{{.*}} amdgpu_kernel void @test_call_kernel(i32 addrspace(1)* nocapture writeonly %out)
 // CHECK: store i32 4, i32 addrspace(1)* %out, align 4
 
 kernel void test_kernel(global int *out)
 {
   out[0] = 4;
 }
 
 __kernel void test_call_kernel(__global int *out)
 {
   test_kernel(out);
 }

clang/test/CodeGenOpenCL/kernels-have-spir-cc-by-default.cl

typedef struct test_struct {
     short elementG;
     double elementH;
 } test_struct;
 
 kernel void test_single(int_single input, global int* output) {
 // CHECK: spir_kernel
 // AMDGCN: define{{.*}} amdgpu_kernel void @test_single
 // CHECK: struct.int_single* nocapture {{.*}} byval(%struct.int_single)
-// CHECK: i32* nocapture %output
+// CHECK: i32* nocapture writeonly %output
  output[0] = input.a;
 }
 
 kernel void test_pair(int_pair input, global int* output) {
 // CHECK: spir_kernel
 // AMDGCN: define{{.*}} amdgpu_kernel void @test_pair
 // CHECK: struct.int_pair* nocapture {{.*}} byval(%struct.int_pair)
-// CHECK: i32* nocapture %output
+// CHECK: i32* nocapture writeonly %output
  output[0] = (int)input.a;
  output[1] = (int)input.b;
 }
 
 kernel void test_kernel(test_struct input, global int* output) {
 // CHECK: spir_kernel
 // AMDGCN: define{{.*}} amdgpu_kernel void @test_kernel
 // CHECK: struct.test_struct* nocapture {{.*}} byval(%struct.test_struct)
-// CHECK: i32* nocapture %output
+// CHECK: i32* nocapture writeonly %output
  output[0] = input.elementA;
  output[1] = input.elementB;
  output[2] = (int)input.elementC;
  output[3] = (int)input.elementD;
  output[4] = (int)input.elementE;
  output[5] = (int)input.elementF;
  output[6] = (int)input.elementG;
  output[7] = (int)input.elementH;
 };
 
 void test_function(int_pair input, global int* output) {
 // CHECK-NOT: spir_kernel
 // AMDGCN-NOT: define{{.*}} amdgpu_kernel void @test_function
-// CHECK: i64 %input.coerce0, i64 %input.coerce1, i32* nocapture %output
+// CHECK: i64 %input.coerce0, i64 %input.coerce1, i32* nocapture writeonly %output
  output[0] = (int)input.a;
  output[1] = (int)input.b;
 }

llvm/lib/Transforms/IPO/FunctionAttrs.cpp

 
 STATISTIC(NumReadNone, "Number of functions marked readnone");
 STATISTIC(NumReadOnly, "Number of functions marked readonly");
 STATISTIC(NumWriteOnly, "Number of functions marked writeonly");
 STATISTIC(NumNoCapture, "Number of arguments marked nocapture");
 STATISTIC(NumReturned, "Number of arguments marked returned");
 STATISTIC(NumReadNoneArg, "Number of arguments marked readnone");
 STATISTIC(NumReadOnlyArg, "Number of arguments marked readonly");
+STATISTIC(NumWriteOnlyArg, "Number of arguments marked writeonly");
 STATISTIC(NumNoAlias, "Number of function returns marked noalias");
 STATISTIC(NumNonNullReturn, "Number of function returns marked nonnull");
 STATISTIC(NumNoRecurse, "Number of functions marked as norecurse");
 STATISTIC(NumNoUnwind, "Number of functions marked as nounwind");
 STATISTIC(NumNoFree, "Number of functions marked as nofree");
 STATISTIC(NumWillReturn, "Number of functions marked as willreturn");
 STATISTIC(NumNoSync, "Number of functions marked as nosync");
 
struct GraphTraits<ArgumentGraph *> : public GraphTraits<ArgumentGraphNode *> {
 
   static ChildIteratorType nodes_end(ArgumentGraph *AG) { return AG->end(); }
 };
 
 } // end namespace llvm
 
 /// Returns Attribute::None, Attribute::ReadOnly or Attribute::ReadNone.
 static Attribute::AttrKind
-determinePointerReadAttrs(Argument *A,
+determinePointerAccessAttrs(Argument *A,
                             const SmallPtrSet<Argument *, 8> &SCCNodes) {
   SmallVector<Use *, 32> Worklist;
   SmallPtrSet<Use *, 32> Visited;
 
   // inalloca arguments are always clobbered by the call.
   if (A->hasInAllocaAttr() || A->hasPreallocatedAttr())
     return Attribute::None;
 
   bool IsRead = false;
-  // We don't need to track IsWritten. If A is written to, return immediately.
+  bool IsWrite = false;
 
   for (Use &U : A->uses()) {
     Visited.insert(&U);
     Worklist.push_back(&U);
   }
 
   while (!Worklist.empty()) {
+    if (IsWrite && IsRead)
+      // No point in searching further..
+      return Attribute::None;
+
     Use *U = Worklist.pop_back_val();
     Instruction *I = cast<Instruction>(U->getUser());
 
     switch (I->getOpcode()) {
     case Instruction::BitCast:
     case Instruction::GetElementPtr:
     case Instruction::PHI:
     case Instruction::Select:
case Instruction::Load:
       // A volatile load has side effects beyond what readonly can be relied
       // upon.
       if (cast<LoadInst>(I)->isVolatile())
         return Attribute::None;
 
       IsRead = true;
       break;
 
+    case Instruction::Store:
+      // A volatile store has side effects beyond what writeonly can be relied
+      // upon.
+      if (cast<StoreInst>(I)->isVolatile())
+        return Attribute::None;
+
+      IsWrite = true;
+      break;
+
     case Instruction::ICmp:
     case Instruction::Ret:
       break;
 
     default:
       return Attribute::None;
     }
   }
 
-  return IsRead ? Attribute::ReadOnly : Attribute::ReadNone;
+  if (IsWrite && IsRead)
+    return Attribute::None;
+  else if (IsRead)
+    return Attribute::ReadOnly;
+  else if (IsWrite)
+    return Attribute::WriteOnly;
+  else
+    return Attribute::ReadNone;
 }
 
 /// Deduce returned attributes for the SCC.
 static void addArgumentReturnedAttrs(const SCCNodeSet &SCCNodes,
                                      SmallSet<Function *, 8> &Changed) {
   // Check each function in turn, determining if an argument is always returned.
   for (Function *F : SCCNodes) {
     // We can infer and propagate function attributes only when we know that the
for (Instruction &I : Entry) {
     }
     if (!isGuaranteedToTransferExecutionToSuccessor(&I))
       break;
   }
 
   return Changed;
 }
 
-static bool addReadAttr(Argument *A, Attribute::AttrKind R) {
-  assert((R == Attribute::ReadOnly || R == Attribute::ReadNone)
-         && "Must be a Read attribute.");
+static bool addAccessAttr(Argument *A, Attribute::AttrKind R) {
+  assert((R == Attribute::ReadOnly || R == Attribute::ReadNone ||
+          R == Attribute::WriteOnly)
+         && "Must be an access attribute.");
   assert(A && "Argument must not be null.");
 
   // If the argument already has the attribute, nothing needs to be done.
   if (A->hasAttribute(R))
       return false;
 
   // Otherwise, remove potentially conflicting attribute, add the new one,
   // and update statistics.
   A->removeAttr(Attribute::WriteOnly);
   A->removeAttr(Attribute::ReadOnly);
   A->removeAttr(Attribute::ReadNone);
   A->addAttr(R);
-  R == Attribute::ReadOnly ? ++NumReadOnlyArg : ++NumReadNoneArg;
+  if (R == Attribute::ReadOnly)
+    ++NumReadOnlyArg;
+  else if (R == Attribute::WriteOnly)
+    ++NumWriteOnlyArg;
+  else
+    ++NumReadNoneArg;
   return true;
 }
 
 /// Deduce nocapture attributes for the SCC.
 static void addArgumentAttrs(const SCCNodeSet &SCCNodes,
                              SmallSet<Function *, 8> &Changed) {
   ArgumentGraph AG;
 
for (Function::arg_iterator A = F->arg_begin(), E = F->arg_end(); A != E;
               if (Use != &*A)
                 HasNonLocalUses = true;
             }
           }
         }
         // Otherwise, it's captured. Don't bother doing SCC analysis on it.
       }
       if (!HasNonLocalUses && !A->onlyReadsMemory()) {
-        // Can we determine that it's readonly/readnone without doing an SCC?
-        // Note that we don't allow any calls at all here, or else our result
-        // will be dependent on the iteration order through the functions in the
-        // SCC.
+        // Can we determine that it's readonly/readnone/writeonly without doing
+        // an SCC? Note that we don't allow any calls at all here, or else our
+        // result will be dependent on the iteration order through the
+        // functions in the SCC.
         SmallPtrSet<Argument *, 8> Self;
         Self.insert(&*A);
-        Attribute::AttrKind R = determinePointerReadAttrs(&*A, Self);
+        Attribute::AttrKind R = determinePointerAccessAttrs(&*A, Self);
         if (R != Attribute::None)
-          if (addReadAttr(A, R))
+          if (addAccessAttr(A, R))
             Changed.insert(F);
       }
     }
   }
 
   // The graph we've collected is partial because we stopped scanning for
   // argument uses once we solved the argument trivially. These partial nodes
   // show up as ArgumentGraphNode objects with an empty Uses list, and for
for (scc_iterator<ArgumentGraph *> I = scc_begin(&AG); !I.isAtEnd(); ++I) {
 
     for (unsigned i = 0, e = ArgumentSCC.size(); i != e; ++i) {
       Argument *A = ArgumentSCC[i]->Definition;
       A->addAttr(Attribute::NoCapture);
       ++NumNoCapture;
       Changed.insert(A->getParent());
     }
 
-    // We also want to compute readonly/readnone. With a small number of false
-    // negatives, we can assume that any pointer which is captured isn't going
-    // to be provably readonly or readnone, since by definition we can't
-    // analyze all uses of a captured pointer.
+    // We also want to compute readonly/readnone/writeonly. With a small number
+    // of false negatives, we can assume that any pointer which is captured
+    // isn't going to be provably readonly or readnone, since by definition
+    // we can't analyze all uses of a captured pointer.
     //
     // The false negatives happen when the pointer is captured by a function
     // that promises readonly/readnone behaviour on the pointer, then the
     // pointer's lifetime ends before anything that writes to arbitrary memory.
     // Also, a readonly/readnone pointer may be returned, but returning a
     // pointer is capturing it.
 
-    Attribute::AttrKind ReadAttr = Attribute::ReadNone;
-    for (unsigned i = 0, e = ArgumentSCC.size(); i != e; ++i) {
+    auto meetAccessAttr = [](Attribute::AttrKind A, Attribute::AttrKind B) {
+      if (A == B)
+        return A;
+      if (A == Attribute::ReadNone)
+        return B;
+      if (B == Attribute::ReadNone)
+        return A;
+      return Attribute::None;
+    };
+
+    Attribute::AttrKind AccessAttr = Attribute::ReadNone;
+    for (unsigned i = 0, e = ArgumentSCC.size();
+         i != e && AccessAttr != Attribute::None; ++i) {
       Argument *A = ArgumentSCC[i]->Definition;
-      Attribute::AttrKind K = determinePointerReadAttrs(A, ArgumentSCCNodes);
-      if (K == Attribute::ReadNone)
-        continue;
-      if (K == Attribute::ReadOnly) {
-        ReadAttr = Attribute::ReadOnly;
-        continue;
-      }
-      ReadAttr = K;
-      break;
+      Attribute::AttrKind K = determinePointerAccessAttrs(A, ArgumentSCCNodes);
+      AccessAttr = meetAccessAttr(AccessAttr, K);
     }
 
-    if (ReadAttr != Attribute::None) {
+    if (AccessAttr != Attribute::None) {
       for (unsigned i = 0, e = ArgumentSCC.size(); i != e; ++i) {
         Argument *A = ArgumentSCC[i]->Definition;
-        if (addReadAttr(A, ReadAttr))
+        if (addAccessAttr(A, AccessAttr))
           Changed.insert(A->getParent());
       }
     }
   }
 }
 
 /// Tests whether a function is "malloc-like".
 ///

llvm/test/Analysis/TypeBasedAliasAnalysis/functionattrs.ll

 ; code path which isn't ever executed.
 
 ; CHECK: define void @test0_yes(i32* nocapture %p) #0 {
 define void @test0_yes(i32* %p) nounwind {
   store i32 0, i32* %p, !tbaa !1
   ret void
 }
 
-; CHECK: define void @test0_no(i32* nocapture %p) #1 {
+; CHECK: define void @test0_no(i32* nocapture writeonly %p) #1 {
 define void @test0_no(i32* %p) nounwind {
   store i32 0, i32* %p, !tbaa !2
   ret void
 }
 
 ; Add the readonly attribute, since there's just a call to a function which 
 ; TBAA says doesn't modify any memory.
 

llvm/test/Feature/OperandBundles/pr26510.ll

 ; RUN: opt -S -globals-aa -function-attrs < %s | FileCheck %s
 ; RUN: opt -S -O3 < %s | FileCheck %s
 
 ; Apart from checking for the direct cause of the bug, we also check
 ; if any problematic aliasing rules have accidentally snuck into -O3.
 ;
 ; Since the "abc" operand bundle is not a special operand bundle that
 ; LLVM knows about, all of the stores and loads in @test below have to
 ; stay.
 
 declare void @foo() readnone
 
-; CHECK-LABEL: define i8* @test(i8* %p)
+; CHECK-LABEL: define i8* @test(i8* writeonly %p)
 ; CHECK:   %a = alloca i8*, align 8
 ; CHECK:   store i8* %p, i8** %a, align 8
 ; CHECK:   call void @foo() [ "abc"(i8** %a) ]
 ; CHECK:   %reload = load i8*, i8** %a, align 8
 ; CHECK:   ret i8* %reload
 ; CHECK: }
 
 define i8* @test(i8* %p) {
   %a = alloca i8*, align 8
   store i8* %p, i8** %a, align 8
   call void @foo() ["abc" (i8** %a)]
   %reload = load i8*, i8** %a, align 8
   ret i8* %reload
 }

llvm/test/Transforms/Coroutines/coro-async.ll

entry:
   %continuation_actor_arg = extractvalue {i8*, i8*, i8*} %res.2, 1
   %actor.2 =  bitcast i8* %continuation_actor_arg to %async.actor*
 
   tail call swiftcc void @asyncReturn(i8* %async.ctxt, %async.task* %task.2, %async.actor* %actor.2)
   call i1 @llvm.coro.end(i8* %hdl, i1 0)
   unreachable
 }
 
-; CHECK-LABEL: define swiftcc void @my_async_function2(%async.task* %task, %async.actor* %actor, i8* %async.ctxt)
+; CHECK-LABEL: define swiftcc void @my_async_function2(%async.task* %task, %async.actor* %actor, i8* writeonly %async.ctxt)
 ; CHECK-SAME: #[[FRAMEPOINTER:[0-9]+]]
 ; CHECK-SAME: !dbg ![[SP3:[0-9]+]]
 ; CHECK: store i8* %async.ctxt,
 ; CHECK: store %async.actor* %actor,
 ; CHECK: store %async.task* %task,
 ; CHECK: [[CALLEE_CTXT:%.*]] =  tail call i8* @llvm.coro.async.context.alloc(
 ; CHECK: store i8* [[CALLEE_CTXT]],
 ; CHECK: store i8* bitcast (void (i8*, i8*, i8*)* @my_async_function2.resume.0 to i8*),
 ; CHECK: store i8* %async.ctxt,
 ; CHECK: tail call swiftcc void @asyncSuspend(i8* [[CALLEE_CTXT]], %async.task* %task, %async.actor* %actor)
 ; CHECK: ret void
 
-; CHECK-LABEL: define internal swiftcc void @my_async_function2.resume.0(i8* %0, i8* nocapture readnone %1, i8* nocapture readonly %2)
+; CHECK-LABEL: define internal swiftcc void @my_async_function2.resume.0(i8* writeonly %0, i8* nocapture readnone %1, i8* nocapture readonly %2)
 ; CHECK-SAME: #[[FRAMEPOINTER]]
 ; CHECK-SAME: !dbg ![[SP4:[0-9]+]]
 ; CHECK: [[CALLEE_CTXT_ADDR:%.*]] = bitcast i8* %2 to i8**
 ; CHECK: [[CALLEE_CTXT:%.*]] = load i8*, i8** [[CALLEE_CTXT_ADDR]]
 ; CHECK: [[CALLEE_CTXT_SPILL_ADDR:%.*]] = getelementptr inbounds i8, i8* [[CALLEE_CTXT]], i64 152
 ; CHECK: [[CALLEE_CTXT_SPILL_ADDR2:%.*]] = bitcast i8* [[CALLEE_CTXT_SPILL_ADDR]] to i8**
 ; CHECK: store i8* bitcast (void (i8*, i8*, i8*)* @my_async_function2.resume.1 to i8*),
 ; CHECK: [[CALLLE_CTXT_RELOAD:%.*]] = load i8*, i8** [[CALLEE_CTXT_SPILL_ADDR2]]

llvm/test/Transforms/FunctionAttrs/2009-01-02-LocalStores.ll

 ; RUN: opt < %s -function-attrs -S | FileCheck %s
 ; RUN: opt < %s -passes=function-attrs -S | FileCheck %s
 
 ; CHECK: define i32* @a(i32** nocapture readonly %p)
 define i32* @a(i32** %p) {
   %tmp = load i32*, i32** %p
   ret i32* %tmp
 }
 
-; CHECK: define i32* @b(i32* %q)
+; CHECK: define i32* @b(i32* writeonly %q)
 define i32* @b(i32 *%q) {
   %mem = alloca i32*
   store i32* %q, i32** %mem
   %tmp = call i32* @a(i32** %mem)
   ret i32* %tmp
 }
 
 ; CHECK: define i32* @c(i32* readnone returned %r)
 @g = global i32 0
 define i32* @c(i32 *%r) {
   %a = icmp eq i32* %r, null
   store i32 1, i32* @g
   ret i32* %r
 }

llvm/test/Transforms/FunctionAttrs/nocapture.ll

 ; RUN: opt -function-attrs -S < %s | FileCheck %s --check-prefixes=FNATTR
 ; RUN: opt -passes=function-attrs -S < %s | FileCheck %s --check-prefixes=FNATTR
 
 @g = global i32* null   ; <i32**> [#uses=1]
 
 ; FNATTR: define i32* @c1(i32* readnone returned %q)
 define i32* @c1(i32* %q) {
   ret i32* %q
 }
 
-; FNATTR: define void @c2(i32* %q)
+; FNATTR: define void @c2(i32* writeonly %q)
 ; It would also be acceptable to mark %q as readnone. Update @c3 too.
 define void @c2(i32* %q) {
   store i32* %q, i32** @g
   ret void
 }
 
 ; FNATTR: define void @c3(i32* %q)
 define void @c3(i32* %q) {

 @g2 = global i8* null
 ; FNATTR: define void @captureLaunder(i8* %p)
 define void @captureLaunder(i8* %p) {
   %b = call i8* @llvm.launder.invariant.group.p0i8(i8* %p)
   store i8* %b, i8** @g2
   ret void
 }
 
-; FNATTR: @nocaptureStrip(i8* nocapture %p)
+; FNATTR: @nocaptureStrip(i8* nocapture writeonly %p)
 define void @nocaptureStrip(i8* %p) {
 entry:
   %b = call i8* @llvm.strip.invariant.group.p0i8(i8* %p)
   store i8 42, i8* %b
   ret void
 }
 
 @g3 = global i8* null
-; FNATTR: define void @captureStrip(i8* %p)
+; FNATTR: define void @captureStrip(i8* writeonly %p)
 define void @captureStrip(i8* %p) {
   %b = call i8* @llvm.strip.invariant.group.p0i8(i8* %p)
   store i8* %b, i8** @g3
   ret void
 }
 
 ; FNATTR: define i1 @captureICmp(i32* readnone %x)
 define i1 @captureICmp(i32* %x) {

llvm/test/Transforms/FunctionAttrs/readattrs.ll

 declare void @test4_1(i8* nocapture) readonly
 
 ; CHECK: define void @test4_2(i8* nocapture readonly %p)
 define void @test4_2(i8* %p) {
   call void @test4_1(i8* %p)
   ret void
 }
 
-; CHECK: define void @test5(i8** nocapture %p, i8* %q)
+; CHECK: define void @test5(i8** nocapture writeonly %p, i8* writeonly %q)
 ; Missed optz'n: we could make %q readnone, but don't break test6!
 define void @test5(i8** %p, i8* %q) {
   store i8* %q, i8** %p
   ret void
 }
 
 declare void @test6_1()
-; CHECK: define void @test6_2(i8** nocapture %p, i8* %q)
+; CHECK: define void @test6_2(i8** nocapture writeonly %p, i8* writeonly %q)
 ; This is not a missed optz'n.
 define void @test6_2(i8** %p, i8* %q) {
   store i8* %q, i8** %p
   call void @test6_1()
   ret void
 }
 
 ; CHECK: define void @test7_1(i32* nocapture inalloca(i32) %a)

 }
 
 ; CHECK: define i32* @test8_1(i32* readnone returned %p)
 define i32* @test8_1(i32* %p) {
 entry:
   ret i32* %p
 }
 
-; CHECK: define void @test8_2(i32* %p)
+; CHECK: define void @test8_2(i32* writeonly %p)
 define void @test8_2(i32* %p) {
 entry:
   %call = call i32* @test8_1(i32* %p)
   store i32 10, i32* %call, align 4
   ret void
 }
 
 ; CHECK: declare void @llvm.masked.scatter

llvm/test/Transforms/FunctionAttrs/writeonly.ll

 
 ; CHECK: define void @nouses-argworn-funwo(i32* nocapture readnone %.aaa) #2 {
 define void @nouses-argworn-funwo(i32* writeonly %.aaa) {
 nouses-argworn-funwo_entry:
   store i8 0, i8* getelementptr inbounds (%_type_of_d-ccc, %_type_of_d-ccc* @d-ccc, i32 0, i32 3)
   ret void
 }
 
-; CHECK: define void @test_store(i8* nocapture %p)
+; CHECK: define void @test_store(i8* nocapture writeonly %p)
 define void @test_store(i8* %p) {
   store i8 0, i8* %p
   ret void
 }
 
-; CHECK: define void @test_addressing(i8* nocapture %p)
+; CHECK: define void @test_addressing(i8* nocapture writeonly %p)
 define void @test_addressing(i8* %p) {
   %gep = getelementptr i8, i8* %p, i64 8
   %bitcast = bitcast i8* %gep to i32*
   store i32 0, i32* %bitcast
   ret void
 }
 
 ; CHECK: define void @test_readwrite(i8* nocapture %p)