<div dir="ltr">Hi Eric, <div><br></div><div>In the IR, besides "target datalayout" and "target triple", we have a special "target cpu" string which is set by the Clang front-end according to its -target-cpu flag. We also write a Module::getTargetCPU() method to retrieve this string from the IR. </div>
<div><br></div><div>Jingyue</div></div><div class="gmail_extra"><br><br><div class="gmail_quote">On Tue, Jun 17, 2014 at 2:22 PM, Eric Christopher <span dir="ltr"><<a href="mailto:echristo@gmail.com" target="_blank">echristo@gmail.com</a>></span> wrote:<br>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">Eh? How do you envision this?<br>
<span class="HOEnZb"><font color="#888888"><br>
-eric<br>
</font></span><div class="HOEnZb"><div class="h5"><br>
On Tue, Jun 17, 2014 at 2:09 PM, Jingyue Wu <<a href="mailto:jingyue@google.com">jingyue@google.com</a>> wrote:<br>
> Hi Nick,<br>
><br>
> That makes sense. I think a main issue here is that the ranges of these PTX<br>
> special registers (e.g., threadIdx.x) depend on -target-cpu which is only<br>
> visible to clang and llc. Would you mind we specify "target cpu" in the IR<br>
> similar to what we did for "target triple"?<br>
><br>
> Thanks,<br>
> Jingyue<br>
><br>
><br>
> On Tue, Jun 17, 2014 at 12:19 PM, Nick Lewycky <<a href="mailto:nlewycky@google.com">nlewycky@google.com</a>> wrote:<br>
>><br>
>> On 17 June 2014 06:41, Eli Bendersky <<a href="mailto:eliben@google.com">eliben@google.com</a>> wrote:<br>
>>><br>
>>> On Tue, Jun 17, 2014 at 1:38 AM, Nick Lewycky <<a href="mailto:nicholas@mxc.ca">nicholas@mxc.ca</a>> wrote:<br>
>>>><br>
>>>> Chandler Carruth wrote:<br>
>>>>><br>
>>>>> This seems fine to me, but I'd like to make sure it looks OK to Nick as<br>
>>>>> well.<br>
>>>><br>
>>>><br>
>>>> I strongly prefer baking in knowledge about the intrinsics themselves<br>
>>>> into the passes if possible. Metadata will always be secondary.<br>
>>><br>
>>><br>
>>> So you're saying that in this particular case you'd prefer LLVM passes to<br>
>>> know about the range of these PTX intrinsics, rather than Clang adding them<br>
>>> as metadata?<br>
>><br>
>><br>
>> Yep.<br>
>><br>
>>> ValueTracking.cpp already has some iffy target knowledge (someone sneaked<br>
>>> a direct  Intrinsic::x86_sse42_crc32_64_64 check in there), but extending it<br>
>>> to other intrinsics in other targets seems like too much...<br>
>><br>
>><br>
>> That's not iffy. That's exactly how it should work, and we should have<br>
>> more of that. There is a major gotcha and that's dealing with the case where<br>
>> the intrinsics don't exist because the backend wasn't compiled in. If<br>
>> x86_sse42_crc32_64_64 is in there (and also in instcombine btw), presumably<br>
>> that problem is solved somehow? Or does llvm actually not build if you don't<br>
>> enable the x86 target? I feel like we would've heard about that.<br>
>><br>
>> Nick<br>
>><br>
>>> So should target info be passed into it in some way? Any suggestions<br>
>>> where to put it? TargetLibraryInfo? TargetTransformInfo? In any case this<br>
>>> seems like the target interface will have to be augmented, and we'll have to<br>
>>> carry an object around into ValueTracking's compute* functions. If this is<br>
>>> the right way, then this is the way it will be done - design ideas are<br>
>>> appreciated.<br>
>>><br>
>>> Eli<br>
>>><br>
>>><br>
>>><br>
>>>><br>
>>>><br>
>>>> Separately, should value tracking look use range metadata when it's<br>
>>>> available? Absolutely.<br>
>>>><br>
>>>> I think it should apply to all CallInst not just IntrinsicInst (which is<br>
>>>> derived from CallInst).<br>
>>>><br>
>>>> Nick<br>
>>>><br>
>>>>> On Tue, Jun 17, 2014 at 12:37 AM, Jingyue Wu <<a href="mailto:jingyue@google.com">jingyue@google.com</a><br>
>>>>> <mailto:<a href="mailto:jingyue@google.com">jingyue@google.com</a>>> wrote:<br>
>>>>><br>
>>>>>     Hi,<br>
>>>>><br>
>>>>>     The range metadata can only be attached to LoadInst for now. I am<br>
>>>>>     considering extending its usage to IntrinsicInst so that the<br>
>>>>>     frontend can annotate the range of the return value of an intrinsic<br>
>>>>>     call. e.g.,<br>
>>>>>     %a = call i32 @llvm.xxx(), !range !0<br>
>>>>>     !0 = metadata !{ i32 0, i23 1024 }<br>
>>>>><br>
>>>>>     The motivation behind this extension is some optimizations we are<br>
>>>>>     working on for CUDA programs. Some special registers in CUDA (e.g.,<br>
>>>>>     threadIdx.x) are bounded per CUDA programming guide, and knowing<br>
>>>>>     their ranges can improve the precision of ValueTracking and benefit<br>
>>>>>     optimizations such as InstCombine.<br>
>>>>><br>
>>>>>     To implement this idea, we need ValueTracking to be aware of the<br>
>>>>>     ranges of these special variables. These special registers are so<br>
>>>>>     far read-only and accessed using intrinsics. e.g.,<br>
>>>>>     %threadIdx.x = call i32 @llvm.nvvm.read.ptx.sreg.tid.x().<br>
>>>>><br>
>>>>>     One possible approach is to have ValueTracking compute the known<br>
>>>>>     bits of these intrinsics as special cases. This approach is already<br>
>>>>>     taken for the x86_sse42_crc32_64_64 intrinsic. However, this<br>
>>>>>     approach may not be elegant because the ranges of these CUDA<br>
>>>>> special<br>
>>>>>     registers depend on the GPU compute capability specified by<br>
>>>>>     -target-cpu. For instance, blockIdx.x is bounded by 65535 in sm_20<br>
>>>>>     but 2^31-1 in sm_30. Exposing -target-cpu to ValueTracking is<br>
>>>>>     probably discouraged.<br>
>>>>><br>
>>>>>     Therefore, the approach I am considering is to have clang annotate<br>
>>>>>     the ranges of these CUDA special registers according to the<br>
>>>>>     -target-cpu flag, and have ValueTracking pick the range metadata<br>
>>>>> for<br>
>>>>>     optimization. By doing so, we hide the target-specific info from<br>
>>>>>     ValueTracking.<br>
>>>>><br>
>>>>>     The code change in llvm minus clang won't be large. The core change<br>
>>>>>     is only a few lines:<br>
>>>>>     <a href="http://reviews.llvm.org/differential/diff/10464/" target="_blank">http://reviews.llvm.org/differential/diff/10464/</a>. If this extension<br>
>>>>>     sounds good to you, I'll definitely add more tests and revise the<br>
>>>>>     documents on range metadata.<br>
>>>>><br>
>>>>>     Best,<br>
>>>>>     Jingyue<br>
>>>>><br>
>>>>>     _______________________________________________<br>
>>>>>     LLVM Developers mailing list<br>
>>>>>     <a href="mailto:LLVMdev@cs.uiuc.edu">LLVMdev@cs.uiuc.edu</a> <mailto:<a href="mailto:LLVMdev@cs.uiuc.edu">LLVMdev@cs.uiuc.edu</a>><br>
>>>>> <a href="http://llvm.cs.uiuc.edu" target="_blank">http://llvm.cs.uiuc.edu</a><br>
>>>>><br>
>>>>>     <a href="http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev" target="_blank">http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev</a><br>
>>>>><br>
>>>>><br>
>>>>><br>
>>>>><br>
>>>>> _______________________________________________<br>
>>>>> LLVM Developers mailing list<br>
>>>>> <a href="mailto:LLVMdev@cs.uiuc.edu">LLVMdev@cs.uiuc.edu</a>         <a href="http://llvm.cs.uiuc.edu" target="_blank">http://llvm.cs.uiuc.edu</a><br>
>>>>> <a href="http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev" target="_blank">http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev</a><br>
>>>><br>
>>>><br>
>>>> _______________________________________________<br>
>>>> LLVM Developers mailing list<br>
>>>> <a href="mailto:LLVMdev@cs.uiuc.edu">LLVMdev@cs.uiuc.edu</a>         <a href="http://llvm.cs.uiuc.edu" target="_blank">http://llvm.cs.uiuc.edu</a><br>
>>>> <a href="http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev" target="_blank">http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev</a><br>
>>><br>
>>><br>
>>><br>
>>> _______________________________________________<br>
>>> LLVM Developers mailing list<br>
>>> <a href="mailto:LLVMdev@cs.uiuc.edu">LLVMdev@cs.uiuc.edu</a>         <a href="http://llvm.cs.uiuc.edu" target="_blank">http://llvm.cs.uiuc.edu</a><br>
>>> <a href="http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev" target="_blank">http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev</a><br>
>>><br>
>><br>
>><br>
>> _______________________________________________<br>
>> LLVM Developers mailing list<br>
>> <a href="mailto:LLVMdev@cs.uiuc.edu">LLVMdev@cs.uiuc.edu</a>         <a href="http://llvm.cs.uiuc.edu" target="_blank">http://llvm.cs.uiuc.edu</a><br>
>> <a href="http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev" target="_blank">http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev</a><br>
>><br>
><br>
><br>
> _______________________________________________<br>
> LLVM Developers mailing list<br>
> <a href="mailto:LLVMdev@cs.uiuc.edu">LLVMdev@cs.uiuc.edu</a>         <a href="http://llvm.cs.uiuc.edu" target="_blank">http://llvm.cs.uiuc.edu</a><br>
> <a href="http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev" target="_blank">http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev</a><br>
><br>
</div></div></blockquote></div><br></div>