# Loop limit in CUDA kernel ? Too large loop =\> loop not launched

**URL:** <https://forums.developer.nvidia.com/t/loop-limit-in-cuda-kernel-too-large-loop-loop-not-launched/20035>\
**Category:** CUDA Programming and Performance\
**Created:** [November 30, 2010, 7:43pm UTC](https://forums.developer.nvidia.com/t/loop-limit-in-cuda-kernel-too-large-loop-loop-not-launched/20035 "2010-11-30T19:43:04Z")\
**Posts on this page:** 13\
**Page:** 1

<div class="post-metadata">

**Author:** ![nezix](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/nezix/32/28228_2.png) [@nezix](https://forums.developer.nvidia.com/u/nezix)\
**Post date:** [November 30, 2010, 7:43pm UTC](https://forums.developer.nvidia.com/t/loop-limit-in-cuda-kernel-too-large-loop-loop-not-launched/20035/1 "2010-11-30T19:43:04Z")

</div>

Hello,

I am wondering if my GPU is not lazy : I have to do a pretty big loop in my kernel and even if I split it into several loops, it just does not start it/them.

```auto
float dt=0.005;

...

__global__ void kernel(float array,...){

   float it;...

   for(it=5000;it<50000;it+=dt)

      new=do(stuff);

   array[idx]=new; //retrieve_stuff

}

```

When I cudaMemcpy back to my RAM nothing has changed between my host initialization and the result. But kernel takes effective time to run even if cuda timer says otherwise. I checked every small detail like memcpys, initialisations … I retrieved my loop indice and it does NOT change after my loop.

Whereas if I change 50 000 to 8000 the loop runs.

I hope this is not a too simple question but is there a limit for loops in CUDA ? What is the thing I am doing wrong ?

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:** [November 30, 2010, 8:33pm UTC](https://forums.developer.nvidia.com/t/loop-limit-in-cuda-kernel-too-large-loop-loop-not-launched/20035/2 "2010-11-30T20:33:34Z")

</div>

Yes, there is a limit for loops in CUDA. It is the same as in C: You cannot loop over more iterations than the loop counter has different states.  
Increase [font=“Courier New”]dt[/font], or switch to double precision.

---

<div class="post-metadata">

**Author:** ![nezix](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/nezix/32/28228_2.png) [@nezix](https://forums.developer.nvidia.com/u/nezix)\
**Post date:** [November 30, 2010, 8:45pm UTC](https://forums.developer.nvidia.com/t/loop-limit-in-cuda-kernel-too-large-loop-loop-not-launched/20035/3 "2010-11-30T20:45:21Z")

</div>

Actually, the loop does not cause troubles in C. And as I said even if I split my loop in 2 different ones it does not the trick, and double precision doesn’t help too (I have a 470 gtx and I am compiling with -arch=sm\_20)

---

<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:** [November 30, 2010, 8:57pm UTC](https://forums.developer.nvidia.com/t/loop-limit-in-cuda-kernel-too-large-loop-loop-not-launched/20035/4 "2010-11-30T20:57:07Z")

</div>

Ok, so you may have a second problem. Do you check return codes?

---

<div class="post-metadata">

**Author:** ![nezix](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/nezix/32/28228_2.png) [@nezix](https://forums.developer.nvidia.com/u/nezix)\
**Post date:** [December 1, 2010, 8:40am UTC](https://forums.developer.nvidia.com/t/loop-limit-in-cuda-kernel-too-large-loop-loop-not-launched/20035/5 "2010-12-01T08:40:28Z")

</div>

Yes but no error occures.

---

<div class="post-metadata">

**Author:** ![LSChien](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@LSChien](https://forums.developer.nvidia.com/u/LSChien)\
**Post date:** [December 1, 2010, 9:47am UTC](https://forums.developer.nvidia.com/t/loop-limit-in-cuda-kernel-too-large-loop-loop-not-launched/20035/6 "2010-12-01T09:47:23Z")

</div>

> [@](#):
>
> Yes but no error occures.

Could you post the code which checks error code?

> [@](#):
>
> Whereas if I change 50 000 to 8000 the loop runs.

I wonder if it is watchdog problem, do you connect your GTX470 to display?

---

<div class="post-metadata">

**Author:** ![nezix](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/nezix/32/28228_2.png) [@nezix](https://forums.developer.nvidia.com/u/nezix)\
**Post date:** [December 1, 2010, 12:57pm UTC](https://forums.developer.nvidia.com/t/loop-limit-in-cuda-kernel-too-large-loop-loop-not-launched/20035/7 "2010-12-01T12:57:23Z")

</div>

> [@](#):
>
> Could you post the code which checks error code?

Taken from dr dobb’s site.

```auto
void checkCUDAError(const char *msg){

	cudaError_t err = cudaGetLastError();

	if( cudaSuccess != err){

		fprintf(stderr, "Cuda error: %s: %s.\n", msg,

		cudaGetErrorString( err) );

		exit(EXIT_FAILURE);

	}

}

```

> [@](#):
>
> I wonder if it is watchdog problem, do you connect your GTX470 to display?

Yes I do. I also time my all program with the “time” unix command and it takes more time when I call my kernel than when I don’t.

---

<div class="post-metadata">

**Author:** ![nezix](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/nezix/32/28228_2.png) [@nezix](https://forums.developer.nvidia.com/u/nezix)\
**Post date:** [December 1, 2010, 2:33pm UTC](https://forums.developer.nvidia.com/t/loop-limit-in-cuda-kernel-too-large-loop-loop-not-launched/20035/8 "2010-12-01T14:33:17Z")

</div>

You are right my program is working (my results are wrong but maybe it is an other story) when I stop gnome. Do you know what causes that ?

---

<div class="post-metadata">

**Author:** ![seibert](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/seibert/32/520688_2.png) [@seibert](https://forums.developer.nvidia.com/u/seibert)\
**Post date:** [December 1, 2010, 3:03pm UTC](https://forums.developer.nvidia.com/t/loop-limit-in-cuda-kernel-too-large-loop-loop-not-launched/20035/9 "2010-12-01T15:03:09Z")

</div>

> [@](#):
>
> You are right my program is working (my results are wrong but maybe it is an other story) when I stop gnome. Do you know what causes that ?

If you mean that you have stopped the X server entirely, then you are probably seeing the effect of the watchdog timer. To keep the GUI responsive, the driver will terminate any kernel that takes longer than ~5 seconds. If no GUI is running, then there is no time limit. When the kernel is terminated this way, you should get some error return value from the next CUDA call.

---

<div class="post-metadata">

**Author:** ![Daniel\_nyu](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Daniel\_nyu](https://forums.developer.nvidia.com/u/Daniel_nyu)\
**Post date:** [June 10, 2011, 6:57pm UTC](https://forums.developer.nvidia.com/t/loop-limit-in-cuda-kernel-too-large-loop-loop-not-launched/20035/10 "2011-06-10T18:57:19Z")

</div>

I have the same exact problem but I don’t think it’s a watchdog issue. My code is structured very similarly to the example code:

**global** void advanceSystem(float \*cx, unsigned long int N)  
{  
int idx = blockIdx.x \* blockDim.x + threadIdx.x;  
unsigned long int i;  
float temp;

```
unsigned int lrng_state;

temp = cx[idx];

i=1;
while (i <= N) {
	temp = do stuff
            i++;
}

cx[idx] = temp;

```

}

I have even tried running it by saving N in **constant** memory with cudaMemcpy instead of passign it as an argument to the kernel and I get the same issues. I managet to run the kernel with N=10000 for the most part but sometimes it works with N=1000000. Sometimes (more rarely) it fails even with N=10000.  
If it is of any importance, in the instances where it fails even for N=10000, I am calling the kernel in a loop from my C++ code:

while(k\<something\_big) {  
advanceSystem\<\<num\_threads/64, 64\>\>(cx, N);

}

---

<div class="post-metadata">

**Author:** ![Daniel\_nyu](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Daniel\_nyu](https://forums.developer.nvidia.com/u/Daniel_nyu)\
**Post date:** [June 10, 2011, 7:03pm UTC](https://forums.developer.nvidia.com/t/loop-limit-in-cuda-kernel-too-large-loop-loop-not-launched/20035/11 "2011-06-10T19:03:08Z")

</div>

I have the same exact problem but I don’t think it’s a watchdog issue. My code is structured very similarly to the example code:

**global** void advanceSystem(float \*cx, unsigned long int N)  
{  
int idx = blockIdx.x \* blockDim.x + threadIdx.x;  
unsigned long int i;  
float temp;  
//here I declare other stuff that i need on this thread//

```
unsigned int lrng_state;

temp = cx[idx];

i=1;
while (i <= N) {
	temp = do stuff
            i++;
}

cx[idx] = temp;

```

}

I have even tried running it by saving N in **constant** memory with cudaMemcpy instead of passign it as an argument to the kernel and I get the same issues. I managet to run the kernel with N=10000 for the most part but sometimes it works with N=1000000. Sometimes (more rarely) it fails even with N=10000.  
If it is of any importance, in the instances where it fails even for N=10000, I am calling the kernel in a loop from my C++ code:

while(k\<something\_big) {  
advanceSystem\<\<num\_threads/64, 64\>\>(cx, N);  
cudaThreadSynchronize();  
k++;  
}

I’m starting to think that it must be a memory leak or something.  
For further clarification, the thread\_dependent stuff I declare in the kernel constitutes of 8 vectors of type float (2 floats for each vector) and two additional floats.

If someone could help me I would be immensely grateful.  
I’m goign crazy over this.

Daniele

---

<div class="post-metadata">

**Author:** ![nezix](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/nezix/32/28228_2.png) [@nezix](https://forums.developer.nvidia.com/u/nezix)\
**Post date:** [June 11, 2011, 1:35pm UTC](https://forums.developer.nvidia.com/t/loop-limit-in-cuda-kernel-too-large-loop-loop-not-launched/20035/12 "2011-06-11T13:35:05Z")

</div>

To be sure the kernel is not interrupted by the wachdog timer, just run it with your X server shut down.

An easy way to see if your loop is running is to do :

```auto
__global__ void advanceSystem(float *cx, unsigned long int N)

{

int idx = blockIdx.x * blockDim.x + threadIdx.x;

unsigned long int i;

//here I declare other stuff that i need on this thread//

i=1;

while (i <= N) {

i++;

}

cx[idx] =(float)i;

}

```

Then you check if all the cx array is full of your N value.

---

<div class="post-metadata">

**Author:** ![Skybuck](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/skybuck/32/9717_2.png) [@Skybuck](https://forums.developer.nvidia.com/u/Skybuck)\
**Post date:** [June 12, 2011, 2:34pm UTC](https://forums.developer.nvidia.com/t/loop-limit-in-cuda-kernel-too-large-loop-loop-not-launched/20035/13 "2011-06-12T14:34:11Z")

</div>

What is number of threads in the kernel call ?