[LLVMdev] LLVM CreateStructGEP type assert error

[Lowell, Daniel]

Hi all,

Running LLVM 3.4 to create a custom pass for OpenCL transformations. I am attempting to GEP into a struct using IRBuilder's CreateStructGEP, but I keep getting this assert:

aoc: ../../../../../../compiler/llvm/include/llvm/Instructions.h:703: llvm::Type* llvm::checkGEPType(llvm::Type*): Assertion `Ty && "Invalid GetElementPtrInst indices for type!"' failed.

Which I've decoded as it doesn't recognize my struct as a built in type. This is confusing since it shouldn't care what type I'm indexing into, only that I give the correct indices. Otherwise what is the point of user defined structs?

//This part is atop my cl kernel file:

#define BUFFER_LEN 0x100000
typedef struct RB{
unsigned int x;
unsigned int y;
int z[BUFFER_LEN];
unsigned int xx[BUFFER_LEN];
unsigned int yy[BUFFER_LEN];
float zz[BUFFER_LEN];
} RB_t;

//The kernel sig. with the struct as the last argument:

__kernel void ht( int iteration, global RB_t *cB){ ... }

//My LLVM code:

void MPASS::exPointers(Module& M, Function& F){
BasicBlock *first = F.begin();
IRBuilder<> builder(first->begin());
Value *cB = --(F.arg_end());
Value *x_idx = builder.CreateStructGEP(cB, 0);
...
... < more of the same>
...
return
}

Simple enough, but I can't get past the first CreateStructGEP without the assert. I thought it was a version issue, so I changed the call to:

Value *x_idx = builder.CreateConstGEP2_32(cB, 0, 0);

However I got the same assert. It looks like this line in Instruction.h from LLVM is producing the assert further down the stack:

815 Type *PtrTy = PointerType::get(checkGEPType(
816 getIndexedType(Ptr->getType(), IdxList)),
817 Ptr->getType()->getPointerAddressSpace());

Resulting in the debugger producing:

(gdb) p PtrTy
$19 = (llvm::Type *) 0x7fffffff63d0
(gdb) p PtrTy->dump()
<unrecognized-type> $20 = void

Is there something special which needs to be done in order to index into a user defined struct in LLVM for OpenCL?

-----------------------------------------------







________________________________________
From: llvmdev...@cs.uiuc.edu [llvmdev...@cs.uiuc.edu] on behalf of llvmdev...@cs.uiuc.edu [llvmdev...@cs.uiuc.edu]
Sent: Wednesday, August 20, 2014 11:23 AM
To: llv...@cs.uiuc.edu
Subject: LLVMdev Digest, Vol 122, Issue 56

Send LLVMdev mailing list submissions to
llv...@cs.uiuc.edu

To subscribe or unsubscribe via the World Wide Web, visit
http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
or, via email, send a message with subject or body 'help' to
llvmdev...@cs.uiuc.edu

You can reach the person managing the list at
llvmde...@cs.uiuc.edu

When replying, please edit your Subject line so it is more specific
than "Re: Contents of LLVMdev digest..."


Today's Topics:

1. Re: [RFC] Removing static initializers for command line
options (Pete Cooper)
2. Project DiscoPoP (Entry provided) (Zhen Li)
3. Re: Functions with unnamed parameters in LLVM IR (Dan Liew)
4. Re: Functions with unnamed parameters in LLVM IR (Tim Northover)
5. Re: Functions with unnamed parameters in LLVM IR (Dan Liew)
6. Re: Proposal for ""llvm.mem.vectorize.safelen" (Robison, Arch)
7. Re: Functions with unnamed parameters in LLVM IR (Dan Liew)
8. Re: Proposal for ""llvm.mem.vectorize.safelen" (Robison, Arch)
9. Re: Addressing const reference in ArrayRef (David Blaikie)
10. Re: Proposal for ""llvm.mem.vectorize.safelen"
(Arnold Schwaighofer)
11. Re: [RFC] Removing static initializers for command line
options (Rafael Esp?ndola)
12. Re: Proposal for ""llvm.mem.vectorize.safelen" (Renato Golin)
13. Re: [RFC] Removing static initializers for command line
options (Pete Cooper)
14. Re: Proposal for ""llvm.mem.vectorize.safelen" (Robison, Arch)
15. Re: Proposal for ""llvm.mem.vectorize.safelen" (Johannes Doerfert)


----------------------------------------------------------------------

Message: 1
Date: Tue, 19 Aug 2014 23:21:14 -0700
From: Pete Cooper <peter_...@apple.com>
To: Chandler Carruth <chan...@google.com>
Cc: LLVM Dev <llv...@cs.uiuc.edu>
Subject: Re: [LLVMdev] [RFC] Removing static initializers for command
line options
Message-ID: <272A2B77-7B80-44DE...@apple.com>
Content-Type: text/plain; charset="utf-8"


> On Aug 19, 2014, at 10:24 PM, Chandler Carruth <chan...@google.com> wrote:
>
>
> On Tue, Aug 19, 2014 at 10:10 PM, Pete Cooper <peter_...@apple.com> wrote:
>> On Aug 19, 2014, at 6:45 PM, Chandler Carruth <chan...@google.com> wrote:
>
>
>> FWIW, I largely agree with Rafael's position here, at least in the near term.
>>
>> The nice thing about it is that even if we don't stay there forever, it is a nice incremental improvement over the current state of the world, and we can actually be mindful going forward of whether the restriction it imposes (an inability to use "internal" knobs from within a library context that have multiple different library users in a single address space) proves to be a significant on-going burden.
> I actually disagree with this for two reasons.
>
> The first is that if there are going to be changes to the code anyway to remove static initializers, and we can move the storage to the pass at the same time, the we should make just one change and not two.
>
> No one is suggesting otherwise that I have seen? At least, my interpretation (perhaps incorrect, I've not had time to read all of this thread in 100% detail) was that the removal of static initializers wouldn't require changing every cl::opt variable.
It won?t no. The majority of static initializers are globals in passes, but cl::opt?s and other globals in tools for example are out of scope for this work right now (there?s no opt.cpp in a dylib so we don?t care about its globals for example).
>
>
> The second reason is that in most cases these knobs should not be globals. If I had a piece of data (not a CL::opt) in global scope, only used by one pass, then I'm sure people would question why it's a global at all and move it inside the class. We're treating cl::opt as special here when there's no reason for the opt or the storage to be global.
>
> We frown upon the use of globals, otherwise LLVM would be littered with them like many other C++ code bases. I don't think cl::opts should be special at all in this respect.
>
> Sure, you're arguing against the core design of cl::opt. However, we have it, and it wasn't an accident or for lack of other patterns that we chose it.
Oh yeah, its a fine solution for a tricky problem, but now that we?re having this discussion its clear that its use of static initializers is itself a problem.
>
> For example, we don't require all constants to be per-pass, and instead have per-file constants. Rafael is suggesting that one use case for cl::opt global variables is, in essence, a constant that is somewhat easier for a developer of LLVM (*not* a user) to change during debugging and active development. I don't think the desire for convenience and only supporting the default value in production contexts are completely invalid.
I can see what you?re saying here, but i?m not convinced that changing the value of a constant in global scope is any easier than changing its value in the pass constructor.
>
> Once you factor those in, we have a tradeoff. Historically, the tradeoff was made in the direction of convenience, relying on a very narrow interpretation of the use cases. It isn't clear to me that we should, today, make a different tradeoff. It certainly doesn't make sense why you would gate removing static initializers (a clear win) on forcing a change on a core design pattern within LLVM which not all of the developers are really supportive of (at this point).
I agree here. There?s not a strict need to move the option storage for something like an unsigned in to a pass using it (there may be for a std::string for which we?d again have a static initializer). However, I do think its the right thing to do in terms of coding guidelines. If the code to get and set an option is already in the pass initializer/constructor respectively, then I don?t think the storage should have to live outside the pass just to minimize a patch.

Now saying this, I don?t think if the community agreed to this proposal as is, that it would mean blanket approval to change all cl::opts in all cases. But I think where its an easy change, and obviously the right choice, that options as well as their storage should be allowed to be moved in to the pass. If it makes sense to move only the option but not the storage then that should be done as a first step, and a discussion started on the right thing for the storage.

Thanks,
Pete

-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.cs.uiuc.edu/pipermail/llvmdev/attachments/20140819/d25e9982/attachment-0001.html>

------------------------------

Message: 2
Date: Wed, 20 Aug 2014 12:42:48 +0200
From: Zhen Li <z....@grs-sim.de>
To: <llv...@cs.uiuc.edu>
Cc: disc...@grs-sim.de
Subject: [LLVMdev] Project DiscoPoP (Entry provided)
Message-ID: <1c9515d7-71de-45e0...@HUB1.rwth-ad.de>
Content-Type: text/plain; charset="utf-8"; Format="flowed"

Dear LLVM Devs,

We are a team from German Research School for Simulation Sciences. We
are currently working on DiscoPoP (_Disco_very of _Po_tential
_P_arallelism), which is an LLVM-based tool that assists in identifying
potential parallelism in sequential programs. The tool uses LLVM to
instrument the code for finding both control and data dependences. A
series of analyses are built on top of dependence to explore potential
parallelism and parallel patterns. We have a modified version of Clang
with a new option "-dp" to invoke our tool.

We have a web page about DiscoPoP: http://www.grs-sim.de/discopop, and
we would like to know if it is possible to have our project listed on
"Projects built with LLVM" page of LLVM website. Our plan is to make
DiscoPoP an open-source tool in the near future.

----------------------------------
The entry for our project:

DiscoPoP: A parallelism discovery tool

By the <a
href="http://www.grs-sim.de/research/parallel-programming/multicore-programming/">DiscoPoP
Dev Team</a>

<a href="http://www.grs-sim.de/discopop">DiscoPoP (_Disco_very of
_Po_tential _P_arallelism)</a> is a tool that assists in identifying
potential parallelism in sequential C/C++ programs. It instruments the
code for finding both control and data dependences. A series of analyses
are built on top of dependence to explore potential parallelism and
parallel patterns.

The instrumentation is done with the help of LLVM. A modified version of
Clang with a new option "-dp" is also provided to invoke DiscoPoP.
----------------------------------

Thank you for your attention and we are looking forward to hearing from you.

Best regards,
DiscoPoP Dev Team
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.cs.uiuc.edu/pipermail/llvmdev/attachments/20140820/e6108f0c/attachment-0001.html>

------------------------------

Message: 3
Date: Wed, 20 Aug 2014 15:12:44 +0100
From: Dan Liew <d...@su-root.co.uk>
To: Tim Northover <t.p.no...@gmail.com>
Cc: LLVM Developers Mailing List <llv...@cs.uiuc.edu>
Subject: Re: [LLVMdev] Functions with unnamed parameters in LLVM IR
Message-ID:
<CAJ7DczHn2WQkQow2yw_CuQiz...@mail.gmail.com>
Content-Type: text/plain; charset="utf-8"

Sorry for the delay on this. Is the attached patch any better?

I've modified the note in the Identifiers sections and trimmed down my
addition to the Functions section.

I've been a little inconsistent with using argument/parameter but the
entire document seems to be like this so I figured it would be okay.

Thanks,
Dan.
-------------- next part --------------
A non-text attachment was scrubbed...
Name: [v2]0001-Add-note-to-LangRef-about-how-function-arguments-can.patch
Type: text/x-patch
Size: 1713 bytes
Desc: not available
URL: <http://lists.cs.uiuc.edu/pipermail/llvmdev/attachments/20140820/9241d580/attachment-0001.bin>

------------------------------

Message: 4
Date: Wed, 20 Aug 2014 15:17:08 +0100
From: Tim Northover <t.p.no...@gmail.com>
To: Dan Liew <d...@su-root.co.uk>
Cc: LLVM Developers Mailing List <llv...@cs.uiuc.edu>
Subject: Re: [LLVMdev] Functions with unnamed parameters in LLVM IR
Message-ID:
<CAFHTzfLJEuj4jAS1JZSgbvoa...@mail.gmail.com>
Content-Type: text/plain; charset=UTF-8

Hi Dan,

On 20 August 2014 15:12, Dan Liew <d...@su-root.co.uk> wrote:
> Sorry for the delay on this. Is the attached patch any better?

Yep, that looks fine to me. Would you like me to commit it or do you
have access?

Cheers.

Tim.


------------------------------

Message: 5
Date: Wed, 20 Aug 2014 15:51:04 +0100
From: Dan Liew <d...@su-root.co.uk>
To: Tim Northover <t.p.no...@gmail.com>
Cc: LLVM Developers Mailing List <llv...@cs.uiuc.edu>
Subject: Re: [LLVMdev] Functions with unnamed parameters in LLVM IR
Message-ID:
<CAJ7DczF_VbSnUwhS-vWNf_DW...@mail.gmail.com>
Content-Type: text/plain; charset=UTF-8

> On 20 August 2014 15:12, Dan Liew <d...@su-root.co.uk> wrote:
>> Sorry for the delay on this. Is the attached patch any better?
>
> Yep, that looks fine to me. Would you like me to commit it or do you
> have access?

Thanks for taking a look. I have commit access so I'll commit it later on today.

Thanks,
Dan.


------------------------------

Message: 6
Date: Wed, 20 Aug 2014 15:08:15 +0000
From: "Robison, Arch" <arch.r...@intel.com>
To: Hal Finkel <hfi...@anl.gov>, Roel Jordans <r.jo...@tue.nl>
Cc: "llv...@cs.uiuc.edu" <llv...@cs.uiuc.edu>
Subject: Re: [LLVMdev] Proposal for ""llvm.mem.vectorize.safelen"
Message-ID:
<51C2718BFE68DA4698E4...@fmsmsx116.amr.corp.intel.com>

Content-Type: text/plain; charset="utf-8"

> I recommend that you send patches for an implementation
> (including the Loop::GetAnnotatedVectorSafelen function
> and associated updates to the vectorizer) for review.

I expect to send the patches for review later this week.

> Make sure to remember LangRef updates!

Thanks for the reminder.

> Also, looking at the original proposal again, I see no reason
> to restrict this to an i32: we might as well allow it to be an
> i64 obviously we don't have billion-lane vectors, but the
> metadata can also be useful for other kinds of analysis).

I doubt there is much advantage between knowing the minimum loop-carried
dependence difference is 1u<<31 or 1u<<63. Unless there is a clear
need for the higher values less than infinity, It would seem simpler
for now to adopt the same conventions as llvm.loop.vectorize.width
so that some processing infrastructure can be shared easily.

> I don't like using the name 'safelen' however. I can forgive OpenMP
> for choosing such a short name because people need to type it,
> but I'd prefer that the metadata have a more-descriptive name.
> minimum_dependency_distance is perhaps better.

I'm open to naming suggestions, though I'm wondering if sticking with
what is now a term of art in OpenMP might be the least of all evils,
because the semantics turn out to be a little more subtle than
just a minimum dependence distance. My current wording of the semantics
for safelen of k are:

/// The loop can be assumed to have no loop-carried
/// lexically backward dependencies with distance less than k.

- Arch D. Robison
Intel Corporation




------------------------------

Message: 7
Date: Wed, 20 Aug 2014 16:16:31 +0100
From: Dan Liew <d...@su-root.co.uk>
To: Tim Northover <t.p.no...@gmail.com>
Cc: LLVM Developers Mailing List <llv...@cs.uiuc.edu>
Subject: Re: [LLVMdev] Functions with unnamed parameters in LLVM IR
Message-ID:
<CAJ7DczHz_bb__ATrFOB8wgKD...@mail.gmail.com>
Content-Type: text/plain; charset=UTF-8

Committed in r216070


------------------------------

Message: 8
Date: Wed, 20 Aug 2014 15:17:01 +0000
From: "Robison, Arch" <arch.r...@intel.com>
To: Hal Finkel <hfi...@anl.gov>, Roel Jordans <r.jo...@tue.nl>
Cc: "llv...@cs.uiuc.edu" <llv...@cs.uiuc.edu>
Subject: Re: [LLVMdev] Proposal for ""llvm.mem.vectorize.safelen"
Message-ID:
<51C2718BFE68DA4698E4...@fmsmsx116.amr.corp.intel.com>

Content-Type: text/plain; charset="utf-8"

Now that I'm looking at editing LangRef, I have a question. The current
llvm.loop.vectorize metadata are hints, and have this constraint:

The``llvm.loop.vectorize`` and ``llvm.loop.interleave`` metadata are only
optimization hints and the optimizer will only interleave and vectorize loops
if it believes it is safe to do so.

But safelen is different. It's an assertion about safety, so prefixing it with
llvm.loop.vectorize seems inappropriate. Does is sound reasonable to drop
the "vectorize" portion. Maybe spell it something like this?

llvm.loop.carried_dependence_distance.min

- Arch




------------------------------

Message: 9
Date: Wed, 20 Aug 2014 08:28:37 -0700
From: David Blaikie <dbla...@gmail.com>
To: Joey Ye <joey....@gmail.com>
Cc: LLVM Developers Mailing List <llv...@cs.uiuc.edu>
Subject: Re: [LLVMdev] Addressing const reference in ArrayRef
Message-ID:
<CAENS6EveGNodi4ok3dJNzOnm...@mail.gmail.com>
Content-Type: text/plain; charset="utf-8"

I seem to recall discussing this before - is there an existing llvm bug
filed, another email thread or something (or perhaps it was just an IRC
conversation)? It would be good to keep all the discussion together or at
least reference the prior (llvm community) discussion.

Have you tried applying your suggested patch? I assume you meant to suggest
replacing/modifying the existing ctor (to take non-const ref) rather than
adding another?

I'd assume that causes a lot of build failures as we probably rely on
binding temporaries to ArrayRef's in many places correctly (most ArrayRef's
are temporaries, not local variables).

I think in the previous discussion I suggested we might just want to make a
practice of treating named/local (not parameter) ArrayRef's with greater
suspicion, the same way we do for twine, but I'm not sure.

We could move this ctor into a factory function (and/or just make the CTor
private and friend the makeArrayRef helper for this case) to make it more
obvious/easier to find bad call sites. But it would be helpful to have the
context of the prior discussion to continue from there.
On Aug 19, 2014 11:18 PM, "Joey Ye" <joey....@gmail.com> wrote:

> Analyzing why GCC failed to build LLVM recently, one root cause lies
> in definition of ArrayRef:
> // ArrayRef.h:
> ArrayRef(const T &OneElt) : Data(&OneElt), Length(1) {}
>
> Here address of const reference is taken and stored to an object. It
> is believed that live range of const reference is only at the function
> call site, escaping of its address to an object with a longer live
> range is invalid. Referring to the case and discussion here:
> https://gcc.gnu.org/ml/gcc/2014-08/msg00173.html
>
> So I would suggest to fix ArrayRef. Adding a non-const version of
> constructor should work, but it still leaves the vulnerability in
> const version, which I'd expect on people in the community to work out
> a solution.
> ArrayRef(T &OneElt) : Data(&OneElt), Length(1) {}
>
>
> Thanks,
> Joey
> _______________________________________________
> LLVM Developers mailing list
> LLV...@cs.uiuc.edu http://llvm.cs.uiuc.edu
> http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.cs.uiuc.edu/pipermail/llvmdev/attachments/20140820/55b6ad6e/attachment-0001.html>

------------------------------

Message: 10
Date: Wed, 20 Aug 2014 08:29:48 -0700
From: Arnold Schwaighofer <aschwa...@apple.com>
To: "Robison, Arch" <arch.r...@intel.com>
Cc: "llv...@cs.uiuc.edu" <llv...@cs.uiuc.edu>
Subject: Re: [LLVMdev] Proposal for ""llvm.mem.vectorize.safelen"
Message-ID: <6EDAADA1-F702-4AAD...@apple.com>
Content-Type: text/plain; charset=utf-8


> On Aug 20, 2014, at 8:08 AM, Robison, Arch <arch.r...@intel.com> wrote:
>
>> I recommend that you send patches for an implementation
>> (including the Loop::GetAnnotatedVectorSafelen function
>> and associated updates to the vectorizer) for review.
>
> I expect to send the patches for review later this week.
>
>> Make sure to remember LangRef updates!
>
> Thanks for the reminder.
>
>> Also, looking at the original proposal again, I see no reason
>> to restrict this to an i32: we might as well allow it to be an
>> i64 obviously we don't have billion-lane vectors, but the
>> metadata can also be useful for other kinds of analysis).
>
> I doubt there is much advantage between knowing the minimum loop-carried
> dependence difference is 1u<<31 or 1u<<63. Unless there is a clear
> need for the higher values less than infinity, It would seem simpler
> for now to adopt the same conventions as llvm.loop.vectorize.width
> so that some processing infrastructure can be shared easily.
>
>> I don't like using the name 'safelen' however. I can forgive OpenMP
>> for choosing such a short name because people need to type it,
>> but I'd prefer that the metadata have a more-descriptive name.
>> minimum_dependency_distance is perhaps better.
>
> I'm open to naming suggestions, though I'm wondering if sticking with
> what is now a term of art in OpenMP might be the least of all evils,
> because the semantics turn out to be a little more subtle than
> just a minimum dependence distance. My current wording of the semantics
> for safelen of k are:
>
> /// The loop can be assumed to have no loop-carried
> /// lexically backward dependencies with distance less than k.

This means you would allow lexically forward dependencies which complicates things (this would need more infrastructure in clang). You need to carry over ?source order? from the front-end. Because the dependency is loop carried llvm would be allowed to move loads and stores within the loop:

Lexical forward dependency:

#pragma vectorize safelen(4)
for (i = ?) {
a[i] = b[i] + z[i];
c[i] = a[i-1] + 1;
}

LLVM might tranform this loop to:

for (i = ?) {
c[i] = a[i-1] + 1;
a[i] = b[i] + z[i];
}

if we now vectorize this (without knowledge of the orignal source order) we get the wrong answer:

for (i = ?) {
c[i] = a[i-1:i+2] + 1;
a[i:i+3] = b[i] + z[i];
}

Alternatively, the metadata loop.vectorize.safelen would have to prevent any such reordering of accesses i.e. any pass that reorders memory accesses would have to be aware of it which is fragile.

>
> - Arch D. Robison
> Intel Corporation
>
>
> _______________________________________________
> LLVM Developers mailing list
> LLV...@cs.uiuc.edu http://llvm.cs.uiuc.edu
> http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev




------------------------------

Message: 11
Date: Wed, 20 Aug 2014 11:51:56 -0400
From: Rafael Esp?ndola <rafael.e...@gmail.com>
To: Pete Cooper <peter_...@apple.com>
Cc: LLVM Dev <llv...@cs.uiuc.edu>
Subject: Re: [LLVMdev] [RFC] Removing static initializers for command
line options
Message-ID:
<CAG3jReKcN9niyhGsfZrHg2qF...@mail.gmail.com>
Content-Type: text/plain; charset=UTF-8

> FWIW, I largely agree with Rafael's position here, at least in the near
> term.
>
> The nice thing about it is that even if we don't stay there forever, it is a
> nice incremental improvement over the current state of the world, and we can
> actually be mindful going forward of whether the restriction it imposes (an
> inability to use "internal" knobs from within a library context that have
> multiple different library users in a single address space) proves to be a
> significant on-going burden.
>
> I actually disagree with this for two reasons.
>
> The first is that if there are going to be changes to the code anyway to
> remove static initializers, and we can move the storage to the pass at the
> same time, the we should make just one change and not two.
>
> The second reason is that in most cases these knobs should not be globals.
> If I had a piece of data (not a CL::opt) in global scope, only used by one
> pass, then I'm sure people would question why it's a global at all and move
> it inside the class. We're treating cl::opt as special here when there's no
> reason for the opt or the storage to be global.
>
> We frown upon the use of globals, otherwise LLVM would be littered with them
> like many other C++ code bases. I don't think cl::opts should be special at
> all in this respect.


Note that the call

cl::OptionRegistry::CreateOption<bool>(&ScalarizeLoadStore,
"ScalarizeLoadStore",
"scalarize-load-store", cl::Hidden, cl::init(false),
cl::desc("Allow the scalarizer pass to scalarize loads and store"));

ScalarizeLoadStore can actually be a member variable as long as caller
guarantees it is still around when the command line is parsed. I have
no objections to doing this move in the first pass if you want to.

What I would really like to avoid for now is the

cl::OptionRegistry::GetValue<bool>("ScalarizeLoadStore");

Cheers,
Rafael


------------------------------

Message: 12
Date: Wed, 20 Aug 2014 17:00:07 +0100
From: Renato Golin <renato...@linaro.org>
To: "Robison, Arch" <arch.r...@intel.com>
Cc: "llv...@cs.uiuc.edu" <llv...@cs.uiuc.edu>
Subject: Re: [LLVMdev] Proposal for ""llvm.mem.vectorize.safelen"
Message-ID:
<CAMSE1kd9S41Oy3hxXBMz2r0M...@mail.gmail.com>
Content-Type: text/plain; charset=UTF-8

On 20 August 2014 16:17, Robison, Arch <arch.r...@intel.com> wrote:
> But safelen is different. It's an assertion about safety, so prefixing it with
> llvm.loop.vectorize seems inappropriate. Does is sound reasonable to drop
> the "vectorize" portion. Maybe spell it something like this?

This was true to all *current* vectorizer pragmas, but certainly not
safelen. I think we can change that description instead of changing
the namespace.

cheers,
--renato


------------------------------

Message: 13
Date: Wed, 20 Aug 2014 09:15:02 -0700
From: Pete Cooper <peter_...@apple.com>
To: Rafael Esp?ndola <rafael.e...@gmail.com>
Cc: LLVM Dev <llv...@cs.uiuc.edu>
Subject: Re: [LLVMdev] [RFC] Removing static initializers for command
line options
Message-ID: <154097CE-BB97-4746...@apple.com>
Content-Type: text/plain; charset=utf-8


> On Aug 20, 2014, at 8:51 AM, Rafael Esp?ndola <rafael.e...@gmail.com> wrote:
>
>> FWIW, I largely agree with Rafael's position here, at least in the near
>> term.
>>
>> The nice thing about it is that even if we don't stay there forever, it is a
>> nice incremental improvement over the current state of the world, and we can
>> actually be mindful going forward of whether the restriction it imposes (an
>> inability to use "internal" knobs from within a library context that have
>> multiple different library users in a single address space) proves to be a
>> significant on-going burden.
>>
>> I actually disagree with this for two reasons.
>>
>> The first is that if there are going to be changes to the code anyway to
>> remove static initializers, and we can move the storage to the pass at the
>> same time, the we should make just one change and not two.
>>
>> The second reason is that in most cases these knobs should not be globals.
>> If I had a piece of data (not a CL::opt) in global scope, only used by one
>> pass, then I'm sure people would question why it's a global at all and move
>> it inside the class. We're treating cl::opt as special here when there's no
>> reason for the opt or the storage to be global.
>>
>> We frown upon the use of globals, otherwise LLVM would be littered with them
>> like many other C++ code bases. I don't think cl::opts should be special at
>> all in this respect.
>
>
> Note that the call
>
> cl::OptionRegistry::CreateOption<bool>(&ScalarizeLoadStore,
> "ScalarizeLoadStore",
> "scalarize-load-store", cl::Hidden, cl::init(false),
> cl::desc("Allow the scalarizer pass to scalarize loads and store"));
>
> ScalarizeLoadStore can actually be a member variable as long as caller
> guarantees it is still around when the command line is parsed. I have
> no objections to doing this move in the first pass if you want to.
>
> What I would really like to avoid for now is the
>
> cl::OptionRegistry::GetValue<bool>("ScalarizeLoadStore?);
Ah, I totally see what you mean now. Sorry if i had been confused before.

To be honest, I had exactly the same reservations myself, but then i looked in to how we initialize and create passes and there actually isn?t a nice way to do this right now.

The trouble is that INITIALIZE_PASS doesn?t actually construct a pass. It actually constructs a PassInfo which has a pointer to the default constructor of the pass. Later, if the pass manager needs to (say -scalarize was on the commandline), it will get the default constructor from PassInfo and construct the pass.

Unfortunately, this completely detaches the lifetime of the pass instance where we want to store the ScalarizeLoadStore bool, from the option code to hook in to the cl::opt stuff. I originally wanted to do something gross like pass __offset_of(Scalarizer::ScalarizeLoadStore) to the PassInfo and hide the initialization of the option in there. But that doesn?t work either, as there?s no guarantee you?ll even create a instance of Scalarizer, even though you could pass -scalarize-load-store to the command line.

So, we do have 2 solutions: a global variable, and a call to GetValue. As you can see, either way isn?t perfect, but if you can think of another solution i?m sure we can discuss it.

Thanks,
Pete
>
> Cheers,
> Rafael




------------------------------

Message: 14
Date: Wed, 20 Aug 2014 16:16:32 +0000
From: "Robison, Arch" <arch.r...@intel.com>
To: Arnold Schwaighofer <aschwa...@apple.com>
Cc: "llv...@cs.uiuc.edu" <llv...@cs.uiuc.edu>
Subject: Re: [LLVMdev] Proposal for ""llvm.mem.vectorize.safelen"
Message-ID:
<51C2718BFE68DA4698E4...@fmsmsx116.amr.corp.intel.com>

Content-Type: text/plain; charset="utf-8"

Thanks for alerting me to this issue. This is going to take more
effort than I thought :-(. Given that a major motivation behind
the OpenMP #pragma omp simd is to allow lexical forward dependences,
we should find a way to do this right.

I agree that we want to avoid making passes that reorder accesses fragile.
The extra work in Clang (or Julia :-) seems unavoidable, unless those
producers always emit LLVM instructions in lexical order. If the latter
is the case, perhaps we could have a helper routine that, given the IR early
and in lexical order, finishes the annotation work?

It seems that we need metadata on each memory access, distinct from
llvm.mem.parallel_loop_access. Say something like llvm.mem.vector_loop_access
that includes relative lexical position as a third operand? Second operand
could point back to the metadata with the minimum loop-carried distance,
which in turn could point back to the loop id.

- Arch





------------------------------

Message: 15
Date: Wed, 20 Aug 2014 09:18:14 -0700
From: Johannes Doerfert <doer...@cs.uni-saarland.de>
To: Arnold Schwaighofer <aschwa...@apple.com>
Cc: "llv...@cs.uiuc.edu" <llv...@cs.uiuc.edu>, "Robison, Arch"
<arch.r...@intel.com>
Subject: Re: [LLVMdev] Proposal for ""llvm.mem.vectorize.safelen"
Message-ID: <20140820161814.GD4039@JD-Arch>
Content-Type: text/plain; charset="utf-8"

On 08/20, Arnold Schwaighofer wrote:
>
> > On Aug 20, 2014, at 8:08 AM, Robison, Arch <arch.r...@intel.com> wrote:
> >
> >> I recommend that you send patches for an implementation
> >> (including the Loop::GetAnnotatedVectorSafelen function
> >> and associated updates to the vectorizer) for review.
> >
> > I expect to send the patches for review later this week.
> >
> >> Make sure to remember LangRef updates!
> >
> > Thanks for the reminder.
> >
> >> Also, looking at the original proposal again, I see no reason
> >> to restrict this to an i32: we might as well allow it to be an
> >> i64 obviously we don't have billion-lane vectors, but the
> >> metadata can also be useful for other kinds of analysis).
> >
> > I doubt there is much advantage between knowing the minimum loop-carried
> > dependence difference is 1u<<31 or 1u<<63. Unless there is a clear
> > need for the higher values less than infinity, It would seem simpler
> > for now to adopt the same conventions as llvm.loop.vectorize.width
> > so that some processing infrastructure can be shared easily.
> >
> >> I don't like using the name 'safelen' however. I can forgive OpenMP
> >> for choosing such a short name because people need to type it,
> >> but I'd prefer that the metadata have a more-descriptive name.
> >> minimum_dependency_distance is perhaps better.
> >
> > I'm open to naming suggestions, though I'm wondering if sticking with
> > what is now a term of art in OpenMP might be the least of all evils,
> > because the semantics turn out to be a little more subtle than
> > just a minimum dependence distance. My current wording of the semantics
> > for safelen of k are:
> >
> > /// The loop can be assumed to have no loop-carried
> > /// lexically backward dependencies with distance less than k.
>
> This means you would allow lexically forward dependencies which complicates things (this would need more infrastructure in clang). You need to carry over ?source order? from the front-end. Because the dependency is loop carried llvm would be allowed to move loads and stores within the loop:
This should not be that hard (see below).

> Lexical forward dependency:
>
> #pragma vectorize safelen(4)
> for (i = ?) {
> a[i] = b[i] + z[i];
> c[i] = a[i-1] + 1;
> }
>
> LLVM might tranform this loop to:
>
> for (i = ?) {
> c[i] = a[i-1] + 1;
> a[i] = b[i] + z[i];
> }
>
> if we now vectorize this (without knowledge of the orignal source order) we get the wrong answer:
>
> for (i = ?) {
> c[i] = a[i-1:i+2] + 1;
> a[i:i+3] = b[i] + z[i];
> }
>
> Alternatively, the metadata loop.vectorize.safelen would have to prevent any such reordering of accesses i.e. any pass that reorders memory accesses would have to be aware of it which is fragile.
Could we number the memory accesses for such loops (e.g., in clang)?
Adding metadata to each memory access which points to the safelen
annotation and contains an increasing constant would be similarly
fragile as other constructions we use at the moment. However, it would
allow us to determine iff the current order is still the original one or
not. (At least if I did not miss anything).

Best regards,
Johannes

--

Johannes Doerfert
Researcher / PhD Student

Compiler Design Lab (Prof. Hack)
Saarland University, Computer Science
Building E1.3, Room 4.26

Tel. +49 (0)681 302-57521 : doer...@cs.uni-saarland.de
Fax. +49 (0)681 302-3065 : http://www.cdl.uni-saarland.de/people/doerfert
-------------- next part --------------
A non-text attachment was scrubbed...
Name: not available
Type: application/pgp-signature
Size: 213 bytes
Desc: not available
URL: <http://lists.cs.uiuc.edu/pipermail/llvmdev/attachments/20140820/12b3dcf1/attachment.sig>

------------------------------

_______________________________________________
LLVMdev mailing list
LLV...@cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev


End of LLVMdev Digest, Vol 122, Issue 56
****************************************
_______________________________________________
LLVM Developers mailing list
LLV...@cs.uiuc.edu http://llvm.cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev

[Tom Stellard]

On Wed, Aug 20, 2014 at 05:31:26PM +0000, Lowell, Daniel wrote:
>
> Hi all,
>
> Running LLVM 3.4 to create a custom pass for OpenCL transformations. I am attempting to GEP into a struct using IRBuilder's CreateStructGEP, but I keep getting this assert:
>
> aoc: ../../../../../../compiler/llvm/include/llvm/Instructions.h:703: llvm::Type* llvm::checkGEPType(llvm::Type*): Assertion `Ty && "Invalid GetElementPtrInst indices for type!"' failed.
>

I've hit this assertion before, and it was because I was passing
the wrong number of operands to the GEP instruction.

Can you add

cB->getType()->getPointerElementType()->dump();
cB->dump();

to your pass right before you call CreateStructGEP and show us
what the output looks like.

-Tom

[Lowell, Daniel]

cB->getType()->getPointerElementType()->dump();


gives:

%struct.RB = type opaque2189 x_idx = builder.CreateStructGEP(cB, 0);



cB->dump()

gives:

%struct.RB addrspace(1)* %cB
$1 = void

//To undo confusion, the last cB is the function arg name. For this example I unfortunately chose the same name for my argument Value* as the argument name.
//Also one correction. I am running LLVM 3.2 not 3.4



________________________________________
From: Tom Stellard [t...@stellard.net]
Sent: Wednesday, August 20, 2014 2:39 PM
To: Lowell, Daniel
Cc: llv...@cs.uiuc.edu
Subject: Re: [LLVMdev] LLVM CreateStructGEP type assert error

[Tom Stellard]

On Wed, Aug 20, 2014 at 08:12:58PM +0000, Lowell, Daniel wrote:
>
> cB->getType()->getPointerElementType()->dump();
>
>
> gives:
>
> %struct.RB = type opaque2189 x_idx = builder.CreateStructGEP(cB, 0);
>

Ok, sorry. I thought that would dump the full struct type. Can you dump the
struct as it is defined in LLVM IR, usually if you dump the whole module,
this will be at the top.

-Tom

[Lowell, Daniel]

If I do M.dump(), at the top of the output I have:

%struct.RB = type opaque


Further down I have:

@.str18 = internal addrspace(2) constant [13 x i8] c"RB_t*\00"


However nowhere does it dump the full struct type when I call "M.dump()". I have it explicitly defined above the kernel in the kernel file, but LLVM doesn't seem to pick it up.
Opaque is a placeholder until it can resolve the "actual" type, so at my pass level "optimization pass" it hasn't resolved it.

________________________________________
From: Tom Stellard [t...@stellard.net]

Sent: Wednesday, August 20, 2014 3:43 PM

[Tom Stellard]

On Wed, Aug 20, 2014 at 08:55:32PM +0000, Lowell, Daniel wrote:
> If I do M.dump(), at the top of the output I have:
>

Can you post the full .cl file you are compiling?

-Tom

[Lowell, Daniel]

Sure, it is the Rodinia 2.4 Hotspot benchmark OpenCL kernel (not my kernel), with the addition of my struct as the last argument in the kernel function.



//------- kernel file start -------------------------------

#define BLOCK_SIZE 16


//dlowell's type

#define BUFFER_LEN 0x100000
typedef struct RB{
unsigned int x;
unsigned int y;
int z[BUFFER_LEN];
unsigned int xx[BUFFER_LEN];
unsigned int yy[BUFFER_LEN];
float zz[BUFFER_LEN];

} RBr_t;

#define IN_RANGE(x, min, max) ((x)>=(min) && (x)<=(max))

__kernel void hotspot( int iteration, //number of iteration
global float *power, //power input
global float *temp_src, //temperature input/output
global float *temp_dst, //temperature input/output
int grid_cols, //Col of grid
int grid_rows, //Row of grid
int border_cols, // border offset
int border_rows, // border offset
float Cap, //Capacitance
float Rx,
float Ry,
float Rz,
float step,
//dlowell's argument
global RB_t *cB) {

local float temp_on_cuda[BLOCK_SIZE][BLOCK_SIZE];
local float power_on_cuda[BLOCK_SIZE][BLOCK_SIZE];
local float temp_t[BLOCK_SIZE][BLOCK_SIZE]; // saving temporary temperature result

float amb_temp = 80.0f;
float step_div_Cap;
float Rx_1,Ry_1,Rz_1;

int bx = get_group_id(0);
int by = get_group_id(1);

int tx = get_local_id(0);
int ty = get_local_id(1);

step_div_Cap=step/Cap;

Rx_1=1/Rx;
Ry_1=1/Ry;
Rz_1=1/Rz;

// each block finally computes result for a small block
// after N iterations.
// it is the non-overlapping small blocks that cover
// all the input data

// calculate the small block size
int small_block_rows = BLOCK_SIZE-iteration*2;//EXPAND_RATE
int small_block_cols = BLOCK_SIZE-iteration*2;//EXPAND_RATE

// calculate the boundary for the block according to
// the boundary of its small block
int blkY = small_block_rows*by-border_rows;
int blkX = small_block_cols*bx-border_cols;
int blkYmax = blkY+BLOCK_SIZE-1;
int blkXmax = blkX+BLOCK_SIZE-1;

// calculate the global thread coordination
int yidx = blkY+ty;
int xidx = blkX+tx;

// load data if it is within the valid input range
int loadYidx=yidx, loadXidx=xidx;
int index = grid_cols*loadYidx+loadXidx;

if(IN_RANGE(loadYidx, 0, grid_rows-1) && IN_RANGE(loadXidx, 0, grid_cols-1)){
temp_on_cuda[ty][tx] = temp_src[index]; // Load the temperature data from global memory to shared memory
power_on_cuda[ty][tx] = power[index];// Load the power data from global memory to shared memory
}
barrier(CLK_LOCAL_MEM_FENCE);

// effective range within this block that falls within
// the valid range of the input data
// used to rule out computation outside the boundary.
int validYmin = (blkY < 0) ? -blkY : 0;
int validYmax = (blkYmax > grid_rows-1) ? BLOCK_SIZE-1-(blkYmax-grid_rows+1) : BLOCK_SIZE-1;
int validXmin = (blkX < 0) ? -blkX : 0;
int validXmax = (blkXmax > grid_cols-1) ? BLOCK_SIZE-1-(blkXmax-grid_cols+1) : BLOCK_SIZE-1;

int N = ty-1;
int S = ty+1;
int W = tx-1;
int E = tx+1;

N = (N < validYmin) ? validYmin : N;
S = (S > validYmax) ? validYmax : S;
W = (W < validXmin) ? validXmin : W;
E = (E > validXmax) ? validXmax : E;

bool computed;
for (int i=0; i<iteration ; i++){
computed = false;

if( IN_RANGE(tx, i+1, BLOCK_SIZE-i-2) && \
IN_RANGE(ty, i+1, BLOCK_SIZE-i-2) && \
IN_RANGE(tx, validXmin, validXmax) && \
IN_RANGE(ty, validYmin, validYmax) ) {

computed = true;
temp_t[ty][tx] = temp_on_cuda[ty][tx] + step_div_Cap * (power_on_cuda[ty][tx] +
(temp_on_cuda[S][tx] + temp_on_cuda[N][tx] - 2.0f * temp_on_cuda[ty][tx]) * Ry_1 +
(temp_on_cuda[ty][E] + temp_on_cuda[ty][W] - 2.0f * temp_on_cuda[ty][tx]) * Rx_1 +
(amb_temp - temp_on_cuda[ty][tx]) * Rz_1);

}
barrier(CLK_LOCAL_MEM_FENCE);

if(i==iteration-1)
break;
if(computed) //Assign the computation range
temp_on_cuda[ty][tx]= temp_t[ty][tx];

barrier(CLK_LOCAL_MEM_FENCE);
}

if (computed){
temp_dst[index]=temp_t[ty][tx];
}
}

//-------------- end kernel file -----------------------------------

________________________________________
From: Tom Stellard [t...@stellard.net]

Sent: Wednesday, August 20, 2014 4:40 PM

[Tom Stellard]

Hi,

When I compile this with clang, I get:

%struct.RB = type { i32, i32, [1048576 x i32], [1048576 x i32], [1048576 x i32], [1048576 x float] }

so I'm not sure why it is showing up as an opaque type for you.

-Tom

[Lowell, Daniel]

Ok, thanks for your help anyways. Perhaps AMDdev who put together this particular stack have an idea.
What compiler / version were you using?

[Tom Stellard]

On Wed, Aug 20, 2014 at 10:03:18PM +0000, Lowell, Daniel wrote:
> Ok, thanks for your help anyways. Perhaps AMDdev who put together this particular stack have an idea.
> What compiler / version were you using?
>

I tested it with clang 3.2 and also clang 3.6. Here is how I did it:

/clang -emit-llvm -include /usr/local/include/clc/clc.h -I/usr/local//include/ -Dcl_clang_storage_class_specifiers -c test.cl -S -o -

You can get the libclc headers from here:
https://github.com/llvm-mirror/libclc

-Tom