Stopping kernel using shared mem variable What I'm doing wrong ?

Accelerated Computing CUDA CUDA Programming and Performance
1

Here is the pseudo-code. What I’m trying to do is to prevent threads that are not launched yet from doing anything if one of previous threads met the ‘timed out’ condition (worked for more than allowed number of iterations).

shared bool bStopWorking;

global void Kernel(int* pTimedOut)
{
const int nThreadIdx = __mul24(blockDim.x, blockIdx.x) + threadIdx.x;
const int nThreadIdxInBlock = threadIdx.x;
const int nNumThreads = __mul24(blockDim.x, gridDim.x);

if (nThreadIdx == 0)
	bStopWorking = false;
__syncthreads();

for (int nTestCase = nThreadIdx;
	nTestCase < nNumTestCases;
	nTestCase += nNumThreads)
{
	if (bStopWorking)
		break;

	bool bTimedOut = false;
	EvaluateTestCase(&bTimedOut);

	if (bTimedOut)
	{
		bStopWorking = true;
		__syncthreads();
	}
}

if (bStopWorking)
	*pTimedOut = 1;

}

pTimedOut returns 1 when timed out condition is met, however, the workflow is not stopped and all the threads from the grid of blocks do their full job.

Thanks in advance!

2

a shared variable is shared within a block. So the rest of the blocks will never know they should stop. So you should write to a global variable.

3

Does that mean that shared bStopWorking should be raplaced with device bStopWorking ? I’ve tried this out - but something is still wrong …

Could you please correct the pseudo-code ? Looks like I’m missing something obvious …

4

Sorry, that was an obvious fault in my code, device variable bStopWorking works just fine, thank you!