Possible bug in ztranspose-v2.cu

[hsahasra]

Hi,

I'm having a problem with zgetri_gpu or zgesv_gpu when they perform LU factorization at zgetrf_gpu.cpp:167. I checked that all the pointers I pass to zgesv/zgetri are fine. What is more weird is that there is no such error when only the tests for zgesv/zgetri are run. This only happens when other operations are performed before zgetri/zgesv.

Code: Select all

Program received signal CUDA_EXCEPTION_10, Device Illegal Address.
[Switching focus to CUDA kernel 0, grid 94, block (0,124,0), thread (0,6,0), device 0, sm 2, warp 11, lane 0]
0x0000000005cee738 in ztranspose3_32<<<(2,125,1),(16,8,1)>>> (m32=0, n32=20, __val_paramB=0x136e880000, ldb=4000, __val_paramA=0x130fac0000, lda=3980, m=64, n=3980)
    at ztranspose-v2.cu:85
85	    sA[iny+ 8][inx] = A[ 8*lda];

I get this error when the matrix dimensions are not divisible by 32. I looked at the code at ztranspose-v2.cu:84 and found that there is no check for out-of-bounds memory access in that part:

Code: Select all

    sA[iny+ 0][inx] = A[ 0*lda];
    sA[iny+ 8][inx] = A[ 8*lda];
    sA[iny+16][inx] = A[16*lda];
    sA[iny+24][inx] = A[24*lda];

Whereas, a very similar memory access just a few lines back at ztranspose-v2.cu:30 has that check:

Code: Select all

    int t2 = iby+iny;
    if (ibx+inx < m) {
        if (t2   < n) {
            sA[iny+0][inx] = A[0*lda];
            if (t2+ 8 < n) {
                sA[iny+8][inx] = A[8*lda];
                if (t2 + 16 < n) {
                    sA[iny+16][inx] = A[16*lda];
                    if (t2 + 24 < n) {
                        sA[iny+24][inx] = A[24*lda];
                    }
                }
            }
        }
    }

Is this a bug in MAGMA or this is just how it's supposed to be?

Thanks,
Harshad

[hsahasra]

Just doing this probably fixes the error, not sure if it still gives the correct result.

Code: Select all

    if (ibx+inx+ZSIZE_1SHARED < m) {
        if (t2   < n) {
            sA[iny+0][inx] = A[0*lda];
            if (t2+ 8 < n) {
                sA[iny+8][inx] = A[8*lda];
                if (t2 + 16 < n) {
                    sA[iny+16][inx] = A[16*lda];
                    if (t2 + 24 < n) {
                        sA[iny+24][inx] = A[24*lda];
                    }
                }
            }
        }
    }

    /*sA[iny+ 0][inx] = A[ 0*lda];
    sA[iny+ 8][inx] = A[ 8*lda];
    sA[iny+16][inx] = A[16*lda];
    sA[iny+24][inx] = A[24*lda];*/

[mgates3]

Yes, there was a bug in that function. Thank you for reporting it. It has been fixed for the next MAGMA release.

You can test your fix by re-compiling and running testing_ztranspose.
-mark

[hsahasra]

Thank you. Is there a way I can access the code before it is released? Also, when are you planning to release MAGMA 1.5.0?

Thanks,
Harshad