Peter Collingbourne
2010-Dec-16 22:33 UTC
[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
David Neto
2010-Dec-17 21:21 UTC
[LLVMdev] Function-level metadata for OpenCL (was Re: OpenCL support)
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 >
Nick Lewycky
2010-Dec-17 22:16 UTC
[LLVMdev] [cfe-dev] Function-level metadata for OpenCL (was Re: OpenCL support)
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>
David Neto
2010-Dec-20 19:47 UTC
[LLVMdev] [cfe-dev] Function-level metadata for OpenCL (was Re: OpenCL support)
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
Peter Collingbourne
2010-Dec-20 20:11 UTC
[LLVMdev] Function-level metadata for OpenCL (was Re: OpenCL support)
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
Reasonably Related Threads
- [LLVMdev] Function-level metadata for OpenCL (was Re: OpenCL support)
- [LLVMdev] Function-level metadata for OpenCL (was Re: OpenCL support)
- [LLVMdev] Function-level metadata for OpenCL (was Re: OpenCL support)
- [LLVMdev] Function-level metadata for OpenCL (was Re: OpenCL support)
- [LLVMdev] [cfe-dev] Function-level metadata for OpenCL (was Re: OpenCL support)