OpenACC vector length > 256 gives launch failure

Hi,
I have a very simple Jacobi example. However, if I play around with different vector sizes (loop schedules), I get the error “call to cuStreamSynchronize returned error 700: Launch failed” when taking sizes > 256. That is a problem for 14.2, 14.1, 13.10, 13.9. With 13.6 this code is correctly working.
I assume that is a compiler bug, isn’t it? If so, I can send you the whole example if needed. If not, let me know what I am doing wrong. But actually, changing the loop schedule should not effect the output (as long as I am in a valid thread number range).
BTW: I use a NVIDIA Quadro 6000 (Fermi) and the flags -acc -ta:nvidia,cuda5.0

  #pragma acc data copy(A[0:m*n]) create(Anew[0:m*n])
  {
    while ( err > tol && iter < iter_max ) {

        err = 0.0;

        #pragma acc parallel loop present(Anew[0:n*m],A[0:n*m]) reduction(max:err) vector_length(257)
        for( j = 1; j < n-1; j++) {
            for( i = 1; i < m-1; i++ ) {
                Anew[j *m+ i] = 0.25 * ( A[j     *m+ (i+1)] + A[j     *m+ (i-1)]
                                     +   A[(j-1) *m+ i]     + A[(j+1) *m+ i]);
                err = fmax(err,fabs(Anew[j*m+i]-A[j*m+i]));
            }
        }

        #pragma acc parallel loop present(Anew[0:n*m],A[0:n*m])
        for( j = 1; j < n-1; j++) {
            for( i = 1; i < m-1; i++ ) {
                A[j *m+ i] = Anew[j *m+ i];
            }
        }

        if(iter % 10 == 0) {
            printf("%5d, %0.6f\n", iter, err);
        }

        iter++;
    } // end while
  }

The code uses dynamically allocated arrays.
Sandra

I assume that is a compiler bug, isn’t it? If so, I can send you the whole example if needed.

It’s possibly a compiler issue (maybe the reduction?) but it could be something else as well. Please do send me the program and I’ll see what I can determine.

Thanks,
Mat

I just did. Let me know if there is a workaround or a fix.
Thanks, Sandra

Hi Sandra,

Looks like this is the same issue as a known problem (TPR#18947) which occurs when the threads used in the kernel are greater than the number used in the compiler generated reduction kernel. There’s no work around other than to either not do the reduction or limit the number of threads to 256.

I’ll add your information to the TPR.

Thanks,
Mat