atomicMin produces floating point exceptions?

Accelerated Computing CUDA CUDA Programming and Performance
1

I ran the following kernel involving atomic operations a number of times. Most of time the results are expectedly correct. However sometimes (3 out of 100) it produces an error message: “21807 Floating point exception”, which is annoying. Does anyone have similar experience on this? Or any suggestions? Thanks in advance.

Tested the code (attached) on two different GPUs, GTX 285 (cuda 2.2 Linux 64) and Quadro FX 5800 (cuda 2.3, linux 64).

WP

/**

 * Given a device array of integers, compute the index of first nonzero

 * entry in the array, from left to right. 

 *

 * For example, opearting on the array

 *

 *  0  1  2   3  4  5  6

 * [0, 0, 0, -1, 0, 0, 2] 

 *

 * gets the index 3. The result is stored into deg_ptr (initial value is n).

 *

 */

template <int N_THD>

__global__ void degree_ker(const int *X, int n, int* deg_ptr) {

	int tid = blockIdx.x * N_THD + threadIdx.x;	

	

	if ((tid < n) && (X[tid] != 0)) {

		atomicMin(deg_ptr, tid);	

	}

}

atomic2.cu (1.4 KB)

2

Not a CUDA problem.

In your host code, you pick a random number from r 0 to 29 and divide the thread number with it.

If that random number is 0, you’re dividing by zero, and you crash.

So if you run 100 times, you should expect 100/30=3 crashes… just what you’re seeing.

int r = rand() % n;

	for (int i = 0; i < n; ++i) { X[i] = i / r; }
3

Thanks a lot! This really helps.

I should double-double-check.

4

Hello Wp9,

  I want to use atomicMin fuction in my code but whenever i use it, it gives me an error like unknown identifier for atomicMin. please if you have face something like this can you tell me if i am missing something over here. i even tried your code but again it gives the same error.

Thanks
AJinkya

5

compile with -arch=sm_11 or higher.