[ValueTracking] Consider the bounds of PTX special registers

[jingyue]

Summary

Some PTX special registers are bounded per CUDA programming guide.
Leveraing the bounds of these special registers can lead to more precise
value analysis.

Add two new tests in test/Transforms/InstCombine/intrinsics.ll

Depends on D4144

[eliben]

LGTM

[meheff]

lib/Analysis/ValueTracking.cpp

759

spelling nit: levaraging

778

These values seem to be CUDA version specific. Is there any way of guarding these? And how will we know to update these when CUDA N+1 is supported with different values?

[jingyue]

Mark, I agree with your concern. I just found out we can use -target-cpu to pass the compute capacity (e.g., sm_35) to the clang frontend. I'll send out another diff. Thanks!

[meheff]

I should also mention that I encountered some long compilation times which are superlinear with the unroll count when experimenting with the pragma loop limit. With the current limit (32K) on a simple loop the compilation time is ~7s. Doubling the limit results in a compilation time of ~50s. It seems to be beneath llvm::UnrollLoop -> FoldBlockIntoPredecessor -> llvm::ScalarEvolution::forgetLoop.

[hfinkel]

• Original Message -----

From: "Mark Heffernan" <[e-mail]>
To: [e-mail], "justin holewinski" <[e-mail]>, [e-mail], [e-mail]
Cc: [e-mail]
Sent: Monday, June 16, 2014 2:08:29 PM
Subject: Re: [PATCH] [ValueTracking] Consider the bounds of PTX special registers

I should also mention that I encountered some long compilation times
which are superlinear with the unroll count when experimenting with
the pragma loop limit. With the current limit (32K) on a simple
loop the compilation time is ~7s. Doubling the limit results in a
compilation time of ~50s. It seems to be beneath llvm::UnrollLoop
-> FoldBlockIntoPredecessor -> llvm::ScalarEvolution::forgetLoop.

Can you please file a PR to track this issue?

Thanks,
Hal

http://reviews.llvm.org/D4150

---

llvm-commits mailing list
[e-mail]
http://lists.cs.uiuc.edu/mailman/listinfo/llvm-commits

[jingyue]

Because the ranges of PTX special registers depend on subtarget (-target-cpu), we will have clang attach range metadata to these intrinsics and have the optimizer pick up these metadata. The second part is committed in r211281 (D4187). Will work on the first part.

lib/Analysis/ValueTracking.cpp

if (IntrinsicInst *II = dyn_cast<IntrinsicInst>(I)) {
       case Intrinsic::ctpop: {
         unsigned LowBits = Log2_32(BitWidth)+1;
         KnownZero = APInt::getHighBitsSet(BitWidth, BitWidth - LowBits);
         break;
       }
       case Intrinsic::x86_sse42_crc32_64_64:
         KnownZero = APInt::getHighBitsSet(64, 32);
         break;
+      // Some PTX special registers are bounded per CUDA programming guide
+      // (http://docs.nvidia.com/cuda/cuda-c-programming-guide/
+      // index.html#compute-capabilities).
+      // Leveraing the bounds of these special registers can lead to more
+      // precise value analysis.
+      case Intrinsic::nvvm_read_ptx_sreg_tid_x:
+      case Intrinsic::nvvm_read_ptx_sreg_tid_y:
+        // threadIdx.x, threadIdx.y < 1024
+        KnownZero = APInt::getHighBitsSet(32, 32 - 10);
+        break;
+      case Intrinsic::nvvm_read_ptx_sreg_tid_z:
+        // threadIdx.z < 64
+        KnownZero = APInt::getHighBitsSet(32, 32 - 6);
+        break;
+      case Intrinsic::nvvm_read_ptx_sreg_ctaid_x:
+        // blockIdx.x < 2^31
+        KnownZero = APInt::getHighBitsSet(32, 32 - 31);
+        break;
+      case Intrinsic::nvvm_read_ptx_sreg_ctaid_y:
+      case Intrinsic::nvvm_read_ptx_sreg_ctaid_z:
+        // blockIdx.y, blockIdx.z < 65536
+        KnownZero = APInt::getHighBitsSet(32, 32 - 16);
+        break;
       }
     }
     break;
   case Instruction::ExtractValue:
     if (IntrinsicInst *II = dyn_cast<IntrinsicInst>(I->getOperand(0))) {
       ExtractValueInst *EVI = cast<ExtractValueInst>(I);
       if (EVI->getNumIndices() != 1) break;
       if (EVI->getIndices()[0] == 0) {

test/Transforms/InstCombine/intrinsics.ll

 ; RUN: opt -instcombine -S < %s | FileCheck %s
 
 %overflow.result = type {i8, i1}
 
 declare %overflow.result @llvm.uadd.with.overflow.i8(i8, i8)
 declare %overflow.result @llvm.umul.with.overflow.i8(i8, i8)
 declare double @llvm.powi.f64(double, i32) nounwind readonly
 declare i32 @llvm.cttz.i32(i32, i1) nounwind readnone
 declare i32 @llvm.ctlz.i32(i32, i1) nounwind readnone
 declare i32 @llvm.ctpop.i32(i32) nounwind readnone
 declare i8 @llvm.ctlz.i8(i8, i1) nounwind readnone
+declare i32 @llvm.nvvm.read.ptx.sreg.tid.x()
+declare i32 @llvm.nvvm.read.ptx.sreg.tid.y()
+declare i32 @llvm.nvvm.read.ptx.sreg.tid.z()
+declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
+declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()
+declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.z()
 
 define i8 @uaddtest1(i8 %A, i8 %B) {
   %x = call %overflow.result @llvm.uadd.with.overflow.i8(i8 %A, i8 %B)
   %y = extractvalue %overflow.result %x, 0
   ret i8 %y
 ; CHECK-LABEL: @uaddtest1(
 ; CHECK-NEXT: %y = add i8 %A, %B
 ; CHECK-NEXT: ret i8 %y
define i32 @cttz_select(i32 %Value) nounwind {
   %tobool = icmp ne i32 %Value, 0
   %cttz = call i32 @llvm.cttz.i32(i32 %Value, i1 true)
   %s = select i1 %tobool, i32 %cttz, i32 32
   ret i32 %s
 
 ; CHECK-LABEL: @cttz_select(
 ; CHECK: select i1 %tobool, i32 %cttz, i32 32
 }
+
+define void @nvvm_thread_idx(i32* %output_x, i32* %output_y, i32* %output_z) {
+; CHECK-LABEL: @nvvm_thread_idx(
+  %tid_x = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
+  %tid_y = call i32 @llvm.nvvm.read.ptx.sreg.tid.y()
+  %tid_z = call i32 @llvm.nvvm.read.ptx.sreg.tid.z()
+; 0 <= threadIdx.x, threadIdx.y and threadIdx.z < 2^31.
+; Therefore, add i32 threadIdx.x|y|z, 5 has no unsigned wrap.
+  %x = add i32 %tid_x, 5
+  %y = add i32 %tid_y, 5
+  %z = add i32 %tid_z, 5
+; CHECK: add nuw
+; CHECK: add nuw
+; CHECK: add nuw
+  store i32 %x, i32* %output_x
+  store i32 %y, i32* %output_y
+  store i32 %z, i32* %output_z
+  ret void
+}
+
+define void @nvvm_block_idx(i32* %output_x, i32* %output_y, i32* %output_z) {
+; CHECK-LABEL: @nvvm_block_idx(
+  %bid_x = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
+  %bid_y = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()
+  %bid_z = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.z()
+; 0 <= blockIdx.x, blockIdx.y and blockIdx.z < 2^31.
+; Therefore, add i32 blockIdx.x|y|z, 5 has no unsigned wrap.
+  %x = add i32 %bid_x, 5
+  %y = add i32 %bid_y, 5
+  %z = add i32 %bid_z, 5
+; CHECK: add nuw
+; CHECK: add nuw
+; CHECK: add nuw
+  store i32 %x, i32* %output_x
+  store i32 %y, i32* %output_y
+  store i32 %z, i32* %output_z
+  ret void
+}