while loop in cuda breaks out

Because of kernel launch latency I thought about having one kernel that has a while loop and breaks out of that loop if the host writes “abort flag” in the host’s pinned memory and the kernel reads that using “zero-copy”. I don’t even know if this is a dumb idea because I am not experienced. So I wrote a test program.

Here is a stripped down version that runs expectedly.

[codebox]global void gpuThread(unsigned * flag, unsigned * inB, unsigned * outB)

{

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

int k = 0;

while (1)

{

    if (i < NUM_SAMPLES)

    {

        k++;

    }

    if (k == 100000)

        return;

}

}

//

[/codebox]

However if I remove the return it does not run forever.

[codebox]

global void gpuThread(unsigned * flag, unsigned * inB, unsigned * outB)

{

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

int k = 0;

while (1)

{

    if (i < NUM_SAMPLES)

    {

        k++;

    }

}

}

//

[/codebox]

It quits after some random time between 9 seconds to 15 seconds. I have a GTX 285 and running linux. I read some timeout on Vista. But if that was the case for linux as well it would quit at consistent times right? I would really appreciate if someone could explain what’s going on.

Thanks

No, Linux also has a kernel watchdog timer. It cannot be disabled if the device is used for display.

The watchdog is not a precise timer. It’s likely some low priority polling in the driver so indeed sometimes kernels can run different times from run to run.

I think in linux, watchdog plays a role if you start “X”

Thanks for the replies guys. I will look into disabling the watchdog timer. The card is not used for display. Any idea how I can disable it?

OK, GPU is not used for display but still same problem exists. Something kills the GPU kernel around 10 seconds or so. This is really frustrating. I appreciate any suggestions. Thanks…