similar to: [LLVMdev] [PATCH][RFC] Add llvm.codegen Intrinsic To Support Embedded LLVM IR Code Generation

Dear LLVM NVPTX maintainers, Just to have the issue recorded, I don't know how important it is: clang generates linkonce_odr out of __inline__, and NVPTX generates .weak out of linkonce_odr (how it happens - a big question, btw, because I can't find anything related in NVPTX asm printer - does it chain to some other printer?), and finally ptxas (both 4.2 and 5) fails to compile it to

On Tue, Jun 12, 2012 at 6:11 PM, Dmitry N. Mikushin <maemarcus at gmail.com>wrote: > Dear LLVM NVPTX maintainers, > > Just to have the issue recorded, I don't know how important it is: > > clang generates linkonce_odr out of __inline__, and NVPTX generates .weak > out of linkonce_odr (how it happens - a big question, btw, because I can't > find anything related

On 05/08/2012 12:14 AM, dag at cray.com wrote: > Tobias Grosser<tobias at grosser.es> writes: > >>> I forgot to address this one. With current OpenCL and CUDA >>> specifications, there's no need to do multiple .o files. In my mind, >>> llc should output one .o (one .s, etc.). Anything else wreaks havoc on >>> build systems. >> >>

On 05/07/2012 06:07 PM, dag at cray.com wrote: > Tobias Grosser<tobias at grosser.es> writes: > >>> Doesn't LLVM support taking the address of a function in another address >>> space? If not it probably should. >> >> Hi Dave, >> The llvm.codegen intrinsic seems the perfect match to build up such >> experience. It requires no changes to

Tobias Grosser <tobias at grosser.es> writes: >> Doesn't LLVM support taking the address of a function in another address >> space? If not it probably should. > > Hi Dave, > > I highly appreciate your idea of integrating heterogeneous computing > features directly into LLVM-IR. I believe this can be a way worth > going, but I doubt now is the right moment

Hi LLVMers, The attached patch adds a new Intrinsic named "llvm.codegen" to support embedded LLVM IR code generation. **The 'llvm.codegen' intrinsic uses the LLVM back ends to generate code for embedded LLVM IR strings. The code generation target can be same or different to the one of the parent module. The original motivation inspiring us to add this intrinsic, is to

On 04/29/2012 01:21 AM, Justin Holewinski wrote: > > > On Sat, Apr 28, 2012 at 8:27 AM, Tobias Grosser <tobias at grosser.es > <mailto:tobias at grosser.es>> wrote: > regalloc= is different. It is global and consequently influences > both host and device code generation. However, to me it is rather a > debugging option. It is never set by clang and

Justin Holewinski <justin.holewinski at gmail.com> writes: > For something like PTX, runtime calls take care of the call semantics so > it is either up to the user or the frontend to set up the runtime calls > correctly. We don't need to completely solve this problem. Yet. :) > > But there has to be some interface that allows an LLVM IR function > from

Hi, Looks like "{" and "}" are lost when trying to use the combination of Clang and NVPTX, which may result into clash of definitions of the function-scope and asm-scope. Here is an example: > cat test.cu __attribute__((device)) __attribute__((nv_linkonce_odr)) __inline__ int __any(int a) { int result; asm __volatile__ ("{ \n\t" ".reg .pred

On Fri, Apr 27, 2012 at 7:40 PM, Yabin Hu <yabin.hwu at gmail.com> wrote: > Hi LLVMers, > > The attached patch adds a new Intrinsic named "llvm.codegen" to support > embedded LLVM IR code generation. **The 'llvm.codegen' intrinsic uses > the LLVM back ends to generate code for embedded LLVM IR strings. The code > generation target can be same or

Dmitry, You might be better served by filing this as a bug (http://llvm.org/bugs/). Please include a test case and the steps to reproduce (i.e., what you've provided below). Chad On Jul 10, 2012, at 3:15 PM, Dmitry N. Mikushin wrote: > Hi, > > Looks like "{" and "}" are lost when trying to use the combination of Clang and NVPTX, which may result into clash of

On Apr 29, 2012, at 6:37 AM, Tobias Grosser wrote: > > OK, I get what you mean. The intrinsic is currently targeted at the > OpenCL/CUDA model. It is the most widely used. Stuff like cell sounds > interesting, but probably needs further thoughts. Even with OpenCL/CUDA, > this intrinsic works currently only for PTX code generation, but I hope > we can gain support for other

Tobias Grosser <tobias at grosser.es> writes: >> I forgot to address this one. With current OpenCL and CUDA >> specifications, there's no need to do multiple .o files. In my mind, >> llc should output one .o (one .s, etc.). Anything else wreaks havoc on >> build systems. > > Yes, that's what I am advocating for. There is no need for all this >

Yes, sure, good idea, because might be also Clang-related. http://llvm.org/bugs/show_bug.cgi?id=13322 2012/7/11 Chad Rosier <mcrosier at apple.com> > Dmitry, > You might be better served by filing this as a bug (http://llvm.org/bugs/). > Please include a test case and the steps to reproduce (i.e., what you've > provided below). > > Chad > > On Jul 10, 2012,

On 05/01/2012 11:21 PM, dag at cray.com wrote: > Justin Holewinski<justin.holewinski at gmail.com> writes: > >> For something like PTX, runtime calls take care of the call semantics so >> it is either up to the user or the frontend to set up the runtime calls >> correctly. We don't need to completely solve this problem. Yet. :) >> >>

On 05/01/2012 11:21 PM, dag at cray.com wrote: > Justin Holewinski<justin.holewinski at gmail.com> writes: >> Do we allow more than one Module per file? If not, that seems like an >> arbitrary limitation. If we allowed that we could have each module >> specify a different target. >> >> That could work. > > Given your questions about

On 04/29/2012 08:28 PM, Evan Cheng wrote: > > On Apr 29, 2012, at 6:37 AM, Tobias Grosser wrote: > >> >> OK, I get what you mean. The intrinsic is currently targeted at the >> OpenCL/CUDA model. It is the most widely used. Stuff like cell sounds >> interesting, but probably needs further thoughts. Even with OpenCL/CUDA, >> this intrinsic works currently only

Justin Holewinski <justin.holewinski at gmail.com> writes: > I believe the point Tobias is trying to make is that he wants to > retain the ability to pipe modules between tools and not worry about > the modules ever hitting disk, e.g. > > opt -load GPUOptimizer.so -gpu-opt | llc -march=x86 > where the module coming in to opt is just unoptimized host code, and the module

On Sat, Apr 28, 2012 at 8:27 AM, Tobias Grosser <tobias at grosser.es> wrote: > On 04/28/2012 04:30 PM, Justin Holewinski wrote: > >> We can handle this by provide a new argument (e.g. a string of >> properly-configured Target Machine) instead of or in addition to >> the >> Arch type string argument. >> >> >> I think we

On Tue, May 1, 2012 at 8:22 AM, <dag at cray.com> wrote: > Justin Holewinski <justin.holewinski at gmail.com> writes: > > > I don't think the code base changes are all that bad. We have a > number > > of them to support generating code one function at a time rather > than a > > whole module together. They've been sitting around