llvm dev - Sep 2020 - Information about the number of indices in memory accesses

Hi all,

For loads and stores i want to extract information about the number of 
indices accessed. For instance:

struct S {int X, int *Y};

__global__ void kernel(int *A, int **B, struct S) {
   int x = A[..][..]; // -> L: A[..][..]
   int y = *B[2];   // -> L: B[0][2]
   int z = S.y[..];  // -> L: S.1[..]

   // etc..
}

I am performing some preprocessing on IR to:
1. Move constant inline GEPs into instructions
2. For loads and stores without a GEP operand, explicitly create a 
(trivial) GEP with index 0

So now the operand of every load and store is a GEP instruction.

For simple stuff i am getting the right answer but when the index 
expression becomes more complex multiple GEPs are introduced. For instance:

*(A+2*(blockDim.x*blockIdx.x+threadIdx.x+1)+2+3) = 5;

produces:

   %6 = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
   %7 = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
   %8 = mul i32 %6, %7,
   %9 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
   %10 = add i32 %8, %9,
   %11 = add i32 %10, 1,
   %12 = mul i32 2, %11,
   %13 = zext i32 %12 to i64,
   %14 = getelementptr inbounds i32, i32* %0, i64 %13
   %15 = getelementptr inbounds i32, i32* %14, i64 2
   %16 = getelementptr inbounds i32, i32* %15, i64 3
   store i32 5, i32* %16, align 4,

So i guess relying on the number of GEPs to figure the number of indices 
is only a heuristic. Is there a more robust way to go on about it? Or 
some example i can look at?

PS: I'm only interested about CUDA kernels.

Ees

Anyone? I'd really appreciate any hints to look up as i'm somewhat stuck
with this.

Cheers.

On 23-09-2020 12:27, Ees wrote:
> Hi all,
>
> For loads and stores i want to extract information about the number of 
> indices accessed. For instance:
>
> struct S {int X, int *Y};
>
> __global__ void kernel(int *A, int **B, struct S) {
>   int x = A[..][..]; // -> L: A[..][..]
>   int y = *B[2];   // -> L: B[0][2]
>   int z = S.y[..];  // -> L: S.1[..]
>
>   // etc..
> }
>
> I am performing some preprocessing on IR to:
> 1. Move constant inline GEPs into instructions
> 2. For loads and stores without a GEP operand, explicitly create a 
> (trivial) GEP with index 0
>
> So now the operand of every load and store is a GEP instruction.
>
> For simple stuff i am getting the right answer but when the index 
> expression becomes more complex multiple GEPs are introduced. For 
> instance:
>
> *(A+2*(blockDim.x*blockIdx.x+threadIdx.x+1)+2+3) = 5;
>
> produces:
>
>   %6 = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
>   %7 = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
>   %8 = mul i32 %6, %7,
>   %9 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
>   %10 = add i32 %8, %9,
>   %11 = add i32 %10, 1,
>   %12 = mul i32 2, %11,
>   %13 = zext i32 %12 to i64,
>   %14 = getelementptr inbounds i32, i32* %0, i64 %13
>   %15 = getelementptr inbounds i32, i32* %14, i64 2
>   %16 = getelementptr inbounds i32, i32* %15, i64 3
>   store i32 5, i32* %16, align 4,
>
> So i guess relying on the number of GEPs to figure the number of 
> indices is only a heuristic. Is there a more robust way to go on about 
> it? Or some example i can look at?
>
> PS: I'm only interested about CUDA kernels.
>
> Ees
>

Hi Ees,

SCEV Delinearization is the closest I know. But it has its problems. Well
for one your expression should be SCEVable.

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.

And so since you are too trying to deduce a high-level thing, I believe
that any solution will be imperfect.

Best,
Stefanos

On Sat, Oct 3, 2020, 02:55 Ees via llvm-dev <llvm-dev at lists.llvm.org>
wrote:
> Anyone? I'd really appreciate any hints to look up as i'm somewhat
stuck
> with this.
>
> Cheers.
>
> On 23-09-2020 12:27, Ees wrote:
> > Hi all,
> >
> > For loads and stores i want to extract information about the number of
> > indices accessed. For instance:
> >
> > struct S {int X, int *Y};
> >
> > __global__ void kernel(int *A, int **B, struct S) {
> >   int x = A[..][..]; // -> L: A[..][..]
> >   int y = *B[2];   // -> L: B[0][2]
> >   int z = S.y[..];  // -> L: S.1[..]
> >
> >   // etc..
> > }
> >
> > I am performing some preprocessing on IR to:
> > 1. Move constant inline GEPs into instructions
> > 2. For loads and stores without a GEP operand, explicitly create a
> > (trivial) GEP with index 0
> >
> > So now the operand of every load and store is a GEP instruction.
> >
> > For simple stuff i am getting the right answer but when the index
> > expression becomes more complex multiple GEPs are introduced. For
> > instance:
> >
> > *(A+2*(blockDim.x*blockIdx.x+threadIdx.x+1)+2+3) = 5;
> >
> > produces:
> >
> >   %6 = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
> >   %7 = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
> >   %8 = mul i32 %6, %7,
> >   %9 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
> >   %10 = add i32 %8, %9,
> >   %11 = add i32 %10, 1,
> >   %12 = mul i32 2, %11,
> >   %13 = zext i32 %12 to i64,
> >   %14 = getelementptr inbounds i32, i32* %0, i64 %13
> >   %15 = getelementptr inbounds i32, i32* %14, i64 2
> >   %16 = getelementptr inbounds i32, i32* %15, i64 3
> >   store i32 5, i32* %16, align 4,
> >
> > So i guess relying on the number of GEPs to figure the number of
> > indices is only a heuristic. Is there a more robust way to go on about
> > it? Or some example i can look at?
> >
> > PS: I'm only interested about CUDA kernels.
> >
> > Ees
> >
> _______________________________________________
> LLVM Developers mailing list
> llvm-dev at lists.llvm.org
> https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL:
<http://lists.llvm.org/pipermail/llvm-dev/attachments/20201003/1c3576db/attachment.html>