yongwww commented on PR #16652:
URL: https://github.com/apache/tvm/pull/16652#issuecomment-1969972962

   <p>With this example:</p>
   <pre><code class="language-python">    @I.ir_module
       class Mod:
           @R.function
           def foo(x: R.Tensor((20, 32000), &quot;float32&quot;)):
               m = T.int64()
               n = T.int64()
               with R.dataflow():
                   gv = R.cumsum(x, axis=1)
                   R.output(gv)
               return gv
   
       target = tvm.target.Target(&quot;cuda -libs=thrust&quot;, 
host=&quot;llvm&quot;)
       dev = tvm.cuda(0)
       ex = relax.build(Mod, target)
       vm = relax.VirtualMachine(ex, device=dev)
   </code></pre>
   <p>The 'cuda_gpu_kern_sum' stats report</p>
   <p>w/ this change</p>
   <p>Time (%)  Total Time (ns)  Instances  Avg (ns)  Med (ns)  Min (ns)  Max 
(ns)  StdDev (ns)                                                  Name</p>
   <hr>
   <pre><code> 92.8       18,099,352      1,101  16,439.0  16,384.0    14,944   
 18,336        433.8  void 
cub::CUB_200200_750_NS::DeviceScanByKeyKernel&lt;cub::CUB_200200_750_NS::DeviceScanByKeyPolicy&lt;th…
     7.2        1,409,748      1,101   1,280.4   1,280.0     1,151     1,568    
    120.1  void 
cub::CUB_200200_750_NS::DeviceScanByKeyInitKernel&lt;cub::CUB_200200_750_NS::ReduceByKeyScanTileS…
   
   </code></pre>
   <p>w/o this change</p>
   <hr>
   <pre><code> 51.6       22,603,888      1,101  20,530.3  20,545.0    19,744   
 21,120        293.0  cumsum_kernel
    44.3       19,442,725      1,101  17,659.2  17,665.0    16,641    18,752    
    294.6  void 
cub::CUB_200200_750_NS::DeviceScanByKeyKernel&lt;cub::CUB_200200_750_NS::DeviceScanByKeyPolicy&lt;th…
     4.1        1,800,850      1,101   1,635.6   1,632.0     1,440     1,792    
     35.9  void 
cub::CUB_200200_750_NS::DeviceScanByKeyInitKernel&lt;cub::CUB_200200_750_NS::ReduceByKeyScanTileS…
   
   </code></pre>
   <p>The execution perf numbers I got on NVIDIA GeForce RTX 3070:</p>
   <strong>R.cumsum(x, axis=1)</strong>
   
   Description | Performance
   -- | --
   With this change (w/ thrust) | 0.018 ms
   Without this change (w/ thrust) | 0.039 ms
   Without thrust | 0.460 ms
   
   <strong>R.cumsum(x, axis=0)</strong>
   
   Description | Performance
   -- | --
   With this change (w/ thrust) | 0.073 ms
   Without this change (w/ thrust) | 0.074 ms
   Without thrust | 44.018 ms
   
   <!-- notionvc: bfa52bd0-8dc8-4ed2-b286-54b234125fff -->


-- 
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.

To unsubscribe, e-mail: [email protected]

For queries about this service, please contact Infrastructure at:
[email protected]

Reply via email to