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.cuint 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 memoryfor (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/

GPU 内存的分级综述(gpu memory hierarchy)相关推荐

  1. 服务器开虚拟机总是gpu满载,vSphere 环境机器学习 GPU 加速方案选型

    GPU 已经成为支撑 AI 应用的一种关键计算加速设备,GPU 的多处理器架构非常适合用来加快深度神经网络应用中的大量矩阵运算过程.大量实测数据表明,跟通用处理器相比,GPU 在运行深度神经网络时具有 ...

  2. PyTorch可视化工具:GPU内存分配

    本文来源  机器之心  编辑:杜伟 想要了解自己的 PyTorch 项目在哪些地方分配 GPU 内存以及为什么用完吗?不妨试试这个可视化工具. 近日,PyTorch 核心开发者和 FAIR 研究者 Z ...

  3. tensorflow GPU 内存不够

    tensorflow GPU 内存不够 from tensorflow as tf gpus = tf.config.list_physical_devices('GPU') if gpus:try: ...

  4. 深度学习 占用gpu内存 使用率为0_深度学习的完整硬件指南

    原标题 | A Full Hardware Guide to Deep Learning 作者 | Tim Dettmers 译者 | linlh.呀啦呼(Tufts University).Ryan ...

  5. 腾讯游戏学院专家分析:Unity在移动设备的GPU内存机制

    导语CPU和GPU是共享一份内存的吗?腾讯游戏学院专家Donald将在本文尝试以一张贴图纹理的虚拟内存占用为例,解答一些内存方面的问题.本篇主要分析iOS系统,后续会更新安卓篇. 开发手机游戏时,常听 ...

  6. PostgreSQL GPU 加速(HeteroDB pg_strom) (GPU计算, GPU-DIO-Nvme SSD, 列存, GPU内存缓存)

    标签 PostgreSQL , GPU , heteroDB , pg_strom , CUDA , nvidia , DIO , Nvme , SSD , 列存 , GPU内存缓存 背景 Heter ...

  7. [人工智能-深度学习-39]:环境搭建 - 训练主机硬件选择全指南(CPU/GPU/内存/硬盘/电源)

    作者主页(文火冰糖的硅基工坊):文火冰糖(王文兵)的博客_文火冰糖的硅基工坊_CSDN博客 本文网址:https://blog.csdn.net/HiWangWenBing/article/detai ...

  8. MICCAI 2021 FLARE 挑战:快速和低 GPU 内存腹部器官分割-附代码

    MICCAI 2021 FLARE 挑战:快速和低 GPU 内存腹部器官分割 简介 腹部器官分割在临床实践中发挥着重要作用,在某种程度上,这似乎是一个已解决的问题,因为最先进的方法已经在几个基准数据集 ...

  9. AMD GPU内存管理(1):概览

    参考内核版本:Linux-6.1.8 HMM 待更新...... dumb buffer create/map 在AMDGPU的Graphics业务中,用到了GEM(Graphics Executio ...

最新文章

  1. Vue.js slots: 为什么你需要它们?
  2. Android状态栏颜色修改
  3. 展望2021年:智能机器人可监督工业机器人干活,效率提升30%
  4. docker-compose运行sentry
  5. Bytes int java_Java Bytes.readInt方法代码示例
  6. 基于cookies的小型购物车程序
  7. 【计算机网络】OSI参考模型与TCP/IP分层模型详解(超级详细,三张图完整说明)
  8. JQuery Tree 树形结构插件 zTree
  9. html中函数传递多个值,JavaScript 实战开发经验!函数多参数传参技巧
  10. catti二级笔译综合能力真题_CATTI 二级口笔译教材、真题和模拟试题汇总!
  11. Dijkstra最短路径算法
  12. 使用clonezilla(在生龙)克隆系统
  13. Firefox 尝试与您指定的代理服务器连接时被拒绝
  14. 【Python】只需2行代码,轻松将PDF转换成Word(含示范案例)
  15. qrcode增加二维码中心图片
  16. 【追剧达人JAVA修复版】苹果CMS10原生APP修改版 影视APP源码附安装教程
  17. Temu拼多多跨境店铺如何快速上货?
  18. git 提交时报错 error: failed to push some refs to ‘https://github.com/xxx/demo.git 解决方法
  19. 安装Properties Editor
  20. CS/BS架构是什么?以及他们的区别

热门文章

  1. makefile的命令包定义及使用
  2. python什么时候用类_python中什么时候使用自定义类
  3. HTML5 ArrayBufferView之DataView
  4. 训练指南第二章-基础问题
  5. 【沟通之道】头脑风暴-女人的心思你别猜
  6. ImportError: No module named arcpy
  7. IE, FireFox, Opera 浏览器支持CSS实现Alpha半透明的方法
  8. mysql数据源找不到_mysql报出找不到目标数据源为啥? 同一个方法在其他地方调用正常...
  9. 【数据平台】基于pyhs2库Python作为client driver连接HiveServer
  10. 搭建Java ME 开发环境