zoukankan      html  css  js  c++  java
  • cuda GPU 编程之共享内存的使用

      原理上来说,共享内存是GPU上可受用户控制的一级缓存。在一个SM中,存在着若干cuda core + DP(双精度计算单元) + SFU(特殊函数计算单元)+共享内存+常量内存+纹理内存。相对于全局内存,共享内存的方寸延迟较低,可以达到惊人的1.5TB/s。而全局内存大约只有150GB/s。(最新的NVLINK技术没有考虑在内)。因而共享内存的使用时性能提高的一个重要的因素。但是注意到,将数据拷贝到共享内存中也消耗了部分时间。因而,共享内存仅仅适合存在着数据的重复利用,全局的内存合并或者是线程之间有共享数据的时候,否则直接使用全局内存会更好一些。

       下面介绍两种使用共享内存的方法。

     1. 创建固定大小的共享内存。(在kernel函数内存定义)

    __shared__ float a_in[34];
    

      注意这里的34必须在编译之前指定大小。可以使用宏定义的方式进行。下面的方式是一种错误的示范。

    __shared__ float s_in[blockDim.x+2*RAD];
    

      

    2. 动态申请共享内存数组,声明时需要加上 extern 前缀。

    extern __shared__ float a[];
    

      并且,在调用内核函数的时候,需要在<<<>>>内加上第三个参数来指明所需分配的共享内存的字节大小。

    const size_t smemSize=(TPB+ 2*RAD)*sizeof(float);
    ddkernel<<<Grids, Blocks,smemSize>>>(paramenter);
    

      

    分配好共享内存之后,就可以将全局内存拷贝到共享内存之中。基本的方案是每个线程从全局索引位置读取元素,将它存储到共享内存之中。在使用共享内存的时候,还应该注意数据存在着交叉,应该将边界上的数据拷贝进来。

    __global__ 
    void ddkernel(paramenter)
    {
        const int i=threadIdx.x+blockDim.x*BlockIdx.x;
        if(i.size)return;
    
        const int s_idx=threadIdx.x+RAD;
        extern __shared__ float s_in[];
        
        s_in[s_idx]=d_in[i];
        if(threadIdx.x<RAD){
            s_in[s_idx-RAD]=d_in[i-RAD];
            s_in[s_idx+blockDim.x]=d_in[i+blockDim.x]; 
        }            
        __syncthread();
    
    
    
    } 
    

      

  • 相关阅读:
    结对编程作业——毕设导师智能匹配
    结对项目之需求分析与原型设计
    Excel绘制之甘特图
    Excel绘图之数据波动条形图
    Excel绘图之漏斗图
    Excel绘图之四象限散点图
    软件工程实践总结
    发送手机验证码
    个人作业——软件产品案例分析
    用例图
  • 原文地址:https://www.cnblogs.com/cofludy/p/7622254.html
Copyright © 2011-2022 走看看