# limits on threads per kernel launch

**URL:** <https://forums.developer.nvidia.com/t/limits-on-threads-per-kernel-launch/26038>\
**Category:** CUDA Programming and Performance\
**Created:** [March 11, 2012, 11:04pm UTC](https://forums.developer.nvidia.com/t/limits-on-threads-per-kernel-launch/26038 "2012-03-11T23:04:38Z")\
**Posts on this page:** 4\
**Page:** 1

<div class="post-metadata">

**Author:** ![LHickey](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@LHickey](https://forums.developer.nvidia.com/u/LHickey)\
**Post date:** [March 11, 2012, 11:04pm UTC](https://forums.developer.nvidia.com/t/limits-on-threads-per-kernel-launch/26038/1 "2012-03-11T23:04:38Z")

</div>

I am running a C2050

I have the simple cuda program  
cat slag.cu  
1 #include \<stdlib.h\>  
2 #include \<string.h\>  
3 #include \<assert.h\>  
4 #include “book.h”  
5 **global** void gpuKernel( int nBLks, int nThreadsPerBlk)  
6 {  
7 printf(“Bid=%06ld, Tid=%06ld\n”, blockIdx.x, threadIdx.x);  
8 };  
9  
10 int main(int argc,char \*_argv)  
11 {  
12 if (argc == 1)  
13 {  
14 struct cudaDeviceProp P;  
15 cudaError\_t e = cudaGetDeviceProperties(&P,0);  
16 assert( e == cudaSuccess);  
17 fprintf(stdout,  
18 “name=%s\ntotalGlobalMem=%lu\nsharedMemPerBlock=%lu\nregsPerBlock=%ld\n”  
19 “warpSize=%ld\nmemPitch=%ld\nmaxThreadsPerBlock=%ld\n”  
20 “maxThreadsDim=(%ld,%ld%ld)\n” “maxGridSize=(%ld,%ld%ld)\n”  
21 “totalConstMem=%ld\nmajor=%ld minor=%ld\nclockRate=%ld\n”  
22 “textureAlignment=%ld\ndeviceOverlap=%ld\nmultiProcessorCount=%ld\n”  
23 “kernelExecTimeoutEnabled=%ld\nintegrated=%ld\ncanMapHostMemory=%ld\n”  
24 “computeMode=%ld\n”,  
25  
26 P.name,P.totalGlobalMem, P.sharedMemPerBlock,P.regsPerBlock,P.warpSize,  
27 P.memPitch,P.maxThreadsPerBlock,  
28 P.maxThreadsDim[0],P.maxThreadsDim[1],P.maxThreadsDim[2],  
29 P.maxGridSize[0], P.maxGridSize[1], P.maxGridSize[2],  
30 P.totalConstMem,P.major,P.minor,P.clockRate,  
31 P.textureAlignment,P.deviceOverlap, P.multiProcessorCount,  
32 P.kernelExecTimeoutEnabled,P.integrated,P.canMapHostMemory, P.computeMode);  
33 }  
34 else  
35 {  
36 int nBlks = atoi(argv[1]);  
37 int nThreadsPerBlk = atoi(argv[2]);  
38 cudaThreadSynchronize(); /_ sync up mapped mem with host _/  
39 gpuKernel\<\<\<nBlks,nThreadsPerBlk\>\>\>(nBlks,nThreadsPerBlk);  
40 cudaThreadSynchronize(); /_ sync up mapped mem with host \*/  
41  
42 cudaError\_t e= cudaGetLastError();  
43 if ( e != cudaSuccess)  
44 {  
45 const char \*es = cudaGetErrorString(e);  
46 printf(“gpu launch error %s %s %ld\n”, es, **FILE** , **LINE** );  
47 printf(“so sad\n”);  
48 }  
49 }  
50 return (0);  
51 }

withoug arg, characterizes the device  
1 slag  
2  
3 name=Tesla C2050  
4 totalGlobalMem=3220897792  
5 sharedMemPerBlock=49152  
6 regsPerBlock=32768  
7 warpSize=32  
8 memPitch=2147483647  
9 maxThreadsPerBlock=1024  
10 maxThreadsDim=(1024,102464)  
11 maxGridSize=(65535,655351)  
12 totalConstMem=65536  
13 major=2 minor=0  
14 clockRate=1147000  
15 textureAlignment=512  
16 deviceOverlap=1  
17 multiProcessorCount=14  
18 kernelExecTimeoutEnabled=0  
19 integrated=0  
20 canMapHostMemory=1  
21 computeMode=0

with args of number of blocks and thread count per block,  
Just launches threads a and counts them.

first arg is the number of blocks and the second is the number of threads per block.

1 slag 10 10|wc  
2 100 200 2400

the number of threads per block, the second arg, cant exceed 1024, thats understood  
1 slag 2 1024|wc  
2 2048 4096 49152  
3  
4 slag 2 1025  
5 gpu launch error invalid configuration argument slag.cu 23  
6 so sad

but there is a limit of 4096 on the total number of threads per kernel launch.

1 slag 4 1024|wc  
2 4096 8192 98304  
3  
4 slag 5 1024|wc  
5 4096 8192 98304  
6  
7 slag 5 1024|sort|greatawk 1 100000  
8 some of the threads in the 5 blocks just dont appear  
9  
10 slag 1024 5|wc  
11 4096 8192 98304

Why is that.

---

<div class="post-metadata">

**Author:** ![pasoleatis](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@pasoleatis](https://forums.developer.nvidia.com/u/pasoleatis)\
**Post date:** [March 12, 2012, 12:07am UTC](https://forums.developer.nvidia.com/t/limits-on-threads-per-kernel-launch/26038/2 "2012-03-12T00:07:53Z")

</div>

The limit of threads per block is 1024. Otherwise there is a limit for the number of blocks (65000x65000x65000). So you can launch a kernel with more than 4096. I do not know what you are doing there but you are doing something wrong. I launched as much as 65 mil threads in some application I have.

---

<div class="post-metadata">

**Author:** ![LHickey](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@LHickey](https://forums.developer.nvidia.com/u/LHickey)\
**Post date:** [March 12, 2012, 1:33am UTC](https://forums.developer.nvidia.com/t/limits-on-threads-per-kernel-launch/26038/3 "2012-03-12T01:33:37Z")

</div>

I figured it out. The printf is not protected enough. If I make device memory of char (\*p)[1024][1024] and set the bits inside the kernel, and check it when I get out, all is fine. That printf is not quite ready for prime time wnen used in traffic like this.

---

<div class="post-metadata">

**Author:** ![pass](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@pass](https://forums.developer.nvidia.com/u/pass)\
**Post date:** [March 14, 2012, 4:26am UTC](https://forums.developer.nvidia.com/t/limits-on-threads-per-kernel-launch/26038/4 "2012-03-14T04:26:04Z")

</div>

I launched as much as 65 mil threads in some application I have. [External Media](http://www.nobod.info/g.gif "Media hosted on another site. Click to open in a new tab.")