Displaying 20 results from an estimated 5000 matches similar to: "[LLVMdev] Status of PTX Backend"
2010 Oct 07
1
[LLVMdev] Status of PTX Backend
Hi,
The PTX backend we developed (CBackend approach, does not use the target
independent code generator) is already more advanced.
An older version is published here:
http://sourceforge.net/projects/llvmptxbackend/
We recently eliminated a bug which increased the number of required
registers per thread. Surprisingly, without that bug the generated code
is already comparable to code generated
2010 Oct 06
0
[LLVMdev] Status of PTX Backend
Hi Justin,
I am upstreaming the PTX backend.
My plan is to have a working prototype (that mean you may compile
non-trivial code with some workarounds) by the end of this year or by
January 2011.
I hope I could catch up next release of LLVM (version 2.9), so I will
adjust my plan according to the release schedule once it is announced.
Regards,
Che-Liang
On Wed, Oct 6, 2010 at 11:16 PM, Justin
2013 Dec 09
1
[LLVMdev] PTX generation examples?
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?
2011 Mar 28
5
[LLVMdev] GSoC: PTX Back-End for LLVM
Hi All,
I am going to submit a GSoC proposal for LLVM this year, and I would like to
first post it here to get constructive feedback before I submit it before
the April 8 deadline. This is the first time I have submitted a GSoC
proposal, so please be brutal with the feedback. :)
Additionally, Che-Liang Chiou (the code owner of the PTX back-end) has
agreed to be my mentor if this is accepted.
2011 May 13
2
[LLVMdev] [ptx] Propose a register class naming convention change
On Fri, May 13, 2011 at 5:11 AM, Dan Bailey <drb at dneg.com> wrote:
> That's fine with me. Unless there's a particular reason for it I would
> suggest perhaps changing the immediate syntax as well to swap it round, so
> it would be Immi32, Immi64, Immf32, etc. It doesn't bother me that much the
> way it currently is, but when there are lots of operations taking a
2011 May 13
1
[LLVMdev] [ptx] Propose a register class naming convention change
2011/5/13 Dan Bailey <drb at dneg.com>
> Justin Holewinski wrote:
>
> On Fri, May 13, 2011 at 5:11 AM, Dan Bailey <drb at dneg.com> wrote:
>
>> That's fine with me. Unless there's a particular reason for it I would
>> suggest perhaps changing the immediate syntax as well to swap it round, so
>> it would be Immi32, Immi64, Immf32, etc. It
2011 Oct 24
1
[LLVMdev] Function pointer parameters in PTX backend
Hi everybody,
I am trying to produce ptx code starting from OpenCL C.
I am experiencing a problem concerning pointer parameters.
Here follows an example:
kernel void function(__global float* parameter1) {}
NVIDIA NVCC Compiler:
.entry function(
.param .u32 *.ptr* .global .align 4 function_param_0
)
{
ret;
}
CLANG + LLVM PTX backend
// (skipping builtin functions definitions)
.entry
2012 Nov 08
3
[LLVMdev] translating from OpenMP to CUDA
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/
2012 Nov 09
0
[LLVMdev] translating from OpenMP to CUDA
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:
2011 May 13
0
[LLVMdev] [ptx] Propose a register class naming convention change
<!DOCTYPE html PUBLIC "-//W3C//DTD HTML 4.01 Transitional//EN">
<html>
<head>
<meta content="text/html;charset=UTF-8" http-equiv="Content-Type">
</head>
<body bgcolor="#ffffff" text="#000000">
Justin Holewinski wrote:
<blockquote
cite="mid:BANLkTi=Y9EFmWRu-9dQxydq8zTyF7tEbJw@mail.gmail.com"
2015 Aug 21
2
[CUDA/NVPTX] is inlining __syncthreads allowed?
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
2011 May 13
0
[LLVMdev] [ptx] Propose a register class naming convention change
That's fine with me. Unless there's a particular reason for it I would
suggest perhaps changing the immediate syntax as well to swap it round,
so it would be Immi32, Immi64, Immf32, etc. It doesn't bother me that
much the way it currently is, but when there are lots of operations
taking a register and an immediate, representing them in the same way
might be a little more
2011 May 13
3
[LLVMdev] [ptx] Propose a register class naming convention change
Hi,
Current register class naming has a confusing prefix letter 'R' (it is
my bad), such as the first 'R' of RRegu32 (for unsigned 32-bit
registers).
I propose a 'Reg' + type name naming convention for register classes; such as:
Regu16, Regu32, Regf32, Regf64
With one exception for predicate registers (capitalized first letter of 'pred'):
RegPred
Since
2012 Sep 04
2
[LLVMdev] [NVPTX] Backend cannot handle array-of-arrays constant
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
2011 May 17
1
[LLVMdev] [LLVMDev] Add not instruction to PTX backend
> Though, I have to agree with Dan on assessing whether the selection logic is
> needed. Do you have an example where the PTX back-end cannot generate code
> for some piece of LLVM IR because of the lack of 'not' selection?
Honestly, I don't have such example yet. Just want to try to implement
some instructions by myself. :p
Regards,
chenwj
--
Wei-Ren Chen (陳韋任)
2015 Aug 21
3
[CUDA/NVPTX] is inlining __syncthreads allowed?
Hi Justin,
Is a compiler allowed to inline a function that calls __syncthreads? I saw
nvcc does that, but not sure it's valid though. For example,
void foo() {
__syncthreads();
}
if (threadIdx.x % 2 == 0) {
...
foo();
} else {
...
foo();
}
Before inlining, all threads meet at one __syncthreads(). After inlining
if (threadIdx.x % 2 == 0) {
...
__syncthreads();
} else {
...
2013 Jan 11
4
[LLVMdev] Update PTX section in CodeGenerator.html
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
2012 May 01
2
[LLVMdev] [llvm-commits] [PATCH][RFC] NVPTX Backend
> -----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
2011 Nov 16
4
[LLVMdev] PTX builtin functions.
Dear Justin,
I am trying to add the support for some OpenCL builtin functions to
the PTX backend.
The attached file represent the first stub of a patch for the fmax
builtin function.
The test case I am trying is the following:
define ptx_device float @f(float %x, float %y) {
entry:
%z = call float @fmax(float %x, float %y)
ret float %z
}
declare float @fmax(float, float)
But at the moment
2013 Dec 06
2
[LLVMdev] PTX generation examples?
OK, fine -- an example of MCJIT that sets up for PTX JIT would also be helpful.
On Dec 6, 2013, at 12:32 PM, Eli Bendersky <eliben at google.com> wrote:
>
> You'll have to switch to MCJIT for this purpose. Legacy JIT doesn't emit PTX.
>
> Eli
--
Larry Gritz
lg at larrygritz.com
-------------- next part --------------
An HTML attachment was scrubbed...
URL: