What causes division error in this CUDA kernel?

Question

I recently noticed an error when running our program on different hardware. I could track things down to the initial part of the kernel, where position within a grid ist calculated, using modulo (%) and division (/) operator. This is the minimal working example that reproduces the error:

#include <stdio.h>

__global__ void div_issue( int blocks_x, int* block_offset)
{
    int blks_x = blocks_x;


    //number of block 2d
    int block_id_2d = block_offset[0]; 

    //x-coordinate of block in absolute grid
    int block_idx = block_id_2d % blks_x;

    //y-coordinate of block in absolute grid    
    int block_idy = (block_id_2d - block_idx) / blks_x;

    printf("%d mod %d = %d \n", block_id_2d, blks_x, block_idx);
    printf("%d / %d = %d \n", block_id_2d - block_idx, blks_x, block_idy);

}    

int main(int argc, char *argv[])
{
    int dev_count;
    cudaGetDeviceCount(&dev_count);
    for (unsigned int i=0; i < dev_count; i++)
    {
        cudaSetDevice(i);
        cudaDeviceProp prop;
        cudaGetDeviceProperties(&prop, i);
        printf("using device %s :\n\n", prop.name);

        int block_offset_host[1];
        block_offset_host[0] = 753;
        int* block_offset_dev;
        cudaMalloc(&block_offset_dev, sizeof(int));
        cudaMemcpy(block_offset_dev, block_offset_host, sizeof(int), cudaMemcpyHostToDevice);

        div_issue <<<1, 1 >>>( 251 , block_offset_dev);

        cudaDeviceSynchronize();

        printf("\n");

    }
}

Result (on my machine having 2 GPUs):

using device GeForce GTX 980 Ti :
753 mod 251 = 0 
753 / 251 = 4 
using device GeForce GTX TITAN Black :
753 mod 251 = 0 
753 / 251 = 3 

I am using CUDA 7.0, Visual Studio 2012, 9.18.13.5306 WHQL on Windows 8.

I have no other hint, that the 980Ti could be damaged hardware. Can somebody confirm the issue on their own hardware?

So far it seems it does not happen in debug mode. No additional parameter such as -use_fast_math -prec-div=false -prec-sqrt=false was used for compilation.

Answer 1

I had the same error with cuda 7.5 rc compiler (cuda 7.5.7rc, linux ubuntu 14.04, titan X), however it seems to be fixed for cuda 7.5.18 (final release).