llvm dev - Dec 2010 - [LLVMdev] Function-level metadata for OpenCL (was Re: OpenCL support)

On Thu, Dec 16, 2010 at 06:16:25PM -0000, Anton Lokhmotov
wrote:
> Mike Gist wrote:
> > You could also consider placing all kernel functions in a
'kernel'
> > section, or adding a function attribute for kernels.
> Unlike in Clang, function attribute bit-fields in LLVM are pretty crowded
> (only couple of bits are unused?).  Besides, we do not want to represent
> differently the kernel qualifier and optional kernel qualifiers, which
> require storing <typen> and X, Y, Z values. (I don't even want to
think how
> one would mangle the optional qualifiers into a kernel name and then
> demangle.)
>
> That's why we propose to use metadata for this purpose.  Does anyone
have a
> better idea?
I agree that metadata should be used for function qualifiers;
a prerequisite being support for non-discardable function-level
metadata, which would need to be added to LLVM.  I'm undecided on
whether __kernel should also be represented by metadata; there is
precedent (PTX backend) for using the calling convention.

I do have a concern though with the semantics of the inliner when it
needs to inline a function with metadata.  One possibility would be to
discard the callee's metadata, or somehow merge it with the caller's.
Discarding seems like the right solution for OpenCL and a good starting
point (in future we may wish to add attributes to metadata nodes like
the 'appending' linkage for globals) but sounds like something that
should be discussed first.

Thanks,
-- 
Peter

However we record the fact that a function is a kernel, the mechanism
should handle the case of a kernel calling another kernel.
Recall that a kernel called by another kernel behaves more like a
regular function.  For example it doesn't have workspace iteration
automatically applied to it; rather it just adopts the work item of
the caller.

About using a calling convention to mark a function as a kernel.  It
seems a handy place to hang it, but is it really exclusive of the
other calling conventions?
In particular, does that approach nicely in the case where a CPU is
running the kernels?  Does that lead to special casing or duplication
in the code generator?  For example, you still have to know what
"real" calling convention to use when a kernel is running on a CPU.
(Forgive my ignorance.)


thanks,
david


On Thu, Dec 16, 2010 at 5:33 PM, Peter Collingbourne <peter at pcc.me.uk>
wrote:
> On Thu, Dec 16, 2010 at 06:16:25PM -0000, Anton Lokhmotov wrote:
>> Mike Gist wrote:
>> > You could also consider placing all kernel functions in a
'kernel'
>> > section, or adding a function attribute for kernels.
>> Unlike in Clang, function attribute bit-fields in LLVM are pretty
crowded
>> (only couple of bits are unused?).  Besides, we do not want to
represent
>> differently the kernel qualifier and optional kernel qualifiers, which
>> require storing <typen> and X, Y, Z values. (I don't even
want to think how
>> one would mangle the optional qualifiers into a kernel name and then
>> demangle.)
>>
>> That's why we propose to use metadata for this purpose.  Does
anyone have a
>> better idea?
>
> I agree that metadata should be used for function qualifiers;
> a prerequisite being support for non-discardable function-level
> metadata, which would need to be added to LLVM.  I'm undecided on
> whether __kernel should also be represented by metadata; there is
> precedent (PTX backend) for using the calling convention.
>
> I do have a concern though with the semantics of the inliner when it
> needs to inline a function with metadata.  One possibility would be to
> discard the callee's metadata, or somehow merge it with the
caller's.
> Discarding seems like the right solution for OpenCL and a good starting
> point (in future we may wish to add attributes to metadata nodes like
> the 'appending' linkage for globals) but sounds like something that
> should be discussed first.
>
> Thanks,
> --
> Peter
> _______________________________________________
> LLVM Developers mailing list
> LLVMdev at cs.uiuc.edu         http://llvm.cs.uiuc.edu
> http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
>

On 16 December 2010 14:33, Peter Collingbourne <peter at pcc.me.uk> wrote:
> On Thu, Dec 16, 2010 at 06:16:25PM -0000, Anton Lokhmotov wrote:
> > Mike Gist wrote:
> > > You could also consider placing all kernel functions in a
'kernel'
> > > section, or adding a function attribute for kernels.
> > Unlike in Clang, function attribute bit-fields in LLVM are pretty
crowded
> > (only couple of bits are unused?).  Besides, we do not want to
represent
> > differently the kernel qualifier and optional kernel qualifiers, which
> > require storing <typen> and X, Y, Z values. (I don't even
want to think
> how
> > one would mangle the optional qualifiers into a kernel name and then
> > demangle.)
> >
> > That's why we propose to use metadata for this purpose.  Does
anyone have
> a
> > better idea?
>
> I agree that metadata should be used for function qualifiers;
> a prerequisite being support for non-discardable function-level
> metadata, which would need to be added to LLVM.  I'm undecided on
> whether __kernel should also be represented by metadata; there is
> precedent (PTX backend) for using the calling convention.
>
Being discardable is a design point of metadata. You might add something
else to support this, but it won't be metadata.

Why are you trying to preserve "kernel"-ness into the LLVM IR? What
semantics does it have? What does __kernel actually mean to the optimizers
and code generator?

Could you just make __kernel mean "externally visible" and undecorated
functions be "linkonce_odr"? If that's not enough, could you swing
it around
and maintain single named metadata node with a list of functions that are
marked __kernel?

Nick

I do have a concern though with the semantics of the inliner when
it
> needs to inline a function with metadata.  One possibility would be to
> discard the callee's metadata, or somehow merge it with the
caller's.
> Discarding seems like the right solution for OpenCL and a good starting
> point (in future we may wish to add attributes to metadata nodes like
> the 'appending' linkage for globals) but sounds like something that
> should be discussed first.
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL:
<http://lists.llvm.org/pipermail/llvm-dev/attachments/20101217/a009d416/attachment.html>

On Fri, Dec 17, 2010 at 5:16 PM, Nick Lewycky <nlewycky at google.com>
wrote:
> Being discardable is a design point of metadata. You might add something
> else to support this, but it won't be metadata.
> Why are you trying to preserve "kernel"-ness into the LLVM IR?
What
> semantics does it have? What does __kernel actually mean to the optimizers
> and code generator?
> Could you just make __kernel mean "externally visible" and
undecorated
> functions be "linkonce_odr"? If that's not enough, could you
swing it around
> and maintain single named metadata node with a list of functions that are
> marked __kernel?
> Nick
>>
>> I do have a concern though with the semantics of the inliner when it
>> needs to inline a function with metadata.  One possibility would be to
>> discard the callee's metadata, or somehow merge it with the
caller's.
>> Discarding seems like the right solution for OpenCL and a good starting
>> point (in future we may wish to add attributes to metadata nodes like
>> the 'appending' linkage for globals) but sounds like something
that
>> should be discussed first.
>
Regarding linkage:

A __kernel function is externally visible.  It is callable from the
user program which is logically a separate compilation unit; and from
other functions in its own compilation unit.
The non-kernel functions have private linkage, I believe: they are
only callable by other functions in the same compilation unit.

However, a __kernel behaves differently when called from the user
program vs. another function in the compilation unit.  In OpenCL the
user program can invoke a kernel as an NDRange, i.e. with an implied
loop around it to iterate over an index space of 1 to 3 dimensions.
(This is the "big idea" of OpenCL).  (The index values are available
in the function body from intrinsic functions get_work_dim() and
get_global_id(uint workdim).)
But that implied loop is only applied when directly called from the
user program.  When a kernel is called from another kernel, it behaves
as a regular function call and just adopts the caller's index point.

The spec does not specify whether or how that implied loop is
represented in the IR.  I expect most implementations don't represent
the loop explicitly.

I would be happy to see an OpenCL-specific patch that always marked
non-kernel functions with internal linkage.  Then you could
distinguish the kernel/non-kernel case just by the linkage attribute.
It might be a little unclean / unorthogonal, but I think it would be
ok.

(There are also other minor differences, e.g. the behaviour of a
function-scope-local-addr-space variable in a nested kernel is
implementation-defined.  See the Notes in the functionQualifiers
reference.)


david

References
http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/functionQualifiers.html
http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html
http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/local.html

On Fri, Dec 17, 2010 at 04:21:18PM -0500, David Neto
wrote:
> However we record the fact that a function is a kernel, the mechanism
> should handle the case of a kernel calling another kernel.
> Recall that a kernel called by another kernel behaves more like a
> regular function.  For example it doesn't have workspace iteration
> automatically applied to it; rather it just adopts the work item of
> the caller.
> 
> About using a calling convention to mark a function as a kernel.  It
> seems a handy place to hang it, but is it really exclusive of the
> other calling conventions?
> In particular, does that approach nicely in the case where a CPU is
> running the kernels?  Does that lead to special casing or duplication
> in the code generator?  For example, you still have to know what
> "real" calling convention to use when a kernel is running on a
CPU.
> (Forgive my ignorance.)
As with __local variables, it may be that "kernelness" cannot be
represented in a standard form in LLVM.  For example on a CPU a
kernel function may have an additional parameter which is a pointer to
__local memory space, which would not be necessary on GPUs.  Then in
fact you would use a standard calling convention on a CPU.

But for GPUs, I think using the calling convention is appropriate.
If we standardise the calling convention number, this can be the
default behaviour.
> I would be happy to see an OpenCL-specific patch that always marked
> non-kernel functions with internal linkage.  Then you could
> distinguish the kernel/non-kernel case just by the linkage attribute.
> It might be a little unclean / unorthogonal, but I think it would be
> ok.
Some OpenCL implementations (including my own) may use runtime library
functions which live in a separate compilation unit.  These would
need to be marked external but of course would not be kernel functions.

Nick Lewycky wrote:
> Being discardable is a design point of metadata. You might add something
> else to support this, but it won't be metadata.
There's nothing intrinsic about the concept of metadata which requires
it to be discardable.  In particular, if the metadata is attached to
a function, the only case I can think of where an optimiser needs
to touch the metadata is if a function with metadata is inlined.
And as I mentioned in my previous mail I don't think this will be
any trouble for OpenCL.

The __kernel attribute isn't the only attribute we need to preserve.
There are also:

__attribute__((vec_type_hint(type)))
__attribute__((work_group_size_hint(X, Y, Z)))
__attribute__((reqd_work_group_size(X, Y, Z))) 

which provide hints to the code generator regarding the specific work
load of a particular kernel.
> Why are you trying to preserve "kernel"-ness into the LLVM IR?
What
> semantics does it have? What does __kernel actually mean to the optimizers
> and code generator?
For PTX, if __kernel is set on a function it needs to be codegen'd
with a specific directive which marks it as a kernel entry point.
What this actually means at a lower level I don't know (the low level
machine code representation is undocumented).  I believe there is
also something similar in the AMD Stream IL.

As for the other attributes mentioned above, I don't know off-hand,
but I believe there are PTX directives for at least some of them.
> Could you just make __kernel mean "externally visible" and
undecorated
> functions be "linkonce_odr"?
I think the semantics of undecorated functions is closer to "internal"
than "linkonce_odr" here (kernel programs shouldn't be able to
provide a definition for functions in another module, such as a
runtime library module).
> If that's not enough, could you swing it around
> and maintain single named metadata node with a list of functions that are
> marked __kernel?
Are you saying that named metadata nodes are non-discardable?  Even
if this were true, it would still be difficult to represent the other
attributes unless the metadata were attached to the function.

Thanks,
-- 
Peter