[llvm] r290294 - [NVVMIntrRange] Only set range metadata if none is already present

[Philip Reames via llvm-commits]

 From the description, I wonder if intersecting the two ranges might be 
better?  Note that I have no clue what this code is being used for, so 
take my suggestion with a bag of salt.

Philip

On 12/21/2016 04:52 PM, David Majnemer via llvm-commits wrote:
> Author: majnemer
> Date: Wed Dec 21 18:51:59 2016
> New Revision: 290294
>
> URL: http://llvm.org/viewvc/llvm-project?rev=290294&view=rev
> Log:
> [NVVMIntrRange] Only set range metadata if none is already present
>
> The range metadata inserted by NVVMIntrRange is pessimistic, range
> metadata already present could be more precise.
>
> Modified:
>      llvm/trunk/lib/Target/NVPTX/NVVMIntrRange.cpp
>      llvm/trunk/test/CodeGen/NVPTX/intrinsic-old.ll
>
> Modified: llvm/trunk/lib/Target/NVPTX/NVVMIntrRange.cpp
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVVMIntrRange.cpp?rev=290294&r1=290293&r2=290294&view=diff
> ==============================================================================
> --- llvm/trunk/lib/Target/NVPTX/NVVMIntrRange.cpp (original)
> +++ llvm/trunk/lib/Target/NVPTX/NVVMIntrRange.cpp Wed Dec 21 18:51:59 2016
> @@ -65,6 +65,10 @@ INITIALIZE_PASS(NVVMIntrRange, "nvvm-int
>   // Adds the passed-in [Low,High) range information as metadata to the
>   // passed-in call instruction.
>   static bool addRangeMetadata(uint64_t Low, uint64_t High, CallInst *C) {
> +  // This call already has range metadata, nothing to do.
> +  if (C->getMetadata(LLVMContext::MD_range))
> +    return false;
> +
>     LLVMContext &Context = C->getParent()->getContext();
>     IntegerType *Int32Ty = Type::getInt32Ty(Context);
>     Metadata *LowAndHigh[] = {
>
> Modified: llvm/trunk/test/CodeGen/NVPTX/intrinsic-old.ll
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/NVPTX/intrinsic-old.ll?rev=290294&r1=290293&r2=290294&view=diff
> ==============================================================================
> --- llvm/trunk/test/CodeGen/NVPTX/intrinsic-old.ll (original)
> +++ llvm/trunk/test/CodeGen/NVPTX/intrinsic-old.ll Wed Dec 21 18:51:59 2016
> @@ -155,6 +155,13 @@ define ptx_device i32 @test_nctaid_x() {
>   	ret i32 %x
>   }
>   
> +define ptx_device i32 @test_already_has_range_md() {
> +; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.x;
> +; RANGE: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.x(), !range ![[ALREADY:[0-9]+]]
> +	%x = call i32 @llvm.nvvm.read.ptx.sreg.nctaid.x(), !range !0
> +	ret i32 %x
> +}
> +
>   
>   define ptx_device i32 @test_nctaid_w() {
>   ; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.w;
> @@ -311,6 +318,9 @@ declare i32 @llvm.nvvm.read.ptx.sreg.pm3
>   
>   declare void @llvm.nvvm.bar.sync(i32 %i)
>   
> +!0 = !{i32 0, i32 19}
> +; RANGE-DAG: ![[ALREADY]] = !{i32 0, i32 19}
> +; RANGE-DAG: ![[BLK_IDX_XY]] = !{i32 0, i32 1024}
>   ; RANGE-DAG: ![[BLK_IDX_XY]] = !{i32 0, i32 1024}
>   ; RANGE-DAG: ![[BLK_IDX_Z]] = !{i32 0, i32 64}
>   ; RANGE-DAG: ![[BLK_SIZE_XY]] = !{i32 1, i32 1025}
>
>
> _______________________________________________
> llvm-commits mailing list
> llvm-commits at lists.llvm.org
> http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-commits