<div dir="ltr">Hi Ees,<br><br>You may want to have a look at GetUnderlyingObjects (<a href="https://llvm.org/doxygen/namespacellvm.html#ad3428471506e6c03e9395a697a897a83">https://llvm.org/doxygen/namespacellvm.html#ad3428471506e6c03e9395a697a897a83</a>).<br>The page takes some time to load. Alternatively, you can see the implementation at ValueTracking.cpp (<a href="https://llvm.org/doxygen/ValueTracking_8cpp_source.html">https://llvm.org/doxygen/ValueTracking_8cpp_source.html</a>)<br>and use it by including `llvm/Analysis/ValueTracking.h`.<div><br></div><div>Note that you probably don't want to call this frequently.<br><br>Best,<br>Stefanos</div></div><br><div class="gmail_quote"><div dir="ltr" class="gmail_attr">Στις Σάβ, 3 Οκτ 2020 στις 3:26 μ.μ., ο/η Ees <<a href="mailto:kayesg42@gmail.com">kayesg42@gmail.com</a>> έγραψε:<br></div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">
<div>
<p>Stefanos, Michael, thank you for your responses. I've been toying
with this for some days now and i kind of figured/feared this
might not be possible, i just thought, that maybe there is
something. I'm new to LLVM so i'm not aware of most things. In any
case i've decided to take another direction with this.<br>
<br>
May i ask one more thing? Say i want for each load/store to get a
name for the base pointer. For instance the "A", "B", "S.1" in the
example above. <br>
<br>
In this example:<br>
<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>
I can follow the GEP chain until i hit an argument/global/alloca.
But what if there are temporaries in between them? Is there a more
generic/systematic way to go on about it?<br>
Maybe some project you are aware of that does something similar to
have a look at?<br>
<br>
Cheers,<br>
Ees<br>
</p>
<div>On 03-10-2020 12:43, Stefanos Baziotis
wrote:<br>
</div>
<blockquote type="cite">
<div dir="ltr">Michael makes a great point about aliasing here
and different indexing that accesses the same element!
<div><br>
</div>
<div>Another note: x = A[0][2] is fundamentally different
depending on the type of `A`. If e.g. A was declared: int
A[10][20], there's only _one_ load. A is a (and is treated as)
a linear buffer,<br>
and GEPs only pinpoint the specific position of A[0][2] in
this buffer (i.e. 0*10 + 2). But if A was e.g. this: int **A,
there _two_ loads. One load to get the "pointer of the zeroth
row" and another load to get the 2nd element off of that row.
<div><br>
</div>
<div>So, you see, all these things make any deduction method
very very imprecise.</div>
</div>
<div><br>
</div>
<div>Best,</div>
<div>Stefanos</div>
</div>
<br>
<div class="gmail_quote">
<div dir="ltr" class="gmail_attr">Στις Σάβ, 3 Οκτ 2020 στις 5:13
π.μ., ο/η Michael Kruse <<a href="mailto:llvmdev@meinersbur.de" target="_blank">llvmdev@meinersbur.de</a>>
έγραψε:<br>
</div>
<blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">
<div dir="ltr">As Stefanos mentioned, LLVM-IR is generally too
low-level for this. C/C++ semantics also don't make it
easier due to possible aliasing.
<div><br>
</div>
<div>The 3 cases are very different to analyze.</div>
<div><br>
</div>
<div>1. A[..][..] is syntactically invalid. A is a single
pointer.</div>
<div>2. *B[2] Is not equivalent to
B[0][2], but to B[2][0]. This
jagged/indirect/pointer-chasing arrays accesses are
expressed as chains of getelementptr and load
instructions. However, they cannot be usefully optimized
because pointers on each level can actually point to the
same thing. Eg:</div>
<div><br>
</div>
<div>int data[] = { 41, 42, 43 };</div>
<div>B[0] = data;</div>
<div>B[1] = data;</div>
<div>B[2] = data;</div>
<div><br>
</div>
<div>therefore</div>
<div><br>
</div>
<div>B[0][1] == B[1][1] == B[2][1] == 42;</div>
<div><br>
</div>
<div>3. S.y[..] can be one or multiple getelementptr
instructions and is the most analyzable since all indices
are static at compile-time. The number of subscripts and
their values can directly be taken from the getelementptr
instruction(s). The pointers-to-int can still alias with
other ints in the program.</div>
<div><br>
</div>
<div>4. These case of VLA or manually linearized array: </div>
<div><br>
</div>
<div>int *C = alloca(sizeof(int) * m * n);</div>
<div>C[x + m*y];</div>
<div><br>
</div>
<div>This is intended to have two dimensions, but appear as
just one in the LLVM-IR. The following (C99 VLA) is
compiled to approximately the same LLVM-IR</div>
<div><br>
</div>
<div>int C[n][m];</div>
<div>C[y][x];</div>
<div><br>
</div>
<div>Delinearization as mentioned by Stefanos tries to
recover the two subscripts x and y, but can only do so
heuristically. Also keep in mind that </div>
<div>C[1][-1] appears as the same IR as C[0][m-1], so there
is no unique way to delinerarize. In particular, one
cannot just assume that if all indices are different, that
the memory locations being accessed are different (again,
a pointer aliasing problem)</div>
<div><br>
</div>
<div>Michael</div>
<div><br>
</div>
<div><br>
</div>
<div><br>
</div>
<div><br>
</div>
<div><br>
</div>
<div><br>
</div>
<div><br>
</div>
<div><br>
</div>
<div><br>
</div>
<div><br>
</div>
<div><br>
</div>
<div><br>
</div>
<div><br>
</div>
<div><br>
</div>
<div><br>
</div>
</div>
<br>
<div class="gmail_quote">
<div dir="ltr" class="gmail_attr">Am Fr., 2. Okt. 2020 um
19:25 Uhr schrieb Stefanos Baziotis via llvm-dev <<a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a>>:<br>
</div>
<blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">
<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:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);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>
_______________________________________________<br>
LLVM Developers mailing list<br>
<a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a><br>
<a href="https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev" rel="noreferrer" target="_blank">https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev</a><br>
</blockquote>
</div>
</blockquote>
</div>
</blockquote>
</div>
</blockquote></div>