前往小程序,Get更优阅读体验!
立即前往
首页
学习
活动
专区
工具
TVP
发布
社区首页 >专栏 >GPU编程(五): 利用好shared memory

GPU编程(五): 利用好shared memory

作者头像
sean_yang
发布2019-03-06 11:37:20
1.2K0
发布2019-03-06 11:37:20
举报
文章被收录于专栏:Sorrower的专栏

前言

之前在第三章对比过CPU和GPU, 差距非常大. 这一次来看看GPU自身的优化, 主要是shared memory的用法.


CPU矩阵转置

矩阵转置不是什么复杂的事情. 用CPU实现是很简单的:

代码语言:javascript
复制
#include <stdio.h>
#include <stdlib.h>
#include <sys/time.h>

#define LOG_
#define N 1024

/* 转置 */
void transposeCPU( float in[], float out[] )
{
    for ( int j = 0; j < N; j++ )
    {
        for ( int i = 0; i < N; i++ )
        {
            out[j * N + i] = in[i * N + j];
        }
    }
}


/* 打印矩阵 */
void logM( float m[] )
{
    for ( int i = 0; i < N; i++ )
    {
        for ( int j = 0; j < N; j++ )
        {
            printf( "%.1f ", m[i * N + j] );
        }
        printf( "\n" );
    }
}


int main()
{
    int size    = N * N * sizeof(float);
    float   *in = (float *) malloc( size );
    float   *out    = (float *) malloc( size );

    /* 矩阵赋值 */
    for ( int i = 0; i < N; ++i )
    {
        for ( int j = 0; j < N; ++j )
        {
            in[i * N + j] = i * N + j;
        }
    }

    struct timeval  start, end;
    double      timeuse;
    int     sum = 0;
    gettimeofday( &start, NULL );

    transposeCPU( in, out );

    gettimeofday( &end, NULL );
    timeuse = end.tv_sec - start.tv_sec + (end.tv_usec - start.tv_usec) / 1000000.0;
    printf( "Use Time: %fs\n", timeuse );

#ifdef LOG
    logM( in );
    printf( "\n" );
    logM( out );
#endif

    free( in );
    free( out );
    return(0);
}

GPU实现

简单移植

如果什么都不考虑, 只是把代码移植到GPU:

代码语言:javascript
复制
#include <stdio.h>
#include <stdlib.h>
#include <sys/time.h>

#define N 1024
#define LOG_

/* 转置 */
__global__ void transposeSerial( float in[], float out[] )
{
    for ( int j = 0; j < N; j++ )
        for ( int i = 0; i < N; i++ )
            out[j * N + i] = in[i * N + j];
}

/* 打印矩阵 */
void logM( float m[] ){...}

int main()
{
    int size = N * N * sizeof(float);

    float *in, *out;

    cudaMallocManaged( &in, size );
    cudaMallocManaged( &out, size );

    for ( int i = 0; i < N; ++i )
        for ( int j = 0; j < N; ++j )
            in[i * N + j] = i * N + j;

    struct timeval  start, end;
    double      timeuse;
    gettimeofday( &start, NULL );

    transposeSerial << < 1, 1 >> > (in, out);

    cudaDeviceSynchronize();

    gettimeofday( &end, NULL );
    timeuse = end.tv_sec - start.tv_sec + (end.tv_usec - start.tv_usec) / 1000000.0;
    printf( "Use Time: %fs\n", timeuse );


#ifdef LOG
    logM( in );
    printf( "\n" );
    logM( out );
#endif

    cudaFree( in );
    cudaFree( out );
}

不用想, 这里肯定是还不如单线程的CPU的, 真的是完完全全的资源浪费. 实测下来, 耗时是CPU的20多倍, 大写的丢人.

耗时

单block

单block最多可以开1024线程, 这里就开1024线程跑下.

代码语言:javascript
复制
/* 转置 */
__global__ void transposeParallelPerRow( float in[], float out[] )
{
    int i = threadIdx.x;
    for ( int j = 0; j < N; j++ )
        out[j * N + i] = in[i * N + j];
}

int main()
{
    ...
    transposeParallelPerRow << < 1, N >> > (in, out);
    ...
}

效率一下就提升了, 耗时大幅下降.

耗时

tile

但是的话, 如果可以利用多个block, 把矩阵切成更多的tile, 效率还会进一步提升.

代码语言:javascript
复制
/* 转置 */
__global__ void transposeParallelPerElement( float in[], float out[] )
{
    int i = blockIdx.x * K + threadIdx.x;
    /* column */
    int j = blockIdx.y * K + threadIdx.y;
    /* row */
    out[j * N + i] = in[i * N + j];
}

int main()
{
    ...
    dim3 blocks( N / K, N / K );
    dim3 threads( K, K );

    ...
    
    transposeParallelPerElement << < blocks, threads >> > (in, out);
    ...
}

这些都是GPU的常规操作, 但其实利用率依旧是有限的.

耗时


利用率计算

利用率是可以粗略计算的, 比方说, 这里的Memory Clock rateMemory Bus Width是900Mhz和128-bit, 所以峰值就是14.4GB/s.

GPU参数

之前的最短耗时是0.001681s. 数据量是1024*1024*4(Byte)*2(读写). 所以是4.65GB/s. 利用率就是32%. 如果40%算及格, 这个利用率还是不及格的.


shared memory

那该如何提升呢? 问题在于读数据的时候是连着读的, 一个warp读32个数据, 可以同步操作, 但是写的时候就是散开来写的, 有一个很大的步长. 这就导致了效率下降. 所以需要借助shared memory, 由他转置数据, 这样, 写入的时候也是连续高效的了.

代码语言:javascript
复制
/* 转置 */
__global__ void transposeParallelPerElementTiled( float in[], float out[] )
{
    int in_corner_i = blockIdx.x * K, in_corner_j = blockIdx.y * K;
    int out_corner_i    = blockIdx.y * K, out_corner_j = blockIdx.x * K;

    int x = threadIdx.x, y = threadIdx.y;

    __shared__ float tile[K][K];

    tile[y][x] = in[(in_corner_i + x) + (in_corner_j + y) * N];
    __syncthreads();
    out[(out_corner_i + x) + (out_corner_j + y) * N] = tile[x][y];
}

int main()
{

    ...
    dim3 blocks( N / K, N / K );
    dim3 threads( K, K );

    struct timeval  start, end;
    double      timeuse;
    gettimeofday( &start, NULL );

    transposeParallelPerElementTiled << < blocks, threads >> > (in, out);
    ...

}

这样利用率就来到了44%, 及格了.

耗时

所以这就是依据架构来设计算法, 回顾一下架构图:

GPU存储架构


最后

但是44%也就是达到了及格线, 也就是说, 还有更深层次的优化工作需要做. 这些内容也就放在后续文章中了, 有意见或者建议评论区见~


本文参与 腾讯云自媒体同步曝光计划,分享自作者个人站点/博客。
原始发表:2019.02.18 ,如有侵权请联系 cloudcommunity@tencent.com 删除

本文分享自 作者个人站点/博客 前往查看

如有侵权,请联系 cloudcommunity@tencent.com 删除。

本文参与 腾讯云自媒体同步曝光计划  ,欢迎热爱写作的你一起参与!

评论
登录后参与评论
0 条评论
热度
最新
推荐阅读
目录
  • 前言
  • CPU矩阵转置
  • GPU实现
    • 简单移植
      • 单block
        • tile
        • 利用率计算
        • shared memory
        • 最后
        领券
        问题归档专栏文章快讯文章归档关键词归档开发者手册归档开发者手册 Section 归档