title: 【CUDA 基础】5.2 共享内存的数据布局
categories:
- CUDA
- Freshman
tags:
- 行主序
- 列主序
toc: true
date: 2018-06-02 21:01:03


Abstract: 本文主要研究几个关于共享内存的例子，以此来了解共享内存的性质，为我们的核函数加速

Keywords: 行主序，列主序，填充与无填充，从线程索引体映射数据元素

开篇废话

同一个东西，A花大工夫做到极致，成本100，售价200；C模仿A的做法快速的通过仿制，节省了研发试验的所有开销，但是没有做到A那么完美，成本25，售价140。A虽然好但是不见得销量有C高，并且A的利润并没有C那么高，所以，作为商人，选择C是没错的，商人的目的就是盈利，但是问题来了，如果不是商人呢？而是一个科学家呢？
本文我们主要研究共享内存的数据布局，通过代码实现，来观察运行数据，换句话说，我们主要研究上一篇中的放西瓜，取西瓜，以及放冬瓜等的一些列操作对性能的影响，以及如何才能使效率最大化。
几个例子包括以下几个主题：

  • 方阵与矩阵数组
  • 行主序与列主序
  • 静态与动态共享内存的声明
  • 文件范围与内核范围的共享内存
  • 内存填充与无内存填充

当使用共享内存设计核函数的时候下面两个概念是非常重要的：

  1. 跨内存存储体映射数据元素
  2. 从线程索引到共享内存偏移的映射

当上面这些主题和概念都得到很好地理解，设计一个高效的使用共享内存的核函数就没什么问题了，其可以避免存储体冲突并充分利用共享内存的优势。
注意，从几何上讲，方形属于矩形，这里我们说的矩形时指长方形。

方形共享内存

我们前面说过我们的线程块可以是一维二维和三维的，对应的线程编号是threadIdx.x, threadIdx.y以及threadIdx.z，为了对应一个二维的共享内存，我们假设我们使用二维的线程块，那么对于一个二维的共享内存

#define N 32
...
__shared__ int x[N][N];
...

当我们使用二维块的时候，很有可能会使用下面这种方式来索引x的数据：

#define N 32
...
__shared__ int x[N][N];
...
int a=x[threadIdx.y][threadIdx.x];

当然这个索引就是 gpu共享内存和专用内存区别 gpu共享内存有什么用_数据 对应的，我们也可以用 gpu共享内存和专用内存区别 gpu共享内存有什么用_数据_02 来索引。
在CPU中，如果用循环遍历二维数组，尤其是双层循环的方式，我们倾向于内层循环对应x，因为这样的访问方式在内存中是连续的，因为CPU的内存是线性存储的，但是GPU的共享内存并不是线性的，而是二维的，分成不同存储体的，并且，并行也不是循环，那么这时候，问题完全不同，没有任何可比性。
回顾放西瓜的例子以及存储体冲突的特性，容易想到，我们最应该避免的是存储体冲突，那么对应的问题就来了，我们每次执行一个线程束，对于二维线程块，一个线程束是按什么划分的呢？是按照threadIdx.x 维进行划分还是按照threadIdx.y维进行划分的呢？
这句话有点迷糊？那我再啰嗦一遍，因为这个很关键，我们每次执行的是一个线程束，线程束里面有很多线程，对于一个二维的块，切割线程束有两种方法，顺着y切，那么就是threadIdx.x固定（变化慢），而threadIdx.y是连续的变化，顺着x切相反；CUDA明确的告诉你，我们是顺着x切的，也就是一个线程束中的threadIdx.x 连续变化。
我们的数据是按照行放进存储体中的这是固定的，所以我们希望，这个线程束中取数据是按照行来进行的，所以

x[threadIdx.y][threadIdx.x];

这种访问方式是最优的，threadIdx.x在线程束中体现为连续变化的，而对应到共享内存中也是遍历共享内存的同一行的不同列

上面这个确实有点绕，我们可以画画图，多想象一下CUDA的运行原理，这个就好理解了，说白了就是不要一个线程束中访问一列共享内存，而是要访问一行。

对照上图，我们把一个int类型（四字节）的1024个元素的数组放到共享内存A中，每个int的索引对应到蓝框中，假设我们的块大小是 gpu共享内存和专用内存区别 gpu共享内存有什么用_gpu共享内存和专用内存区别_03

A[threadIdx.x][threadIdx.y];

的索引方式，就会得到绿框的数据，可想而知，这冲突达到了最大，效率最低、
果我们使用

A[threadIdx.y][threadIdx.x];

我们就会得到红色框中的数据，无冲突，一个事务完成。

本文全部代码在GitHub上可下载使用:https://github.com/Tony-Tan/CUDA_Freshman

行主序访问和列主序访问

行主序访问和列主序访问我们上面已经把原理基本介绍清楚了，我们下面看实现后的试验，这里我们研究的访问，包括读和写，也就是加载和存储。

我们定义块的尺寸为

#define BDIMX 32
#define BDIMY 32

核函数只完成简单的两个操作：

  • 将全局线程索引值存入二维共享内存
  • 从共享内存中按照行主序读取这些值并存到全局内存中

项目完整的代码在24_shared_memory_read_data这个文件夹下，下文我们只贴部分代码。
核函数如下

__global__ void setRowReadRow(int * out)
{
    __shared__ int tile[BDIMY][BDIMX];
    unsigned int idx=threadIdx.y*blockDim.x+threadIdx.x;

    tile[threadIdx.y][threadIdx.x]=idx;
    __syncthreads();
    out[idx]=tile[threadIdx.y][threadIdx.x];
}
  • 定义一个共享内存，大小为 gpu共享内存和专用内存区别 gpu共享内存有什么用_共享内存_04
  • 计算当前线程的全局位置的值idx
  • 将idx这个无符号整数值写入二维共享内存tile[threadIdx.y][threadIdx.x]中
  • 同步
  • 将共享内存tile[threadIdx.y][threadIdx.x]中的值写入全局内存对应的idx位置处

核函数的内存工作：

  1. 共享内存的写入
  2. 共享内存的读取
  3. 全局内存的写入

这个核函数按照行主序读和写，所以对于共享内存没有读写冲突
另一种方法就是按照列主序访问了，核函数代码如下：

__global__ void setColReadCol(int * out)
{
    __shared__ int tile[BDIMY][BDIMX];
    unsigned int idx=threadIdx.y*blockDim.x+threadIdx.x;

    tile[threadIdx.x][threadIdx.y]=idx;
    __syncthreads();
    out[idx]=tile[threadIdx.x][threadIdx.y];
}

原理不再赘述，我们直接看运行结果：
对于使用nvprof如果出现 ======== Error: unified memory profiling failed.错误，是因为系统的保护机制，所以使用sudo权限来执行即可，如果sudo找不到你的nvprof，你可以用完整路径，或则添加到环境变量：

可见行主序的平均时间是 gpu共享内存和专用内存区别 gpu共享内存有什么用_数据_05 而列主序是 gpu共享内存和专用内存区别 gpu共享内存有什么用_二维_06

接下来我们看看检测存储体冲突的指标，会是什么数据：

shared_load_transactions_per_request
shared_store_transactions_per_request
  • shared_load_transactions_per_request 结果：
nvprof  --metrics shared_load_transactions_per_request ./shared_memory_read_data


可以看到load过程行主序1个事务，而列主序32个

  • shared_store_transactions_per_request 结果：
nvprof  --metrics shared_store_transactions_per_request ./shared_memory_read_data

同样行主序的事务是1，而列主序的事务是32

注意，我们这个设备是4-byte宽的，上面第二张图中有相关信息。

按行主序写和按列主序读

完整内容在 https://face2ai.com/CUDA-F-5-2-共享内存的数据布局/