【参加CUDA线上训练营】CUDA进阶之路 - Chapter 5 - 多种CUDA存储单元详解

最新推荐文章于 2023-02-11 17:05:30 发布

Chiao.

最新推荐文章于 2023-02-11 17:05:30 发布

阅读量614

点赞数 1

文章标签： c语言人工智能 linux 硬件架构

本文链接：https://blog.csdn.net/weixin_43908429/article/details/128984369

版权

5.1 寄存器

不同存储的位置不同，意味着不同存储的特性（大小、读写速度等）亦有所不同。前面GPU存储单元概述中已经提到过，GPU的存储单元主要分布于两处地方：片上（On chip）和板上（On board）。

寄存器是速度最快的存储单元，位于GPU芯片的SM上，每个SM（SMX）上有成千上万的32位寄存器，当kernel函数启动后，这些寄存器被分配给指定的线程来使用。kernel中没有什么特殊声明的自动变量都是存放在寄存器中的。当数组的索引是const类型且在编译期能被确定的话，就是内置类型，数组也是放在寄存器中的。

寄存器变量是每个线程私有的，一旦thread执行结束，寄存器变量就会失效。
寄存器是稀有资源，可用--maxrregcount指令设置其大小。省着用可以让更多的Block驻留在SM上，增加Occupancy。
不同设备架构，数量不同。

5.2 Shared Memory

Shared Memory位于GPU芯片上，访问延迟仅次于寄存器。Shared Memory是可以被一个Block中的所有Thread来进行访问的，可以实现Block内的\textsf{线程间的低开销通信}。用__shared__修饰符修饰的变量存放在shared memory；

on-chip；
拥有高得多的带宽和低很多的延迟（Latency）；
__syncthreads()同步，因为shared memory涉及多个线程之间的协作，同一个存储地址可能被多个线程进行读写操作，加入同步操作就可保证读写不会冲突；
比较小，省着用，否则会限制活动warp的数量。

使用了shared memory的Kernel函数通常要按如下三个阶段进行编写：

Load shared memory and __syncthreads();
Process shared memory and __syncthreads();
Write results

在nvcc编译阶段可以加入-Xptxas –v,abi=no选项来观察Kernel函数对shared memory的使用情况。在运行时，可以通过cuFuncGetAttribute(CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES)获得Kernel函数对shared memory的占用大小。

5.3 Local Memory

Local Memory本身在硬件中没有特定的存储单元，而是从Global Memory虚拟出来的地址空间。Local Memory是为寄存器无法满足存储需求的情况而设计的}，主要是用于存放单线程的大型数组和变量。Local Memory是线程私有的，线程之间是不可见的。由于GPU硬件单位没有Local Memory的存储单元，所以，针对它的访问是比较慢的。从上面的表格中，也可以看到跟Global Memory的访问速度是接近的。

Local Memory不需要手动指定，更多在以下情况当中CUDA会申请Local Memory：

无法确定其索引是否为常量的数组；
会消耗太多寄存器空间的大型结构或数组；
如果内核使用了多于可用寄存器的任何变量（寄存器溢出）；

在nvcc编译阶段可以加入-Xptxas –v,abi=no选项来观察Kernel函数对shared memory的使用情况。在运行时，可以通过cuFuncGetAttribute(CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES)获得Kernel函数对shared memory的占用大小。

5.4 Constant Memory

Constant Memory类似于Local Memory，也是没有特定的存储单元的，只是Global Memory的虚拟地址。因为它是只读的，所以简化了缓存管理，硬件无需管理复杂的回写策略。Constant Memory同时是对所有Kernel可见的，它启动的条件是同一个warp所有的线程同时访问同样的常量数据。

声明时使用__constant__标记符。
由于Kernel只能从Constant Memory中读取数据，因此其初始化必须在host端使用下面的api调用：

cudaError_t cudaMemcpyToSymbol(const void* symbol, const void* scr, size_t count)

当一个warp中所有thread都从同一个Memory地址读取数据时，constant memory表现会非常好，会触发广播机制。

5.4.1 常量内存运用实例：光追

光线追踪(Ray Tracing)希望能模拟光线在真实环境下的折射，反射，漫反射，间接反射等物理现象。如果我们能够模拟整个场景光线的传播，就能得到非常真实的图像，但如果直接从光源开始发射的光线经过传播后并不一定会进入摄像机，造成了大量计算浪费。使用逆向追踪——即从摄像机开始发射光线进行追踪，将有效减少无效计算，即使如此，计算量依然是巨大的，如果每一个像素点发射一条光线进行追踪，并且追踪10次折射或反射的话.生成一张1080P的图，相交的计算量将是 1920 * 1080 * 10 = 20 736 000，约两千万次光线追踪运算。

请添加图片描述

单独看一个像素上的追踪过程的话:

从视线方向发射一条射线。
取得射线与场景最近的交点。
取得交点的材质颜色。
如果材质包含反射或折射，则改变光线方向。
寻找下一个交点，并重复(2)。直到在场景内找不到交点，或达到最大追踪次数。

常量内存在其中起到的作用就是用来存储场景中所有物体的坐标、形状、颜色、材质等各种信息，共每个线程在进行光线追踪时进行访问。

struct Sphere{
    float r, g, b;
    float radius;
    float x, y, z;

    __device__ float hit(float ox, float oy, float *n){
        // hit方法，计算光线是否与球面相交，若相交则返回光线到命中球面处的距离。
        float dx = ox - x;
        float dy = oy - y;

        if(dx * dx + dy * dy < radius * radius){
            float dz = sqrt(radius * radius - dx * dx - dy * dy);
            *n = dz / sqrt(radius * radius);
            return dz + z;
        }

        return -INF;
    }

    __constant__ Sphere s[SPHERES];

    __global__ void rayTracing(unsigned char* ptr){
        // 将threadIdx及blockIdx映射到像素位置。
        int x =threadIdx.x + blockIdx.x * blockDim.x;
        int y =threadIdx.y + blockIdx.y * blockDim.y;

        int offset = x + y * blockDim.x * gridDim.x;

        // 让图像坐标偏移DIM/2，使z轴穿过图像中心
        float ox = (x - DIM/2);
        float oy = (y - DIM/2);

        float r=0, g=0, b=0; // 初始化背景颜色为黑色
        float maxz = -INF;
        for(int i=0; i<SPHERES; i++){
            float n;
            float t=s[i].hit(ox, oy, &n);
            if(t > maxz){  // 如果比上一次命中距离更近，则将这个距离保存为最近距离，并且保存球面颜色。
                float fscale = n;
                r = s[i].r * fscale;
                g = s[i].g * fscale;
                b = s[i].b * fscale;
                maxz = t;
            }
        }
        // 判断完球面相交情况后，将当前颜色保存到输出图像中。
        ptr[offset*4 + 0] = (int)(r * 255);
        ptr[offset*4 + 1] = (int)(g * 255);
        ptr[offset*4 + 2] = (int)(b * 255);
        ptr[offset*4 + 3] = 255;
    }
}

5.5 Texture Memory

Texture Memory驻留在device memory中，并且只使用一个只读cache。Texture Memory是专门为那些在内存访问模式中存在大量空间局限性（Spacial Locality）的图形应用程序而设计的。假如在某个计算应用程序中，一个线程读取的位置与邻近线程读取的位置"非常接近"，使用纹理内存将会减少内存通信量，从而提高性能。

Texture Memory 实际上也是global memory的一块，但它有自己专有的只读cache。
纹理内存也是缓存在片上的，因此一些情况下比从芯片外的DRAM上获取数据，纹理内存可以通过减少内存请求来提高带宽。
从数学的角度，下图的四个地址并非连续的，在一般的CPU缓存中，这些地址将不会缓存。但由于GPU纹理缓存是专门为了加速这种访问模式而设计的，因此如果在这种情况中使用纹理内存而不是全局内存，将会获得性能的提升。

纹理内存

5.5.1 纹理内存运用实例：热传导

假设一个矩形网格，热源分布如下图所示。假设热源单元本身的温度保持不变，热量可以在相邻单元之间流动，从高温单元传导到低温单元。
请添加图片描述
温度更新的计算方法：将单元与其邻接单元的温差加起来，然后加上原有温度，等于新时刻的温度
$T_{NEW} = T_{OLD}+\sum_{NEIGHBOR}(k(T_{NEIGHBOR}-T_{OLD}))$
由于邻接单元只有上下左右四个，将上式化简，得到：
$T_{NEW}=T_{OLD}+k(T_{TOP}+T_{BOTTOM}+T_{LEFT}+T_{RIGHT}-4T_{OLD})$

代码中温度更新的计算流程：

给定包含初始温度的网格，它的大部分单元都是0，少部分是初始温度值。将其中作为热源的初始温度值复制到当前时间的网格对应单元中。确保“加热单元将保持恒温”，这个复制操作是copy_const_kernal()中执行的。
给定一个输入温度网格，根据温度更新的公式计算输出温度网格。这个更新操作是在blend_kernal()中执行的。
将输入温度网格和输出温度网格交换，为下一个步骤的计算做好准备。当模拟下一个时间步时，步骤2中计算得到的输出温度网格将成为步骤1中的输入温度网格。

Texture Memory的申请是通过绑定完成的，将申请到的内存空间绑定到Texture Memory：cudaBindTexture()。访问Texture Memory：tex1Dfetch()

5.6 Global Memory

Global Memory在某种意义上等同于GPU显存，kernel函数通过Global Memory来读写显存。Global Memory是kernel函数输入数据和写入结果的唯一来源。

空间最大，latency最高，GPU最基础的memory；
驻留在Device memory中
memory transaction对齐，合并访存：当我们从global memory中读取一个数据时，实际返回的一块连续的数据。如果当前时间步骤内，连续的另一个线程刚好要访问的是连续的数据，则可以在一个时间步骤内同时完成。

在GPU实现矩阵乘法的过程中，每一个线程既要从global memory中读取连续的一行，也要从中读取连续的一列？总的来看，按行读更快还是按列读更快？
按列读更快。如下图所示，每一个线程T在A矩阵中按行读取，在B矩阵中按列读取。一个线程按列读取时，将触发广播机制（即合并访存），由于是行主序，连续相邻的数据将同时被返回给相邻的线程。而按行读取时，每一个线程访问global memory后多余的相邻数据并不能为其他线程所用。

5.7 基于ARM平台的Jetson Nano存储单元

5.7.1 统一内存的基本概念

统一内存是可从系统中的任何处理器访问的单个内存地址空间。这种硬件/软件技术允许应用程序分配可以从CPUs或者GPUs上运行的代码读取或写入的数据。分配统一内存非常简单，只需要对malloc()或new的调用替换为对cudaMallocManaged()的调用，这是一个分配函数，返回可从任何处理器访问的指针。

内存空间的统一意味着主机和设备之间不再需要显式内存传输。在托管内存空间中创建的任何分配都会有CUDA的SDK自动迁移到需要的位置。统一内存的出现使得GPU编程显得更加简便。

可直接访问CPU内存、GPU显存，不需要手动拷贝数据；
CUDA在现有的内存结构上增加了一个统一内存系统，程序员可以直接访问任何内存/显存资源，或者在合法的内存空间内寻址，而不用管涉及到的到底是内存还是显存；
CUDA的数据拷贝由程序员的手动转移变成自动执行，但它仍然受制于PCI-E的带宽和延迟。也就是说所谓的统一内存，实际上就是CUDA帮你完成了手动数据传输的过程罢了。

一个没有使用统一内存的简单CPU程序：

void sortfile(FILE *fp, int N){
    char *data;
    data = (char *)malloc(N)};

    fread(data, 1, N, fp);

    qsort(data, N, 1, cmp);

    use_data(data);

    free(data);
}

一个使用了统一内存的简单GPU程序：

void sortfile(FILE *fp, int N){
    char *data;
    cudaMallocManaged(&data, N);

    fread(data, 1, N, fp);

    qsort<<<...>>>(data, N, 1, cmp);
    cudaDeviceSynchronize();

    use_data(data);

    cudaFree(data);
}

Unified Memory有两种实现方法：

cudaError_t cudaMallocManaged(void **devPtr, size_t size, unsigned int flags=0)；
标识符：__managed__；

__global__ void printme(char *str){
    printf(str);
}

int main(){
    char *s;
    cudaMallocManaged(&s, 100);
    strncpy(s, "Hello Unified Memory\n", 99);
    printme<<< 1, 1 >>>(s);
    cudaDeviceSynchronize();
    cudaFree(s)
}

原来没有使用统一内存时，通常要在执行完核函数之后进行cudaMemcpy()，类似的函数已经自带多线程同步功能了。而使用统一内存后，由于不再进行类似的数据拷贝命令，故需要在执行完核函数之后，手动加上cudaDeviceSynchronize()。

__device__ __managed__ int x[2];
__device__ __managed__ int y;
__global__ void kernel(){
    x[1] = x[0] + y;
}

int main(){
    x[0] = 3;
    y = 5;
    kernel<<< 1, 1 >>>();
    cudaDeviceSynchronize();
    printf("result=%d\n", x[1]);
    return 0;
}

从上面的代码实例中可以看出，尽管__managed__前面是__device__，Host仍然是能够访问其申请到的内存地址的。

还有值得提到的一点，对于复杂的结构体，使用统一内存可以省掉非常多的麻烦。例如，当一个结构体内部有一个字符串成员时，手动数据转移不仅要将结构体拷贝过去，还需要另外专门将其字符串所在的存储地址也拷贝过去。使用统一内存则避免了如此繁琐的操作。

5.7.2 基于ARM平台的Jetson Nano存储单元特点

在Jetson Nano上当然可以使用统一内存的方法，然而Jetson Nano还可以有更高级的方法。在Jetson Nano所使用的Tegra SoC芯片上专门设计了一个可供CPU和GPU共享的存储空间。Jetson平台上的GPU是整合的（integrated GPU，iGPU），不同于独立的GPU（discrete GPU, dGPU），iGPU与CPU共享着存储空间。对于SoC DRAM上的同一个内存地址，CPU和iGPU可以同时进行访问（如下图）。

jetson nano 存储单元特点

也就是说在Jetson平台进行CUDA编程时，不需要像在dGPU里那样，两边倒腾数据了。但是，在Jetson平台上进行数据拷贝的操作也是允许的，这样做的好处是具有更好的跨平台支持。

5.8 总结

目前异构计算当中存在的很多瓶颈都是卡在数据传输，因为硬件物理带宽是很难优化的。各种Memory要灵活运用，自定义方法的上限更高。

请添加图片描述

Chiao.

关注

1
点赞
踩
0

收藏

觉得还不错? 一键收藏
0
评论
复制链接

分享到 QQ

分享到新浪微博

扫一扫