找回密码
 立即注册

QQ登录

只需一步,快速开始

查看: 254|回复: 4

GPU图像插值,计算速度太慢,解决办法

[复制链接]
发表于 2018-4-10 20:19:22 | 显示全部楼层 |阅读模式
ESC4000G3
    最近进行GPU的图像插值的时候,发现速度太慢,不知道怎么优化才好。程序只要是对图像进行1/4像素插值,程序如下:

__constant__ int8_t INTPL_FILTERS[4][8] = {
    { 0, 0,   0, 64,  0,  0,  0,  0 }, /* for full-pixel, no use */
    { -1, 4, -10, 57, 19, -7,  3, -1 },
    { -1, 4, -11, 40, 40, -11, 4, -1 },
    { -1, 3,  -7, 19, 57, -10, 4, -1 }
};


#define FLT_4TAP_HOR(src, i, coef) (\
    __mul24((src)[i - 1], (coef)[0]) + \
    __mul24((src)[i    ], (coef)[1]) + \
    __mul24((src)[i + 1], (coef)[2]) + \
    __mul24((src)[i + 2], (coef)[3]))


#define FLT_4TAP_VER(src, i, i_src, coef) (\
    __mul24((src)[i -     i_src], (coef)[0]) + \
    __mul24((src)[i            ], (coef)[1]) + \
    __mul24((src)[i +     i_src], (coef)[2]) + \
    __mul24((src)[i + 2 * i_src], (coef)[3]))


template <int bit_depth>
__device__ __inline__
void cud_intpl_luma_block_hor(gpel_t *dst, int i_dst, gpel_t *__restrict__ src, int i_src, int width, int height, int8_t const *coeff, int tid)
{
    const int max_sample_value = (1 << bit_depth) - 1;
    int x, y, v;
    gpel_t *dst_tmp;
    gpel_t *src_tmp;

    int y_step = 31 / width + 1;
    int x_step = 31 % width + 1;
    for(y = tid / width; y < height; y += y_step){
        src_tmp = src + i_src * y;
        dst_tmp = dst + i_dst * y;
        for(x = tid % width; x < width; x += x_step){
            v = (FLT_8TAP_HOR(src_tmp, x, coeff) + 32) >> 6;
            dst_tmp[x] = (gpel_t)GVC_CLIP1(v);
        }
    }
}

/* ---------------------------------------------------------------------------
*/
template <int bit_depth>
__device__ __inline__
void cud_intpl_luma_block_ver(gpel_t *dst, int i_dst, gpel_t *__restrict__ src, int i_src, int width, int height, int8_t const *coeff, int tid)
{
    const int max_sample_value = (1 << bit_depth) - 1;
    int x, y, v;
    gpel_t *dst_tmp;
    gpel_t *src_tmp;

    int y_step = 31 / width + 1;
    int x_step = 31 % width + 1;
    for(y = tid / width; y < height; y += y_step){
        src_tmp = src + i_src * y;
        dst_tmp = dst + i_dst * y;
        for(x = tid % width; x < width; x += x_step){
            v = (FLT_8TAP_VER(src_tmp, x, i_src, coeff) + 32) >> 6;
            dst_tmp[x] = (gpel_t)GVC_CLIP1(v);
        }
    }
}

template <int bit_depth>
__device__ __inline__
void cud_intpl_luma_block_ext(gpel_t *dst, int i_dst, gpel_t *__restrict__ src, int i_src, int width, int height, mct_t *tmp_buf, const int8_t *coeff_h, const int8_t *coeff_v, int tid)
{
    const int max_sample_value = (1 << bit_depth) - 1;
#define TMP_STRIDE      64

    const int shift1 = bit_depth - 8;
    const int add1 = (1 << shift1) >> 1;
    const int shift2 = 20 - bit_depth;
    const int add2 = 1 << (shift2 - 1);//1<<(19-bit_depth)

   // ALIGN16(mct_t tmp_buf[(64 + 7) * TMP_STRIDE]);
    mct_t *tmp;
    int x, y, v;
    gpel_t *dst_tmp;
    gpel_t *src_tmp;

    tmp_buf += 3 * TMP_STRIDE;
    int y_step = 31 / width + 1;
    int x_step = 31 % width + 1;
    for(y = (tid / width) - 3; y < height + 4; y += y_step){
        src_tmp = src + y * i_src;
        tmp = tmp_buf + y * TMP_STRIDE;
        for(x = tid % width; x < width; x += x_step){
            v = FLT_8TAP_HOR(src_tmp, x, coeff_h);
            tmp[x] = (mct_t)((v + add1) >> shift1);
        }
    }

    for(y = tid / width; y < height; y += y_step){
        dst_tmp = dst + y * i_dst;
        tmp = tmp_buf + y * TMP_STRIDE;
        for(x = tid % width; x < width; x += x_step){
            v = (FLT_8TAP_VER(tmp, x, TMP_STRIDE, coeff_v) + add2) >> shift2;
            dst_tmp[x] = (gpel_t)GVC_CLIP1(v);
        }
    }

#undef TMP_STRIDE
}


__device__ __inline__
void cud_copy_pp(gpel_t *dst, int dstStride, const gpel_t *__restrict__ src, int srcStride, int width, int height, int tid)
{
    int x, y;

    int y_step = 31 / width + 1;
    int x_step = 31 % width + 1;
    for(y = tid / width; y < height; y += y_step){
        for(x = tid % width; x < width; x += x_step){
            dst[y * dstStride + x] = src[y * srcStride + x];
        }
    }
}


template <int bit_depth>
__device__ __inline__
void cud_mc_luma(gpel_t *dst, int i_dst, int posx, int posy, int width, int height, mct_t *pBuf, gpel_t *p_fref, int i_fref, int tid)
{
    const int dx = posx & 3;
    const int dy = posy & 3;
    gpel_t *fref;

    posx >>= 2;
    posy >>= 2;

    uint32_t *fref_32 = (uint32_t *)p_fref;
    uint32_t *dst_tmp = (uint32_t *)dst;
    fref = p_fref;
    p_fref += posy * i_fref + posx;
    fref_32 += posy * (i_fref >> 1) + (posx >> 1);
    if(dx == 0 && dy == 0){
        cud_copy_pp(dst, i_dst, p_fref, i_fref, width, height, tid);
    }else if(dx == 0){
        cud_intpl_luma_block_ver<bit_depth>(dst, i_dst, p_fref, i_fref, width, height, INTPL_FILTERS[dy], tid);
    }else if(dy == 0){
        cud_intpl_luma_block_hor<bit_depth>(dst, i_dst, p_fref, i_fref, width, height, INTPL_FILTERS[dx], tid);
    }else{
        cud_intpl_luma_block_ext<bit_depth>(dst, i_dst, p_fref, i_fref, width, height, pBuf, INTPL_FILTERS[dx], INTPL_FILTERS[dy], tid);
    }
}


分像素位置示例图,分为H、V、B三种插值点,
(1)A为整像素点位置;
(2)H为水平分像素点位置;
(3)V为垂直分像素点位置;
(4)采用INTPL_FILTERS中的系数进行滤波得到。

A00  | H01  H02  H03  |  A04'
-----+----------------+------
V10  | B11  B12  B13  |  V14'
V20  | B21  B22  B23  |  V24'
V30  | B31  B32  B33  |  V34'
-----+----------------+------
A40' | H41' H42' H43' |  A44'


回复

使用道具 举报

发表于 2018-4-11 13:20:30 | 显示全部楼层
Jetson TX2
Hi, 楼主,

麻烦提供一下完整代码。光有几个被调用的宏和__device__部分,无法全面分析。

如果提供完整代码有困难,请尽量至少提供kernel代码(含有__global__开头的部分)。

请配合。

Regards,
屠戮人神
回复 支持 反对

使用道具 举报

 楼主| 发表于 2018-4-11 21:35:15 | 显示全部楼层
Tesla P100
本帖最后由 jiangbo 于 2018-4-11 21:37 编辑
屠戮人神 发表于 2018-4-11 13:20
Hi, 楼主,

麻烦提供一下完整代码。光有几个被调用的宏和__device__部分,无法全面分析。
以下为完整代码,每个block有64个线程,希望您有好的优化方法
__constant__ int8_t INTPL_FILTERS[4][8] = {
    { 0, 0,   0, 64,  0,  0,  0,  0 }, /* for full-pixel, no use */
    { -1, 4, -10, 57, 19, -7,  3, -1 },
    { -1, 4, -11, 40, 40, -11, 4, -1 },
    { -1, 3,  -7, 19, 57, -10, 4, -1 }
};


#define FLT_4TAP_HOR(src, i, coef) (\
    __mul24((src)[i - 1], (coef)[0]) + \
    __mul24((src)[i    ], (coef)[1]) + \
    __mul24((src)[i + 1], (coef)[2]) + \
    __mul24((src)[i + 2], (coef)[3]))


#define FLT_4TAP_VER(src, i, i_src, coef) (\
    __mul24((src)[i -     i_src], (coef)[0]) + \
    __mul24((src)[i            ], (coef)[1]) + \
    __mul24((src)[i +     i_src], (coef)[2]) + \
    __mul24((src)[i + 2 * i_src], (coef)[3]))


template <int bit_depth>
__device__ __inline__
void cud_intpl_luma_block_hor(gpel_t *dst, int i_dst, gpel_t *__restrict__ src, int i_src, int width, int height, int8_t const *coeff, int tid)
{
    const int max_sample_value = (1 << bit_depth) - 1;
    int x, y, v;
    gpel_t *dst_tmp;
    gpel_t *src_tmp;

    int y_step = 31 / width + 1;
    int x_step = 31 % width + 1;
    for(y = tid / width; y < height; y += y_step){
        src_tmp = src + i_src * y;
        dst_tmp = dst + i_dst * y;
        for(x = tid % width; x < width; x += x_step){
            v = (FLT_8TAP_HOR(src_tmp, x, coeff) + 32) >> 6;
            dst_tmp[x] = (gpel_t)GVC_CLIP1(v);
        }
    }
}

/* ---------------------------------------------------------------------------
*/
template <int bit_depth>
__device__ __inline__
void cud_intpl_luma_block_ver(gpel_t *dst, int i_dst, gpel_t *__restrict__ src, int i_src, int width, int height, int8_t const *coeff, int tid)
{
    const int max_sample_value = (1 << bit_depth) - 1;
    int x, y, v;
    gpel_t *dst_tmp;
    gpel_t *src_tmp;

    int y_step = 31 / width + 1;
    int x_step = 31 % width + 1;
    for(y = tid / width; y < height; y += y_step){
        src_tmp = src + i_src * y;
        dst_tmp = dst + i_dst * y;
        for(x = tid % width; x < width; x += x_step){
            v = (FLT_8TAP_VER(src_tmp, x, i_src, coeff) + 32) >> 6;
            dst_tmp[x] = (gpel_t)GVC_CLIP1(v);
        }
    }
}

template <int bit_depth>
__device__ __inline__
void cud_intpl_luma_block_ext(gpel_t *dst, int i_dst, gpel_t *__restrict__ src, int i_src, int width, int height, mct_t *tmp_buf, const int8_t *coeff_h, const int8_t *coeff_v, int tid)
{
    const int max_sample_value = (1 << bit_depth) - 1;
#define TMP_STRIDE      64

    const int shift1 = bit_depth - 8;
    const int add1 = (1 << shift1) >> 1;
    const int shift2 = 20 - bit_depth;
    const int add2 = 1 << (shift2 - 1);//1<<(19-bit_depth)

   // ALIGN16(mct_t tmp_buf[(64 + 7) * TMP_STRIDE]);
    mct_t *tmp;
    int x, y, v;
    gpel_t *dst_tmp;
    gpel_t *src_tmp;

    tmp_buf += 3 * TMP_STRIDE;
    int y_step = 31 / width + 1;
    int x_step = 31 % width + 1;
    for(y = (tid / width) - 3; y < height + 4; y += y_step){
        src_tmp = src + y * i_src;
        tmp = tmp_buf + y * TMP_STRIDE;
        for(x = tid % width; x < width; x += x_step){
            v = FLT_8TAP_HOR(src_tmp, x, coeff_h);
            tmp[x] = (mct_t)((v + add1) >> shift1);
        }
    }

    for(y = tid / width; y < height; y += y_step){
        dst_tmp = dst + y * i_dst;
        tmp = tmp_buf + y * TMP_STRIDE;
        for(x = tid % width; x < width; x += x_step){
            v = (FLT_8TAP_VER(tmp, x, TMP_STRIDE, coeff_v) + add2) >> shift2;
            dst_tmp[x] = (gpel_t)GVC_CLIP1(v);
        }
    }

#undef TMP_STRIDE
}


__device__ __inline__
void cud_copy_pp(gpel_t *dst, int dstStride, const gpel_t *__restrict__ src, int srcStride, int width, int height, int tid)
{
    int x, y;

    int y_step = 31 / width + 1;
    int x_step = 31 % width + 1;
    for(y = tid / width; y < height; y += y_step){
        for(x = tid % width; x < width; x += x_step){
            dst[y * dstStride + x] = src[y * srcStride + x];
        }
    }
}


template <int bit_depth>
__device__ __inline__
void cud_mc_luma(gpel_t *dst, int i_dst, int posx, int posy, int width, int height, mct_t *pBuf, gpel_t *p_fref, int i_fref, int tid)
{
    const int dx = posx & 3;
    const int dy = posy & 3;
    gpel_t *fref;

    posx >>= 2;
    posy >>= 2;

    uint32_t *fref_32 = (uint32_t *)p_fref;
    uint32_t *dst_tmp = (uint32_t *)dst;
    fref = p_fref;
    p_fref += posy * i_fref + posx;
    fref_32 += posy * (i_fref >> 1) + (posx >> 1);
    if(dx == 0 && dy == 0){
        cud_copy_pp(dst, i_dst, p_fref, i_fref, width, height, tid);
    }else if(dx == 0){
        cud_intpl_luma_block_ver<bit_depth>(dst, i_dst, p_fref, i_fref, width, height, INTPL_FILTERS[dy], tid);
    }else if(dy == 0){
        cud_intpl_luma_block_hor<bit_depth>(dst, i_dst, p_fref, i_fref, width, height, INTPL_FILTERS[dx], tid);
    }else{
        cud_intpl_luma_block_ext<bit_depth>(dst, i_dst, p_fref, i_fref, width, height, pBuf, INTPL_FILTERS[dx], INTPL_FILTERS[dy], tid);
    }
}

template <int bit_depth>
__device__ __inline__
void cud_inter_dec_lcu_y(lcu_rec_info_t *p_rec_info,
                         gpel_t *pred_sh, mct_t *pBuf,
                         gpel_t *rec_y, Ref_Pic ref_pics,
                         int stride_y, gcoeff_t *p_gpu_resi_y,
                         int i_resi, int tid)
{
    const int max_sample_value = (1 << bit_depth) - 1;
    int i, x, y;
    gpel_t *p_dst_y;
    gpel_t *p_fref1;
    gpel_t *p_fref2;

    int thrd_idx = tid & 31;
    /* luma prediction */
    for(i = tid / 32; i < p_rec_info->num_inter_pu; i += 2){
        const inter_pred_unit_t *p_inter_pu = &p_rec_info->pred_luma[255 - i].inter_blk;
        int pix_x, pix_y, width, height;
        int vec1_x, vec1_y, vec2_x, vec2_y;
        mv_t mv_1st, mv_2nd;

        int8_t r0 = p_inter_pu->ref_idxs.r[0];
        int8_t r1 = p_inter_pu->ref_idxs.r[1];
        if(r0 != INVALID_REF){
            p_fref1 = ref_pics.ref_y[p_inter_pu->ref_idxs.r[0]];
        }
        if(r1 != INVALID_REF){
            p_fref2 = ref_pics.ref_y[p_inter_pu->ref_idxs.r[1]];
        }

        pix_x = p_inter_pu->x;
        pix_y = p_inter_pu->y;
        width = p_inter_pu->w;
        height = p_inter_pu->h;
        mv_1st.v = p_inter_pu->real_mvs[0].v;
        mv_2nd.v = p_inter_pu->real_mvs[1].v;

        vec1_x = mv_1st.x;
        vec1_y = mv_1st.y;
        vec2_x = mv_2nd.x;
        vec2_y = mv_2nd.y;

        int i_pred = 65;
        int i_fref = stride_y;
        p_dst_y = rec_y + pix_y * stride_y + pix_x;
        gpel_t *p_pred = pred_sh + pix_y * i_pred + pix_x;
        mct_t *tmp_buf = pBuf + (i & 3) * 71 * 64;
        if(r0 != INVALID_REF){
            cud_mc_luma<bit_depth>(p_dst_y, stride_y, vec1_x, vec1_y, width, height, tmp_buf, p_fref1, i_fref, thrd_idx, 0);
            if(r1 != INVALID_REF){
                cud_mc_luma<bit_depth>(p_pred, i_pred, vec2_x, vec2_y, width, height, tmp_buf, p_fref2, i_fref, thrd_idx, 0);

                cud_block_avg(p_dst_y, stride_y, p_dst_y, stride_y, p_pred, i_pred, width, height, thrd_idx);
            }
        }else{
            return;
        }
    }
    __syncthreads();
    /* inter reconstruction: luma */
    for(i = tid / 32; i < p_rec_info->num_inter_tb_luma; i += 2){
        const transform_block_t *p_inter_tb = &p_rec_info->trans_luma[255 - i];
        if(p_inter_tb->is_nonzero_block[0]){
            int bsx = p_inter_tb->w;
            int bsy = p_inter_tb->h;
            int img_x = p_inter_tb->x;
            int img_y = p_inter_tb->y;

            int i_coeff = i_resi;
            gcoeff_t *p_resi_y = p_gpu_resi_y + img_y * i_resi + img_x;
            p_dst_y = rec_y + img_y * stride_y + img_x;

            int y_step = 31 / bsx + 1;
            int x_step = 31 % bsx + 1;
            for(y = thrd_idx / bsx; y < bsy; y += y_step){
                for(x = thrd_idx % bsx; x < bsx; x += x_step){
                    p_dst_y[y * stride_y + x] = GVC_CLIP1(p_dst_y[y * stride_y + x] + p_resi_y[y * i_resi + x]);
                }
            }
        }
    }
}

template <int bit_depth>
__global__
void gvcd_inter_dec_lcu_y(gpuvc_lcu_info_t *p_dev_lcus, Ref_Pic ref_pics, lcu_pred_buffer_t *lcu_pred_buffer, gpel_t *rec_y,
                        int stride_y, gcoeff_t **p_gpu_resi, int *i_gpu_resi, int log2_lcu_size)
{
    gpel_t *pred_sh;
    mct_t *buf_tmp;
    int lcu_x = gridDim.x;
    gpel_t *rec_y_tmp;
    gpel_t *rec_u_tmp;
    gcoeff_t *p_gpu_resi_y;
    gcoeff_t *p_gpu_resi_u;
    int x, y, num_lcu;

    int tid = threadIdx.x;
    x = blockIdx.x;
    y = blockIdx.y;
    /* Luma component */
    num_lcu = y*lcu_x + x;
    lcu_rec_info_t *p_rec_info = &p_dev_lcus[num_lcu].rec_info;
    pred_sh = lcu_pred_buffer[num_lcu].pred_buf;
    buf_tmp = lcu_pred_buffer[num_lcu].tmp_buf;
    rec_y_tmp = rec_y + (y << log2_lcu_size) * stride_y + (x << log2_lcu_size);
    p_gpu_resi_y = p_gpu_resi[0] + (y << log2_lcu_size) * i_gpu_resi[0] + (x << log2_lcu_size);
    cud_inter_dec_lcu_y<bit_depth>(p_rec_info, pred_sh, buf_tmp, rec_y_tmp, ref_pics, stride_y, p_gpu_resi_y, i_gpu_resi[0], tid);
}
回复 支持 反对

使用道具 举报

发表于 2018-4-12 13:01:43 | 显示全部楼层
jiangbo 发表于 2018-4-11 21:35
以下为完整代码,每个block有64个线程,希望您有好的优化方法
__constant__ int8_t INTPL_FILTERS[4][8] = ...

工程师要看你的代码需要1-2天时间
回复 支持 反对

使用道具 举报

发表于 2018-4-13 07:11:17 | 显示全部楼层
本帖最后由 屠戮人神 于 2018-4-13 07:32 编辑
jiangbo 发表于 2018-4-11 21:35
以下为完整代码,每个block有64个线程,希望您有好的优化方法
__constant__ int8_t INTPL_FILTERS[4][8] = ...

楼主你好,

你这代码我真心看不了.

通篇无任何注释, 大量的无规则无说明的用户定义类型(光有一个名字, 里面有啥全靠猜, 例如gpel_t, mct_t, lcu_rec_info_t, transform_block_t, gpuvc_lcu_info_t, Ref_Pic, lcu_Pred_buffer_t, gcoeff_t....), 大量的多级指针/地址使用/未知结构体成员, 例如pred_sh/pred_buf, p_gpu_resi, p_dev_lcus, rec_y, i_gpu_resi, p_rec_info, p_inter_tb, p_inter_pu, p_fref....),大量的嵌套的自定义函数套用.  绕到头来, 我连哪些访存行为是block进行的, 哪些是warp和线程进行的, 它们之间的跨步有多大, 是否能合并的访问, 完全看不出来.

请原谅我无法提供任何建议. 已经给你看了很多小时了, 现在还在头晕在各种调用关系和缩写命名中.
你这代码也许使用多核的CPU较好(无惧你的多级指针推导, 多次依赖性前后访存, 各种看不出关系的无规则访存).

抱歉无任何建议.
回复 支持 反对

使用道具 举报

您需要登录后才可以回帖 登录 | 立即注册

本版积分规则

关闭

站长推荐上一条 /1 下一条

快速回复 返回顶部 返回列表