i noticed error when running our program on different hardware. track things down initial part of kernel, position within grid ist calculated, using modulo (%) , division (/) operator. minimal working example reproduces 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); (unsigned int i=0; < 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 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 using cuda 7.0, visual studio 2012, 9.18.13.5306 whql on windows 8.
i have no other hint, 980ti damaged hardware. can confirm issue on own hardware?
so far seems not happen in debug mode. no additional parameter such -use_fast_math -prec-div=false -prec-sqrt=false used compilation.
i had same error cuda 7.5 rc compiler (cuda 7.5.7rc, linux ubuntu 14.04, titan x), seems fixed cuda 7.5.18 (final release).
Comments
Post a Comment