<html xmlns:v="urn:schemas-microsoft-com:vml" xmlns:o="urn:schemas-microsoft-com:office:office" xmlns:w="urn:schemas-microsoft-com:office:word" xmlns:m="http://schemas.microsoft.com/office/2004/12/omml" xmlns="http://www.w3.org/TR/REC-html40"><head><meta http-equiv=Content-Type content="text/html; charset=utf-8"><meta name=Generator content="Microsoft Word 15 (filtered medium)"><style><!--
/* Font Definitions */
@font-face
{font-family:Wingdings;
panose-1:5 0 0 0 0 0 0 0 0 0;}
@font-face
{font-family:"Cambria Math";
panose-1:0 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;}
/* Style Definitions */
p.MsoNormal, li.MsoNormal, div.MsoNormal
{margin:0cm;
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.EmailStyle17
{mso-style-type:personal-reply;
font-family:"Calibri",sans-serif;
color:#1F497D;}
.MsoChpDefault
{mso-style-type:export-only;
font-family:"Calibri",sans-serif;
mso-fareast-language:EN-US;}
@page WordSection1
{size:612.0pt 792.0pt;
margin:70.85pt 70.85pt 70.85pt 70.85pt;}
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]--></head><body lang=FR link=blue vlink=purple><div class=WordSection1><p class=MsoNormal><span style='font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1F497D;mso-fareast-language:EN-US'>Thanks<o:p></o:p></span></p><p class=MsoNormal><span style='font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1F497D;mso-fareast-language:EN-US'><o:p> </o:p></span></p><p class=MsoNormal><span lang=EN-US style='font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1F497D;mso-fareast-language:EN-US'>The thing is I’m actually writing those headers (because I have to). <o:p></o:p></span></p><p class=MsoNormal><span lang=EN-US style='font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1F497D;mso-fareast-language:EN-US'>So builtins are not really an issue. <o:p></o:p></span></p><p class=MsoNormal><span lang=EN-US style='font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1F497D;mso-fareast-language:EN-US'><o:p> </o:p></span></p><p class=MsoNormal><span lang=EN-US style='font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1F497D;mso-fareast-language:EN-US'>My problem was<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-US style='font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1F497D;mso-fareast-language:EN-US'>__global__<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-US style='font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1F497D;mso-fareast-language:EN-US'>__shared__ …<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-US style='font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1F497D;mso-fareast-language:EN-US'><o:p> </o:p></span></p><p class=MsoNormal><span lang=EN-US style='font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1F497D;mso-fareast-language:EN-US'>Which I defined as custom attributes<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-US style='font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1F497D;mso-fareast-language:EN-US'><o:p> </o:p></span></p><p class=MsoNormal><span lang=EN-US style='font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1F497D;mso-fareast-language:EN-US'>And the more than painful kernel<<<32, 128>>> which I can’t handle without modifying the front end myself (and I doesn’t look easy </span><span lang=EN-US style='font-size:11.0pt;font-family:Wingdings;color:#1F497D;mso-fareast-language:EN-US'>J</span><span lang=EN-US style='font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1F497D;mso-fareast-language:EN-US'> )<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-US style='font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1F497D;mso-fareast-language:EN-US'><o:p> </o:p></span></p><p class=MsoNormal><span lang=EN-US style='font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1F497D;mso-fareast-language:EN-US'>I’ll try the </span><span lang=EN-US>-fcuda-is-device</span><span lang=EN-US style='font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1F497D;mso-fareast-language:EN-US'> and keep you posted<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-US style='font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1F497D;mso-fareast-language:EN-US'><o:p> </o:p></span></p><p class=MsoNormal><span lang=EN-US style='font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1F497D;mso-fareast-language:EN-US'>thanks<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-US style='font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1F497D;mso-fareast-language:EN-US'><o:p> </o:p></span></p><p class=MsoNormal><b><span lang=EN-US style='font-size:11.0pt;font-family:"Calibri",sans-serif'>From:</span></b><span lang=EN-US style='font-size:11.0pt;font-family:"Calibri",sans-serif'> Eli Bendersky [mailto:eliben@google.com] <br><b>Sent:</b> mardi 7 avril 2015 00:44<br><b>To:</b> C Bergström<br><b>Cc:</b> Régis Portalez; clang-dev Developers<br><b>Subject:</b> Re: [cfe-dev] Parallelism TS implementation and feasibility of GPUexecution policies<o:p></o:p></span></p><p class=MsoNormal><o:p> </o:p></p><div><p class=MsoNormal><o:p> </o:p></p><div><p class=MsoNormal><o:p> </o:p></p><div><p class=MsoNormal>On Mon, Apr 6, 2015 at 3:22 PM, C Bergström <<a href="mailto:cbergstrom@pathscale.com" target="_blank">cbergstrom@pathscale.com</a>> wrote:<o:p></o:p></p><blockquote style='border:none;border-left:solid #CCCCCC 1.0pt;padding:0cm 0cm 0cm 6.0pt;margin-left:4.8pt;margin-right:0cm'><p class=MsoNormal>On Tue, Apr 7, 2015 at 4:47 AM, Eli Bendersky <<a href="mailto:eliben@google.com">eliben@google.com</a>> wrote:<br>><br>><br>> On Mon, Apr 6, 2015 at 11:48 AM, Régis Portalez<br>> <<a href="mailto:regis.portalez@altimesh.com">regis.portalez@altimesh.com</a>> wrote:<br>>><br>>> Hi.<br>>><br>>> Sorry to interrupt, but I understood there is way to emit llvm ir from<br>>> cuda code?<br>>><br>><br>> In general, the Clang frontend (-cc1) can generate LLVM IR for the nvptx<br>> triples/targets, when passed -fcuda-is-device. To use this in practice,<br>> you'll need to supply a bunch of things in headers (definitions of builtins,<br>> CUDA types and such), and no such headers exist in the open yet. Clang won't<br>> be able to parse the NVIDIA headers as these collide with the standard C++<br>> headers in some ways.<br><br>The old cuda headers used to be permissively licensed. (below)<br><br>Those headers are probably sufficient to get things "rolling", but I<br>don't know if they are really a good start. (Not to mention you'd be<br>missing a runtime)<br><br>The problem is that you'd have to do 2 passes with different<br>(conflicting) defines. Once for host and once for device. To get<br>host+device to play nice together is a b*. We have this resolved in<br>our clang, but it's really specific to our compilation flow. A general<br>solution would most likely involve extensive changes to the headers or<br>a rewrite. :/<o:p></o:p></p></blockquote><div><p class=MsoNormal><o:p> </o:p></p></div><div><p class=MsoNormal>Yep, this (2 passes with different defines) is the path taken in our approach, and the one we're pushing in <a href="http://reviews.llvm.org/D8463">http://reviews.llvm.org/D8463</a><o:p></o:p></p></div><div><p class=MsoNormal><o:p> </o:p></p></div><div><p class=MsoNormal>Note that, headers notwithstanding, the 2 pass compilation flow is enforced by the definition of the CUDA language, because __CUDA_ARCH__ is defined only for device code, and undefined for host code, even though these two can live in the same TU. So you *have* to compile the code twice.<o:p></o:p></p></div><div><p class=MsoNormal><o:p> </o:p></p></div><div><p class=MsoNormal>Back to the headers - these will remain a problem. We're now looking at different options.<o:p></o:p></p></div><div><p class=MsoNormal><o:p> </o:p></p></div><div><p class=MsoNormal><o:p> </o:p></p></div><div><p class=MsoNormal>Eli<o:p></o:p></p></div><div><p class=MsoNormal><o:p> </o:p></p></div><div><p class=MsoNormal><o:p> </o:p></p></div><div><p class=MsoNormal><o:p> </o:p></p></div><div><p class=MsoNormal><o:p> </o:p></p></div><div><p class=MsoNormal><o:p> </o:p></p></div><div><p class=MsoNormal> <o:p></o:p></p></div><blockquote style='border:none;border-left:solid #CCCCCC 1.0pt;padding:0cm 0cm 0cm 6.0pt;margin-left:4.8pt;margin-right:0cm'><p class=MsoNormal><br>I've fought with this for the past 5 years. I'll try to help in a<br>general way if/where I can.<br><br>Old CUDA header license<br>----------<br>/*<br> * Copyright 1993-2010 NVIDIA Corporation. All rights reserved.<br> *<br> * NOTICE TO USER:<br> *<br> * This source code is subject to NVIDIA ownership rights under U.S. and<br> * international Copyright laws. Users and possessors of this source code<br> * are hereby granted a nonexclusive, royalty-free license to use this code<br> * in individual and commercial software.<br> *<br> * NVIDIA MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THIS SOURCE<br> * CODE FOR ANY PURPOSE. IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR<br> * IMPLIED WARRANTY OF ANY KIND. NVIDIA DISCLAIMS ALL WARRANTIES WITH<br> * REGARD TO THIS SOURCE CODE, INCLUDING ALL IMPLIED WARRANTIES OF<br> * MERCHANTABILITY, NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE.<br> * IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL,<br> * OR CONSEQUENTIAL DAMAGES, OR ANY DAMAGES WHATSOEVER RESULTING FROM LOSS<br> * OF USE, DATA OR PROFITS, WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE<br> * OR OTHER TORTIOUS ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE<br> * OR PERFORMANCE OF THIS SOURCE CODE.<br> *<br> * U.S. Government End Users. This source code is a "commercial item" as<br> * that term is defined at 48 C.F.R. 2.101 (OCT 1995), consisting of<br> * "commercial computer software" and "commercial computer software<br> * documentation" as such terms are used in 48 C.F.R. 12.212 (SEPT 1995)<br> * and is provided to the U.S. Government only as a commercial end item.<br> * Consistent with 48 C.F.R.12.212 and 48 C.F.R. 227.7202-1 through<br> * 227.7202-4 (JUNE 1995), all U.S. Government End Users acquire the<br> * source code with only those rights set forth herein.<br> *<br> * Any use of this source code in individual and commercial software must<br> * include, in the user documentation and internal comments to the code,<br> * the above Disclaimer and U.S. Government End Users Notice.<br> */<o:p></o:p></p></blockquote></div><p class=MsoNormal><o:p> </o:p></p></div></div></div></body></html>