コード例 #1
0
ファイル: symv.c プロジェクト: AndreasMiller/clBLAS
// global memory based kernel generator
static ssize_t
generator(
   char *buf,
   size_t buflen,
   const struct SubproblemDim *subdims,
   const struct PGranularity *pgran,
   void *extra)
{
    struct KgenContext *ctx;
    CLBLASKernExtra *kextra = (CLBLASKernExtra*)extra;
    KernelExtraFlags kflags = kextra->flags;
    bool upper = ((kflags & KEXTRA_UPPER_TRIANG) != 0) ^
                  ((kflags & KEXTRA_COLUMN_MAJOR) != 0);
    char tmp[2048];
    const char *typeName;
    DataType dtype = kextra->dtype;
    BlasGenSettings gset, tgset, lset, gset1;
    CLBLASKernExtra kextraTmp;
    TileMulOpts mulOpts, tmulOpts;
    KernelVarNames *vnames = &gset.varNames;
    ssize_t ret;
    size_t vecLen = kextra->vecLen;
    const char *outTypeName;
    bool b;
    TilePostFetchPrivate pfPriv;
    struct symvPrivate priv;
    size_t wgSize;
    bool tailM = (kflags & KEXTRA_TAILS_M) != 0;
    bool tailK = (kflags & KEXTRA_TAILS_K) != 0;
    bool tra = (kflags & KEXTRA_COLUMN_MAJOR) != 0;
    bool rowMaj = !isMatrixAccessColMaj(CLBLAS_SYMV, kflags, MATRIX_A);
    bool isComplex = isComplexType(dtype);
    Tile tileb;
    const char *gid = "get_group_id(0)";
    const char *lid = "get_local_id(0)";
    bool isHoriz = subdims[1].bwidth >= subdims[1].y;
    unsigned int bStep = subdims[0].bwidth / subdims[1].bwidth;
    unsigned int cLocal;
    unsigned int nPlans;

    wgSize = (subdims[0].y / subdims[1].y) *
            (subdims[0].bwidth / subdims[1].bwidth);
    assert(pgran->wgSize[0] == wgSize);
    assert(subdims[0].x == 1);
    assert(subdims[1].x == 1);

    memset(&gset, 0, sizeof(gset));
    memset(&mulOpts, 0, sizeof(mulOpts));
    memset(&pfPriv, 0, sizeof(pfPriv));
    memset(&priv, 0, sizeof(priv));
    ctx = createKgenContext(buf, buflen, true);
    if (ctx == NULL) {
        return -ENOMEM;
    }

    // at first, generate needed declarations
    b = isDoubleBasedType(dtype);
    kgenDeclareUptrs(ctx, b);

    typeName = dtypeBuiltinType(dtype);

    declareSymvKernel(ctx, dtype, pgran, kflags);

    ret = kgenBeginFuncBody(ctx);
    /* 1D work space. Matrix is divided among wi, each calculates it's own
     * part of vector y */

    kgenAddStmt(ctx, "#define M actualN\n");
    memcpy(gset.subdims, subdims, sizeof(gset.subdims));
    gset.subdims[0].itemX = gset.subdims[0].x = 1;
    gset.subdims[1].itemX = gset.subdims[1].x = 1;
    gset.subdims[0].bwidth = gset.subdims[1].bwidth;
    gset.flags |= BGF_WHOLE_A | BGF_UPTRS;

    gset.kextra = kextra;
    gset.pgran = pgran;

    initDefaultTiles(&gset, CLBLAS_SYMV, 0, PRIV_STORAGE_VARIABLE_SET);
    gset.tileA.vecLen = umin(8u, tra ? gset.tileA.nrCols : gset.tileA.nrRows);

    if (isComplex) {
         gset.tileCY.vecLen = 1;
    }
    declareTileStorages(ctx, &gset);
    genZeroTile(ctx, &gset.tileCY);
    getVectorTypeName(dtype, gset.tileCY.vecLen, &outTypeName, NULL);
    cLocal = wgSize / bStep;
    nPlans = gset.tileCY.nrRows / gset.tileCY.vecLen;

    sprintf(tmp, "__local %s localRes[%u][%u];\n",
                outTypeName, pgran->wgSize[0], nPlans);
    kgenAddStmt(ctx, tmp);
    sprintf(tmp, "uint coordA = (%s * %u + %s / %u) * %lu + startN;\n",
                 gid, cLocal, lid, bStep, subdims[1].y);
    kgenAddStmt(ctx, tmp);
    sprintf(tmp, "uint n = coordA;\n");
    kgenAddStmt(ctx, tmp);
    sprintf(tmp, "uint k0 = (%s %% %u) * %lu;\n",
                 lid,  bStep, subdims[1].bwidth);
    kgenAddStmt(ctx, tmp);
    kgenAddStmt(ctx, "actualN += startN;\n");

    kgenAddBlankLine(ctx);

    kgenBeginBranch(ctx,"if (coordA < actualN && k0 < N)");

    genIncPointers(ctx, kflags);
    sprintf(tmp,
            "const GPtr Ag = {(__global %s*)A};\n"
            "const GPtr Xg = {(__global %s*)X};\n",
            typeName, typeName);
    kgenAddStmt(ctx, tmp);

    kgenAddBlankLine(ctx);

    kgenAddStmt(ctx, "uint k = k0;\n");

    if (tailK) {
        sprintf(tmp, "uint Ntail = N %% %lu;\n", subdims[1].bwidth);
        kgenAddStmt(ctx, tmp);
        sprintf(tmp, "uint Ktail = N %% %lu;\n\n", subdims[1].y);
        kgenAddStmt(ctx, tmp);
        kgenBeginBranch(ctx, "if (n + Ktail < N)");
        kgenAddStmt(ctx, "N -= Ntail;\n");
        kgenAddBlankLine(ctx);
    }

    mulOpts.flags |= TILEMUL_OPTIMIZE_COORD_CALC;
    if (tailM) {
        vnames->sizeM = "N";
    }

    vnames->A = "Ag";
    vnames->B = "Xg";
    vnames->coordA = "coordA";
    vnames->coordB = ""; //should not be used for vector
    vnames->k = "k";
    vnames->lda = "lda";
    vnames->sizeK = "N";
    vnames->sizeM = "N";

    mulOpts.flags |= TILEMUL_NOT_FETCH_B | TILEMUL_TRB | TILEMUL_NOT_INC_K;
    if ((kflags & KEXTRA_CONJUGATE_A) != 0) {
        mulOpts.flags |= TILEMUL_CONJA;
    }
    if ((kflags & KEXTRA_ENABLE_MAD) != 0) {
        mulOpts.core = TILEMUL_MAD;
    }
    else {
        mulOpts.core = TILEMUL_MULADD;
    }
    mulOpts.memA = CLMEM_GLOBAL_MEMORY;
    mulOpts.memB = CLMEM_GLOBAL_MEMORY;

    if (rowMaj) {
        mulOpts.flags |= TILEMUL_BW_STRIDE;
    }

    if (upper) {
        kgenAddStmt(ctx, "// k loop over column from the beginning of the column till the diagonal\n");
    }
    else {
        kgenAddStmt(ctx, "// k loop over row from the beginning of the row till the diagonal\n");
    }
    sprintf(tmp, "for (; k < n/%lu*%lu; k += %lu)",
        subdims[1].bwidth, subdims[1].bwidth, bStep*subdims[1].bwidth);
    kgenBeginBranch(ctx, tmp);

    genFetchX(ctx, &gset.tileBX, gset.kextra->vecLen, dtype, vnames,
            mulOpts.flags, kflags);

    upper ^= rowMaj;
    tra ^= rowMaj;
    if (upper ^ rowMaj && tra) {
        mulOpts.flags |= TILEMUL_TRA;
    }
    gset.tileA.trans ^= !upper;
    tgset = gset;
    tmulOpts = mulOpts;

    ret = tileMulGen(ctx, &gset, &mulOpts);
    if (ret != 0) {
        return ret;
    }
    kgenEndBranch(ctx, NULL); /* k loop */

    if (tailK)
    {
            kextraTmp = *kextra;
            gset1 = gset;

            kextraTmp.vecLen = 1;
            gset1.kextra = &kextraTmp;

            gset1.subdims[0].bwidth = gset1.subdims[1].bwidth = 1;

            gset1.tileBX.nrRows = 1;
            gset1.tileA.nrCols = 1;
            kextraTmp.vecLenA = 1;
    }


    if (isHoriz)
    {
        lset = gset;
        lset.subdims[0].bwidth = lset.subdims[1].bwidth =
            lset.subdims[1].y = umin(subdims[1].bwidth, subdims[1].y);
        lset.tileA.nrCols = lset.tileA.nrRows =
            lset.tileBX.nrRows = lset.subdims[1].y;

        kgenAddStmt(ctx, "// the diagonal\n");
        kgenBeginBranch(ctx, "if (k <= n)");
        kgenAddStmt(ctx, "uint k1 = k;\n");

        if (subdims[1].bwidth != subdims[1].y) {
            kgenAddStmt(ctx, "// the pred diagonal\n");
            sprintf(tmp, "for (; k < n; k += %lu)", lset.subdims[1].bwidth);
            kgenBeginBranch(ctx, tmp);

            genFetchX(ctx, &lset.tileBX, lset.subdims[1].bwidth, dtype, vnames,
                    mulOpts.flags, kflags);

            ret = tileMulGen(ctx, &lset, &mulOpts);
            if (ret != 0) {
                return ret;
            }
            kgenEndBranch(ctx, NULL); /* k loop */
        }

        initTile(&tileb, "b", lset.subdims[1].bwidth, lset.subdims[1].bwidth,
            lset.subdims[1].bwidth, lset.tileA.dtype, PRIV_STORAGE_VARIABLE_SET,
            lset.tileA.trans, lset.tileA.packed);
        declareOneTileStorage(ctx, &tileb);

        genFetchX(ctx, &lset.tileBX, lset.subdims[1].bwidth, dtype, vnames,
                mulOpts.flags, kflags);

        priv.mulOpts = &mulOpts;
        priv.pfPriv = &pfPriv;
        priv.tilea = lset.tileA;
        priv.diag = false;

        pfPriv.funcID = CLBLAS_SYMV;
        pfPriv.gset = &lset;
        lset.tileA = tileb;
        mulOpts.postFetch = genPostFetchMirror;
        mulOpts.postFetchPriv = &priv;

        ret = tileMulGen(ctx, &lset, &mulOpts);
        if (ret != 0) {
            return ret;
        }

        if (upper ^ rowMaj && tra) {
            mulOpts.flags &= ~TILEMUL_TRA;
        }
        else {
            mulOpts.flags |= TILEMUL_TRA;
        }
        gset.tileA.trans = lset.tileA.trans ^= true;
        mulOpts.postFetch = NULL;
        mulOpts.postFetchPriv = NULL;

        if (subdims[1].bwidth != subdims[1].y) {
            size_t width = umax(subdims[1].bwidth, subdims[1].y);
            kgenAddStmt(ctx, "// the post diagonal\n");
            if (tailK) {
                kgenBeginBranch(ctx, "if(k < N)");
            }
            sprintf(tmp, "for (k += %lu; k < n/%lu*%lu+%lu; k += %lu)",
                    lset.subdims[1].bwidth,
                    width, width, width,
                    lset.subdims[1].bwidth);
            kgenBeginBranch(ctx, tmp);

            genFetchX(ctx, &lset.tileBX, lset.subdims[1].bwidth, dtype, vnames,
                    mulOpts.flags, kflags);

            ret = tileMulGen(ctx, &lset, &mulOpts);
            if (ret != 0) {
                return ret;
            }
            kgenEndBranch(ctx, NULL); /* k loop */

            if (tailK) {
                kgenEndBranch(ctx, NULL);
                kgenBeginBranch(ctx, "else");
                /* Handle tail along vector X */

                kgenAddStmt(ctx, "N += Ntail;\n");

                mulOpts.flags |= TILEMUL_GLOBAL_CYCLIC_A;
#if 1
                sprintf(tmp, "for (k += %lu; k < actualN; k++)",
                    lset.subdims[1].bwidth);
                kgenBeginBranch(ctx, tmp);

                gset1.tileA.trans = gset.tileA.trans;

                genFetchX(ctx, &gset1.tileBX, gset1.kextra->vecLen, dtype, vnames,
                          mulOpts.flags, kflags);
                ret = tileMulGen(ctx, &gset1, &mulOpts);
                if (ret != 0) {
                    return ret;
                }
                kgenEndBranch(ctx, NULL); /* k loop for tails along vector X */
#else
                mulOpts.flags |= TILEMUL_SKEW_B | TILEMUL_NOT_INC_K;
                genFetchX(ctx, &gset.tileBX, gset.kextra->vecLen, dtype, vnames,
                          mulOpts.flags, kflags);
                ret = tileMulGen(ctx, &gset, &mulOpts);
                if (ret != 0) {
                    return ret;
                }
#endif

                mulOpts.flags &= ~TILEMUL_GLOBAL_CYCLIC_A;
                kgenEndBranch(ctx, NULL);
            }
        }

        sprintf(tmp, "k = k1 + %lu;\n", bStep*subdims[1].bwidth);
        kgenAddStmt(ctx, tmp);
        kgenEndBranch(ctx, NULL);
    }
    else
    {

        kgenAddStmt(ctx, "// the diagonal\n");
        sprintf(tmp, "if (k <= (n  + (get_local_id(0)%%%lu)*%lu))",
            subdims[1].y/subdims[1].bwidth, subdims[1].bwidth);
        kgenBeginBranch(ctx, tmp);

        genFetchX(ctx, &gset.tileBX, gset.subdims[1].bwidth, dtype, vnames,
                    mulOpts.flags, kflags);

        kgenBeginBranch(ctx, NULL);

        priv.mulOpts = &mulOpts;
        priv.pfPriv = &pfPriv;
        priv.diag = true;

        pfPriv.funcID = CLBLAS_SYMV;
        pfPriv.gset = &gset;
        mulOpts.postFetch = genPostFetchVertDiag;
        mulOpts.postFetchPriv = &priv;

        ret = tileMulGen(ctx, &gset, &mulOpts);
        if (ret != 0) {
            return ret;
        }
        kgenEndBranch(ctx, NULL);

        if (upper ^ rowMaj && tra) {
            mulOpts.flags &= ~TILEMUL_TRA;
        }
        else {
            mulOpts.flags |= TILEMUL_TRA;
        }
        gset.tileA.trans ^= true;
        lset = gset;

        sprintf(tmp, "n += (get_local_id(0)%%%lu)*%lu;\n",
            subdims[1].y/subdims[1].bwidth, subdims[1].bwidth);
        kgenAddStmt(ctx, tmp);
        kgenBeginBranch(ctx, NULL);

        priv.diag = false;
        ret = tileMulGen(ctx, &gset, &mulOpts);
        if (ret != 0) {
            return ret;
        }
        kgenEndBranch(ctx, NULL);

        mulOpts.postFetch = NULL;
        mulOpts.postFetchPriv = NULL;

        sprintf(tmp, "k += %lu;\n", bStep*subdims[1].bwidth);
        kgenAddStmt(ctx, tmp);
        kgenEndBranch(ctx, NULL); /* if */
    }

    if (upper) {
        kgenAddStmt(ctx, "// k loop over row from the diagonal till the right\n");
    }
    else {
        kgenAddStmt(ctx, "// k loop over column from the diagonal till the bottom\n");
    }
    sprintf(tmp, "for (; k < N; k += %lu)", bStep*subdims[1].bwidth);
    kgenBeginBranch(ctx, tmp);

    genFetchX(ctx, &gset.tileBX, gset.kextra->vecLen, dtype, vnames,
            mulOpts.flags, kflags);

    ret = tileMulGen(ctx, &gset, &mulOpts);
    if (ret != 0) {
        return ret;
    }
    kgenEndBranch(ctx, NULL); /* k loop */

    if (tailK) {
        /* Handle tail along vector X */
        kgenAddStmt(ctx, "N += Ntail;\n");

        mulOpts.flags |= TILEMUL_GLOBAL_CYCLIC_A;
#if 1
        sprintf(tmp, "for (; k < N; k++)");
        kgenBeginBranch(ctx, tmp);

        gset1.tileA.trans = gset.tileA.trans;

        genFetchX(ctx, &gset1.tileBX, gset1.kextra->vecLen, dtype, vnames,
                  mulOpts.flags, kflags);
        ret = tileMulGen(ctx, &gset1, &mulOpts);
        if (ret != 0) {
            return ret;
        }
        kgenEndBranch(ctx, NULL); /* k loop for tails along vector X */
#else
        mulOpts.flags |= TILEMUL_SKEW_B | TILEMUL_NOT_INC_K;
        genFetchX(ctx, &gset.tileBX, gset.kextra->vecLen, dtype, vnames,
                  mulOpts.flags, kflags);
        ret = tileMulGen(ctx, &gset, &mulOpts);
        if (ret != 0) {
            return ret;
        }
#endif

        kgenEndBranch(ctx, NULL);

        kgenBeginBranch(ctx, "else");

        sprintf(tmp, "for (; k < N; k += %lu)", bStep*subdims[1].bwidth);
        kgenBeginBranch(ctx, tmp);

        tmulOpts.flags |= TILEMUL_SKEW_B | TILEMUL_GLOBAL_CYCLIC_A;
        genFetchX(ctx, &tgset.tileBX, tgset.kextra->vecLen, dtype, vnames,
                tmulOpts.flags, kflags);

        priv.mulOpts = &tmulOpts;
        priv.pfPriv = &pfPriv;
        pfPriv.gset = &tgset;
        priv.diag = false;

        pfPriv.funcID = CLBLAS_SYMV;
        tmulOpts.postFetch = genPostFetchDiag;
        tmulOpts.postFetchPriv = &priv;

        ret = tileMulGen(ctx, &tgset, &tmulOpts);
        if (ret != 0) {
            return ret;
        }

        if (isHoriz) {
            sprintf(tmp, "if (k + %lu > N) break;\n", subdims[1].bwidth);
        }
        else {
            sprintf(tmp, "if (k + %lu > N + (get_local_id(0)%%%lu)*%lu) break;\n",
                subdims[1].y, subdims[1].y/subdims[1].bwidth, subdims[1].bwidth);
        }
        kgenAddStmt(ctx, tmp);

        kgenEndBranch(ctx, NULL); /* k loop */

        kgenBeginBranch(ctx, "if (k < N)");
        if (isHoriz) {
            kgenAddStmt(ctx, "k = n;\n");
        }
        else {
            sprintf(tmp, "n += (get_local_id(0)%%%lu)*%lu;\n",
                subdims[1].y/subdims[1].bwidth, subdims[1].bwidth);
            kgenAddStmt(ctx, tmp);
        }

        genFetchX(ctx, &lset.tileBX, lset.kextra->vecLen, dtype, vnames,
                tmulOpts.flags, kflags);

        priv.mulOpts = &tmulOpts;
        priv.pfPriv = &pfPriv;
        priv.diag = true;

        pfPriv.funcID = CLBLAS_SYMV;
        pfPriv.gset = &lset;
        tmulOpts.postFetch = genPostFetchDiag;
        tmulOpts.postFetchPriv = &priv;

        if (!isHoriz) {
            if (upper ^ rowMaj && tra) {
                tmulOpts.flags &= ~TILEMUL_TRA;
            }
            else {
                tmulOpts.flags |= TILEMUL_TRA;
            }
            kgenAddStmt(ctx, "Ktail = N - n;\n");
            priv.coord = true;
        }
        else {
            priv.coord = false;
        }
        tmulOpts.flags |= TILEMUL_SKEW_B | TILEMUL_GLOBAL_CYCLIC_A | TILEMUL_GLOBAL_CYCLIC_K;


        ret = tileMulGen(ctx, &lset, &tmulOpts);
        if (ret != 0) {
            return ret;
        }

        kgenEndBranch(ctx, NULL);

        kgenEndBranch(ctx, NULL);
    }


    if (!isMatrixAccessColMaj(CLBLAS_GEMV, kflags, MATRIX_A)) {
        mulOpts.flags &= ~TILEMUL_BW_STRIDE;
    }

    kgenEndBranch(ctx,NULL);

    genStoreLocalResult(ctx, &gset.tileCY, lid);

    kgenAddBarrier(ctx, CLK_LOCAL_MEM_FENCE);
    kgenAddBlankLine(ctx);

    sprintf(tmp, "if ((%s %% %u) == 0 && coordA < actualN && k0 < N)", lid, bStep);
    kgenBeginBranch(ctx, tmp);

    genAddLocalResult(ctx, &gset.tileCY, lid, bStep, 1);

    /* write back the results */
    /* y := alpha*A*x + beta*y */
    sprintf(tmp,"(%s - startN)", vnames->coordA);
    setResultPos(ctx, kflags, tmp);

    updateResultVectorTiled(ctx, kflags, vecLen, &gset.tileCY);

    kgenEndBranch(ctx, NULL);

    kgenEndFuncBody(ctx);
    ret = kgenAddBlankLine(ctx);
    if (!ret) {
        ret = (ssize_t)kgenSourceSize(ctx) + 1;
    }

    destroyKgenContext(ctx);
    return (ret < 0) ? -EOVERFLOW : ret;
}
コード例 #2
0
ファイル: t_tilemul.c プロジェクト: BenjaminCoquelle/clBLAS
void
genTest(
    struct KgenContext *ctx,
    BlasGenSettings *gset,
    TileMulOpts *mulOpts,
    bool separateFetch)
{
    char s[1024];
    Kstring kstr;
    char *tName, tVect[64], *ptrName;
    KernelVarNames *vnames = &gset->varNames;
    DataType dtype = gset->kextra->dtype;
    const SubproblemDim *subdims = gset->subdims;
    unsigned int vecLen = gset->kextra->vecLen;
    size_t m, n, k;
    unsigned int i, j;
    bool tra, trb, localA, localB, vecCoords;
    int ret;
    TileMulFlags flags = mulOpts->flags;
    FetchOpts fetchOpts;

    m = gset->subdims[1].y;
    n = gset->subdims[1].x;
    k = gset->subdims[1].bwidth;

    tra = ((flags & TILEMUL_TRA) != 0);
    trb = ((flags & TILEMUL_TRB) != 0);
    localA = (mulOpts->memA == CLMEM_LOCAL_MEMORY);
    localB = (mulOpts->memB == CLMEM_LOCAL_MEMORY);

    vecCoords = ((flags & TILEMUL_OPTIMIZE_VEC_COORDS) != 0);

    tVect[0] = '\0';

    if (vecCoords && vecLen != 1) {
        sprintf(tVect, "%u", vecLen);
    }

    switch (dtype) {
    case TYPE_FLOAT:
        tName = "float";
        ptrName = "f";
        break;
    case TYPE_DOUBLE:
        tName = "double";
        ptrName = "d";
        break;
    case TYPE_COMPLEX_FLOAT:
        tName = "float2";
        ptrName = "f2v";
        break;
    case TYPE_COMPLEX_DOUBLE:
        tName = "double2";
        ptrName = "d2v";
        break;
    default:
        return;
    }

    if (vecCoords) {
        //Do not use GPtrs in fetching
        vnames->A = "A";
        vnames->B = "B";
    }
    else {
        vnames->A = localA ? "LAptr" : "((GPtr)A)";
        vnames->B = localB ? "LBptr" : "((GPtr)B)";
    }
    if (!localA) {
        vnames->lda = "lda";

    }
    if (!localB) {
        vnames->ldb = "ldb";
    }
    vnames->sizeM = "M";
    vnames->sizeN = "N";
    vnames->sizeK = "K";
    vnames->skewA = "skewA";
    vnames->skewB = "skewB";
    vnames->skewK = "skewK";
    vnames->coordA = "workItemM";
    vnames->coordB = "workItemN";
    vnames->k = "k";

    kgenAddBlankLine(ctx);
    sprintf(s, "__attribute__((reqd_work_group_size(%i, %i, 1)))\n",
            ITEM_WORK_M, ITEM_WORK_N);
    kgenAddStmt(ctx, s);
    kgenAddStmt(ctx, "__kernel void\n");
    sprintf(s, "%s(\n", kernelName);
    kgenAddStmt(ctx, s);
    sprintf(s,"    %s alpha,\n", tName);
    kgenAddStmt(ctx, s);
    sprintf(s,"    __global %s%s *A,\n", tName, tVect);
    kgenAddStmt(ctx, s);
    sprintf(s,"    __global %s%s *B,\n", tName, tVect);
    kgenAddStmt(ctx, s);
    kgenAddStmt(ctx, "    uint M,\n"
                     "    uint N,\n"
                     "    uint K,\n");
    sprintf(s,
            "    __global %s *C,\n"
            "    const uint iter)\n", tName);
    kgenAddStmt(ctx, s);
    kgenBeginFuncBody(ctx);
    sprintf(s, "uint workItemM = %lu * get_global_id(0);\n"
               "uint workItemN = %lu * get_global_id(1);\n",
            m, n);
    kgenAddStmt(ctx, s);
    if ((flags & TILEMUL_SKEW_A) != 0) {
        kgenAddStmt(ctx, "uint skewA = 0u;\n");
    }
    if ((flags & TILEMUL_SKEW_B) != 0) {
        kgenAddStmt(ctx, "uint skewB = 0u;\n");
    }
    if ((flags & TILEMUL_SKEW_K) != 0) {
        kgenAddStmt(ctx, "uint skewK = 0u;\n");
    }

    if (localA) {
        sprintf(s, "__local %s LA[%lu];\n",
                tName, subdims[0].bwidth * subdims[0].y);
        kgenAddStmt(ctx, s);
    }
    else { //global A
        sprintf(s, "uint lda = %s;\n", tra ? "M" : "K");
        kgenAddStmt(ctx, s);
    }
    if (localB) {
        sprintf(s, "__local %s LB[%lu];\n",
                tName, subdims[0].bwidth * subdims[0].x);
        kgenAddStmt(ctx, s);
    }
    else { //global B
        sprintf(s, "uint ldb = %s;\n", trb ? "K" : "N");
        kgenAddStmt(ctx, s);
    }

    initDefaultTiles(gset, CLBLAS_GEMM, TILE_PACKED, PRIV_STORAGE_ARRAY);
    declareTileStorages(ctx, gset);

    if (vecCoords) {
        size_t ha, hb;
        char *str;

        ha = tra ? k : m;
        hb = trb ? n : k;

        if (ha > 1) {
            str = s;
            str += sprintf(str, "uint%lu ca = {0", ha);
            for (i = 1; i < ha; i++) {
                str += sprintf(str, ", %s * %u / %u", vnames->lda, i, vecLen);
            }
            str += sprintf(str, "};\n");
            kgenAddStmt(ctx, s);
        }
        else {
            kgenAddStmt(ctx, "uint ca = 0;\n");
        }
        vnames->vectCoordA = "ca";

        if (hb > 1) {
            str = s;
            str += sprintf(str, "uint%lu cb = {0", hb);
            for (i = 1; i < hb; i++) {
                str += sprintf(str, ", %s * %u / %u", vnames->ldb, i, vecLen);
            }
            str += sprintf(str, "};\n");
            kgenAddStmt(ctx, s);
        }
        else {
            kgenAddStmt(ctx, "uint cb = 0;\n");
        }
        vnames->vectCoordB = "cb";

//        uint4 ca = {0, vecLDA, vecLDA * 2, vecLDA * 3};
//        uint4 cb = {0, vecLDB, vecLDB * 2, vecLDB * 3};
    }

    kgenAddBlankLine(ctx);

    sprintf(s, "for (int it = 0; it < iter; it++)");
    kgenBeginBranch(ctx, s);

    if (!(localA && localB)) {
        kgenAddStmt(ctx, "uint k = 0;\n");
    }

    genZeroTile(ctx, &gset->tileCY);

    if (vecCoords) {
        char *coordsA[2] = {"workItemM", "k"};
        char *coordsB[2] = {"k", "workItemN"};
        sprintf(s, "A += %s * (lda / %u) + %s / %u;\n",
                coordsA[tra], vecLen, coordsA[1 - tra], vecLen);
        kgenAddStmt(ctx, s);
        sprintf(s, "B += %s * (ldb / %u) + %s / %u;\n",
                coordsB[trb], vecLen, coordsB[1 - trb], vecLen);
        kgenAddStmt(ctx, s);
    }

    sprintf(s, "for (int k0 = 0; k0 < K; k0 += %lu)", subdims[0].bwidth);
    kgenBeginBranch(ctx, s);

    /* Copy data to local memory. We know that the size of matrix is the same
     * that the size of one block and use that.
     */
    if (localA) {
        sprintf(s,
                "event_t evA = async_work_group_copy(LA, A, %lu, 0);\n"
                "wait_group_events(1, &evA);\n"
                "barrier(CLK_LOCAL_MEM_FENCE);\n",
                subdims[0].y * subdims[0].bwidth);
        kgenAddStmt(ctx, s);
        kgenAddStmt(ctx, "LPtr LAptr;\n");
        if (tra) {
            sprintf(s,
                    "LAptr.%s = LA + workItemM;\n", ptrName);
        }
        else {
            sprintf(s,
                    "LAptr.%s = LA + workItemM * %lu;\n",
                    ptrName, subdims[0].bwidth);
        }
        kgenAddStmt(ctx, s);
    }
    if (localB) {
        sprintf(s,
                "event_t evB = async_work_group_copy(LB, B, %lu, 0);\n"
                "wait_group_events(1, &evB);\n"
                "barrier(CLK_LOCAL_MEM_FENCE);\n",
                subdims[0].x * subdims[0].bwidth);
        kgenAddStmt(ctx, s);
        kgenAddStmt(ctx, "LPtr LBptr;\n");
        if (trb) {
            sprintf(s, "LBptr.%s = LB + workItemN * %lu;\n",
                    ptrName, subdims[0].bwidth);
        }
        else {
            sprintf(s, "LBptr.%s = LB + workItemN;\n", ptrName);
        }
        kgenAddStmt(ctx, s);
    }

    if (!separateFetch) {
        ret = tileMulGen(ctx, gset, mulOpts);
        checkRet(ret, "Multiplier");
    }
    else {
        Tile *tileA = &gset->tileA;
        Tile *tileB = &gset->tileBX;

        memset(&fetchOpts, 0, sizeof(fetchOpts));
        if (localA) {
            fetchOpts.memA = CLMEM_LOCAL_MEMORY;
        }
        if (localB) {
            fetchOpts.memB = CLMEM_LOCAL_MEMORY;
        }

        genFillTileWithNAN(ctx, tileA);
        genFillTileWithNAN(ctx, tileB);

        if (subdims[0].bwidth != subdims[1].bwidth) {
            sprintf(s, "for (int k1 = 0; k1 < %lu; k1 += %lu)",
                    subdims[0].bwidth, k);
            kgenBeginBranch(ctx, s);
        }

#if JUST_MULTIPLICATION
        for (i = 0; i < tileA->nrRows; i++) {
            for(j = 0; j < tileA->nrCols; j++) {
                sprintfTileElement(&kstr, tileA, i, j, 1);
                sprintf(s, "%s = %u;\n", kstr.buf, i * tileA->nrCols + j);
                kgenAddStmt(ctx, s);
            }
        }

        for (i = 0; i < tileB->nrRows; i++) {
            for(j = 0; j < tileB->nrCols; j++) {
                sprintfTileElement(&kstr, tileB, i, j, 1);
                sprintf(s, "%s = %u;\n", kstr.buf, i * tileB->nrCols + j);
                kgenAddStmt(ctx, s);
            }
        }
#else
        fetchOpts.mrole = MATRIX_B;
        fetchOpts.lineOffset = 0;
        fetchOpts.linesNum = (tileB->trans) ? tileB->nrCols : tileB->nrRows;
        ret = genFetchInputTile(ctx, NULL, gset, &fetchOpts);
        checkRet(ret, "Fetching tile b");

        fetchOpts.mrole = MATRIX_A;
        fetchOpts.linesNum = (tileA->trans) ? tileA->nrCols : tileA->nrRows;
        kgenAddBlankLine(ctx);
        fetchOpts.lineOffset = 0;
        ret = genFetchInputTile(ctx, NULL, gset, &fetchOpts);
        checkRet(ret, "Fetching tile a");
#endif
        ret = genMulTiles(ctx, gset, mulOpts);
        checkRet(ret, "Multiplier");
#if ! JUST_MULTIPLICATION
        sprintf(s, "k += %lu;\n", k);
        kgenAddStmt(ctx, s);
#endif
        if (subdims[0].bwidth != subdims[1].bwidth) {
            kgenEndBranch(ctx, NULL);
        }
    }
    kgenEndBranch(ctx, NULL); // K loop
    kgenEndBranch(ctx, NULL); // iterations loop

    kgenAddBlankLine(ctx);

    for (i = 0; i < m; i++) {
        for (j = 0; j < n; j++) {
            sprintfTileElement(&kstr, &gset->tileCY, i, j, 1);
                sprintf(s,
                        "((GPtr)C).%s"
                    "[(%d + workItemM) * N  + %d + workItemN] = %s;\n",
                    ptrName, i, j, kstr.buf);
                kgenAddStmt(ctx, s);
            }
                }

    kgenEndFuncBody(ctx);
}
コード例 #3
0
ファイル: gemv.c プロジェクト: AndreasMiller/clBLAS
// global memory based kernel generator
static ssize_t
generator(
   char *buf,
   size_t buflen,
   const struct SubproblemDim *subdims,
   const struct PGranularity *pgran,
   void *extra)
{
    struct KgenContext *ctx;
    CLBLASKernExtra *kextra = (CLBLASKernExtra*)extra;
    KernelExtraFlags kflags = kextra->flags;
    size_t staggered = ((extraData_t*)&kextra->solverPriv)->staggered;
    //yes, KEXTRA_TAILS_K because it is set if N % bw != 0
    bool tailN = ((kflags & KEXTRA_TAILS_K) != 0);
    bool tailM = ((kflags & KEXTRA_TAILS_M) != 0);
    char tmp[4096];
    DataType dtype = kextra->dtype;
    bool doubleBased = isDoubleBasedType(dtype);
    BlasGenSettings gset;
    TileMulOpts mulOpts;
    KernelVarNames *vnames = &gset.varNames;
    ssize_t ret;
    TilePostFetchPrivate pfPriv;
    unsigned int vecLen = kextra->vecLen;
    const char *outTypeName;
    const char *gid = "get_group_id(0)";
    const char *lid = "get_local_id(0)";
    const char *typeName;
    size_t wgSize;
    //unsigned int nStep = 32;
    unsigned int bStep = subdims[0].bwidth / subdims[1].bwidth; //8;
    unsigned int cLocal;
    bool isComplex = isComplexType(dtype);
    unsigned int nPlans;

    typeName = dtypeBuiltinType(dtype);
    memset(&gset, 0, sizeof(gset));
    memset(&mulOpts, 0, sizeof(mulOpts));
    ctx = createKgenContext(buf, buflen, true);
    if (ctx == NULL) {
        return -ENOMEM;
    }

    // at first, generate needed declarations
    kgenDeclareUptrs(ctx, doubleBased);

    // now, generate the kernel
    declareGemvKernel(ctx, dtype, pgran, kflags);
    ret = kgenBeginFuncBody(ctx);
    kgenAddStmt(ctx, "// M always denotes length of Y "
                     "and N denotes length of X in the kernel\n");
    /* 1D work space. Matrix is divided among wi, each calculates it's own
     * part of vector y */

    wgSize = (subdims[0].y / subdims[1].y) *
            (subdims[0].bwidth / subdims[1].bwidth);
    assert(pgran->wgSize[0] == wgSize);
    assert(subdims[0].x == 1);
    assert(subdims[1].x == 1);
    cLocal = wgSize/bStep;

    memcpy(gset.subdims, subdims, sizeof(gset.subdims));
    gset.subdims[0].itemX = gset.subdims[0].x = 1;
    gset.subdims[1].itemX = gset.subdims[1].x = 1;
    gset.subdims[0].bwidth = gset.subdims[1].bwidth;

    gset.pgran = pgran;
    gset.kextra = kextra;
    gset.flags = BGF_UPTRS;

    initDefaultTiles(&gset, CLBLAS_GEMV, 0, PRIV_STORAGE_VARIABLE_SET);
    if (isComplex) {
         gset.tileCY.vecLen = 1;
    }
    declareTileStorages(ctx, &gset);
    genZeroTile(ctx, &gset.tileCY);
    getVectorTypeName(dtype, gset.tileCY.vecLen, &outTypeName, NULL);
    nPlans = gset.tileCY.nrRows / gset.tileCY.vecLen;

    sprintf(tmp, "__local %s localRes[%u][%u];\n",
                outTypeName, pgran->wgSize[0], nPlans);
    kgenAddStmt(ctx, tmp);
    sprintf(tmp, "uint coordA = (%s * %u + %s %% %u) * %lu;\n",
                 gid, bStep, lid, bStep, subdims[1].y);
    kgenAddStmt(ctx, tmp);
    sprintf(tmp, "uint k0 = (%s / %u) * %lu;\n",
                 lid,  bStep, subdims[1].bwidth);
    kgenAddStmt(ctx, tmp);

    kgenAddBlankLine(ctx);

    kgenBeginBranch(ctx,"if (coordA < M && k0 < N)");

    genIncPointers(ctx, kflags);
    sprintf(tmp,
            "const GPtr Ag = {(__global %s*)A};\n"
            "const GPtr Xg = {(__global %s*)X};\n",
            typeName, typeName);
    kgenAddStmt(ctx, tmp);

    kgenAddBlankLine(ctx);

    if (tailN) {
        sprintf(tmp, "uint Ntail = N %% %lu;\n", subdims[1].bwidth);
        kgenAddStmt(ctx, tmp);
        kgenAddStmt(ctx, "N -= Ntail;\n");
        kgenAddBlankLine(ctx);
    }

    mulOpts.flags |= TILEMUL_OPTIMIZE_COORD_CALC;
    if (tailM) {
        mulOpts.flags |= TILEMUL_GLOBAL_CYCLIC_A;
    }

    vnames->A = "Ag";
    vnames->B = "Xg";
    vnames->coordA = "coordA";
    vnames->coordB = ""; //should not be used for vector
    vnames->k = "k";
    vnames->lda = "lda";
    vnames->sizeK = "N";
    vnames->sizeM = "M";

    mulOpts.flags |= TILEMUL_NOT_FETCH_B | TILEMUL_TRB | TILEMUL_C_COLUMN_MAJOR | TILEMUL_NOT_INC_K;
    if ((kflags & KEXTRA_CONJUGATE_A) != 0) {
        mulOpts.flags |= TILEMUL_CONJA;
    }
    if (isMatrixAccessColMaj(CLBLAS_GEMV, kflags, MATRIX_A)) {
        mulOpts.flags |= TILEMUL_TRA;
    }
    if ((kflags & KEXTRA_ENABLE_MAD) != 0) {
        mulOpts.core = TILEMUL_MAD;
    }
    else {
        mulOpts.core = TILEMUL_MULADD;
    }
    mulOpts.memA = CLMEM_GLOBAL_MEMORY;
    mulOpts.memB = CLMEM_GLOBAL_MEMORY;

    if (!isMatrixAccessColMaj(CLBLAS_GEMV, kflags, MATRIX_A)) {
        gset.subdims[0].bwidth = pgran->wgSize[0] * subdims[1].bwidth;
        mulOpts.flags |= TILEMUL_BW_STRIDE;
    }

    sprintf(tmp, "uint k = k0;\nfor (; k < N; k += %lu)", cLocal*subdims[1].bwidth);
    kgenBeginBranch(ctx, tmp);

    if (staggered) {
        vnames->k = "k1";
        sprintf(tmp, "const uint k1 = (k + get_group_id(0)*%lu)%%N;\n",staggered);
        kgenAddStmt(ctx, tmp);
    }

    genFetchX(ctx, &gset.tileBX, gset.kextra->vecLen, dtype, vnames,
            mulOpts.flags, kflags);

    ret = tileMulGen(ctx, &gset, &mulOpts);
    if (ret != 0) {
        return ret;
    }
    vnames->k = "k";
    kgenEndBranch(ctx, NULL); /* k loop */

    if (tailN) {
        /* Handle tail along vector X */
        kgenAddStmt(ctx, "N += Ntail;\n");
        kgenBeginBranch(ctx, "if (k < N)");

        mulOpts.flags |= TILEMUL_SKEW_B;
        genFetchX(ctx, &gset.tileBX, gset.kextra->vecLen, dtype, vnames,
                  mulOpts.flags, kflags);
        mulOpts.flags |= TILEMUL_GLOBAL_CYCLIC_K|TILEMUL_WRAP_AROUND_TAIL;
        setFetchHandler(&mulOpts, &gset, defaultTilePostFetch, &pfPriv);
        ret = tileMulGen(ctx, &gset, &mulOpts);
        if (ret != 0) {
            return ret;
        }
        kgenEndBranch(ctx, NULL);
    }

    if (!isMatrixAccessColMaj(CLBLAS_GEMV, kflags, MATRIX_A)) {
        gset.subdims[0].bwidth = subdims[1].bwidth;
        mulOpts.flags &= ~TILEMUL_BW_STRIDE;
    }

    kgenEndBranch(ctx,NULL);

    genStoreLocalResult(ctx, &gset.tileCY, lid);

    kgenAddBarrier(ctx, CLK_LOCAL_MEM_FENCE);
    kgenAddBlankLine(ctx);

    sprintf(tmp, "if (%s < %u && coordA < M && k0 < N)", lid, bStep);
    kgenBeginBranch(ctx, tmp);

    genAddLocalResult(ctx, &gset.tileCY, lid, cLocal, bStep);

    /* write back the results */
    /* y := alpha*A*x + beta*y */
    setResultPos(ctx, kflags, vnames->coordA);

    updateResultVectorTiled(ctx, kflags, vecLen, &gset.tileCY);

    kgenEndBranch(ctx, NULL);

    kgenEndFuncBody(ctx);
    ret = kgenAddBlankLine(ctx);

    if (!ret) {
        ret = (ssize_t)kgenSourceSize(ctx) + 1;
    }

    destroyKgenContext(ctx);
    return (ret < 0) ? -EOVERFLOW : ret;
}
コード例 #4
0
ファイル: trmm.c プロジェクト: clMathLibraries/clBLAS
static ssize_t
generator(
   char *buf,
   size_t buflen,
   const struct SubproblemDim *subdims,
   const struct PGranularity *pgran,
   void *extra)
{
    char tmp[4096];
    struct KgenContext *ctx;
    CLBLASKernExtra *kextra = (CLBLASKernExtra*)extra;
    KernelExtraFlags kflags = kextra->flags;
    DataType dtype = kextra->dtype;
    bool doubleBased = isDoubleBasedType(dtype);
    size_t staggered = ((extraData_t*)&kextra->solverPriv)->staggered;
    int ret;
    BlasGenSettings gset;
    TileMulOpts mulOpts;
    int tra = isMatrixAccessColMaj(CLBLAS_TRMM, kflags, MATRIX_A);
    int trb = isMatrixAccessColMaj(CLBLAS_TRMM, kflags, MATRIX_B);
    unsigned int l1Pans;
    TilePostFetchPrivate pfPriv[2];
    UpdateResultFlags upResFlags;
    TailStatus tailStatus;
    bool subgMode = false;
    SubgVarNames subgVNames;

    ctx = createKgenContext(buf, buflen, true);
    if (ctx == NULL) {
        return -ENOMEM;
    }

    // mismatching subdims define case with subgroup decomposition
    subgMode = ( subdims[0].bwidth != subdims[1].bwidth );

    memset(&gset, 0, sizeof(gset));
    memcpy(gset.subdims, subdims, sizeof(gset.subdims));
    gset.flags = BGF_DISTINCT_VECLEN;

    gset.flags |= BGF_WHOLE_A;

    /*FIXME: This used to be a workaround for compilation issues with dtrmm on
     * cpu. Normally BGF_WHOLE_A should be enabled always. But for now,
     * there are wrong results for non-aligned cases on CPU and there is
     * no workaround yet.
    if (kflags & (KEXTRA_TAILS_M | KEXTRA_TAILS_N | KEXTRA_TAILS_K)) {
        gset.flags &= ~BGF_WHOLE_A;
    }*/
    gset.kextra = kextra;
    gset.pgran = pgran;
    //avoid [0].bw loop
    //gset.subdims[0].bwidth = gset.subdims[1].bwidth;

    memset(pfPriv, 0, sizeof(pfPriv));
    pfPriv[0].funcID = CLBLAS_TRMM;
    pfPriv[0].gset = &gset;
    if ((gset.flags & BGF_WHOLE_A) != 0) {
        pfPriv[0].wholeA = 1;
    }

    // at first, generate needed declarations
    kgenDeclareUptrs(ctx, doubleBased);

    // For inner callback, because both callbacks use own fetchNumA
    memcpy(&pfPriv[1], &pfPriv[0], sizeof(pfPriv[0]));

    // if both matrices are accessed row-major - using subgroup pattern
    if ( subgMode ) {

        declareTrxmKernel(ctx,
            dtype,
            pgran,
            kflags,
            CLBLAS_TRMM,
            "Subgroup",
            true,
            true);
        gset.flags |= BGF_UPTRS;
    }
    else {

        declareTrxmKernel(ctx,
            dtype,
            pgran,
            kflags,
            CLBLAS_TRMM,
            "Block",
            true,
            true);

    }
    kgenBeginFuncBody(ctx);

    initDefaultTiles(&gset, CLBLAS_TRMM, 0, PRIV_STORAGE_VARIABLE_SET);
    declareTileStorages(ctx, &gset);

    kgenAddStmt(ctx,
                "uint currM, currN;\n"
                "uint4 coord = 0; /* contains coordB, coordA, k */\n");

    kgenDeclareLocalID(ctx, "lid", pgran);
    kgenDeclareGroupID(ctx, "gid", pgran);

    if ( subgMode ) {

        gset.varNames.LDS = "scratch";

        // declaring variables used by subgroup mode
        subgVNames.itemId = "itemId";
        subgVNames.subgCoord = "subgCoord";

        kgenAddBlankLine( ctx );
        kgenAddBlankLine(ctx);

        kgenPrintf(ctx, "int2 %s;\n", subgVNames.itemId );
        kgenPrintf(ctx, "int2 %s;\n", subgVNames.subgCoord);

        // item ID
        kgenPrintf( ctx,
            "%s.x = get_local_id(0)%%%d;\n",
            subgVNames.itemId,
            subdims[0].bwidth/subdims[1].bwidth);

        // subgroup ID
        kgenPrintf( ctx,
            "%s.y = get_local_id(0)/%d;\n",
            subgVNames.itemId,
            subdims[0].bwidth/subdims[1].bwidth);

        // subgroup coordX
        kgenPrintf( ctx,
            "%s.x = %s.y/%d;\n",
            subgVNames.subgCoord,
            subgVNames.itemId,
            subdims[0].y/subdims[1].y );

        // subgroup coordY
        kgenPrintf( ctx,
            "%s.y = %s.y%%%d;\n",
            subgVNames.subgCoord,
            subgVNames.itemId,
            subdims[0].y/subdims[1].y );
    }

    kgenAddBlankLine(ctx);

    sprintf(tmp, "currN = gid * %lu;\n", subdims->x);
    kgenAddStmt(ctx, tmp);
    genInitCurrM(ctx, subdims, kflags);

    if (kflags & KEXTRA_A_OFF_NOT_ZERO) {
        kgenAddStmt(ctx, "A += offA;\n");
    }
    genTrxmBMatrShift(ctx, kflags, true);

    if ( subgMode ) {
        kgenAddStmt(ctx,
            "GPtr Ag = {A};\n"
            "GPtr Bg = {B};\n");
    }

    l1Pans = (unsigned int)subdims[0].x / (unsigned int)subdims[1].x;

    memset(&mulOpts, 0, sizeof(mulOpts));
    mulOpts.core = ((kflags & KEXTRA_ENABLE_MAD) != 0)
            ? TILEMUL_MAD
            : TILEMUL_MULADD;
    mulOpts.memA = CLMEM_GLOBAL_MEMORY;
    mulOpts.memB = CLMEM_GLOBAL_MEMORY;
    mulOpts.postFetch = NULL;
    mulOpts.postFetchPriv = &pfPriv;
    mulOpts.flags = TILEMUL_NO_FLAGS;
    mulOpts.flags |= TILEMUL_EXTERN_RDECL;

    if ( subgMode ) {

        mulOpts.flags |= TILEMUL_NOT_INC_K;
        mulOpts.flags |= TILEMUL_BW_STRIDE;
    }

    if (kflags & KEXTRA_TAILS_M_LOWER) {
        mulOpts.flags |= TILEMUL_GLOBAL_CYCLIC_A;
    }
    if (kflags & KEXTRA_TAILS_N_LOWER) {
        mulOpts.flags |= TILEMUL_GLOBAL_CYCLIC_B;
    }
    if (kflags & KEXTRA_TAILS_K_LOWER) {
        mulOpts.flags |= TILEMUL_GLOBAL_CYCLIC_K;
        mulOpts.flags |= TILEMUL_WRAP_AROUND_TAIL;
    }

    if (tra) {
        mulOpts.flags |= TILEMUL_TRA;
    }
    if (!trb) {
        mulOpts.flags |= TILEMUL_TRB;
    }
    if (isMatrixConj(kflags, MATRIX_A)) {
        mulOpts.flags |= TILEMUL_CONJA;
    }
    if (isMatrixConj(kflags, MATRIX_B)) {
        mulOpts.flags |= TILEMUL_CONJB;
    }

    initKernelVarNames(&gset.varNames);

    if ( subgMode ) {

        kgenPrintf( ctx,
            "coord.x = currN + %s.x*%d;\n",
            subgVNames.subgCoord,
            subdims[1].x );
    }
    else {

        sprintf(tmp, "coord.x = currN + lid %% %u * %lu;\n", l1Pans, subdims[1].x);
        kgenAddStmt(ctx, tmp);
    }

    // loop over M
    sprintf(tmp, "for (uint m0 = 0; m0 < M; m0 += %lu)", subdims[0].y);
    kgenBeginBranch(ctx, tmp);

    genStartPosK( ctx, subdims, kflags, subgMode );

    sprintf(tmp, "coord.z = kBegin;\n");
    kgenAddStmt(ctx, tmp);

    if ( subgMode ) {

        kgenPrintf(ctx,
            "coord.y = currM + %s.y*%d;\n",
            subgVNames.subgCoord,
            subdims[1].y);
    }
    else {

        sprintf( tmp,
            "coord.y = currM + lid / %u * %lu;\n",
            l1Pans,
            subdims[1].y );
        kgenAddStmt(ctx, tmp);
    }

    genZeroTile(ctx, &gset.tileCY);

    checkGenBeginHitMatrixBlock(ctx, kflags);
    tailStatus = checkGenAdjustTailCoords(ctx, CLBLAS_TRMM, &gset, NULL);

    // loops along 'K'
    if ( subgMode ) {
        ret = genSubgLoopsK( ctx, &gset, &mulOpts, &subgVNames, staggered);
    }
    else {
        ret = genLoopsK( ctx, &gset, &mulOpts, tmp );
    }

    if (ret != 0) {
        printf("%s", buf);
        return ret;
    }

    checkGenEndHitMatrixBlock(ctx, kflags);
    kgenAddBarrier(ctx, CLK_GLOBAL_MEM_FENCE);

    // store results
    // for result update - x coordinate is in elements, not in vectors

    checkGenRestoreTailCoords(ctx, &gset, tailStatus);
    upResFlags = kextraToUpresFlags(CLBLAS_TRMM, kflags);
    upResFlags |= tailStatusToUpresFlags(tailStatus);
    upResFlags |= UPRES_INDEXING_WITH_CONSTANTS;
    upResFlags |= UPRES_TRIANG_WRITE_C;
    upResFlags |= UPRES_EXCEED_PROBLEM_CONDITION;

    if ( subgMode ) {

        mergeUpdateResult( ctx,
            CLBLAS_TRMM,
            &gset,
            &subgVNames,
            upResFlags,
            genResultUpdateWithFlags );
    }
    else {

        //checkGenBeginHitMatrixBlock(ctx, kflags);
        genResultUpdateWithFlags( ctx,
            CLBLAS_TRMM,
            &gset,
            upResFlags,
            NULL,
            NULL,
            NULL );
        //checkGenEndHitMatrixBlock(ctx, kflags);
    }

    if (isMatrixUpper(kflags)) {
        sprintf(tmp, "currM += %lu;\n", subdims[0].y);
    }
    else {
        sprintf(tmp, "currM -= %lu;\n", subdims[0].y);
    }
    kgenAddStmt(ctx, tmp);

    kgenEndBranch(ctx, NULL);

    kgenEndFuncBody(ctx);
    ret = kgenAddBlankLine(ctx);

    if (!ret) {
        ret = (ssize_t)kgenSourceSize(ctx) + 1;
    }

    destroyKgenContext(ctx);

    return (ret < 0) ? -EOVERFLOW : ret;
}