【参加CUDA线上训练营】第2天课程知识梳理和实验记录（上午知识梳理）

实验前置知识

错误检测和事件

等待事件完成，设立flag：

cudaError_t cudaEventSynchronize(cudaEvent_t event);//阻塞（事件完成才记录）
cudaError_t cudaEventQuery(cudaEvent_t event);//非阻塞（事件没完成也会记录）

注意cudaEventSynchronize是阻塞的，需要等待时间完成，而cudaEventQuery是非阻塞的，即使事件未完成也会立即执行。

CUDA编程模型中的错误检测中常用cuda error的四个函数：

__host__​__device__​const char*     cudaGetErrorName ( cudaError_t error )
Returns the string representation of an error code enum name.  

__host__​__device__​const char*     cudaGetErrorString ( cudaError_t error )
Returns the description string for an error code.  

__host__​__device__​cudaError_t     cudaGetLastError ( void )
Returns the last error from a runtime call.  

__host__​__device__​cudaError_t     cudaPeekAtLastError ( void )

为了CUDA程序的debug方便，我们可以采用这里的cudaGetErrorString函数，将其封装在error.cuh中：

#pragma once
#include <stdio.h>

#define CHECK(call)                                   \
do                                                    \
{                                                     \
    const cudaError_t error_code = call;              \
    if (error_code != cudaSuccess)                    \
    {                                                 \
        printf("CUDA Error:\n");                      \
        printf("    File:       %s\n", __FILE__);     \
        printf("    Line:       %d\n", __LINE__);     \
        printf("    Error code: %d\n", error_code);   \
        printf("    Error text: %s\n",                \
            cudaGetErrorString(error_code));          \
        exit(1);                                      \
    }                                                 \
} while (0)

然后我们就可以在自己的.cu文件中引用该检错头文件，#include “error.cuh”，利用其定义的CHECK()函数对我们使用的CUDA api进行检错，形如：
CHECK(cudaMallocHost((void **) &h_a, sizeof(int)*m*n));
CHECK(cudaEventCreate(&start));
CHECK(cudaEventCreate(&stop));
CHECK(cudaMemcpy(d_a, h_a, sizeof(int)*m*n, cudaMemcpyHostToDevice));
CHECK(cudaEventRecord(stop));
CHECK(cudaEventSynchronize(stop));
CHECK(cudaEventElapsedTime(&elapsed_time, start, stop));
下面的实验有检错回显，这个封装思路很巧妙，用来Debug确定出错的位置很好用，具体可以参考樊老师的《CUDA编程基础与实践》一书的git repo：
CUDA-Programming/src/04-error-check at master · brucefan1983/CUDA-Programming · GitHub

CUDA存储单元

右边单向箭头表示是可读的，双向箭头表示是可读可写的（据说这里考试会有小陷阱）

下面我们大体上从由快到慢的顺序介绍GPU的各种存储单元，最后讨论主机端的存储器内存。

Register

寄存器最快，我们要尽量让更多的block主流在SM中，以增加Occupancy，省着点使用寄存器资源。

Shared Memory

• 比Register稍慢一点
• On-chip
• 拥有高的多带宽和低很多的延迟
• 同一个Block中的线程共享一块Shared Memory
• 用__syncthreads()进行同步
• 比较小，要节省使用，否则会限制活动warp的数量。
• SMem被分成32个逻辑块（banks）
  

注意这里的bank conflict，只有在warp中的线程都访问同一个bank的资源的时候才不会存在bank conflict

Local Memory

属于On-board（而不是on chip），但却是GPU线程私有的，空间比较大，而读写比较慢。（这一点需要注意）
Register不够的时候就会用Local Mem来替代，但更多地是在以下情况使用Local Memory：

1. 无法确定其索引是否为常量的数组
2. 会消耗太多寄存器孔家你的大型结构或数组
3. 如果内核使用了多于可用寄存器的任何变量（这也称为寄存器溢出）
4. --ptxas-options=-v

Constant Memory

• 固定内存空间驻留在设备内存中，并缓存在固定缓存中（constant cache）。
• 范围是全局（对所有kernel可见）
• kernel从CM只能读而不能写，因此初始化必须在host端使用cudaError_t cudaMemcpyToSymbol(const void* symbol, const void* src, size_t count);
• 当一个warp中所有的thread都从同一个Memory地址读取数据时，constant Memory表现会非常好，会触发广播机制。

常量内存应用举例-光线追踪

这里是效果图，我们要在某个位置显示球的颜色或者黑色，其中球之间可能存在遮盖，所以需要经过计算距离来确定哪个球在最前面，然后显示这个球的颜色。

这里采用的是hit方法，计算光线是否与球面相交，若相交则返回光线到命中球面的距离。

这里我们将需要大量访问的内容放到常量内存中，也就是说将球体的位置数据部分放到__constant__ Sphere s[SPHERES];中。实现光线追踪部分的代码，首先将threadIdx映射到像素的位置，每个线程都干自己的事情，然后让图像坐标偏移DIM/2，使z轴穿过图像中心，初始化背景颜色为黑色，距离初始化为负无穷-INF，然后开始计算距离：遍历每一个球体，调用上面的hit方法计算光线和球面的距离，如果距离更近则将距离更新为此值，否则不用修改距离值。完成对球面相交的判断后，将当前的颜色等信息保存到我们输出的图像中，ptr[offset*4 + 0]这里的4表示每个点的r，g，b和透明度信息共需要四个存储单元，这些信息存储在一个一维数组中，所以组织的时候需要引入offset偏移量进行索引。

最后生成球面的中心坐标颜色和半径，通过球面数据生成bitmap（这里讲的比较粗略）。

Texture Memory

• 驻留在device Mem中，属于On-board，并使用一个只读cache。
• 与global Memory在一块，但是有自己专有的只读cache
• on-chip，所以比DRAM上取数据减少了内存请求和提高带宽
• 专门为那些在内存访问模式中存在大量空间局部性的图形应用程序而设计的。（也就是说，一个thread读的位置可能与临近的thread读的位置非常接近，如下）。
  

举个texture memory的应用实例-热传导模型：

Global Memory

• 空间最大，latency最高，是GPU中最基础的memory。
• On-board，驻留在Device memory中
• memory transction 对齐，合并访存。
• 合并访存机制，如下图的矩阵乘法：
  

我们可以让线程按照行或者列进行读取，那么哪一种更快呢？答案是按列读快，如下左边是按行读（例如thread0用四次迭代分别读取A00，A01，A02，A03），右边是按照列读（例如thread0用四次迭代分别读取B00，B10，B20，B30）。白色空格部分指的是访存取出来的数据空间部分，显然按行取效率低下，而这里按照列读取的话，我们每一次迭代只需要一次访存即可满足四个线程的取数据操作。

Host Memory

主机端存储器主要是内存可以分为两类：可分页内存（Pageable）和页面 （Page-Locked 或 Pinned）内存。
可分页内存通过操作系统 API(malloc/free) 分配存储器空间，该内存是可以换页的，即内存页可以被置换到磁盘中。可分页内存是不可用使用DMA（Direct Memory Acess)来进行访问的，普通的C程序使用的内存就是这个内存。

总结

不同的应用可能适用于不同的存储单元或他们的各种组合，我们要掌握好每种存储单元的特点并合理架构，所有这些常用的GPU存储单元的特性汇总如下表：