similar to: [LLVMdev] Emitting IR in older formats (for NVVM)

Since SPIR can be (easily) transformed to NVVM IR at least for me this helps a lot. Thank you Tobias. -MH On January 12, 2015, Tobias Grosser <tgrosser at inf.ethz.ch> wrote: > On 12.01.2015 05:48, Jonathan Ragan-Kelley wrote: > > This question is specifically motivated by the practical constraints of > > NVVM, but I don't know anywhere better to ask (hopefully, e.g.,

Thanks, all. I didn’t realize a 7.0 RC was public and changed to 3.4—I will go down that road for now, though I’ll probably also look into integrating variants of the SPIR converter in the future. Another possibility is to skip libnvvm altogether and use LLVM's NVPTX target.  This is of course harder since you have to configure the passes yourself instead of just calling a few C

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

I can't seem to find a direct answer. If I want to convert NVVM IR to PTX code, do I have to use libNVVM or can I use a publicly released version of LLVM? -------------- next part -------------- An HTML attachment was scrubbed... URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/20130131/2067682b/attachment.html>

After reading about Intel's 'Shevlin Park' project to implement C++AMP in llvm/clang, and failing to find any code for it, I decided to try to implement something similar. I did it as an excuse to explore and hack on llvm/clang, which I hadn't done before, but it's now at the point where it will run the simplest matrix multiplication sample from MSDN, so I thought I might

I have written test.ll as below and ran 'opt' on it as " opt -std-compile-opts test.ll -S -o -" . But the output shows that there is code motion around the barrier intrinsics. test.ll ------- ; ModuleID = 'test.bc' define void @test(i16* %I_0, i16* %I_1, i16* %I_2, i16* %I_3, i16* %O_0) { entry: %T_0 = load volatile i16* %I_0 %T_1 = load volatile i16* %I_1 %T_2 =

> On Jul 31, 2015, at 1:01 PM, Eric Christopher <echristo at gmail.com> wrote: > > > > On Fri, Jul 31, 2015 at 12:40 PM Eric Christopher <echristo at gmail.com <mailto:echristo at gmail.com>> wrote: > On Fri, Jul 31, 2015 at 12:22 PM Owen Anderson <resistor at mac.com <mailto:resistor at mac.com>> wrote: > >> On Jul 31, 2015, at 11:28

Hey Jingyue, Though I tried `opt -nvvm-reflect` on both bc files, the nvvm reflect anchor didn't go away; ptxas is still complaining about the duplicate definition of of function '_ZL21__nvvm_reflect_anchorv' . Did I misused the nvvm-reflect pass? Thanks! yuanfeng On Fri, Mar 11, 2016 at 10:10 AM, Jingyue Wu <jingyue at google.com> wrote: > According to the examples you

> -----Original Message----- > From: Dan Bailey [mailto:dan at dneg.com] > Sent: Sunday, April 29, 2012 8:46 AM > To: Justin Holewinski > Cc: Jim Grosbach; llvm-commits at cs.uiuc.edu; Vinod Grover; > llvmdev at cs.uiuc.edu > Subject: Re: [llvm-commits] [PATCH][RFC] NVPTX Backend > > Justin, > > Firstly, this is great! It seems to be so much further forward in

Hi everybody, I have noticed that the new NVPTX backend requires new metadata to identify the kernels in the module: define void @metadata_kernel(float* %a) { ret void } !nvvm.annotations = !{!1} !1 = metadata !{void (float*)* @metadata_kernel, metadata !"kernel", i32 1} Is clang going to support the emission of this metadata soon ? Or do I have to write it on my own ? :) Thanks,

is there any guarantee that the nvptx intrinsic "llvm.nvvm.barrier0" will not be moved around by opt ? In other words, can I expect all the instructions above "llvm.nvvm.barrier0" to remain above it and those below it to remain below, after all the opt passes are run ? If that is not the case, is there a way to define such an intrinsic ? Thanks. -------------- next part

Hey Jingyue, Thanks for being so responsive! I finally figured out a way to resolve the issue: all I have to do is to use `-only-needed` when merging the device bitcodes with llvm-link. However, since we actually need to instrument the host code as well, I encountered another issue when I tried to glue the instrumented host code and fatbin together. When I only instrumented the device code, I

Hi Jingyue, Sorry to ask again, but how exactly could I glue the fatbin with the instrumented host code? Or does it mean we actually cannot instrument both the host & device code at the same time? Thanks! yuanfeng On Tue, Mar 15, 2016 at 10:09 AM, Jingyue Wu <jingyue at google.com> wrote: > Including fatbin into host code should be done in frontend. > > On Mon, Mar 14, 2016

Hi all, For loads and stores i want to extract information about the number of indices accessed. For instance: struct S {int X, int *Y}; __global__ void kernel(int *A, int **B, struct S) {   int x = A[..][..]; // -> L: A[..][..]   int y = *B[2];   // -> L: B[0][2]   int z = S.y[..];  // -> L: S.1[..]   // etc.. } I am performing some preprocessing on IR to: 1. Move constant

Hello, FYI, this is a bug http://llvm.org/bugs/show_bug.cgi?id=13324 When compiling the following code for sm_20, func params are by some reason given with .align 0, which is invalid. Problem does not occur if compiled for sm_10. > 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 =

----- Original Message ----- > From: corngood at gmail.com > To: llvmdev at cs.uiuc.edu > Sent: Saturday, April 13, 2013 9:13:57 PM > Subject: [LLVMdev] C++AMP -> OpenCL (NVPTX) prototype > > After reading about Intel's 'Shevlin Park' project to implement > C++AMP in > llvm/clang, and failing to find any code for it, I decided to try to > implement >

<!DOCTYPE html PUBLIC "-//W3C//DTD HTML 4.01 Transitional//EN"> <html> <head> <meta content="text/html;charset=ISO-8859-1" http-equiv="Content-Type"> </head> <body bgcolor="#ffffff" text="#000000"> Justin,<br> <br> Firstly, this is great! It seems to be so much further forward in terms of features

This new metadata format is currently optional. The old ptx_kernel calling convention should still work. The only thing you should have to change when converting from PTX -> NVPTX is the address space map. The calling conventions and intrinsics should be compatible with both. > -----Original Message----- > From: llvmdev-bounces at cs.uiuc.edu [mailto:llvmdev-bounces at cs.uiuc.edu]

Hi, I wonder whether the current version of LLVM supports separate compilation and linking of device code, i.e., is there a flag analogous to nvcc's --relocatable-device-code flag? If not, is there any plan to support this? Thanks! Yuanfeng Peng -------------- next part -------------- An HTML attachment was scrubbed... URL:

<!DOCTYPE html PUBLIC "-//W3C//DTD HTML 4.01 Transitional//EN"> <html> <head> <meta content="text/html;charset=ISO-8859-1" http-equiv="Content-Type"> </head> <body bgcolor="#ffffff" text="#000000"> Justin Holewinski wrote: <blockquote