Index: magma_morse/compute/pzgetrf_incpiv.c
===================================================================
--- magma_morse/compute/pzgetrf_incpiv.c	(révision 1174)
+++ magma_morse/compute/pzgetrf_incpiv.c	(copie de travail)
@@ -44,8 +44,13 @@
     RUNTIME_options_init(&options, morse, sequence, request);
 
     ib = MORSE_IB;
-    h_work_size  = sizeof(MORSE_Complex64_t)*( ib*L->nb );
-    d_work_size  = 0;
+#if defined(MORSE_USE_CUDA)
+    h_work_size = sizeof(PLASMA_Complex64_t)*( 2*ib + 2*dL->nb )*2*dA->mb;
+    d_work_size = sizeof(PLASMA_Complex64_t)*(   ib            )*2*dA->mb;
+#else
+    h_work_size = sizeof(MORSE_Complex64_t)*( ib*L->nb );
+    d_work_size = 0;
+#endif
     RUNTIME_options_ws_alloc( &options, h_work_size, d_work_size );
 
     for (k = 0; k < min(A->mt, A->nt); k++) {
@@ -55,7 +60,9 @@
         MORSE_TASK_zgetrf_incpiv(
             &options,
             tempkm, tempkn, ib, L->nb,
-            A(k, k), ldak, IPIV(k, k),
+            A(k, k), ldak,
+            L(k, k), L->mb,
+            IPIV(k, k),
             k == A->mt-1, A->nb*k);
 
         for (n = k+1; n < A->nt; n++) {
@@ -65,6 +72,7 @@
                 tempkm, tempnn, tempkm, ib, L->nb,
                 IPIV(k, k),
                 A(k, k), ldak,
+                L(k, k), L->mb,
                 A(k, n), ldak);
         }
         for (m = k+1; m < A->mt; m++) {
Index: magma_morse/include/runtime_z.h
===================================================================
--- magma_morse/include/runtime_z.h	(révision 1174)
+++ magma_morse/include/runtime_z.h	(copie de travail)
@@ -101,6 +101,7 @@
 void MORSE_TASK_zgessm(MORSE_option_t *options,
                        int m, int n, int k, int ib, int nb,
                        int *IPIV,
+                       MORSE_desc_t *D, int Dm, int Dn, int ldd,
                        MORSE_desc_t *L, int Lm, int Ln, int ldl,
                        MORSE_desc_t *A, int Am, int An, int lda);
 void MORSE_TASK_zgessq_f1( MORSE_option_t *options,
@@ -111,25 +112,24 @@
                        int m, int n, int nb,
                        MORSE_desc_t *A, int Am, int An, int lda,
                        int *IPIV,
-                        
                        MORSE_bool check_info, int iinfo);
 void MORSE_TASK_zgetrf_incpiv(MORSE_option_t *options,
                               int m, int n, int ib, int nb,
                               MORSE_desc_t *A, int Am, int An, int lda,
+                              MORSE_desc_t *L, int Al, int Al, int ldl,
                               int *IPIV,
-                               
                               MORSE_bool check_info, int iinfo);
 void MORSE_TASK_zgetrf_reclap(MORSE_option_t *options,
                               int m, int n, int nb,
                               MORSE_desc_t *A, int Am, int An, int lda,
                               int *IPIV,
-                               
+
                               MORSE_bool check_info, int iinfo,
                               int nbthread);
 void MORSE_TASK_zgetrf_rectil(MORSE_option_t *options,
                               MORSE_desc_t A, MORSE_desc_t *Amn, int Amnm, int Amnn, int size,
                               int *IPIV,
-                               
+
                               MORSE_bool check_info, int iinfo,
                               int nbthread);
 void MORSE_TASK_zgetrip(MORSE_option_t *options,
@@ -151,7 +151,7 @@
                        int itype, MORSE_enum uplo, int N,
                        MORSE_desc_t *A, int Am, int An, int LDA,
                        MORSE_desc_t *B, int Bm, int Bn, int LDB,
-                        
+
                        int iinfo);
 void MORSE_TASK_zherk(MORSE_option_t *options,
                       MORSE_enum uplo, MORSE_enum trans,
@@ -251,7 +251,7 @@
 void MORSE_TASK_zpotrf(MORSE_option_t *options,
                        MORSE_enum uplo, int n, int nb,
                        MORSE_desc_t *A, int Am, int An, int lda,
-                        
+
                        int iinfo);
 void MORSE_TASK_zshift( MORSE_option_t *options,
                         int s, int m, int n, int L,
@@ -320,7 +320,7 @@
 void MORSE_TASK_ztrtri(MORSE_option_t *options,
                        MORSE_enum uplo, MORSE_enum diag, int n, int nb,
                        MORSE_desc_t *A, int Am, int An, int lda,
-                        
+
                        int iinfo);
 void MORSE_TASK_ztslqt(MORSE_option_t *options,
                        int m, int n, int ib, int nb,
@@ -380,7 +380,7 @@
                        MORSE_desc_t *A, int Am, int An, int lda,
                        MORSE_desc_t *L, int Lm, int Ln, int ldl,
                        int *IPIV,
-                        
+
                        MORSE_bool check_info, int iinfo);
 void MORSE_TASK_zttmqr(MORSE_option_t *options,
                        MORSE_enum side, MORSE_enum trans,
Index: magma_morse/runtime/starpu/codelets/codelet_zgessm.c
===================================================================
--- magma_morse/runtime/starpu/codelets/codelet_zgessm.c	(révision 1174)
+++ magma_morse/runtime/starpu/codelets/codelet_zgessm.c	(copie de travail)
@@ -70,6 +70,7 @@
 void MORSE_TASK_zgessm(MORSE_option_t *options,
                        int m, int n, int k, int ib, int nb,
                        int *IPIV,
+                       MORSE_desc_t *D, int Dm, int Dn, int ldd,
                        MORSE_desc_t *L, int Lm, int Ln, int ldl,
                        MORSE_desc_t *A, int Am, int An, int lda)
 {
@@ -77,17 +78,19 @@
     struct starpu_codelet *codelet = &cl_zgessm;
     void (*callback)(void*) = options->profiling ? cl_zgessm_callback : NULL;
     starpu_insert_task(codelet,
-        STARPU_VALUE,     &m,                        sizeof(int),
-        STARPU_VALUE,     &n,                        sizeof(int),
-        STARPU_VALUE,     &k,                        sizeof(int),
+        STARPU_VALUE,    &m,                         sizeof(int),
+        STARPU_VALUE,    &n,                         sizeof(int),
+        STARPU_VALUE,    &k,                         sizeof(int),
         STARPU_VALUE,    &ib,                        sizeof(int),
-        STARPU_VALUE,          &IPIV,                      sizeof(int*),
-        STARPU_R,             RTBLKADDR(L, MORSE_Complex64_t, Lm, Ln),
-        STARPU_VALUE,   &ldl,                        sizeof(int),
-        STARPU_RW,             RTBLKADDR(A, MORSE_Complex64_t, Am, An),
-        STARPU_VALUE,   &lda,                        sizeof(int),
-        STARPU_PRIORITY,    options->priority,
-        STARPU_CALLBACK,    callback,
+        STARPU_VALUE,    &IPIV,                      sizeof(int*),
+        STARPU_R,        RTBLKADDR(D, MORSE_Complex64_t, Dm, Dn),
+        STARPU_VALUE,    &ldd,                       sizeof(int),
+        STARPU_R,        RTBLKADDR(L, MORSE_Complex64_t, Lm, Ln),
+        STARPU_VALUE,    &ldl,                       sizeof(int),
+        STARPU_RW,       RTBLKADDR(A, MORSE_Complex64_t, Am, An),
+        STARPU_VALUE,    &lda,                       sizeof(int),
+        STARPU_PRIORITY, options->priority,
+        STARPU_CALLBACK, callback,
         0);
 }
 
@@ -98,17 +101,57 @@
     int k;
     int ib;
     int *IPIV;
-    MORSE_Complex64_t *L;
+    MORSE_Complex64_t *D;
+    int ldd;
     int ldl;
     MORSE_Complex64_t *A;
     int lda;
 
-    L = (MORSE_Complex64_t *)STARPU_MATRIX_GET_PTR(descr[0]);
-    A = (MORSE_Complex64_t *)STARPU_MATRIX_GET_PTR(descr[1]);
-    starpu_codelet_unpack_args(cl_arg, &m, &n, &k, &ib, &IPIV, &ldl, &lda);
-    CORE_zgessm(m, n, k, ib, IPIV, L, ldl, A, lda);
+    D = (MORSE_Complex64_t *)STARPU_MATRIX_GET_PTR(descr[0]);
+    A = (MORSE_Complex64_t *)STARPU_MATRIX_GET_PTR(descr[2]);
+    starpu_codelet_unpack_args(cl_arg, &m, &n, &k, &ib, &IPIV, &ldd, &ldl, &lda);
+    CORE_zgessm(m, n, k, ib, IPIV, D, ldd, A, lda);
 }
+
 /*
+ * Codelet GPU
+ */
+#ifdef MORSE_USE_CUDA
+static void cl_zgessm_cuda_func(void *descr[], void *cl_arg)
+{
+    int m;
+    int n;
+    int k;
+    int ib;
+    int *IPIV;
+    cuDoubleComplex *dD;
+    int ldd;
+    cuDoubleComplex *dL;
+    int ldl;
+    cuDoubleComplex *dA;
+    int lda;
+    int info;
+
+    dD = (cuDoubleComplex *)STARPU_MATRIX_GET_PTR(descr[0]);
+    dL = (cuDoubleComplex *)STARPU_MATRIX_GET_PTR(descr[1]);
+    dA = (cuDoubleComplex *)STARPU_MATRIX_GET_PTR(descr[2]);
+    dL += ib; /* The kernel is just using the inverted part */
+
+    starpu_codelet_unpack_args(cl_arg, &m, &n, &k, &ib, &IPIV, &ldd, &ldl, &lda);
+
+    magma_zgessm_gpu(
+        'C', m, n, k, ib,
+        IPIV,
+        dL, ldl,
+        dD, ldd,
+        dA, lda,
+        &info);
+
+    cudaThreadSynchronize();
+}
+#endif
+
+/*
  * Codelet definition
  */
-CODELETS_CPU(zgessm, 2, cl_zgessm_cpu_func)
+CODELETS(zgessm, 3, cl_zgessm_cpu_func, cl_zgessm_cuda_func)
Index: magma_morse/runtime/starpu/codelets/codelet_zgetrf_incpiv.c
===================================================================
--- magma_morse/runtime/starpu/codelets/codelet_zgetrf_incpiv.c	(révision 1174)
+++ magma_morse/runtime/starpu/codelets/codelet_zgetrf_incpiv.c	(copie de travail)
@@ -78,6 +78,7 @@
 void MORSE_TASK_zgetrf_incpiv(MORSE_option_t *options,
                               int m, int n, int ib, int nb,
                               MORSE_desc_t *A, int Am, int An, int lda,
+                              MORSE_desc_t *L, int Lm, int Ln, int ldl,
                               int *IPIV,
                               MORSE_bool check_info, int iinfo)
 {
@@ -85,16 +86,20 @@
     struct starpu_codelet *codelet = &cl_zgetrf_incpiv;
     void (*callback)(void*) = options->profiling ? cl_zgetrf_incpiv_callback : NULL;
     starpu_insert_task(codelet,
-        STARPU_VALUE,             &m,                        sizeof(int),
-        STARPU_VALUE,             &n,                        sizeof(int),
-        STARPU_VALUE,            &ib,                        sizeof(int),
-        STARPU_RW,                     RTBLKADDR(A, MORSE_Complex64_t, Am, An),
-        STARPU_VALUE,           &lda,                        sizeof(int),
-        STARPU_VALUE,                  &IPIV,                      sizeof(int*),
+        STARPU_VALUE,    &m,                         sizeof(int),
+        STARPU_VALUE,    &n,                         sizeof(int),
+        STARPU_VALUE,    &ib,                        sizeof(int),
+        STARPU_RW,        RTBLKADDR(A, MORSE_Complex64_t, Am, An),
+        STARPU_VALUE,    &lda,                       sizeof(int),
+        STARPU_RW,        RTBLKADDR(L, MORSE_Complex64_t, Lm, Ln),
+        STARPU_VALUE,    &ldl,                       sizeof(int),
+        STARPU_VALUE,    &IPIV,                      sizeof(int*),
+        STARPU_VALUE,    &h_work,                    sizeof(MORSE_starpu_ws_t *),
+        STARPU_VALUE,    &d_work,                    sizeof(MORSE_starpu_ws_t *),
         STARPU_VALUE,    &check_info,                sizeof(MORSE_bool),
-        STARPU_VALUE,         &iinfo,                        sizeof(int),
-        STARPU_PRIORITY,    options->priority,
-        STARPU_CALLBACK,    callback,
+        STARPU_VALUE,    &iinfo,                     sizeof(int),
+        STARPU_PRIORITY,  options->priority,
+        STARPU_CALLBACK,  callback,
         0);
 }
 
@@ -104,7 +109,7 @@
     int n;
     int ib;
     MORSE_Complex64_t *A;
-    int lda;
+    int lda, ldl;
     int *IPIV;
     MORSE_bool check_info;
     int iinfo;
@@ -112,10 +117,85 @@
     int info = 0;
 
     A = (MORSE_Complex64_t *)STARPU_MATRIX_GET_PTR(descr[0]);
-    starpu_codelet_unpack_args(cl_arg, &m, &n, &ib, &lda, &IPIV, &check_info, &iinfo);
+    starpu_codelet_unpack_args(cl_arg, &m, &n, &ib, &lda, &ldl, &IPIV,
+                               &h_work, &d_work, &check_info, &iinfo);
     CORE_zgetrf_incpiv(m, n, ib, A, lda, IPIV, &info);
+
+#if defined(MORSE_USE_CUDA)
+    {
+        PLASMA_Complex64_t *L;
+        L = (PLASMA_Complex64_t *)STARPU_MATRIX_GET_PTR(descr[1]);
+
+        /*
+         * L stores:
+         *      L1     L2    L3     ...
+         *      L1^-1  L2^-1 L3^-1  ...
+         */
+        /* Compute L-1 in lower rectangle of L */
+        L += ib;
+        {
+            int i, sb;
+            for (i=0; i<n; i+=ib) {
+                sb = min( ib, n-i );
+                CORE_zlacpy(PlasmaUpperLower, sb, sb, A+(i*lda+i), lda, L+(i*ldl), ldl );
+
+                CORE_ztrtri( PlasmaLower, PlasmaUnit, sb, L+(i*ldl), ldl, &info );
+                if (info != 0 ) {
+                    fprintf(stderr, "ERROR, trtri returned with info = %d\n", info);
+                }
+            }
+        }
+    }
+#endif
 }
+
+#if defined(MORSE_USE_CUDA)
+static void cl_zgetrl_cuda_func(void *descr[], void *cl_arg)
+{
+    morse_starpu_ws_t *h_work;
+    morse_starpu_ws_t *d_work;
+    cuDoubleComplex *hA, *dA;
+    cuDoubleComplex *hL, *dL;
+    cuDoubleComplex *dwork;
+    int *IPIV;
+    int m, n, ib;
+    int lda, ldl;
+    MORSE_bool check_info;
+    int iinfo;
+    int info = 0;
+
+    dA = (cuDoubleComplex *)STARPU_MATRIX_GET_PTR(descr[0]);
+    dL = (cuDoubleComplex *)STARPU_MATRIX_GET_PTR(descr[1]);
+    starpu_codelet_unpack_args(cl_arg, &m, &n, &ib, &lda, &ldl, &IPIV,
+                               &h_work, &d_work, &check_info, &iinfo);
+
+    /*
+     * hwork => at least (2*IB+NB)*NB contains all hA and hL
+     * dwork => at least IB*NB
+     */
+    hA    = morse_starpu_ws_getlocal(h_work);
+    dwork = morse_starpu_ws_getlocal(d_work);
+
+    hL = hA + lda*n;
+
+    /* Initialize L to 0 */
+    memset(hL, 0, ldl*n*sizeof(cuDoubleComplex));
+
+    /* Copy First panel */
+    cublasGetMatrix( m, min(ib,m), sizeof(cuDoubleComplex), dA, lda, hA, lda );
+
+    magma_zgetrf_incpiv_gpu( 'C', m, n, ib,
+                             hA, lda, dA, lda,
+                             hL, ldl, dL, ldl,
+                             IPIV,
+                             dwork, lda,
+                             &info );
+
+    cudaThreadSynchronize();
+}
+#endif
+
 /*
  * Codelet definition
  */
-CODELETS_CPU(zgetrf_incpiv, 1, cl_zgetrf_incpiv_cpu_func)
+CODELETS(zgetrf_incpiv, 2, cl_zgetrf_incpiv_cpu_func)
Index: magma_morse/runtime/starpu/codelets/codelet_zssssm.c
===================================================================
--- magma_morse/runtime/starpu/codelets/codelet_zssssm.c	(révision 1174)
+++ magma_morse/runtime/starpu/codelets/codelet_zssssm.c	(copie de travail)
@@ -113,11 +113,11 @@
         STARPU_VALUE,  &lda1,                        sizeof(int),
         STARPU_RW,            RTBLKADDR(A2, MORSE_Complex64_t, A2m, A2n),
         STARPU_VALUE,  &lda2,                        sizeof(int),
-        STARPU_R,            RTBLKADDR(L1, MORSE_Complex64_t, L1m, L1n),
+        STARPU_R,             RTBLKADDR(L1, MORSE_Complex64_t, L1m, L1n),
         STARPU_VALUE,  &ldl1,                        sizeof(int),
-        STARPU_R,            RTBLKADDR(L2, MORSE_Complex64_t, L2m, L2n),
+        STARPU_R,             RTBLKADDR(L2, MORSE_Complex64_t, L2m, L2n),
         STARPU_VALUE,  &ldl2,                        sizeof(int),
-        STARPU_VALUE,          &IPIV,                      sizeof(int*),
+        STARPU_VALUE,    &IPIV,                      sizeof(int*),
         STARPU_PRIORITY,    options->priority,
         STARPU_CALLBACK,    callback,
         0);
@@ -148,6 +148,45 @@
     starpu_codelet_unpack_args(cl_arg, &m1, &n1, &m2, &n2, &k, &ib, &lda1, &lda2, &ldl1, &ldl2, &IPIV);
     CORE_zssssm(m1, n1, m2, n2, k, ib, A1, lda1, A2, lda2, L1, ldl1, L2, ldl2, IPIV);
 }
+
+#ifdef MORSE_USE_CUDA
+static void cl_zssssm_cuda_func(void *descr[], void *cl_arg)
+{
+    int m1;
+    int n1;
+    int m2;
+    int n2;
+    int k;
+    int ib;
+    cuDoubleComplex *dA1;
+    int lda1;
+    cuDoubleComplex *dA2;
+    int lda2;
+    cuDoubleComplex *dL1;
+    int ldl1;
+    cuDoubleComplex *dL2;
+    int ldl2;
+    int *IPIV;
+    int info;
+
+    dA1  = (cuDoubleComplex *)STARPU_MATRIX_GET_PTR(descr[0]);
+    dA2  = (cuDoubleComplex *)STARPU_MATRIX_GET_PTR(descr[1]);
+    dL1  = (cuDoubleComplex *)STARPU_MATRIX_GET_PTR(descr[2]);
+    dL2  = (cuDoubleComplex *)STARPU_MATRIX_GET_PTR(descr[3]);
+    dL1 += ib; /* The kernel is just using the inverted part */
+
+    starpu_codelet_unpack_args(cl_arg, &m1, &n1, &m2, &n2, &k, &ib, &lda1, &lda2, &ldl1, &ldl2, &IPIV);
+
+    magma_zssssm_gpu(
+        'C', m1, n1, m2, n2, k, ib,
+        dA1, lda1, dA2, lda2,
+        dL1, ldl1, dL2, ldl2,
+        IPIV, &info);
+
+    cudaThreadSynchronize();
+}
+#endif
+
 /*
  * Codelet definition
  */
Index: magma_morse/runtime/starpu/codelets/codelet_ztstrf.c
===================================================================
--- magma_morse/runtime/starpu/codelets/codelet_ztstrf.c	(révision 1174)
+++ magma_morse/runtime/starpu/codelets/codelet_ztstrf.c	(copie de travail)
@@ -85,8 +85,6 @@
  *              singular, and division by zero will occur if it is used
  *              to solve a system of equations.
  *
- *************************************************************************************************************************************************************
- *
  **/
 void MORSE_TASK_ztstrf(MORSE_option_t *options,
                        int m, int n, int ib, int nb,
@@ -153,7 +151,68 @@
     WORK = (MORSE_Complex64_t*)RUNTIME_starpu_ws_getlocal(h_work);
     CORE_ztstrf(m, n, ib, nb, U, ldu, A, lda, L, ldl, IPIV, WORK, ldwork, &info);
 }
+
+#if defined(MORSE_USE_CUDA)
+static void cl_ztstrf_cuda_func(void *descr[], void *cl_arg)
+{
+    int m;
+    int n;
+    int ib;
+    int nb;
+    cuDoubleComplex *hU, *dU;
+    int ldu;
+    cuDoubleComplex *hA, *dA;
+    int lda;
+    cuDoubleComplex *hL, *dL;
+    int ldl;
+    int *ipiv;
+    cuDoubleComplex *work, *hw, *dw;
+    int ldwork;
+    PLASMA_bool check_info;
+    MORSE_starpu_ws_t *h_work;
+    MORSE_starpu_ws_t *d_work;
+    int iinfo;
+    int info;
+
+    dU = (cuDoubleComplex *)STARPU_MATRIX_GET_PTR(descr[0]);
+    dA = (cuDoubleComplex *)STARPU_MATRIX_GET_PTR(descr[1]);
+    dL = (cuDoubleComplex *)STARPU_MATRIX_GET_PTR(descr[2]);
+
+    starpu_codelet_unpack_args(cl_arg, &m, &n, &ib, &nb, &ldu, &lda, &ldl, &ipiv,
+                               &h_work, &d_work, &ldwork, &check_info, &iinfo);
+
+    /*
+     *  hwork => 2*nb*(2*ib+2nb)
+     *  dwork => 2*nb*ib
+     */
+    hw = (cuDoubleComplex*)morse_starpu_ws_getlocal(h_work);
+    dw = (cuDoubleComplex*)morse_starpu_ws_getlocal(d_work);
+
+    hU = hw;
+    hA = hU + ldu * nb;
+    hL = hA + lda * nb;
+    work = hL + ldl * nb;
+
+    /* Download first panel from A and U */
+    cublasGetMatrix( m, n,  sizeof(cuDoubleComplex), dU, ldu, hU, ldu );
+    cublasGetMatrix( m, ib, sizeof(cuDoubleComplex), dA, lda, hA, lda );
+
+    /* Initialize L to 0 */
+    memset(hL, 0, ldl*nb*sizeof(cuDoubleComplex));
+
+    magma_ztstrf_gpu( 'C', m, n, ib, nb,
+                      hU, ldu, dU, ldu,
+                      hA, lda, dA, lda,
+                      hL, ldl, dL, ldl,
+                      ipiv,
+                      work, ldwork, dw, lda,
+                      &info );
+
+    cudaThreadSynchronize();
+}
+#endif
+
 /*
  * Codelet definition
  */
-CODELETS_CPU(ztstrf, 3, cl_ztstrf_cpu_func)
+CODELETS(ztstrf, 3, cl_ztstrf_cpu_func, cl_ztstrf_cuda_func)
Index: magma_morse/runtime/quark/codelets/codelet_zgetrf_incpiv.c
===================================================================
--- magma_morse/runtime/quark/codelets/codelet_zgetrf_incpiv.c	(révision 1174)
+++ magma_morse/runtime/quark/codelets/codelet_zgetrf_incpiv.c	(copie de travail)
@@ -80,21 +80,22 @@
 void MORSE_TASK_zgetrf_incpiv(MORSE_option_t *options,
                               int m, int n, int ib, int nb,
                               MORSE_desc_t *A, int Am, int An, int lda,
+                              MORSE_desc_t *L, int Lm, int Ln, int ldl,
                               int *IPIV,
                               MORSE_bool check_info, int iinfo)
 {
     DAG_CORE_GETRF;
     QUARK_Insert_Task(options->quark, CORE_zgetrf_incpiv_quark, options->task_flags,
-        sizeof(int),                        &m,             VALUE,
-        sizeof(int),                        &n,             VALUE,
-        sizeof(int),                        &ib,            VALUE,
-        sizeof(MORSE_Complex64_t)*nb*nb,    RTBLKADDR(A, MORSE_Complex64_t, Am, An),                     INOUT,
-        sizeof(int),                        &lda,           VALUE,
-        sizeof(int)*nb,                      IPIV,                  OUTPUT,
-        sizeof(MORSE_sequence_t*),           &(options->sequence),      VALUE,
-        sizeof(MORSE_request_t*),            &(options->request),       VALUE,
-        sizeof(MORSE_bool),                &check_info,    VALUE,
-        sizeof(int),                        &iinfo,         VALUE,
+        sizeof(int),                     &m,             VALUE,
+        sizeof(int),                     &n,             VALUE,
+        sizeof(int),                     &ib,            VALUE,
+        sizeof(MORSE_Complex64_t)*nb*nb,  RTBLKADDR(A, MORSE_Complex64_t, Am, An),      INOUT,
+        sizeof(int),                     &lda,           VALUE,
+        sizeof(int)*nb,                   IPIV,                                         OUTPUT,
+        sizeof(MORSE_sequence_t*),       &(options->sequence),      VALUE,
+        sizeof(MORSE_request_t*),        &(options->request),       VALUE,
+        sizeof(MORSE_bool),              &check_info,    VALUE,
+        sizeof(int),                     &iinfo,         VALUE,
         0);
 }
 
Index: magma_morse/runtime/quark/codelets/codelet_zgessm.c
===================================================================
--- magma_morse/runtime/quark/codelets/codelet_zgessm.c	(révision 1174)
+++ magma_morse/runtime/quark/codelets/codelet_zgessm.c	(copie de travail)
@@ -72,6 +72,7 @@
 void MORSE_TASK_zgessm(MORSE_option_t *options,
                        int m, int n, int k, int ib, int nb,
                        int *IPIV,
+                       MORSE_desc_t *D, int Dm, int Dn, int ldd,
                        MORSE_desc_t *L, int Lm, int Ln, int ldl,
                        MORSE_desc_t *A, int Am, int An, int lda)
 {
@@ -82,8 +83,8 @@
         sizeof(int),                        &k,     VALUE,
         sizeof(int),                        &ib,    VALUE,
         sizeof(int)*nb,                      IPIV,          INPUT,
-        sizeof(MORSE_Complex64_t)*nb*nb,    RTBLKADDR(L, MORSE_Complex64_t, Lm, Ln),             INPUT | QUARK_REGION_L,
-        sizeof(int),                        &ldl,   VALUE,
+        sizeof(MORSE_Complex64_t)*nb*nb,    RTBLKADDR(D, MORSE_Complex64_t, Dm, Dn),             INPUT | QUARK_REGION_L,
+        sizeof(int),                        &ldd,   VALUE,
         sizeof(MORSE_Complex64_t)*nb*nb,    RTBLKADDR(A, MORSE_Complex64_t, Am, An),             INOUT,
         sizeof(int),                        &lda,   VALUE,
         0);
