[PATCH] D11311: Replace hardcoded threshold with an option when analyze read attributes

[Xuetian Weng]

wengxt created this revision.
wengxt added reviewers: nlewycky, jingyue.
wengxt added a subscriber: llvm-commits.

In NVPTX, a specialized instruction ld.global.nc can be used to load memory
with non-coherent texture cache. We notice that in SHOC [1] benchmark, some
function arguments are not marked with readonly because FunctionAttrs reaches
a hardcoded threshold when analysis uses.

This patch changes the hardcoded threshold into an option to enable us to
change the threshold at runtime.

[1] https://github.com/vetter/shoc

http://reviews.llvm.org/D11311

Files:
  lib/Transforms/IPO/FunctionAttrs.cpp
  test/Transforms/FunctionAttrs/readattrs-threshold.ll

Index: test/Transforms/FunctionAttrs/readattrs-threshold.ll
===================================================================
--- /dev/null
+++ test/Transforms/FunctionAttrs/readattrs-threshold.ll
@@ -0,0 +1,13 @@
+; RUN: opt < %s -functionattrs -S | FileCheck %s --check-prefix=CHECK
+; RUN: opt < %s -functionattrs -readattrs-analysis-uses-threshold=1 -S | FileCheck %s --check-prefix=CHECK-NO-READATTR
+
+; %p has two uses
+; CHECK: define i32 @test(i32* nocapture readonly %p)
+; CHECK-NO-READATTR: define i32 @test(i32* nocapture %p)
+define i32 @test(i32* %p) {
+  %1 = load i32, i32* %p
+  %2 = getelementptr i32, i32* %p, i32 1
+  %3 = load i32, i32* %2
+  %4 = add i32 %1, %3
+  ret i32 %4
+}
Index: lib/Transforms/IPO/FunctionAttrs.cpp
===================================================================
--- lib/Transforms/IPO/FunctionAttrs.cpp
+++ lib/Transforms/IPO/FunctionAttrs.cpp
@@ -44,6 +44,12 @@
 STATISTIC(NumNoAlias, "Number of function returns marked noalias");
 STATISTIC(NumAnnotated, "Number of attributes added to library functions");
 
+static cl::opt<int> ReadAttrsAnalysisUsesThreshold(
+    "readattrs-analysis-uses-threshold", cl::Hidden, cl::init(20),
+    cl::ZeroOrMore,
+    cl::desc("Control the maximum number of uses to analyze for "
+             "readonly/readnone attribute (default = 20)"));
+
 namespace {
   struct FunctionAttrs : public CallGraphSCCPass {
     static char ID; // Pass identification, replacement for typeid
@@ -425,7 +431,7 @@
   // We don't need to track IsWritten. If A is written to, return immediately.
 
   for (Use &U : A->uses()) {
-    if (Count++ >= 20)
+    if (Count++ >= ReadAttrsAnalysisUsesThreshold)
       return Attribute::None;
 
     Visited.insert(&U);


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D11311.30032.patch
Type: text/x-patch
Size: 1737 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20150717/ab55b177/attachment.bin>