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

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.

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.

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.

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.

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

ToSendNvidia.zip (8.1 KB)