GPU 内存的分级综述（gpu memory hierarchy）

GPU 内存的分级（gpu memory hierarchy）

小普 中科院化学所在读博士研究生

研究课题，计算机模拟并行软件的开发与应用

Email:  yaopu2019@126.com （欢迎和我讨论问题）

CSDN：

博客园：

摘要(Abstact)

GPU 的存储是多样化的, 其速度和数量并不相同，了解GPU存储对于程序的性能调优有着重要的意义。本文介绍如下几个问题:

1.内存类型有什么？2)查询自己设备的内存大小 3）内存访问速度4）不同级别的存储关系5）使用注意事项。各种存储结构的优缺点。

 

正文

GPU结构图

 

 

①寄存器内存（Register memory）

优点：访问速度的冠军！

缺点：数量有限

使用：在__global__函数 ，或者___device__ 函数内，定义的普通变量，就是寄存器变量。

例子：

//kernel.cu

__global__ void register_test()

{

 int a = 1.0;

double b = 2.0;

}



//main.cu

int main()

{

int nBlock = 100;

register_test <<<nBlock,128>>>();

return 0;

}

 

②共享内存（Shared memory）

优点：

1缓存速度快 比全局内存 快2两个数量级

2 线程块内，所有线程可以读写。

 3 生命周期与线程块同步

缺点：大小有限制

使用：关键词 __shared__  如 __shared__ double A[128];

适用条件：

使用场合，如规约求和 : a = sum A[i]

如果不是频繁修改的变量，比如矢量加法。

是编程优化中的重要手段！

C[i] = A[i] + B[i] 则没有必要将A,B进行缓存到shared memory 中。

 

//kernel.cu

__global__ void shared_test()

{

__shared__ double A[128];

 int a = 1.0;

double b = 2.0;

int tid = threadIdx.x;

A[tid] = a;

}

 

③全局内存 (Global Memory)

优点：

1空间最大（GB级别）

2.可以通过cudaMemcpy 等与Host端，进行交互。

3.生命周期比Kernel函数长

4.所有线程都能访问

缺点：访存最慢

//kernel.cu

__global__ void shared_test(int *B)

{

double b = 2.0;

int tid = threadIdx.x;

int id = blockDim.x*128 + threadIdx.x;

int a = B[id] ;

}

 

④纹理内存

优点，比普通的global memory 快

缺点：使用起来，需要四个步骤，麻烦一点

适用场景：比较大的只需要读取array，采用纹理方式访问，会实现加速

使用的四个步骤（这里以1维float数组为例子），初学者，自己手敲一遍代码！！！

第一步，声明纹理空间,全局变量：

texture<float, 1, cudaReadModeElementType> tex1D_load;

第二步,绑定纹理

 

 

声明语句：

#include <iostream>

#include <time.h>

#include <assert.h>

#include <cuda_runtime.h>

#include "helper_cuda.h"

#include <iostream>

#include <ctime>

#include <stdio.h>



using namespace std;



texture<float, 1, cudaReadModeElementType> tex1D_load;

//第一步，声明纹理空间,全局变量



__global__ void kernel(float *d_out, int size)

{

    //tex1D_load 为全局变量，不在参数表中

    int index;

    index = blockIdx.x * blockDim.x + threadIdx.x;

    if (index < size)

    {

        d_out[index] = tex1Dfetch(tex1D_load, index); //第三步，抓取纹理内存的值

        //从纹理中抓取值

        printf("%f\n", d_out[index]);

    }

}



int main()

{

    int size = 120;

    size_t Size = size * sizeof(float);

    float *harray;

    float *d_in;

    float *d_out;



    harray = new float[size];

    checkCudaErrors(cudaMalloc((void **)&d_out, Size));

    checkCudaErrors(cudaMalloc((void **)&d_in, Size));



    //initial host memory



    for (int m = 0; m < 4; m++)

    {

        printf("m = %d\n", m);

        for (int loop = 0; loop < size; loop++)

        {

            harray[loop] = loop + m * 1000;

        }

        //拷贝到d_in中

        checkCudaErrors(cudaMemcpy(d_in, harray, Size, cudaMemcpyHostToDevice));



        //第二步,绑定纹理

        checkCudaErrors(cudaBindTexture(0, tex1D_load, d_in, Size));

        //0表示没有偏移



        int nBlocks = (Size - 1) / 128 + 1;

        kernel<<<nBlocks, 128>>>(d_out, size); //第三步

        cudaUnbindTexture(tex1D_load);         //第四，解纹理

        getLastCudaError("Kernel execution failed");

        checkCudaErrors(cudaDeviceSynchronize());

    }

    delete[] harray;

    cudaUnbindTexture(&tex1D_load);

    checkCudaErrors(cudaFree(d_in));

    checkCudaErrors(cudaFree(d_out));

    return 0;

}

 

总结如下表

 

结束语

小普 中科院化学所在读博士研究生

研究课题，计算机模拟并行软件的开发与应用

Email:  yaopu2019@126.com （欢迎和我讨论问题，私信和邮件都OK!）

让程序使得更多人受益！

参考文献

1. CUDA专家手册 GPU编程权威指南 [M] 2014
2. CUDA Toolkit Documentation v10.1.168 https://docs.nvidia.com/cuda/