similar to: [LLVMdev] A simple tutorial on generating PTX assembler out of Ada source code using LLVM NVPTX backend

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 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

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,

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

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

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

Dear all, Looks like the NVPTX backend cannot handle array-of-arrays contant (please see the reporocase below). Is it supposed to work? Any ideas how to get it working? Important for our target applications. Thanks, - Dima. $ cat test.ll ; ModuleID = '__kernelgen_main_module' target datalayout =

I think our test case demonstrates that requiring the array item being initialized to be constant is incorrect. NVPTX does not crash anymore and produces correct result with the following change: --- NVPTXAsmPrinter.cpp 2012-09-03 15:14:00.000000000 +0200 +++ NVPTXAsmPrinter.cpp 2012-09-04 15:47:17.859398193 +0200 @@ -1890,17 +1890,15 @@ case Type::ArrayTyID: case Type::VectorTyID: case

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

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 >

NVCC successfully handles the same IR, if we try to process the same .cu file with clang+nvptx and nvcc: CLANG/NVPTX: ============= $ cat dayofweek.cu __attribute__((device)) char yweek[7][4] = { "MON", "TUE", "WED", "THU", "FRI", "SAT", "SUN" }; $ clang -cc1 -emit-llvm -fcuda-is-device dayofweek.cu -o dayofweek.ll $ cat

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: constant one_f = 1.0f; __kernel void test( ...){ ... } Then it is compiled to llvm-ir: @one_f = addrspace(4) const float 1.000000e+00, align 4 define void test(...){ ... } Finally ptx: .visible .global .align 4 .f32

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

On 09/04/2012 09:57 AM, Dmitry N. Mikushin wrote: > I think our test case demonstrates that requiring the array item being > initialized to be constant is incorrect. NVPTX does not crash anymore > and produces correct result with the following change: > > --- NVPTXAsmPrinter.cpp 2012-09-03 15:14:00.000000000 +0200 > +++ NVPTXAsmPrinter.cpp 2012-09-04 15:47:17.859398193 +0200 >

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

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

Dear Yuan, Sorry for delay with reply, Answers on your questions could be different, depending on the math library placement in the code generation pipeline. At KernelGen, we currently have a user-level CUDA math module, adopted from cicc internals [1]. It is intended to be linked with the user LLVM IR module, right before proceeding with the final optimization and backend. Last few months we

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: > >

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