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

[Nick Lewycky]

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