# Launch bounds with 2D Threads dimension (x,y)

**URL:** <https://forums.developer.nvidia.com/t/launch-bounds-with-2d-threads-dimension-x-y/109378>\
**Category:** CUDA Programming and Performance\
**Created:** [December 19, 2019, 3:30pm UTC](https://forums.developer.nvidia.com/t/launch-bounds-with-2d-threads-dimension-x-y/109378 "2019-12-19T15:30:30Z")\
**Posts on this page:** 7\
**Page:** 1

<div class="post-metadata">

**Author:** ![Abdopensky](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Abdopensky](https://forums.developer.nvidia.com/u/Abdopensky)\
**Post date:** [December 19, 2019, 3:30pm UTC](https://forums.developer.nvidia.com/t/launch-bounds-with-2d-threads-dimension-x-y/109378/1 "2019-12-19T15:30:30Z")

</div>

Hello,

I read the **launch\_bounds** function from CUDA Programming Guide and it seems that the example is applied only for one single dimension. As I am using a 2D Kernel function, do you know how to implement the launch bounds into my application?

Thanks

---

<div class="post-metadata">

**Author:** ![tera](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@tera](https://forums.developer.nvidia.com/u/tera)\
**Post date:** [December 19, 2019, 3:44pm UTC](https://forums.developer.nvidia.com/t/launch-bounds-with-2d-threads-dimension-x-y/109378/2 "2019-12-19T15:44:07Z")

</div>

**launch\_bounds** () doesn’t care about the layout, only the total number of threads. So you can safely use blockDim.x_blockDim.y_blockDim.z (where blockDim.z==1 in your case) for the number of trheads per block.

---

<div class="post-metadata">

**Author:** ![Abdopensky](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Abdopensky](https://forums.developer.nvidia.com/u/Abdopensky)\
**Post date:** [December 19, 2019, 3:50pm UTC](https://forums.developer.nvidia.com/t/launch-bounds-with-2d-threads-dimension-x-y/109378/3 "2019-12-19T15:50:40Z")

</div>

Ok. So why I am receiving this compilation warning : CUDACOMPILE : ptxas warning : Value of threads per SM for entry \_Z30MyKernelPfS\_S\_S\_PKfS\_S\_fffbiiS\_b is out of range. .minnctapersm will be ignored?

---

<div class="post-metadata">

**Author:** ![mnicely](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/mnicely/32/14047_2.png) [@mnicely](https://forums.developer.nvidia.com/u/mnicely)\
**Post date:** [December 22, 2019, 11:07pm UTC](https://forums.developer.nvidia.com/t/launch-bounds-with-2d-threads-dimension-x-y/109378/4 "2019-12-22T23:07:03Z")

</div>

[https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html)

- **maxThreadsPerBlock** specifies the maximum number of threads per block with which the application will ever launch MyKernel(); it compiles to the .maxntidPTX directive
- **minBlocksPerMultiprocessor** is optional and specifies the desired minimum number of resident blocks per multiprocessor; it compiles to the .minnctapersmPTX directive

I was able to reproduce the error doing this…

```auto
__launch_bounds__ ( 512, 8 ) __global__ void doMath() {...}

```

You get that error because when you ask for 8 blocks minimum per SM with maximum threads per block equaling 512. So you’re asking for 4096 threads per SM. All architectures allow a maximum of 2048 threads per block, (except 7.5\<1024\>). Therefore, the compiler is ignoring your **minBlocksPerMultiprocessor** parameter.

If you were to do the following, the error should go away.

```auto
__launch_bounds__ ( 512, 4 ) __global__ void doMath() {...}

```

---

<div class="post-metadata">

**Author:** ![Abdopensky](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Abdopensky](https://forums.developer.nvidia.com/u/Abdopensky)\
**Post date:** [December 23, 2019, 1:19am UTC](https://forums.developer.nvidia.com/t/launch-bounds-with-2d-threads-dimension-x-y/109378/5 "2019-12-23T01:19:54Z")

</div>

Thanks !!!

Indeed, I have a RTX 2060 (7.5). So, i get this error when exceeding the maximum 1024.

1. Do you know why the limit was lowered for Turing architecture?
2. For my case, is it still worth to use Launch bounds?

---

<div class="post-metadata">

**Author:** ![mnicely](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/mnicely/32/14047_2.png) [@mnicely](https://forums.developer.nvidia.com/u/mnicely)\
**Post date:** [December 23, 2019, 2:26am UTC](https://forums.developer.nvidia.com/t/launch-bounds-with-2d-threads-dimension-x-y/109378/6 "2019-12-23T02:26:11Z")

</div>

It really depends on the applications. **launch\_bounds** is just an additional way to hand-tune for performance. I would suggest using Nsight Compute to help determine if the number of registers you are using and then go from there.

---

<div class="post-metadata">

**Author:** ![Abdopensky](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Abdopensky](https://forums.developer.nvidia.com/u/Abdopensky)\
**Post date:** [December 23, 2019, 2:41am UTC](https://forums.developer.nvidia.com/t/launch-bounds-with-2d-threads-dimension-x-y/109378/7 "2019-12-23T02:41:30Z")

</div>

Ok thanks