similar to: [LLVMdev] PTX BE status enquiry

Hi Justin, I download llvm-ptx-samples [1] and try to build them. I found it seems lack of a complete document on how to build them with LLVM/Clang/libclc. Do you think it's a good idea to put a complete document/tutorial in _one_ place? Currently, there are your website [2], LLVM [3], Clang and libclc websites [5] over there. I feel people might get lost among those websites. ;-) Here

On Tue, Dec 6, 2011 at 10:17 PM, 陳韋任 <chenwj at iis.sinica.edu.tw> wrote: > Hi Justin, > > I download llvm-ptx-samples [1] and try to build them. I found it seems > lack > of a complete document on how to build them with LLVM/Clang/libclc. Do you > think > it's a good idea to put a complete document/tutorial in _one_ place? > Currently, > there are your

I am trying to generate PTX code for 'nbody' sample program's kernel (nbody_kernel.cu) using clang/LLVM version 3.2. The nbody CUDA program is available in Nvidia's SDK. I am referring to https://github.com/jholewinski/llvm-ptx-samples project. Following are my commands, clang++ -O4 -S -I/usr/local/cuda/include -emit-llvm -target nvptx64 nbody_kernel.cu -o nbody_kernel.ll

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

Hi, Is it possible to translate an OpenMP program to CUDA using LLVM? I read that dragonegg has a OpenMP front-end and LLVM has a PTX back-end. I don't know how mature these tools are. Please let me know. Thanks. -Apala Postdoctoral Scholar Department of Computer Science, University of Chicago Computation Institute, Argonne National Laboratory http://sites.google.com/site/apalaguha/home/

The PTX back-end is robust (it's based on the sources used by nvcc), but I'm not sure about the OpenMP representation in LLVM IR. I believe the OpenMP constructs are already lowered into libgomp calls before leaving DragonEgg. It's been awhile since I've loooked at it though. If you use the PTX back-end and have any issues, please don't hesitate to post to the list and cc:

Hi Hal, Thanks for your comments! I'm inlining my responses below. Jingyue On Sat, Apr 19, 2014 at 6:38 AM, Hal Finkel <hfinkel at anl.gov> wrote: > Jingyue, > > I can't speak for the NVPTX backend, but I think this looks useful as an > (optional) target-independent pass. A few thoughts: > > - Running GVN tends to be pretty expensive; have you tried EarlyCSE

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 disabled due to dependencies on LLVM core changes that we are also in the process of

Hi Justin, Cool stuff, to be sure. Excited to see this. As a pre-cursor to more involved technical feedback, I suggest going through and fixing up the coding style and formatting issues. Just glancing through, I see lots of things like function names starting with capital letters, compound statements with the opening brace on the line following an if/for/while/etc., single-statements after an

Hi, We wrote an optimization that eliminates common sub-expressions in a group of similar GEPs for the NVPTX backend. It speeds up some of our benchmarks by up to 20%, which convinces us to try to upstream it. Here's a brief description of why we wrote this optimization, what we did, and how we did it. Loops in CUDA programs are often extensively unrolled by programmers and compilers,

I'm using 7.0. I am attaching the reduced example. nvcc sync.cu -arch=sm_35 -ptx gives // .globl _Z3foov .visible .entry _Z3foov( ) { .reg .pred %p<2>; .reg .s32 %r<3>; mov.u32 %r1, %tid.x; and.b32 %r2, %r1, 1; setp.eq.b32 %p1, %r2, 1; @!%p1 bra BB7_2; bra.uni

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

Thanks for the feedback! The attached patch addresses the style issues that have been found. From: Jim Grosbach [mailto:grosbach at apple.com] Sent: Wednesday, April 25, 2012 2:22 PM To: Justin Holewinski Cc: llvm-commits at cs.uiuc.edu; llvmdev at cs.uiuc.edu; Vinod Grover Subject: Re: [llvm-commits] [PATCH][RFC] NVPTX Backend Hi Justin, Cool stuff, to be sure. Excited to see this. As a

Hello , How to execute a cuda program using llvm? More specifically, nvcc produces some temporary files during its compilation. I want to convert the .cu.cpp to .ll format and optimize it. The .cu.cpp file contains typedefs and enums used by cuda runtime and also the host part of the code and the ptx file contains the kernel definition. How can i run the program after optimization? Will Rhodin

Hi Adarsh, to my knowledge there is no publicly available CUDA-Frontend for LLVM yet. The work of Helge Rhodin you mentioned is on the backend-side: It allows to generate PTX code from LLVM IR. It is still being maintained, although I think the currently available source code is a little outdated. There is also a PTX backend in the current version of LLVM that makes use of LLVM's

I noticed you're using cuda_runtime.h in the source file. Where are you getting this file? From the CUDA toolkit? Since the error is in the back-end, can you just post the .ll or .bc file you are trying to compile? On Mon, Mar 18, 2013 at 12:42 AM, upit <uday_pitambare at yahoo.com> wrote: > I am trying to generate PTX code for 'nbody' sample program's kernel >

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

I believe it's "Dotty", I simply output "opt -view-regions-only". On Thu, Feb 16, 2012 at 8:22 AM, Schuster, Vince J <vinces at lanl.gov> wrote: > Hi Ryan, > > wrt your llvm email. What do you use to get a visual > display of your .obj file? > > thanks, > Vince Schuster > > > ------------------------------ > > Message: 2 >

Hello, I see some people including myself confused by the different object-oriented approaches in R (S3, S4, OOP, R.oo etc.). Would it be ok to collect examples and solutions for the different OO-packages in one package and add a vignette for documentation? (assuming I find time for this task) I mean in this case the package would not add data or functionality to R or serve as a companion

Justins: I noticed that the intrinsics in IntrinsicsNVVM don't specify a TargetPrefix. This seems like a simple omission, so I was going to simply throw a `let TargetPrefix = "nvvm" ` block around them, but this doesn't quite work. There seem to be three prefixes that are used in this file. About 900 are int_nvvm_*, 30 are int_ptx_*, and 1 is int_cuda. It isn't clear to me