similar to: [LLVMdev] How to generate constant memory for ptx code by NVPTX?

Address space 4 globals should be going into .const. Please open a bug report and provide a repro .ll file. On Fri, Jun 7, 2013 at 11:37 AM, Antony Yu <swpenim at gmail.com> wrote: > Hello, > > I work on compiling OpenCL kernel to PTX code by clang and NVPTX with > libclc. > I have a kernel that contains constant variable declared in file scope like > this: > >

Dear LLVM, FWIW, attached presentation provides a simple jump-start tutorial for newbies to start doing something useful with LLVM, DragonEgg, NVPTX backend and some custom high-level language. Along with short software overview from the user perspective it contains instructions on building components necessary to experiment with NVPTX and an example usecase of generating PTX assembler out of the

Hi Dima, > FWIW, attached presentation provides a simple jump-start tutorial for newbies to > start doing something useful with LLVM, DragonEgg, NVPTX backend and some custom > high-level language. Along with short software overview from the user > perspective it contains instructions on building components necessary to > experiment with NVPTX and an example usecase of generating

On Mon, Feb 4, 2013 at 10:46 AM, <nkavv at physics.auth.gr> wrote: > Hi all, > > I'm trying to use the newly added (in LLVM 3.2) NVPTX backend for > producing PTX (Parallel Thread eXecution) assembly from simple C programs. > > While using llc with -march for mips and x86 works, -march=nvptx doesn't > work. This seems reasonable since I can see that the >

Hi all, I'm trying to use the newly added (in LLVM 3.2) NVPTX backend for producing PTX (Parallel Thread eXecution) assembly from simple C programs. While using llc with -march for mips and x86 works, -march=nvptx doesn't work. This seems reasonable since I can see that the libLLVMNVPTXAsmPrinter.a library is about 500 bytes (thus empty). However, the strange thing is that

I'm sorry if it is too funny, I just thought it could be a simple enough basic tutorial. If you are asking whether many people are interested particularly in Ada, then I don't think so. Ada here is just an illustration of "whatever unusual thing you need - you can do it with LLVM easily by pipelining different components". Anyway, if you think it could be useful for the website,

On Fri, Jun 15, 2012 at 09:06:03AM +0200, Duncan Sands wrote: > Hi Dima, > > > FWIW, attached presentation provides a simple jump-start tutorial for newbies to > > start doing something useful with LLVM, DragonEgg, NVPTX backend and some custom > > high-level language. Along with short software overview from the user > > perspective it contains instructions on

Hi Nikolaos, Following commands work great for me. $ clang -S -emit-llvm -target nvptx -x cl -include clc/clctypes.h ../data-types/scalar.cl $ llc -mcpu=sm_30 scalar.s You can follow Justin's blog [1]. It helped me a lot to understand where to start. [1] http://jholewinski.org/blog/llvm-3-0-ptx-backend/ Best, Ankur On Mon, Feb 4, 2013 at 11:40 PM, Justin Holewinski < justin.holewinski

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

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,

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 Mon, Feb 4, 2013 at 1:09 PM, <nkavv at physics.auth.gr> wrote: > Hi Justin, > > > Has anyone had similar problems with the NVPTX backend? Shouldn't this >>> code be linked to the AsmPrinter library for NVPTX (already)? >>> >> >> What do you mean by "doesn't work"? The AsmPrinter library really houses >> the MCInst

Alright, couple of points here: 1. Address space 0 is invalid for global variables. This is causing a crash in llc where we use llvm_unreachable() on this case. This is most likely why you're seeing llc run forever. The fix for this is to use address space 1 for globals, which puts them into PTX global memory. On our side, we should provide a meaningful error message in this case. 2. The

Hi Justin, >> Has anyone had similar problems with the NVPTX backend? Shouldn't this >> code be linked to the AsmPrinter library for NVPTX (already)? > > What do you mean by "doesn't work"? The AsmPrinter library really houses > the MCInst printer, which isn't implemented for NVPTX yet. The older > assembly printer works just fine. This is

Hi, > Can you post the llc command line you're using? Can you post an LLVM IR > file that causes this behavior? yes: ${LLVM_PATH}/bin/llc -o helloworld.s -march=nvptx helloworld.ll where LLVM_PATH my local installation path for LLVM. Also attaching helloworld.c: #include <stdio.h> int main(void) { printf("Hello World!\n"); return 0; } and helloworld.ll:

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

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

In ptx, variables need to be defined before referenced. NVPTX emits the global variables in the order as in the LLVM IR and does not sort them. It is a bug in the NVPTX backend. Thanks. Yuan From: Dmitry N. Mikushin [mailto:maemarcus at gmail.com] Sent: Wednesday, July 18, 2012 7:44 AM To: LLVM-Dev Cc: Justin Holewinski; Yuan Lin Subject: [NVPTX] PTXAS - Unimplemented feature: labels as

On 4/24/2012 1:50 PM, Justin Holewinski wrote: > > Hi LLVMers, > > We at NVIDIA would like to contribute back to the LLVM open-source > community by up-streaming the NVPTX back-end for LLVM. This back-end > is based on the sources used by NVIDIA, and currently provides > significantly more functionality than the current PTX back-end. Some > functionality is currently

Dear NVPTX community, PTXAS fails to compile the ptx code generated by NVPTX. Is it an issue of backend or an issue of PTXAS or a known reasonable restriction? Thanks, - Dima. > cat test.ll ; ModuleID = '__kernelgen_main_module' target datalayout = "e-p:64:64-i64:64:64-f64:64:64-n1:8:16:32:64" target triple = "ptx64-unknown-unknown" %struct.__st_parameter_dt.0.4