共享内存 共享内存和L1缓存同属为一块物理内存，所以他的访问和写入速度是和L1缓存相当的。并且可以通过CUDA提供的API，设置L1缓存和共享内存的大小，调整他们所占的比例，如果你几乎用不到共享内存，就可以将共享内存的大小调到最小，这样可以获得更大的L1缓存。 共享内存对线程束占用率的影响 和L1缓存一样，共享内存也是每个SM单元都独有一块共享内存块，并且在核函数被调用时，共享内存是在线程块中被共享的。 也就是说： ​ 共享内存使用量 = 每个SM单元中线程块的数量 X 核函数中用到的共享内存大小 如果共享内存资源…

前言 NVIDIA Nsight Compute 和NVIDIA Nsight Systems作为CUDA工具包中的性能分析工具，但是他们的定位有所不同。 NVIDIA Nsight Systems 更强调的是整个程序的性能分析，不仅对于CUDA的运行信息，以及CPU中所执行的任务也是可以进行数据采样的。 NVIDIA NSight Compute 是一个专门针对核函数的性能分析工具，它对核函数的信息采集得非常详细，可以对列如缓存命中、显存吞吐量、计算吞吐量等信息进行追踪。它可以对每个核函数生成报表，这些报表中的信…

前言 NVIDIA Nsight Systems 是一个作用于CUDA的轻度性能分析工具，如果是想知道核函数的执行顺序、消耗时间、资源占用率等信息，那么这个工具很合适。 这里使用的是 2023.2.3版本的软件，软件版本不同，可能使用方法会有所不同，注意甄别。 使用之前 这个软件在CUDA 工具包中，如果安装了CUDA是随带一起安装的。 （windows）在使用之前要注意用户是否为管理员，如果是访客用户，可能会出现无法正确链接设备的情况。 开始使用 NVIDIA Nsight Systems 提供了多种链接设备的方…

前言 这是对官方文档中对占用率解释的部分的翻译。原文： 点击前往 在原文的基础上增加了自己的一些评述，不一定对，看个乐呵。 概述 在CUDA的Profile和Trace模式下记录的所有CUDA核函数启动中的Occupancy（占用率）实验详情窗格显示“理论占用率”，即由核函数启动配置和CUDA设备的能力强加的占用率上限。在Profile模式的Achieved Occupancy实验中，测量了核函数执行期间的占用率，并将实际值添加到Occupancy实验详情窗格中，与理论值一起显示。额外的图表显示每个SM的实际占用率…

前言 核函数不能像普通函数一样可以在IDE中调试，在运行时发生错误时也不会中断，或者抛出错误信息。 但是CUDA提供了一些函数用于捕获这些错误信息。 错误捕捉函数 //返回运行时的最后一个错误代码，并清除掉这个错误 __host__​__device__​cudaError_t cudaGetLastError ( void ) //返回运行时的最后一个错误代码，但不清除掉这个错误 __host__​__device__​cudaError_t cudaPeekAtLastError ( void ) //通过错误…

前言 统一内存是指只需要一次内存申请，就可以在任意的GPU或者CPU中访问这块内存中的数据，不用在手动的在设备与主机之间拷贝数据。 在了解统一内存之前，需要知道一些相关的概念。 页面错误 页面错误是计算机操作系统中的一种异常情况，它发生在程序试图访问虚拟内存中的某个页面（也称为页）时，而该页当前并未加载到物理内存中。 虚拟内存是一种使用硬盘空间模拟扩展内存的技术。在现代操作系统中，每个程序都有其虚拟地址空间，其中包含多个页面。这些页面通常是固定大小的。 当程序访问某个虚拟地址时，操作系统会检查该地址对应的页面是否已…

CUDA GPU编程指南：内存的申请，释放与拷贝 前言 ​ 在CPU上运行的程序，会将数据存放在内存中，在C++中直接使用new delete 或者malloc free就可以管理内存的申请和释放。 ​ 在GPU上运行的程序，会将数据放入显存中，需要CUDA提供的API来管理内存，并且在显存中存放的数据，是不能被CPU直接访问的。 主机内存（host）和设备内存（device） ​ 主机内存指在CPU上运行的代码所使用的内存，正常C/C++代码中的使用的变量，无论是在堆还是栈上，使用的都是主机内存。 ​ 使用设备内…

前言 这篇博文会介绍如何使用CUDA 在屏幕上打出Hello World！ 直接上代码 #include <stdio.h> #include <cuda_runtime.h> __host__ void cpuSayHello() { printf("CPU: hello world!\n"); } __device__ int getThreadID() { return threadIdx.x; } __global__ void sayHello() { prin…

准备工作 ​ 在使用CUDA前，应该在电脑上安装如下工具 CUDA Toolkit 英伟达为CUDA提供的工具包，下载地址：https://developer.nvidia.com/cuda-toolkit。 CMake 版本在3.11以上。 C/C++编译器，如果是在Windows系统上，推荐使用visual studio 20xx。 添加CUDA环境 ​ 有两种方法可以添加CUDA的支持，如果工程默认使用CUDA，那么可以使用这条命令: project(CUDA_PROJECT LANGUAGES CUDA C…