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!)
让程序使得更多人受益!
参考文献
- CUDA专家手册 GPU编程权威指南 [M] 2014
- CUDA Toolkit Documentation v10.1.168 https://docs.nvidia.com/cuda/