Dumb question about if statement inside kernel


I’m CUDA newbie and I’m trying to implement image variance filter with CUDA. The problem that has occured to me is that “if” statement kind of does not work.

__global__ void __cudaVarianceFilter(cv::cuda::PtrStep<int> integralImg, cv::cuda::PtrStep<uint64_t > integralImg_squared,
		int * windows_d, int * d_inWinIndices, int numInWins, float minVar)
		int idx = blockDim.x * blockIdx.x + threadIdx.x;

		if (idx < numInWins) {
			int winIdx = d_inWinIndices[idx];
			int * win = &windows_d[winIdx * TLD_WINDOW_SIZE];

			int x1 = win[0] + 1;
			int y1 = win[1] + 1;
			int w = win[2];
			int h = win[3];
			int x2 = x1 + w - 1;
			int y2 = y1 + h - 1;
			float area = w * h;

			uint64_t mX = __fdividef(integralImg(y2, x2) - integralImg(y1 - 1, x2) - integralImg(y2, x1 - 1) + integralImg(y1 - 1, x1 - 1), area);
			uint64_t  mX2 = __fdividef((integralImg_squared(y2, x2) - integralImg_squared(y1 - 1, x2) - integralImg_squared(y2, x1 - 1) + integralImg_squared(y1 - 1, x1 - 1)), area);

			uint64_t variance = mX2 - __powf(mX, 2);

			#if __CUDA_ARCH__ >= 200
				printf("%d , minvar = %f \n", variance, minVar);
			if(variance <minVar)
				printf("The variance < minVar");
			d_inWinIndices[idx] = (variance < minVar)  ? -1 : d_inWinIndices[idx];

The problem is that the statement “if(variance < minVar)” is never true, and I do not know why, because I’ve even printf some of the values to see if maybe the calculations are wrong, but that is not the point. Here are some results of my code:

Can you please point out my mistake?
Thanks a lot

I’m not positive this is the issue, but won’t you have problems comparing an unsigned 64bit int to a signed float?

I think in this case, type casting minVar to uint64 might help alleviate the issue.

Or I could totally be wrong too. I’m not sure. I haven’t worked enough with uint64 to be positive.

I just tried that, and I think that is not the case.

Instead of using %d for your printf can you try using %u and observe what the output is in that case?
%d is for signed integers whereas %u is for unsigned ints.

%d is not the correct format specifier, neither is %u

You should be using %lu

You can get an indication that your method is not correct by noting that you are printing out negative values for a uint64_t, that could not possibly be correct.

Since your minvar is 1, the only way your if statement could ever be satisfied is if variance is zero. There are no other values in uint64_t space which are less than 1.

It worked! Thank you so much!

I totally missed the l. Thanks txbob. I didn’t realize that printf had a long modifier for ints and doubles.

The long modifier does not apply to double. It applies to int-like quantities.

%f is the correct printf format specifier for both float and double.


Isn’t there a long double modifier? “L”