[LLVMdev] Attaching range metadata to IntrinsicInst

[Hal Finkel]

----- Original Message -----
> From: "Chandler Carruth" <chandlerc at google.com>
> To: "Jingyue Wu" <jingyue at google.com>, "Nick Lewycky" <nlewycky at google.com>
> Cc: "LLVM Developers Mailing List" <llvmdev at cs.uiuc.edu>
> Sent: Tuesday, June 17, 2014 1:44:52 AM
> Subject: Re: [LLVMdev] Attaching range metadata to IntrinsicInst
> 
> 
> 
> This seems fine to me, but I'd like to make sure it looks OK to Nick
> as well.

Is there any reason not to allow these on calls generally (not just intrinsic calls)?

 -Hal

> 
> 
> 
> 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
> 
> 
> 
> _______________________________________________
> LLVM Developers mailing list
> LLVMdev at cs.uiuc.edu         http://llvm.cs.uiuc.edu
> http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
> 

-- 
Hal Finkel
Assistant Computational Scientist
Leadership Computing Facility
Argonne National Laboratory