<div dir="auto">Hi Ees,<div dir="auto"><br></div><div dir="auto">SCEV Delinearization is the closest I know. But it has its problems. Well for one your expression should be SCEVable.</div><div dir="auto"><br></div><div dir="auto">But more importantly, SCEV Delinearization is trying to deduce something that is high-level (actually source-level) from a low-level IR in which a lot of this info has been lost. So, since there's not a 1-1 mapping from high-level code to LLVM IR, going backwards will always be imperfect.</div><div dir="auto"><br></div><div dir="auto">And so since you are too trying to deduce a high-level thing, I believe that any solution will be imperfect.</div><div dir="auto"><br></div><div dir="auto">Best,</div><div dir="auto">Stefanos</div></div><br><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Sat, Oct 3, 2020, 02:55 Ees via llvm-dev <<a href="mailto:llvm-dev@lists.llvm.org" rel="noreferrer noreferrer noreferrer" target="_blank">llvm-dev@lists.llvm.org</a>> wrote:<br></div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">Anyone? I'd really appreciate any hints to look up as i'm somewhat stuck <br>
with this.<br>
<br>
Cheers.<br>
<br>
On 23-09-2020 12:27, Ees wrote:<br>
> Hi all,<br>
><br>
> For loads and stores i want to extract information about the number of <br>
> indices accessed. For instance:<br>
><br>
> struct S {int X, int *Y};<br>
><br>
> __global__ void kernel(int *A, int **B, struct S) {<br>
>   int x = A[..][..]; // -> L: A[..][..]<br>
>   int y = *B[2];   // -> L: B[0][2]<br>
>   int z = S.y[..];  // -> L: S.1[..]<br>
><br>
>   // etc..<br>
> }<br>
><br>
> I am performing some preprocessing on IR to:<br>
> 1. Move constant inline GEPs into instructions<br>
> 2. For loads and stores without a GEP operand, explicitly create a <br>
> (trivial) GEP with index 0<br>
><br>
> So now the operand of every load and store is a GEP instruction.<br>
><br>
> For simple stuff i am getting the right answer but when the index <br>
> expression becomes more complex multiple GEPs are introduced. For <br>
> instance:<br>
><br>
> *(A+2*(blockDim.x*blockIdx.x+threadIdx.x+1)+2+3) = 5;<br>
><br>
> produces:<br>
><br>
>   %6 = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()<br>
>   %7 = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()<br>
>   %8 = mul i32 %6, %7,<br>
>   %9 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()<br>
>   %10 = add i32 %8, %9,<br>
>   %11 = add i32 %10, 1,<br>
>   %12 = mul i32 2, %11,<br>
>   %13 = zext i32 %12 to i64,<br>
>   %14 = getelementptr inbounds i32, i32* %0, i64 %13<br>
>   %15 = getelementptr inbounds i32, i32* %14, i64 2<br>
>   %16 = getelementptr inbounds i32, i32* %15, i64 3<br>
>   store i32 5, i32* %16, align 4,<br>
><br>
> So i guess relying on the number of GEPs to figure the number of <br>
> indices is only a heuristic. Is there a more robust way to go on about <br>
> it? Or some example i can look at?<br>
><br>
> PS: I'm only interested about CUDA kernels.<br>
><br>
> Ees<br>
><br>
_______________________________________________<br>
LLVM Developers mailing list<br>
<a href="mailto:llvm-dev@lists.llvm.org" rel="noreferrer noreferrer noreferrer noreferrer" target="_blank">llvm-dev@lists.llvm.org</a><br>
<a href="https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev" rel="noreferrer noreferrer noreferrer noreferrer noreferrer" target="_blank">https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev</a><br>
</blockquote></div>