[LLVMdev] Attaching range metadata to IntrinsicInst

[Chandler Carruth]

This seems fine to me, but I'd like to make sure it looks OK to Nick as
well.


On Tue, Jun 17, 2014 at 12:37 AM, Jingyue Wu <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         http://llvm.cs.uiuc.edu
> http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
>
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/20140617/a0cbddc3/attachment.html>