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

[Xuetian Weng]

wengxt added a comment.

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


In our case, we intend to transform LoadInst into ld.global.nc during Instruction
selection, not the other direction. It's only safe to use ld.global.nc for LoadInst
if the pointer being accessed is readonly (and noalias, which is irreverent to this
patch).

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


Yeah, I also think this is linear in the number of instructions, and I don't think 
there can be multiple call on this function that make it quadratic or something.
The complexity should be O(num_of_arguments * num_of_instructions)
Maybe simply remove this threshold is good enough. Is there any other reason
to introduce this threshold?

> 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