CUDA编程笔记(5)
创始人
2024-05-16 16:56:28
0

文章目录

  • 前言
  • CUDA的内存组织
    • 全局内存
    • 常量内存
    • 纹理内存和表面内存
    • 寄存器
    • 局部内存
    • 共享内存
    • L1和L2缓存
    • SM的构成
  • API函数查询设备
  • 总结


前言

cuda的内存组织,在使用GPU时尽可能提高性能,合理的使用设备的内存也是十分重要的。

CUDA的内存组织

如表所示:

内存类型物理位置访问权限可见范围生命周期
全局内存在芯片外可读可写所有线程和主机端由主机分配和释放
常量内存在芯片外仅可读所有线程和主机端由主机分配和释放
纹理和表面内存在芯片外一般仅可读所有线程和主机端由主机分配和释放
寄存器内存在芯片内可读可写单个线程所在线程
局部内存在芯片外可读可写单个线程所在线程
共享内存在芯片内可读可写单个线程块所在线程块

全局内存

定义:这里的全局内存,指的是核函数中所有线程都能访问到数据的内存。
作用:保存核函数提供数据,并在主机与设备及设备与设备之间传递数据。
不在GPU芯片上,所以为核函数提供数据时具有较高的延迟和较低的访问速度。
内存容量基本和GPU的显存差不多。
是可读可写的。
动态全局内存变量:前面cuda数组相加的程序中定义的d_x,d_y,d_z就是动态分配的,要先通过cudaMalloc()为其分配设备内存和cudaMemcpy()将主机上的数据传递到设备上,然后在核函数中访问分配的内存和改变其中的数值。
静态全局内存变量:使用cudaMemcpyToSymbol()进行主机与设备之间的数据传输和cudaMemcpyFromSymbol()进行设备与主机之间的数据传输。在核函数中,可直接对静态全局内存变量进行访问,并不需要将它们以参数的形式传给核函数。
由以下方式在函数外部定义

__device__ T x;  // 单个变量
__device__ T y[N];	// 固定长度的数组

例子:
在这里插入图片描述

常量内存

定义:是有常量缓存的全局内存,数量有限,仅有64kb。
作用:和全局内存一样。
仅可读不可写,而且由于有缓存,常量内存的访问速度比全局内存要高。
使用:cuda数组相加的程序里的const int N,就是使用了常量内存的变量。

纹理内存和表面内存

定义:类似于常量内存
一般仅可读,表面内存也可写。对于计算能力不小于3.5的GPU来说,将某些只读全局内存数据用__ldg()函数通过只读数据缓存读取,既可以达到使用纹理内存的加速效果,又可使代码简洁。

寄存器

定义:在核函数中不加任何限定符的变量一般来说就存放于寄存器中(可能在局部内存中)。
寄存器可读可写。寄存器内存在芯片内,是所有内存中访问速度最高的,但其数量有限。
使用:cuda数组相加的程序里的 int n = blockDim.x * blockIdx.x + threadIdx.x;
其中n就是一个寄存器变量。在核函数中使用z[n] = x[n] + y[n],寄存器变量n并将赋值号右边计算出来的赋值给它。
生命周期与所属线程的生命周期一致,从定义它开始到线程结束。寄存器变量仅仅被一个线程可见,不同的线程中该寄存器变量名一样,但是变量的值是不同的,相当于另外创建了副本。

局部内存

定义:和寄存器几乎一样。
寄存器里放不下的变量可能放在局部内存里,这种判断是由编译器自动做。

共享内存

定义:与寄存器类似,但共享内存对整个线程块可见。
作用:减少对全局内存的访问,或者改善对全局内存的访问模式。
其生命周期与整个线程块一致。
使用:在核函数中要将一个变量定义为共享内存变量,就要在定义语句中加上一个限定符__shared__

__shared__ real s_y[128];

L1和L2缓存

从费米架构开始,有了SM层次的L1缓存(一级缓存)和设备层次的L2缓存(二级缓存)。
主要用来缓存全局内存和局部内存的访问,减少延迟。L1和L2缓存是不可编程的缓存(用户最多能引导编译器做一些选择)。

SM的构成

(1)一定数量的寄存器
(2)一定数量的共享内存
(3)常量内存的缓存
(4)纹理和表面内存的缓存
(5)L1缓存
(6)两个线程束调度器,用于在不同线程的上下文之间迅速切换及为准备就绪的线程束发出执行指令。
(7)执行核心:若干整型数运算的核心,若干单精度浮点数运算的核心,若干双精度浮点数运算的核心,若干单精度浮点数超越函数的特殊函数单元,若干混合精度的张量核心。

API函数查询设备

用一些cuda的api程序来查询设备的一些规格。

#include "error.cuh"
#include int main(int argc, char *argv[])
{// 设置查询的设备编号.int device_id = 0; if (argc > 1) device_id = atoi(argv[1]);// cudaSetDevice()函数将对所指定的设备进行初始化CHECK(cudaSetDevice(device_id));// 定义设备输出规格的一些结构体变量cudaDeviceProp prop;CHECK(cudaGetDeviceProperties(&prop, device_id));  // 得到了device_id设备的性质,存放在结构体变量中的prop中.printf("Device id:                                 %d\n", device_id);printf("Device name:                               %s\n",prop.name);printf("Compute capability:                        %d.%d\n",prop.major, prop.minor);printf("Amount of global memory:                   %g GB\n",prop.totalGlobalMem / (1024.0 * 1024 * 1024));printf("Amount of constant memory:                 %g KB\n",prop.totalConstMem  / 1024.0);printf("Maximum grid size:                         %d %d %d\n",prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);printf("Maximum block size:                        %d %d %d\n",prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);printf("Number of SMs:                             %d\n",prop.multiProcessorCount);printf("Maximum amount of shared memory per block: %g KB\n",prop.sharedMemPerBlock / 1024.0);printf("Maximum amount of shared memory per SM:    %g KB\n",prop.sharedMemPerMultiprocessor / 1024.0);printf("Maximum number of registers per block:     %d K\n",prop.regsPerBlock / 1024);printf("Maximum number of registers per SM:        %d K\n",prop.regsPerMultiprocessor / 1024);printf("Maximum number of threads per block:       %d\n",prop.maxThreadsPerBlock);printf("Maximum number of threads per SM:          %d\n",prop.maxThreadsPerMultiProcessor);return 0;
}

查询的一些设备设置:
在这里插入图片描述
从这些输出可以看出GPU的内存组织,和所占各内存的最大容量大小。

总结

cuda程序执行的计时方式和GPU性能加速的分析
参考:
如博客内容有侵权行为,可及时联系删除!
CUDA 编程:基础与实践
https://docs.nvidia.com/cuda/
https://docs.nvidia.com/cuda/cuda-runtime-api
https://github.com/brucefan1983/CUDA-Programming

相关内容

热门资讯

安卓系统苹果手机识别,跨界融合... 你知道吗?在科技飞速发展的今天,手机已经成为了我们生活中不可或缺的一部分。而说到手机,安卓系统和苹果...
harmonyos系统是不是安... 亲爱的读者,你是否曾好奇过HarmonyOS系统与安卓系统之间的关系?是不是安卓的“亲戚”?今天,就...
手机怎么装系统安卓,安卓系统安... 手机卡顿了?想给安卓系统来个大变身?别急,跟着我一步步来,保证让你的手机焕然一新!一、准备工作在开始...
安卓Linux系统内网穿透,A... 你有没有想过,你的安卓手机里那些看似普通的APP,其实可能正在悄悄地帮你打通网络世界的任督二脉呢?没...
win怎么安装安卓系统,Win... 亲爱的读者,你是不是对Win系统上的安卓应用垂涎已久,但又苦于不知道如何安装安卓系统呢?别急,今天我...
升级小米平板安卓系统,畅享全新... 你有没有发现,你的小米平板用久了,是不是感觉有点卡呢?别急,今天就来教你怎么给它来个系统升级,让它焕...
捷豹安卓系统车载,捷豹安卓系统... 哇,你有没有想过,当你的手机和汽车融为一体,会是怎样的体验呢?想象你正驾驶着你的捷豹,车窗外的风景如...
安卓1到10系统,安卓1.0至... 你有没有想过,手机里的安卓系统就像是我们生活中的好朋友,从青涩的少年成长为稳重的青年呢?从安卓1.0...
安卓8.0停用系统应用,提升使... 你知道吗?最近安卓系统又来了一次大动作,那就是安卓8.0系统开始停用一些系统应用了。这可真是让人有点...
安卓系统修改mtu值,轻松提升... 你有没有想过,你的安卓手机其实是个小小的电脑呢?它里面藏着许多可以自定义的秘密功能,就像修改MTU值...
安卓平板改window系统,探... 你有没有想过,你的安卓平板其实可以摇身一变,变成一个Windows系统的电脑呢?没错,就是那种可以运...
时空猎人安卓苹果系统,探索无尽... 你知道吗?最近在手机游戏圈里,有一款叫做《时空猎人》的游戏可是火得一塌糊涂呢!不管是安卓用户还是苹果...
安卓9.0系统的电视,新一代电... 亲爱的读者们,你是否也像我一样,对科技新玩意儿充满好奇?今天,我要和你聊聊一个让人眼前一亮的话题——...
小pc安装安卓系统,轻松安装安... 你有没有想过,你的小PC也能变身成为安卓系统的超级玩家呢?没错,就是那个平时默默无闻的小家伙,现在也...
高通备份安卓系统,全方位数据安... 你知道吗?在这个科技飞速发展的时代,手机备份可是个不得不提的话题。尤其是对于安卓用户来说,选择一个靠...
谷歌安卓系统有多少,从诞生到全... 你有没有想过,那个无处不在的谷歌安卓系统,究竟在全球有多少用户呢?它就像一个神秘的数字,每天都在悄悄...
fc黄金传说安卓系统,畅享复古... 你有没有听说最近安卓系统上的一款超酷的游戏——《FC黄金传说》?这款游戏可是让不少玩家都沉迷其中,今...
变小的我安卓系统,安卓系统演变... 你有没有发现,最近你的手机好像变轻了?没错,说的就是你,那个陪伴你多年的安卓系统。它悄无声息地进行了...
vivo安卓系统小彩蛋,体验科... 你知道吗?在vivo的安卓系统中,竟然隐藏着一些超有趣的小彩蛋!这些小彩蛋就像是在手机里埋下的宝藏,...
安卓系统如何强制重启,安卓系统... 手机突然卡壳了,是不是又该给它来个“大保健”了?没错,今天就来聊聊安卓系统如何强制重启。别小看这个看...