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