这次回顾ECE408 Lecture 7，这次介绍了卷积和Constant Memory。

课程主页：

搬运视频：

本讲目标

  • 学习卷积，一种重要的并行计算模式；
    • 广泛应用于信号、图像和视频处理；
    • 许多科学和工程应用中的计算基础；
    • 卷积神经网络(CNN)的关键组成部分；
  • 重要的GPU技术；
    • 利用高速缓存；

卷积计算公式

连续和离散情形：

卷积应用

  • 一种在信号处理、数字记录、图像处理、视频处理、计算机视觉和机器学习中以各种形式使用的流行操作；
  • 卷积通常作为滤波器执行，以某种上下文感知方式转换输入信号（音频、视频等）；
    • 某些滤波器平滑信号值，以便可以看到大趋势；
    • 或高斯滤波器模糊图像、背景；

卷积运算

  • 数组运算，其中每个输出数据元素都是相邻输入元素集合的加权和；
  • 加权和计算中使用的权重由输入mask数组定义，通常称为卷积核；

1D卷积

基本内容

  • 常用于音频处理
    • MASK_WIDTH通常是奇数（为了对称性，本例中为5）；
    • MASK_RADIUS是计算输出的每一侧的卷积中使用的元素数（本例中为 2）；

计算P[2]：

计算P[3]：

边界值

  • 计算输入数组边界附近的输出元素需要处理“幽灵”元素；
    • 不同的策略（0，边界值的复制等）；

示例：

代码

  • 该内核将有效范围外的所有元素强制为0；
  • 每个线程计算P的一个元素；

代码：

__global__
void convolution_1D_kernel(float *N, float *M, float *P, int Mask_Width, int Width)
{
    int i = blockIdx.x*blockDim.x + threadIdx.x;
    
    float Pvalue = 0;
    int N_start_point = i - (Mask_Width/2);
    
    for (int j = 0; j < Mask_Width; j++) {
        if (((N_start_point + j) >= 0) && ((N_start_point + j) < Width)) {
        	Pvalue += N[N_start_point + j]*M[j];
        }
    }
    
    P[i] = Pvalue;
}

2D卷积

示例

这里通过几张图给出2D卷积的计算方式，边界值处理，以及幽灵单元：

M​的访问模式

  • M的元素称为mask（kernel，filter）系数；
    • 计算P的所有元素需要M；
    • M在grid执行期间不改变；
  • 在计算所有P元素时，以相同顺序访问M元素；
  • 因此M很适合使用Constant Memory；

缓存

CUDA内存的程序员视图

  • 每个线程可以：
    • 读/写每线程寄存器（~1 个周期）；
    • 读/写每块shared memory （~5 个周期）;
    • 读/写每个网格的global memory（~500 个周期）;
    • 读每个网格的constant memory（在使用缓存的情况下~5 个周期）；

内存层次结构

  • 回顾：如果所有数据都在全局内存中，GPU的执行速度将受到全局内存带宽的限制；
  • 我们在tiled矩阵乘法中使用共享内存来减少此限制；
  • 另一个重要的解决方案：缓存和constant memory；

缓存存储内存行

回顾：内存针对突发进行了优化；

  • 突变时会读取多个比特，例如1024位(128B)；
  • 连续的（线性）地址；
  • 让我们称一次突发为行（line）；

什么是缓存？

  • 一组缓存行（和标签）；
  • 内存读取产生一行；
  • 缓存存储该行的副本；
  • 标签记录行的内存地址；

内存访问显示局部性

一个正在执行的程序：

  • 从内存加载和存储数据；
  • 考虑访问地址的顺序；

序列通常显示两种类型的局部性：

  • 空间：访问X意味着很快将访问X+1（和X+2，等等）；
  • 时间：访问X意味着很快再次访问X；

缓存提高了两种类型的性能。

shared memory vs 缓存

  • shared memory和缓存：
    • 都在芯片上，具有相似的性能；
    • 从Nvidia Volta一代开始，两者都使用相同的物理资源，动态分配！
  • 有什么不同？
    • 程序员控制shared memory内容（显式）；
    • 硬件决定高速缓存的内容（隐式）；

GPU同时具有常量缓存和L1缓存

为了支持写入（修改行）：

  • 必须将更改复制回内存；
  • 缓存必须跟踪修改状态；
  • GPU中的L1缓存（用于全局内存访问）支持写入；

缓存常量/纹理内存：

  • 特殊情况：行是只读的；
  • 为常见的GPU内核访问模式启用比L1更高的吞吐量访问；

如何使用constant memory

  • host代码与以前的版本相似；

  • 但是为M（掩码）分配设备内存：

    • 在所有函数之外；
    • 使用__constant__ （告诉GPU缓存是安全的）；

  • 要复制到设备内存，请使用

    cudaMemcpyToSymbol(dest, src, size, offset = 0, kind = cudaMemcpyHostToDevice)

  • 目标地址为__constant__；

代码示例

// global variable, outside any kernel/function
__constant__ float Mc[MASK_WIDTH][MASK_WIDTH];

// Initialize Mask
float Mask[MASK_WIDTH][MASK_WIDTH]
for(unsigned int i = 0; i < MASK_WIDTH * MASK_WIDTH; i++) {
    Mask[i] = (rand() / (float)RAND_MAX);
    if(rand() % 2) Mask[i] = - Mask[i]
}
cudaMemcpyToSymbol(Mc, Mask, MASK_WIDTH*MASK_WIDTH*sizeof(float));

ConvolutionKernel<<<dimGrid, dimBlock>>>(Nd, Pd);