找回密码
 立即注册

QQ登录

只需一步,快速开始

查看: 215|回复: 5

Global memory读取速度太慢

[复制链接]
发表于 2017-7-5 16:04:51 | 显示全部楼层 |阅读模式
ESC4000G3
最近在写cuda的时候发现kernel在读取Global memory的时候速度很慢,占用了大部分的时间,因为程序要去速度要尽量快,所以我想尽量减少这部分内存读取的时间,请问有什么好的办法呢?我用的平台是TX2
回复

使用道具 举报

发表于 2017-7-11 14:44:06 | 显示全部楼层
Jetson TX2
仅从这几句话看不出速度慢的任何原因,不知道内核怎么写的无法回答。比如是因为内存读写已经达到瓶颈了还是因为内存访问优化的不够。
回复 支持 反对

使用道具 举报

 楼主| 发表于 2017-7-12 11:17:07 | 显示全部楼层
Tesla P100
sisiy 发表于 2017-7-11 14:44
仅从这几句话看不出速度慢的任何原因,不知道内核怎么写的无法回答。比如是因为内存读写已经达到瓶颈了还是 ...

核函数主要就是一个查表,我测了一下,对global memmory读写这块挺慢的程序如下:
__device__ uint8_t bilinear_inplt(uint8_t *img, int width, float col, float row)
{
    int int_col = col;
    int int_row = row;
    float data1 = img[int_row * width + int_col];
    float data2 = img[int_row * width + int_col + 1];
    float data3 = img[(int_row + 1) * width + int_col];
    float data4 = img[(int_row + 1) * width + int_col + 1];

    float weight1 = (int_row + 1 - row)*(int_col + 1 - col);
    float weight2 = (int_row + 1 - row)*(col - int_col);
    float weight3 = (row - int_row)*(int_col + 1 - col);
    float weight4 = (row - int_row)*(col - int_col);

    return (uint8_t)(data1*weight1 + data2*weight2 + data3*weight3 + data4*weight4);
}

__global__ void kernel_make_pano(device_t device, uint8_t *in_img)
{
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int j = blockIdx.y * blockDim.y + threadIdx.y;
   
    Seam_mask_data *seam_file = device.m_seam_file;
    uint8_t *out_img_Y = device.out_img_Y;
    uint8_t *out_img_U = device.out_img_U;
    uint8_t *out_img_V = device.out_img_V;

    int64_t pos = j * device.out_width + i;
    int64_t pos_c = device.out_width_c * j / 2 + i / 2;

    uint8_t mask_index = seam_file[pos].index & 0x000F;
    int seam_file_dead_zone = seam_file[pos].index & 0x8000;

    float row = seam_file[pos].row;
    float col = seam_file[pos].col;

    uint64_t offset_Y = device.in_width * device.in_height * 3 / 2;
    uint64_t offset_U = device.in_width * device.in_height;
    uint64_t offset_V = device.in_width * device.in_height * 5 / 4;

    if (i<device.out_width && j<device.out_height){
        if (!seam_file_dead_zone){
            out_img_Y[pos] = bilinear_inplt(in_img + mask_index * offset_Y, device.in_width, col, row);

            if (!((i & 0x01) | (j & 0x01))){  //both even
                out_img_U[pos_c] = bilinear_inplt(in_img + mask_index * offset_Y + offset_U, device.in_width_c, col / 2, row / 2);
                out_img_V[pos_c] = bilinear_inplt(in_img + mask_index * offset_Y + offset_V, device.in_width_c, col / 2, row / 2);
            }            
        }
        else{
            out_img_Y[pos] = 0;
            if (!((i & 0x01) | (j & 0x01))){  //both even
                out_img_U[pos_c] = 128;
                out_img_V[pos_c] = 128;
            }
        }
    }
}
回复 支持 反对

使用道具 举报

发表于 2017-7-12 14:23:11 | 显示全部楼层
jiangbo 发表于 2017-7-12 11:17
核函数主要就是一个查表,我测了一下,对global memmory读写这块挺慢的程序如下:
__device__ uint8_t b ...

访存有优化空间,晚一点再细说
回复 支持 反对

使用道具 举报

 楼主| 发表于 2017-7-12 15:02:23 | 显示全部楼层
sisiy 发表于 2017-7-12 14:23
访存有优化空间,晚一点再细说

好的,感激不尽
回复 支持 反对

使用道具 举报

发表于 2017-7-12 20:02:11 | 显示全部楼层
1 seam_file中的数据结构从AOS改成SOA,也就是index, col, row各一个数组,然后访问时这样:
index[pos], col[pos], row[pos],
2 每个线程计算4各元素(这样读取时可以是uint32_t或是uchar4类型),这样可以提高内存存取的效率;这时候index,in_img,out_img_Y/U/V都要换成uint32_t或是uchar4类型
3 使用纹理缓存:uint8_t* in_imag改成 const uchar4* __restrict__ in_img
回复 支持 反对

使用道具 举报

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

本版积分规则

关闭

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

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