llvm dev - Jun 2014 - [LLVMdev] Attaching range metadata to IntrinsicInst

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
-------------- next part --------------
An HTML attachment was scrubbed...
URL:
<http://lists.llvm.org/pipermail/llvm-dev/attachments/20140616/7850a0af/attachment.html>

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>

----- 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

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.

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