similar to: [LLVMdev] C-to-PTX compilation issues

The issue you're seeing is actually a problem with clang integration. Generally, clang does not understand all of the conventions required by the NVPTX back-end, and compilation from general C may not always work. In this case, it's the local array "b" that is the problem. Clang pulls this out into the global scope, but keeps it in address space 0. In the back-end, address

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

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

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

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

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 >

On Thu, Jan 24, 2013 at 7:20 AM, Ahmed Bougacha <ahmed.bougacha at gmail.com>wrote: > On Thu, Jan 24, 2013 at 12:46 PM, <nkavv at physics.auth.gr> wrote: > > Hi all, > > > > i'm just starting out with LLVM (although i've been observing its > evolution > > since that first release some years ago :) > > > > I would like to develop a

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:

Hi all, i'm just starting out with LLVM (although i've been observing its evolution since that first release some years ago :) I would like to develop a backend for a generic assembly-like language, called NAC (N-Address Code). More info on NAC can be found here: http://www.nkavvadias.com/hercules/nac-refman.html (HTML) http://www.nkavvadias.com/hercules/nac-refman.pdf (PDF) You

On Thu, Jan 24, 2013 at 12:46 PM, <nkavv at physics.auth.gr> wrote: > Hi all, > > i'm just starting out with LLVM (although i've been observing its evolution > since that first release some years ago :) > > I would like to develop a backend for a generic assembly-like language, > called NAC (N-Address Code). More info on NAC can be found here: >

On Mon, Feb 11, 2013 at 6:32 PM, <nkavv at physics.auth.gr> wrote: > Hi Justin and all, > > you've mentioned that you used an easy "trick" for defining an infinite > register file in the backend. > The original PTX back-end just didn't perform register allocation. All registers emitted in the assembly were virtual registers, mapped to a consecutive range.

Hi Justin I just came across your presentation regarding the LLVM 3.0 (?) PTX backend: http://llvm.org/devmtg/2011-11/Holewinski_PTXBackend.pdf Is this the NVPTX backend or the predecessor (in PTX? directory) backend? Best regards Nikolaos Kavvadias

Hi Justin and all, you've mentioned that you used an easy "trick" for defining an infinite register file in the backend. Does this involve defining a single dummy register, and then adding this dummy reg to each register class? Is anything more needed? I refer to the RegisterInfo.td file, as e.g: llvm-3.0.src/lib/Target/PTX/PTXRegisterInfo.td I'm (just starting)

Ah, that's helpful. I knew that I'd need to end up with PTX as text, not a true binary, but I would have figured that it would come out of MCJIT. Thanks for helping to steer me away from the wrong trail. OK, one more question: Can anybody clarify the pros and cons of generating the PTX through the standard LLVM distro, versus using the "libnvvm" that comes with the Cuda SDK?

There is no MCJIT support for PTX at the moment (mainly because PTX does not have a binary format, and is not machine code per se). To generate PTX at run-time, you just set up a standard codegen pass manager like you would like an off-line compiler. The output will be a string buffer that contains the PTX, which you can load into the CUDA runtime. As for determining if PTX support is compiled

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 Justin, I believe the PTX section in the link below need some love, updating "lib/Target/PTX" to "lib/Target/NVPTX" for example. Would you like to take a look? http://llvm.org/docs/CodeGenerator.html#the-ptx-backend Regards, chenwj -- Wei-Ren Chen (陳韋任) Computer Systems Lab, Institute of Information Science, Academia Sinica, Taiwan (R.O.C.) Tel:886-2-2788-3799 #1667

Right, I forgot to update that originally. I also have a separate document in the works that goes into more detail for NVPTX. On Fri, Jan 11, 2013 at 2:57 AM, 陳韋任 (Wei-Ren Chen) < chenwj at iis.sinica.edu.tw> wrote: > Hi Justin, > > I believe the PTX section in the link below need some love, > updating "lib/Target/PTX" to "lib/Target/NVPTX" for example.

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