Possible bug in ztranspose-v2.cu

Open discussion for MAGMA library (Matrix Algebra on GPU and Multicore Architectures)
Post Reply
hsahasra
Posts: 32
Joined: Mon Jun 24, 2013 3:40 pm

Possible bug in ztranspose-v2.cu

Post by hsahasra » Fri Jun 20, 2014 10:03 am

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
Last edited by hsahasra on Fri Jun 27, 2014 10:52 am, edited 1 time in total.

hsahasra
Posts: 32
Joined: Mon Jun 24, 2013 3:40 pm

Re: Possible bug in ztranspose-v2.cu

Post by hsahasra » Fri Jun 20, 2014 10:56 am

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
Posts: 918
Joined: Fri Jan 06, 2012 2:13 pm

Re: Possible bug in ztranspose-v2.cu

Post by mgates3 » Mon Jul 14, 2014 2:35 pm

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
Posts: 32
Joined: Mon Jun 24, 2013 3:40 pm

Re: Possible bug in ztranspose-v2.cu

Post by hsahasra » Mon Jul 14, 2014 4:15 pm

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

Post Reply