for cycles error in nvcc in evaluating second argument in multidimensional array

Accelerated Computing CUDA CUDA Programming and Performance
1

This bug was tested on Tesla С2050 and GTX480,
and two nvcc versions — 3.2 и 4.0.

There is an error in nvcc in for cycle “for” in evaluating second argument in multidimensional array (secondbuf in this sample code):

Sample which force error (this code isn’t work right - gdb was not used, so what number it use for second argument instead of evaluating right value, I can’t…):
for(j=y-1; j<y+2; j++)
{for(k=z-1; k<z+2; k++) {
sharedbuf[j-y+1][k-z+1] = cubufferOld[( ((k+vz)%vz) * vy + ((j+vy)%vy) ) * vx + x];
}
}

while following variant with added temporary variable in second argument of sharedbuf is work fine:
for(j=y-1; j<y+2; j++)
{for(k=z-1; k<z+2; k++) {
int q = j-y+1;
sharedbuf[q][k-z+1] = cubufferOld[( ((k+vz)%vz) * vy + ((j+vy)%vy) ) * vx + x];
}
}

Of course, for manually unrolled cycle all is work fine too).
It is very interesting why second argument only.

2

What is the type of the variables ‘j’ and ‘y’? If either is of type unsigned int, the expression “j-y+1” will be evaluated in unsigned integer arithmetic per C/C++ type promotion rules, which could cause the result to be a very large positive number depending on the values of y and j, causing an out-of-bounds access.

3

all variables: x, y, z, vx, vy, vz, j, q, k are int.

For example we used vx=100, vy=100, vz=100.
x, y, z are evaluated from blockIdx and threadIdx.

If it is need, I can post full kernel code - ask.

4

The reason I brought up the ‘int’ vs ‘unsigned int’ issue is because just two weeks ago I root-caused a supposed compiler bug to this, so I thought it would be worth checking. If after careful checking of the code you think the problem is due to a problem with the compiler (entirely possible, there have been issues with addressing arithmetic before), please file a bug, attaching a self-contained repro case. The smaller the repro case the easier it will be to resolve the issue. Thanks.

5

Here two examples, that must have identical results, but they don’t have the same…:

ToSendNvidia.zip (8.1 KB)