Justin Holewinski
2012-Apr-27 19:50 UTC
[LLVMdev] [llvm-commits] [PATCH][RFC] NVPTX Backend
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 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 'if' enclosed in a compound statement, container.end() being evaluated every iterations of loops, etc, etc.. -Jim On Apr 24, 2012, at 11:50 AM, Justin Holewinski <jholewinski at nvidia.com<mailto:jholewinski at nvidia.com>> wrote: 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 up-streaming, but the back-end is very usable in its current state and would benefit all current and future users of the LLVM PTX back-end. The goal is to phase out the existing PTX back-end, while maintaining compatibility with it. To that end, the NVPTX back-end maintains its own set of intrinsics but also supports the old PTX back-end intrinsics to ensure compatibility with out-of-tree users. We would like to get your feedback on the attached patch to make sure it is up to LLVM commit quality. We would like to commit this as soon as the community is satisfied with it. Also, as the current maintainer of the PTX back-end, I give my consent to deprecate it. :) Thanks, Justin Holewinski ________________________________ This email message is for the sole use of the intended recipient(s) and may contain confidential information. Any unauthorized review, use, disclosure or distribution is prohibited. If you are not the intended recipient, please contact the sender by reply email and destroy all copies of the original message. ________________________________ <nvptx-backend-public.patch>_______________________________________________ llvm-commits mailing list llvm-commits at cs.uiuc.edu<mailto:llvm-commits at cs.uiuc.edu> http://lists.cs.uiuc.edu/mailman/listinfo/llvm-commits -------------- next part -------------- An HTML attachment was scrubbed... URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/20120427/8f219f70/attachment.html> -------------- next part -------------- A non-text attachment was scrubbed... Name: nvptx-backend-public.patch Type: application/octet-stream Size: 1026178 bytes Desc: nvptx-backend-public.patch URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/20120427/8f219f70/attachment.obj>
<!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 compared to the PTX backend and I'm pleased to see all the
support for vector instructions, atomics, debugging information, etc -
much of which was still to be added to the other backend. <br>
<br>
I have a few comments to make, but bear in mind my knowledge comes
solely from working on the PTX backend, so I don't know much about how
this compares to other backends and the feasibility of reusing
components designed to be shared across backends, plus my knowledge of
tablegen is also relatively limited.<br>
<br>
This is in no particular order:<br>
<br>
* Is there any reason why the 32-bit arch name is nvptx and the 64-bit
arch name is nvptx64, especially as the default for the NVCC compiler
is to pass the -m64 flag? Also, all the internal naming
(NVPTXTargetMachine for example) use 32 / 64 suffixes.<br>
<br>
* The register naming seems a little arbitrary as well, using FL
prefixes for 64-bit float and da prefixes for 64-bit float arguments
for example. <br>
<br>
* Something I picked up in the NVVM IR spec - it seems to only be
possible to use the bar.sync 0 instruction. Unless this is being
removed for PTX 3.0, the spec (and the PTX backend) support using
bar.sync {0..15}. The old PTX intrinsic also supports a non-zero
integer operand.<br>
<br>
* I guess this raises the question of whether or not it's actually
worthwhile retaining compatibility with the old backend. I converted my
Jet compiler to use NVVM intrinsics and still found the address spaces
were out of line, so it was largely useless until I'd done the full
conversion anyway. If the plan is to deprecate the old backend (as I
fully support), I'd be tempted to say maybe just leave all these old
intrinsics out too?<br>
<br>
* I know this isn't an easy problem as I remember discussing this at
length with you at the time Justin, but supporting i8 for ld, st, cvt
seems to introduce a load of blocks of string concat code in tablegen.
Is there not a better way of tackling this, it looks like a fairly ugly
solution to me?<br>
<br>
* There are a few places in the code where there are blocks of
commented out code that I presume are waiting for upstreaming of other
additions. Some of these are documented, but many of them aren't. Maybe
these should either be removed for addition later or some explanation
added as to why they're not being used.<br>
<br>
* There seems to be a lot of features that are lacking tests too. I was
trying to find the equivalent flags to use for switching on and off FMA
(a feature we use quite a bit to compare results with x86) and couldn't
find any tests demonstrating this. Considering the amount of code in
the backend committed to handling mad / fma, there really should be a
bunch of tests to go along with this. In fact non of the nvptx-* flags
seem to have any accompanying test cases, nor did atomics, nvvm
intrinsics or many of the other instructions. I imagine NVidia has a
fairly full featured CUDA testing framework, but if this is going to be
made open-source, I'd rather as many of the instructions could have
tests too. Having said that, maybe some of the test cases for the PTX
backend could be migrated across?<br>
<br>
* Finally I spotted a number of typos in the comments in the code that
could be cleaned up.<br>
<br>
There's a lot of code here, so I might have missed some things. On the
whole though, most of this is fairly minor. It looks well designed and
I'm looking forward to this making it in. I'm especially excited to see
the new pass for vectorization and other passes to do code
transformations. I managed to run our fluid solver through this backend
and it generated PTX that did the right thing. I haven't had time to do
any profiling comparing the old backend yet. <br>
<br>
Dan<br>
<br>
Justin Holewinski wrote:
<blockquote
cite="mid:448D0C8E64E8084981082A23DEDAEC70A28C88F926@HQMAIL02.nvidia.com"
type="cite">
<meta http-equiv="Content-Type" content="text/html;
">
<meta name="Generator" content="Microsoft Word 14 (filtered
medium)">
<base href="x-msg://14879/"><!--[if
!mso]><style>v\:* {behavior:url(#default#VML);}
o\:* {behavior:url(#default#VML);}
w\:* {behavior:url(#default#VML);}
.shape {behavior:url(#default#VML);}
</style><![endif]-->
<style><!--
/* Font Definitions */
@font-face
{font-family:Helvetica;
panose-1:2 11 5 4 2 2 2 2 2 4;}
@font-face
{font-family:Wingdings;
panose-1:5 0 0 0 0 0 0 0 0 0;}
@font-face
{font-family:Wingdings;
panose-1:5 0 0 0 0 0 0 0 0 0;}
@font-face
{font-family:Calibri;
panose-1:2 15 5 2 2 2 4 3 2 4;}
@font-face
{font-family:Tahoma;
panose-1:2 11 6 4 3 5 4 4 2 4;}
/* Style Definitions */
p.MsoNormal, li.MsoNormal, div.MsoNormal
{margin:0in;
margin-bottom:.0001pt;
font-size:12.0pt;
font-family:"Times New Roman","serif";}
a:link, span.MsoHyperlink
{mso-style-priority:99;
color:blue;
text-decoration:underline;}
a:visited, span.MsoHyperlinkFollowed
{mso-style-priority:99;
color:purple;
text-decoration:underline;}
span.apple-converted-space
{mso-style-name:apple-converted-space;}
span.EmailStyle18
{mso-style-type:personal-reply;
font-family:"Calibri","sans-serif";
color:#1F497D;}
.MsoChpDefault
{mso-style-type:export-only;
font-size:10.0pt;}
@page WordSection1
{size:8.5in 11.0in;
margin:1.0in 1.0in 1.0in 1.0in;}
div.WordSection1
{page:WordSection1;}
--></style><!--[if gte mso 9]><xml>
<o:shapedefaults v:ext="edit" spidmax="1026" />
</xml><![endif]--><!--[if gte mso 9]><xml>
<o:shapelayout v:ext="edit">
<o:idmap v:ext="edit" data="1" />
</o:shapelayout></xml><![endif]-->
<div class="WordSection1">
<p class="MsoNormal"><span
style="font-size: 11pt; font-family:
"Calibri","sans-serif"; color: rgb(31, 73,
125);">Thanks
for the feedback!<o:p></o:p></span></p>
<p class="MsoNormal"><span
style="font-size: 11pt; font-family:
"Calibri","sans-serif"; color: rgb(31, 73,
125);"><o:p> </o:p></span></p>
<p class="MsoNormal"><span
style="font-size: 11pt; font-family:
"Calibri","sans-serif"; color: rgb(31, 73,
125);">The
attached patch addresses the style issues that have been
found.<o:p></o:p></span></p>
<p class="MsoNormal"><span
style="font-size: 11pt; font-family:
"Calibri","sans-serif"; color: rgb(31, 73,
125);"><o:p> </o:p></span></p>
<div>
<div
style="border-style: solid none none; border-color: rgb(181, 196, 223)
-moz-use-text-color -moz-use-text-color; border-width: 1pt medium medium;
padding: 3pt 0in 0in;">
<p class="MsoNormal" style="margin-left:
0.5in;"><b><span
style="font-size: 10pt; font-family:
"Tahoma","sans-serif";">From:</span></b><span
style="font-size: 10pt; font-family:
"Tahoma","sans-serif";"> Jim
Grosbach [<a class="moz-txt-link-freetext"
href="mailto:grosbach@apple.com">mailto:grosbach@apple.com</a>]
<br>
<b>Sent:</b> Wednesday, April 25, 2012 2:22 PM<br>
<b>To:</b> Justin Holewinski<br>
<b>Cc:</b> <a class="moz-txt-link-abbreviated"
href="mailto:llvm-commits@cs.uiuc.edu">llvm-commits@cs.uiuc.edu</a>;
<a class="moz-txt-link-abbreviated"
href="mailto:llvmdev@cs.uiuc.edu">llvmdev@cs.uiuc.edu</a>;
Vinod Grover<br>
<b>Subject:</b> Re: [llvm-commits] [PATCH][RFC] NVPTX
Backend<o:p></o:p></span></p>
</div>
</div>
<p class="MsoNormal" style="margin-left:
0.5in;"><o:p> </o:p></p>
<p class="MsoNormal" style="margin-left: 0.5in;">Hi
Justin,<o:p></o:p></p>
<div>
<p class="MsoNormal" style="margin-left:
0.5in;"><o:p> </o:p></p>
</div>
<div>
<p class="MsoNormal" style="margin-left:
0.5in;">Cool stuff, to be
sure. Excited to see this.<o:p></o:p></p>
<div>
<p class="MsoNormal" style="margin-left:
0.5in;"><o:p> </o:p></p>
</div>
<div>
<p class="MsoNormal" style="margin-left: 0.5in;">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 'if' enclosed in a
compound statement, container.end() being evaluated every iterations of
loops, etc, etc..<o:p></o:p></p>
</div>
<div>
<p class="MsoNormal" style="margin-left:
0.5in;"><o:p> </o:p></p>
</div>
<div>
<p class="MsoNormal" style="margin-left:
0.5in;">-Jim<o:p></o:p></p>
</div>
<div>
<p class="MsoNormal" style="margin-left:
0.5in;"><o:p> </o:p></p>
<div>
<div>
<p class="MsoNormal" style="margin-left: 0.5in;">On
Apr 24, 2012, at
11:50 AM, Justin Holewinski <<a moz-do-not-send="true"
href="mailto:jholewinski@nvidia.com">jholewinski@nvidia.com</a>>
wrote:<o:p></o:p></p>
</div>
<p class="MsoNormal" style="margin-left:
0.5in;"><br>
<br>
<o:p></o:p></p>
<div>
<div>
<p class="MsoNormal" style="margin-left:
0.5in;"><span
style="font-size: 11pt; font-family:
"Calibri","sans-serif";">Hi
LLVMers,<o:p></o:p></span></p>
</div>
<div>
<p class="MsoNormal" style="margin-left:
0.5in;"><span
style="font-size: 11pt; font-family:
"Calibri","sans-serif";">
<o:p></o:p></span></p>
</div>
<div>
<p class="MsoNormal" style="margin-left:
0.5in;"><span
style="font-size: 11pt; font-family:
"Calibri","sans-serif";">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 up-streaming, but the back-end is very usable in
its current state and would benefit all current and future users of the
LLVM PTX back-end.<o:p></o:p></span></p>
</div>
<div>
<p class="MsoNormal" style="margin-left:
0.5in;"><span
style="font-size: 11pt; font-family:
"Calibri","sans-serif";">
<o:p></o:p></span></p>
</div>
<div>
<p class="MsoNormal" style="margin-left:
0.5in;"><span
style="font-size: 11pt; font-family:
"Calibri","sans-serif";">The goal
is to phase out the existing PTX back-end, while maintaining
compatibility with it. To that end, the NVPTX back-end maintains its
own set of intrinsics but also supports the old PTX back-end intrinsics
to ensure compatibility with out-of-tree
users.<o:p></o:p></span></p>
</div>
<div>
<p class="MsoNormal" style="margin-left:
0.5in;"><span
style="font-size: 11pt; font-family:
"Calibri","sans-serif";">
<o:p></o:p></span></p>
</div>
<div>
<p class="MsoNormal" style="margin-left:
0.5in;"><span
style="font-size: 11pt; font-family:
"Calibri","sans-serif";">We would
like to get your feedback on the attached patch to make sure it is up
to LLVM commit quality. We would like to commit this as soon as the
community is satisfied with it.<o:p></o:p></span></p>
</div>
<div>
<p class="MsoNormal" style="margin-left:
0.5in;"><span
style="font-size: 11pt; font-family:
"Calibri","sans-serif";">
<o:p></o:p></span></p>
</div>
<div>
<p class="MsoNormal" style="margin-left:
0.5in;"><span
style="font-size: 11pt; font-family:
"Calibri","sans-serif";">Also, as
the current maintainer of the PTX back-end, I give my consent to
deprecate it.<span class="apple-converted-space">
</span></span><span
style="font-size: 11pt; font-family:
Wingdings;">J</span><span
style="font-size: 11pt; font-family:
"Calibri","sans-serif";"><o:p></o:p></span></p>
</div>
<div>
<p class="MsoNormal" style="margin-left:
0.5in;"><span
style="font-size: 11pt; font-family:
"Calibri","sans-serif";">
<o:p></o:p></span></p>
</div>
<div>
<p class="MsoNormal" style="margin-left:
0.5in;"><span
style="font-size: 11pt; font-family:
"Calibri","sans-serif";">
<o:p></o:p></span></p>
</div>
<div>
<p class="MsoNormal" style="margin-left:
0.5in;"><span
style="font-size: 11pt; font-family:
"Calibri","sans-serif";">Thanks,<o:p></o:p></span></p>
</div>
<div>
<p class="MsoNormal" style="margin-left:
0.5in;"><span
style="font-size: 11pt; font-family:
"Calibri","sans-serif";">
<o:p></o:p></span></p>
</div>
<div>
<p class="MsoNormal" style="margin-left:
0.5in;"><span
style="font-size: 11pt; font-family:
"Calibri","sans-serif";">Justin
Holewinski<o:p></o:p></span></p>
</div>
<div>
<p class="MsoNormal" style="margin-left:
0.5in;"><span
style="font-size: 11pt; font-family:
"Calibri","sans-serif";">
<o:p></o:p></span></p>
</div>
<div>
<div class="MsoNormal" style="margin-left: 0.5in;
text-align: center;"
align="center"><span
style="font-size: 13.5pt; font-family:
"Helvetica","sans-serif";">
<hr align="center" size="2"
width="100%"></span></div>
</div>
<div>
<p class="MsoNormal" style="margin-left:
0.5in;"><span
style="font-size: 13.5pt; font-family:
"Helvetica","sans-serif";">This
email message is for the sole use of the intended recipient(s) and may
contain confidential information. Any unauthorized review, use,
disclosure or distribution is prohibited. If you are not the intended
recipient, please contact the sender by reply email and destroy all
copies of the original message.<o:p></o:p></span></p>
</div>
<div>
<div class="MsoNormal" style="margin-left: 0.5in;
text-align: center;"
align="center"><span
style="font-size: 13.5pt; font-family:
"Helvetica","sans-serif";">
<hr align="center" size="2"
width="100%"></span></div>
</div>
<div>
<p class="MsoNormal" style="margin-left:
0.5in;"><span
style="font-size: 13.5pt; font-family:
"Helvetica","sans-serif";"><o:p>
</o:p></span></p>
</div>
<p class="MsoNormal" style="margin-left:
0.5in;"><span
style="font-size: 13.5pt; font-family:
"Helvetica","sans-serif";"><nvptx-backend-public.patch>_______________________________________________<br>
llvm-commits mailing list<br>
<a moz-do-not-send="true"
href="mailto:llvm-commits@cs.uiuc.edu"><span
style="color:
purple;">llvm-commits@cs.uiuc.edu</span></a><br>
<a moz-do-not-send="true"
href="http://lists.cs.uiuc.edu/mailman/listinfo/llvm-commits"><span
style="color:
purple;">http://lists.cs.uiuc.edu/mailman/listinfo/llvm-commits</span></a><o:p></o:p></span></p>
</div>
</div>
<p class="MsoNormal" style="margin-left:
0.5in;"><o:p> </o:p></p>
</div>
</div>
</div>
<pre wrap="">
<hr size="4" width="90%">
_______________________________________________
llvm-commits mailing list
<a class="moz-txt-link-abbreviated"
href="mailto:llvm-commits@cs.uiuc.edu">llvm-commits@cs.uiuc.edu</a>
<a class="moz-txt-link-freetext"
href="http://lists.cs.uiuc.edu/mailman/listinfo/llvm-commits">http://lists.cs.uiuc.edu/mailman/listinfo/llvm-commits</a>
</pre>
</blockquote>
</body>
</html>
Justin Holewinski
2012-May-01 14:30 UTC
[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 terms of > features compared to the PTX backend and I'm pleased to see all the support > for vector instructions, atomics, debugging information, etc - much of which > was still to be added to the other backend. > > I have a few comments to make, but bear in mind my knowledge comes > solely from working on the PTX backend, so I don't know much about how > this compares to other backends and the feasibility of reusing components > designed to be shared across backends, plus my knowledge of tablegen is > also relatively limited. > > This is in no particular order: > > * Is there any reason why the 32-bit arch name is nvptx and the 64-bit arch > name is nvptx64, especially as the default for the NVCC compiler is to pass > the -m64 flag? Also, all the internal naming (NVPTXTargetMachine for > example) use 32 / 64 suffixes.As far as I know, there is no fundamental reason for this. If it bugs you too much, I'm sure we could change it. :)> > * The register naming seems a little arbitrary as well, using FL prefixes for 64- > bit float and da prefixes for 64-bit float arguments for example.Really, any choice is going to be arbitrary.> > * Something I picked up in the NVVM IR spec - it seems to only be possible to > use the bar.sync 0 instruction. Unless this is being removed for PTX 3.0, the > spec (and the PTX backend) support using bar.sync {0..15}. The old PTX > intrinsic also supports a non-zero integer operand.The NVVM intrinsic is there to implement CUDA's __syncthreads(). The old intrinsic is still exposed. I'll see about adding a proper NVVM intrinsic.> > * I guess this raises the question of whether or not it's actually worthwhile > retaining compatibility with the old backend. I converted my Jet compiler to > use NVVM intrinsics and still found the address spaces were out of line, so it > was largely useless until I'd done the full conversion anyway. If the plan is to > deprecate the old backend (as I fully support), I'd be tempted to say maybe > just leave all these old intrinsics out too?This is something I've debated back and forth. The original thought was to maintain as much backwards compatibility as possible with the old back-end. I don't think it hurts anything to leave the old intrinsics in there.> > * I know this isn't an easy problem as I remember discussing this at length > with you at the time Justin, but supporting i8 for ld, st, cvt seems to introduce > a load of blocks of string concat code in tablegen. Is there not a better way of > tackling this, it looks like a fairly ugly solution to me?Oh there are quite a few ugly TableGen solutions in there. :) I'm trying to find ways to clean some of this stuff up, but it works so it's hard to justify devoting time to it.> > * There are a few places in the code where there are blocks of commented > out code that I presume are waiting for upstreaming of other additions. > Some of these are documented, but many of them aren't. Maybe these > should either be removed for addition later or some explanation added as to > why they're not being used.Good point. I'll see about adding more documentation to these.> > * There seems to be a lot of features that are lacking tests too. I was trying to > find the equivalent flags to use for switching on and off FMA (a feature we > use quite a bit to compare results with x86) and couldn't find any tests > demonstrating this. Considering the amount of code in the backend > committed to handling mad / fma, there really should be a bunch of tests to > go along with this. In fact non of the nvptx-* flags seem to have any > accompanying test cases, nor did atomics, nvvm intrinsics or many of the > other instructions. I imagine NVidia has a fairly full featured CUDA testing > framework, but if this is going to be made open-source, I'd rather as many of > the instructions could have tests too. Having said that, maybe some of the > test cases for the PTX backend could be migrated across?I'm working on converting the tests now.> > * Finally I spotted a number of typos in the comments in the code that could > be cleaned up. > > There's a lot of code here, so I might have missed some things. On the whole > though, most of this is fairly minor. It looks well designed and I'm looking > forward to this making it in. I'm especially excited to see the new pass for > vectorization and other passes to do code transformations. I managed to run > our fluid solver through this backend and it generated PTX that did the right > thing. I haven't had time to do any profiling comparing the old backend yet.Can I put you down as a technical reviewer of the back-end? :) As the main consumer of the open-source back-end, your seal of approval would go a long way!> > Dan > > Justin Holewinski wrote: > > 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 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 'if' enclosed in a compound > statement, container.end() being evaluated every iterations of loops, etc, > etc.. > > > > -Jim > > > > On Apr 24, 2012, at 11:50 AM, Justin Holewinski > <jholewinski at nvidia.com> wrote: > > > > > > 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 up-streaming, but the back-end is very usable in its > current state and would benefit all current and future users of the LLVM PTX > back-end. > > > > The goal is to phase out the existing PTX back-end, while maintaining > compatibility with it. To that end, the NVPTX back-end maintains its own set > of intrinsics but also supports the old PTX back-end intrinsics to ensure > compatibility with out-of-tree users. > > > > We would like to get your feedback on the attached patch to make > sure it is up to LLVM commit quality. We would like to commit this as soon as > the community is satisfied with it. > > > > Also, as the current maintainer of the PTX back-end, I give my > consent to deprecate it. J > > > > > > Thanks, > > > > Justin Holewinski > > > > > ________________________________ > > > This email message is for the sole use of the intended recipient(s) > and may contain confidential information. Any unauthorized review, use, > disclosure or distribution is prohibited. If you are not the intended recipient, > please contact the sender by reply email and destroy all copies of the original > message. > > > ________________________________ > > > > > <nvptx-backend- > public.patch>_______________________________________________ > llvm-commits mailing list > llvm-commits at cs.uiuc.edu <mailto:llvm-commits at cs.uiuc.edu> > http://lists.cs.uiuc.edu/mailman/listinfo/llvm-commits > <http://lists.cs.uiuc.edu/mailman/listinfo/llvm-commits> > > > > > ________________________________ > > > _______________________________________________ > llvm-commits mailing list > llvm-commits at cs.uiuc.edu > http://lists.cs.uiuc.edu/mailman/listinfo/llvm-commits >