Displaying 20 results from an estimated 10000 matches similar to: "[LLVMdev] LLVMdev Digest, Vol 83, Issue 25"
2011 May 16
0
[LLVMdev] TargetRegisterInfo and "infinite" register files
Justin,
We have the same issue with the AMDIL code generator. We tried #1, but there are passes after register allocator that don't like virtual registers. #3 could be done by having the two spill functions [load|store]Reg[From|To]StackSlot keep track of the FrameIndex to register mapping internally, but again, more of a hack than a proper solution.
My solution was to just create a very large
2011 May 17
0
[LLVMdev] TargetRegisterInfo and "infinite" register files
I have faced this same problem in my backend, and I'm working around it
by providing a large physical register set. There are two problems with
this:
1. There's a chance that the register allocator will run out of
registers to assign, in which case the allocation will fail - making it
necessary to retry with a larger register set
2. The code generator consumes storage proportional to
2011 May 16
0
[LLVMdev] TargetRegisterInfo and "infinite" register files
On May 16, 2011, at 6:52 AM, Justin Holewinski wrote:
> Currently, the TableGen register info files for all of the back-ends define concrete registers and divide them into logical register classes. I would like to get some input from the LLVM experts around here on how best to map this model to an architecture that does *not* have a concrete, pre-defined register file. The architecture is
2011 May 16
6
[LLVMdev] TargetRegisterInfo and "infinite" register files
Currently, the TableGen register info files for all of the back-ends define
concrete registers and divide them into logical register classes. I would
like to get some input from the LLVM experts around here on how best to map
this model to an architecture that does *not* have a concrete, pre-defined
register file. The architecture is PTX, which is more of an intermediate
form than a final
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 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"
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
2012 Jun 13
0
[LLVMdev] [NVPTX] For linkonce_odr NVPTX generates .weak, but even newest PTXAS can't handle it
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
2012 Jul 18
0
[LLVMdev] [NVPTX] PTXAS - Unimplemented feature: labels as initial values
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
2012 Jun 12
2
[LLVMdev] [NVPTX] For linkonce_odr NVPTX generates .weak, but even newest PTXAS can't handle it
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
2012 Jul 18
2
[LLVMdev] [NVPTX] PTXAS - Unimplemented feature: labels as initial values
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
2014 Apr 21
2
[LLVMdev] [NVPTX] Eliminate common sub-expressions in a group of similar GEPs
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
2012 Jul 10
0
[LLVMdev] [NVPTX] CUDA inline PTX asm definitions scoping "{" "}" is broken
Dmitry,
You might be better served by filing this as a bug (http://llvm.org/bugs/). Please include a test case and the steps to reproduce (i.e., what you've provided below).
Chad
On Jul 10, 2012, at 3:15 PM, Dmitry N. Mikushin wrote:
> Hi,
>
> Looks like "{" and "}" are lost when trying to use the combination of Clang and NVPTX, which may result into clash of
2012 Jul 10
1
[LLVMdev] [NVPTX] CUDA inline PTX asm definitions scoping "{" "}" is broken
Yes, sure, good idea, because might be also Clang-related.
http://llvm.org/bugs/show_bug.cgi?id=13322
2012/7/11 Chad Rosier <mcrosier at apple.com>
> Dmitry,
> You might be better served by filing this as a bug (http://llvm.org/bugs/).
> Please include a test case and the steps to reproduce (i.e., what you've
> provided below).
>
> Chad
>
> On Jul 10, 2012,
2012 Jul 10
2
[LLVMdev] [NVPTX] CUDA inline PTX asm definitions scoping "{" "}" is broken
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
2011 Jun 20
2
[LLVMdev] New Configure Option for LLVM Builds
On Jun 20, 2011, at 11:47 AM, John Criswell wrote:
> On 6/20/11 10:39 AM, Justin Holewinski wrote:
>>
>> For the PTX back-end, I would like to introduce a configure-time option to determine the number of architectural registers available to LLVM register allocation during code generation. The motivation for this is PTX is a virtual instruction set and the number of registers is
2017 Nov 08
2
Debug info for Cuda
Nobody blames ptxas. I'm not saying that these are the troubles, I'm just saying that it has some features and we have some problems to be solved.
But lack of labels, label arithmetics in DWARF sections is the real problem, because LLVM actively uses it in DWARF sections
Best regards,
Alexey Bataev
8 нояб. 2017 г., в 5:35, Madhur Amilkanthwar <madhur13490 at
2017 Nov 06
2
Debug info for Cuda
06.11.2017 14:56, Robinson, Paul пишет:
>> Hi everybody,
>> As you know, Cuda/NVPTX target has very limited support of the debug
>> info in Clang/LLVM. Currently, LLVM supports only emission of the line
>> numbers debug info.
>> This is caused by limitations of the Cuda/NVPTX codegen. Clang/LLVM
>> translates the source code to LLVM IR, which is then lowered to
2011 May 17
0
[LLVMdev] TargetRegisterInfo and "infinite" register files
On Tue, May 17, 2011 at 2:52 PM, Jakob Stoklund Olesen <stoklund at 2pi.dk>wrote:
>
> On May 17, 2011, at 11:32 AM, Andrew Clinton wrote:
>
> > On 05/17/2011 12:54 PM, Jakob Stoklund Olesen wrote:
> >> What you can do instead is:
> >>
> >> 1) Just use virtual registers and skip register allocation, or
> >>
> >> 2) Allocate to a
2016 Apr 08
2
[GPUCC] how to remove _ZL21__nvvm_reflect_anchorv() automatically?
Yeah, '.' is the direct reason for the ptxas failure here. I'm curious,
however, about what the purpose of nvvm_reflect_anchorv() is here, and why
does the front-end always generate this function? Since the current PTX
emission doesn't mangle dots, it would be a reasonable workaround for me to
prevent the front-end from generating this function in the first place.
Is there any