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

[Nick Lewycky]

nicholas added subscribers: test, nicholas.
nicholas added a comment.

Xuetian Weng wrote:

> 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.


Question one, what does this have to do with ld.global.nc? It looks like 
you've got a load instruction (wrapped in an intrinsic) and it in turn 
has many uses. The fact it's ld.global.nc doesn't seem to be relevant?

In theory, if you replaced the ld.global.nc with a LoadInst, would you 
have the same problem? Or is there some other optimization that occurs 
on those uses of a LoadInst that doesn't occur on the ld.global.nc?

Question two, can you think of a way to design determinePointerReadAttrs 
so that it doesn't need a use count but doesn't have a bad worst case 
performance problem? For that matter, why isn't this code already linear 
in the number of instructions?

Nick

> [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);

> 

> 

> 

>  _______________________________________________

> 

> llvm-commits mailing list

>  llvm-commits at cs.uiuc.edu

>  http://lists.cs.uiuc.edu/mailman/listinfo/llvm-commits



http://reviews.llvm.org/D11311