CUDA编程：基础与实践-NVIDIA Webinar-樊哲勇-CUDA编程：基础中的基础

听AI的话
460 阅读16分钟

NVIDIA Webinar-樊哲勇-CUDA编程：基础中的基础_哔哩哔哩_bilibili

配置

image.png

单文件CUDA程序的基本框架

image.png

内存的分配与释放

image.png

主机与设备之间的数据传输

image.png

核函数中线程与数据的对应

image.png

核函数的要求

image.png

自定义设备函数 device functions

image.png

检测CUDA Runtime API函数的返回值

image.png

检测核函数调用的错误

image.png

第5章 获得GPU加速的关键

关键1 数据传输比例较小

image.png

关键2 核函数的算数强度 arithmetic intensity 较高

如果数据较多，就关注核函数算数强度 image.png

关键3 并行规模较大

总的线程数要比较大才行，什么意思？

第6章 CUDA的内存组织

image.png

第7章 全局内存

全局内存的合并访问

CUDA中各种内存的用途总结

似乎共享内存效果不大越来越

image.png

第9章 原子函数的合理使用

利用原子函数实现全部归约

image.png

第10章 线程束基本函数与协作组

image.png 线程束内的洗牌函数 warp shuffle functions 协作组 cooperative groups

数组归约计算的性能测试

image.png

问题

gpu的线程块里的线程个数是固定的吗？线程块是硬件决定的还是软件通过<<<2,3>>>任意指定？

找到了第 1 篇资料：GPU体系结构(1)线程的执行 - 知乎

找到了第 2 篇资料：CUDA为什么要分线程块和线程网格？ - 知乎

找到了第 3 篇资料：有人能说一下GPU并行计算和CPU多线程计算有什么区别 ...

找到了第 4 篇资料：理解CUDA中的thread,block,grid和warp - 知乎

找到了第 5 篇资料：SM详解与Warp Scheduler，合理块和线程的数量对GPU利用 ...

找到了第 6 篇资料：如何理解 NVIDIA CUDA 中线程块的三维线程的概念？ - 知乎

一共找到 6 篇资料，正在阅读...

找到了 6 篇资料作为参考：

GPU的线程块（block）中的线程个数不是固定的，它可以通过软件编程来指定。在CUDA编程模型中，线程块的大小（即线程数）是由程序员在调用内核函数（kernel function）时通过<<<gridDim, blockDim>>>这样的语法来指定的。这里的blockDim是一个三维向量，可以指定每个维度上的线程数。

例如，如果你想要创建一个二维的线程块布局，你可以这样指定：

kernel<<<gridDim, blockDim>>>() {
    // 内核代码
}

在这个例子中，gridDim指定了整个网格（grid）的线程块数量，而blockDim指定了每个线程块内的线程数量。线程块的总线程数是blockDim.x * blockDim.y * blockDim.z的乘积。

线程块的大小需要根据具体的GPU架构和硬件特性来合理设置。如果线程块设置得过大或过小，都可能影响GPU的利用率。过大的线程块可能会导致资源浪费，而过小的线程块可能会导致线程调度开销增加。通常，线程块的大小会根据GPU的SM（Streaming Multiprocessor）结构和Warp Scheduler来优化，以确保高效的并行执行。

在Fermi架构的GPU中，每个SM可以同时执行多个线程块，但每个SM上最多可以同时执行8个线程块，且每个SM上最多可以同时执行48个warp（一个warp通常包含32个线程）。这些硬件特性可以帮助程序员在设计线程块大小时做出更合理的决策。

为什么CUDA核函数printf()能打印到终端，不是在GPU执行吗？

找到了第 1 篇资料：cuda printf调试？ - 知乎

找到了第 2 篇资料：自己动手，实现在kernel函数中printf()!-CSDN社区

找到了第 3 篇资料：自己动手，实现在kernel函数中printf()!（转）_内核printf实现 ...

找到了第 4 篇资料：CUDA内核printf()在终端中不产生输出，在分析器中工作-腾讯 ...

找到了第 5 篇资料：printf in CUDA kernel 函数_cuda kernel内部打印执行的数据 ...

找到了第 6 篇资料：用CUDA实现简单的打印操作_cuda printf-CSDN博客

一共找到 6 篇资料，正在阅读...

找到了 6 篇资料作为参考：

CUDA核函数中的printf()实际上并不是直接在GPU上执行的，而是通过特定的机制将打印信息发送到主机（CPU）的控制台。这是因为GPU本身没有直接访问主机标准输出流（如终端）的能力。以下是这个过程的大致步骤：

  1. 缓冲区分配：在主机端（CPU）上，首先会分配一个内存缓冲区，这个缓冲区将用于存储从GPU发送过来的打印信息。
  2. 初始化：在核函数执行之前，需要在主机端调用一个初始化函数（如device_printf_init()），这个函数会设置一个回调函数，当GPU的打印信息准备好后，这个回调函数会被触发。
  3. 核函数执行：在核函数中，你可以像在C语言中一样使用printf()函数。但是，这个printf()实际上是一个特殊的CUDA版本，它会将打印信息写入之前分配的内存缓冲区。
  4. 同步和输出：核函数执行完成后，主机端需要调用一个同步函数（如cudaDeviceSynchronize()），这会确保所有GPU操作都已完成。然后，主机端的回调函数会被触发，它会从内存缓冲区读取信息并将其输出到控制台。
  5. 清理：在所有打印操作完成后，需要释放之前分配的内存缓冲区。

这个过程实际上是一个异步操作，GPU在执行核函数时，打印信息被缓存在内存中，然后由主机端的回调函数处理并输出。这就是为什么你能够在终端看到从CUDA核函数中打印的信息。这种机制允许开发者在不直接访问GPU内部状态的情况下，查看和调试GPU上的执行情况。

显存带宽怎么通过实验测试？

找到了第 1 篇资料：超能课堂(2)：显卡带宽应该怎么算？ - 超能网

找到了第 2 篇资料：GPU 显存带宽 - 知乎

找到了第 3 篇资料：CUDA 修炼笔记(五) -- 性能测量 - 知乎

找到了第 4 篇资料：CUDA工具包中的 bandwidthTest 是什么？ - 知乎

找到了第 5 篇资料：GPU 的性能该如何测试？ - 知乎

一共找到 5 篇资料，正在阅读...

找到了 5 篇资料作为参考：

显存带宽可以通过实验测试来测量，这通常涉及到编写特定的程序来执行大量的内存读写操作，并记录所需的时间。以下是一些常见的方法来测试显存带宽：

  1. 使用CUDA工具包中的示例程序： NVIDIA的CUDA工具包提供了一个名为bandwidthTest的示例程序，它可以测量不同类型的内存传输带宽，包括从CPU到GPU（Host to Device）、从GPU到CPU（Device to Host）以及在GPU内部不同内存区域之间的数据传输（Device to Device）。这个工具可以帮助开发者了解和比较不同GPU硬件和系统配置之间的内存带宽性能。
  2. 编写自定义测试程序： 开发者可以编写自己的CUDA程序，通过在GPU上执行大量的内存读写操作来测试显存带宽。这通常涉及到创建一个大型的内存块，然后进行多次读写操作，并使用CUDA事件（events）来测量操作的开始和结束时间。通过计算这些时间差，可以得到实际的带宽值。
  3. 使用第三方性能测试工具： 有一些第三方工具和库，如GPGPU-Sim、GPGPU-Bench等，专门用于测试和评估GPU的性能，包括显存带宽。这些工具提供了一套标准化的测试流程，可以帮助用户快速获取显存带宽的测量结果。
  4. 利用GPU监控工具： 一些GPU监控工具，如NVIDIA的NSight、AMD的Radeon Software Adrenalin Edition等，可以在运行应用程序时实时监控GPU的性能，包括显存带宽。这些工具通常提供了图形界面，使得性能监控更加直观。

在进行显存带宽测试时，重要的是要确保测试条件尽可能接近实际应用场景，以便获得有意义的性能数据。例如，测试时应该考虑内存访问模式（如连续访问、随机访问）、内存对齐、内存碎片等因素，因为这些都可能影响实际的带宽表现。

核函数是保存在显存上吗？

找到了第 1 篇资料：cuda 如何将二维数组传入显存? - 知乎

找到了第 2 篇资料：GPU编程2--CUDA核函数和线程配置 - 知乎

找到了第 3 篇资料：CUDA（二）：GPU的内存体系及其优化指南 - 知乎

找到了第 4 篇资料：CUDA 编程(一)：CUDA C 编程及 GPU 基本知识 - 知乎

找到了第 5 篇资料：深度学习GPU显存为什么很容易就满了？ - 知乎

一共找到 5 篇资料，正在阅读...

找到了 5 篇资料作为参考：

是的，CUDA核函数（Kernel Function）在执行时会被加载到GPU的显存（VRAM）上。核函数是GPU上并行计算的基本单元，它们被设计为在GPU的多个处理单元（CUDA核心）上并行执行。当您在主机（CPU）上调用一个核函数时，CUDA运行时环境会负责将核函数的代码、必要的数据以及执行配置（如线程块和网格的配置）传输到GPU。

在核函数执行过程中，每个线程都会在GPU的寄存器中保存其线程ID和其他必要的状态信息，以便在执行时能够访问和操作数据。这些线程ID和状态信息通常由CUDA运行时环境自动管理，并且对于开发者来说是透明的。

核函数的执行涉及到的数据（如输入参数、输出结果等）也需要被传输到显存中，以便GPU可以访问和处理这些数据。在核函数执行完成后，结果数据通常需要从显存传输回主机内存，以便CPU可以进一步处理或显示。

总的来说，核函数及其所需的数据在GPU上执行时，都会占用显存空间。这也是为什么在设计CUDA程序时，合理管理显存使用对于优化性能和避免显存溢出非常重要。

GPU全局内存 和 共享内存什么关系？局部内存和共享内存什么关系？

找到了第 1 篇资料：关于共享内存shm和内存映射mmap的区别是什么？ - 知乎

找到了第 2 篇资料：理解GPU中的各种内存类型，为什么SRAM和L1Cache被 ...

找到了第 3 篇资料：GPU共享内存，这里的共享内存指的是共享运行内存条的内存 ...

找到了第 4 篇资料：CUDA（二）：GPU的内存体系及其优化指南 - 知乎

找到了第 5 篇资料：进程间的通信方式（一）：共享内存 - 知乎

找到了第 6 篇资料：共享内存与线程局部存储 - 小张慢慢成了老张 - 博客园

找到了第 7 篇资料：gpu cuda中各种内存的优缺点及调用方法 - 知乎

找到了第 8 篇资料：CUDA（二）：GPU的内存体系及其优化指南 - 知乎

一共找到 8 篇资料，正在阅读...

找到了 8 篇资料作为参考：

在CUDA编程中，全局内存（Global Memory）和共享内存（Shared Memory）是两种不同类型的内存空间，它们在GPU上扮演着不同的角色，并且与局部内存（Local Memory）有着不同的关系。

  1. 全局内存与共享内存的关系：

    • 全局内存是GPU上最大的内存空间，可以被所有线程块（Block）上的线程访问。它是在GPU上分配的，用于存储程序中所有线程共享的数据。全局内存的访问速度相对较慢，因为所有线程都可能访问它，这可能导致内存访问冲突和瓶颈。
    • 共享内存是位于每个线程块内部的内存空间，它对整个线程块内的线程可见。共享内存的访问速度比全局内存快得多，因为它是为单个线程块设计的，减少了内存访问冲突。共享内存的大小有限，通常比全局内存小得多，但它提供了一种快速的数据共享机制，适用于线程块内部的线程间通信。

  2. 局部内存与共享内存的关系：

    • 局部内存（Local Memory）通常指的是线程局部存储（Thread Local Storage），它为每个线程提供了一个私有的内存空间。在CUDA中，局部内存通常指的是线程栈（Thread Stack），它用于存储线程的局部变量和函数调用的参数。局部内存的访问速度非常快，因为它完全是线程私有的，没有并发访问的问题。
    • 共享内存与局部内存的主要区别在于共享性。共享内存允许线程块内的线程共享数据，而局部内存则是每个线程独有的。在某些情况下，局部内存可能会被用作共享内存的辅助，例如，当需要在线程块内部进行数据重组时，可以先将数据从全局内存加载到局部内存，然后在局部内存中进行操作，最后再写回全局内存。

在CUDA编程中，合理地使用全局内存、共享内存和局部内存对于优化程序性能至关重要。开发者需要根据数据访问模式和并行计算的需求来决定如何分配和使用这些内存资源。

add<<<128,32>>>(x,y,z); 的共享内存是32个线程共用一个共享内存？共享内存在硬件的哪个位置？是全局内存的一部分吗？还是靠近核心

找到了第 1 篇资料：【CUDA 基础】5.1 CUDA共享内存概述 - 知乎

找到了第 2 篇资料：GPU加速03:多流和共享内存—让你的CUDA程序如虎添翼的 ...

找到了第 3 篇资料：在 CUDA C / C ++ 中使用共享内存 - 知乎

找到了第 4 篇资料：CUDA 编程入门（6）：共享内存 - 知乎

找到了第 5 篇资料：CUDA 编程入门（6）：共享内存 - 知乎

找到了第 6 篇资料：【CUDA 基础】5.1 CUDA共享内存概述 - 知乎

找到了第 7 篇资料：CUDA共享内存操作(__shared__关键字) - CSDN博客

一共找到 7 篇资料，正在阅读...

找到了 7 篇资料作为参考：

在CUDA中，共享内存（Shared Memory）是位于每个多处理器（multiprocessor）内部的片上内存（on-chip memory），它不是全局内存的一部分，而是专门为一个线程块（thread block）内的线程设计的。共享内存位于每个多处理器的L1缓存中，这意味着它比全局内存（Global Memory）访问速度快得多，但容量有限。

当你使用<<<128, 32>>>这样的启动配置时，你创建了一个128个线程块的网格，每个线程块包含32个线程。在这种情况下，每个线程块内的32个线程会共享一个单独的共享内存空间。这个共享内存空间是为这个线程块动态分配的，当线程块执行完毕后，这个共享内存空间会被释放。

共享内存的生命周期与线程块的生命周期相同。当线程块开始执行时，共享内存被分配；当线程块执行完毕，共享内存被释放。这意味着，共享内存是线程块内的线程之间进行通信和协作的机制，它允许线程之间快速交换数据，而不需要通过全局内存，从而减少了内存访问延迟。

在硬件层面，共享内存位于GPU的核心附近，这使得它比全局内存更接近执行计算的线程，从而提供了更低的访问延迟。这种设计是为了优化性能，因为共享内存的访问速度比全局内存快，可以显著提高并行计算的效率。

A100有多少个 SM？每个有多少个核心？

找到了第 1 篇资料：NVIDIA A100 | Tensor Core GPU

找到了第 2 篇资料：NVIDIA A100 | Tensor Core GPU

找到了第 3 篇资料：Nvidia Ampere 架构深度解析 - 知乎

找到了第 4 篇资料：NVIDIA A100 深度解密（一）：GPU 峰值计算那些事 - 知乎

找到了第 5 篇资料：NVIDIA Ampere架构之SM算力及多线程原理 - 知乎

找到了第 6 篇资料：NVIDIA GPU A100 Ampere(安培) 架构深度解析_a100 架构 ...

一共找到 6 篇资料，正在阅读...

找到了 6 篇资料作为参考：

NVIDIA A100 GPU 拥有108个流式多处理器（Streaming Multiprocessors，简称SM）。每个SM内部包含有多个计算单元，具体来说，A100的每个SM有4个Tensor Core，这些Tensor Core可以执行混合精度（FP16/FP32）的融合乘加（Fused Multiply-Add，FMA）操作。在A100中，每个Tensor Core每个时钟周期可以执行256个FP16/FP32 FMA操作，这意味着每个SM在每个时钟周期内总共可以执行1024个FMA操作。

此外，A100 GPU的每个SM还包含64个FP32 CUDA Cores，这些CUDA Cores是执行基本数学计算的单元。这些CUDA Cores通过共享内存（Shared Memory）和L1数据缓存进行通信和数据交换。A100 GPU的设计旨在提供高效的并行计算能力，特别是在AI、数据分析和高性能计算（HPC）等领域。

eatcosmos的学习方法小结

  1. 重点关注程序执行，不要试图等理解了再写程序，学习就类似在黑屋里走迷宫，你没有全局灯当然无法知道迷宫原理，但是编程本身就是解释，编程语言就是语文。不要把程序理解为结果，程序即理解。
  2. 很多技巧是人脑的直觉，不可否认有很多逻辑细节，但是不需要试图细化，只要觉得合理就行了，重点就是你知道应该怎么让程序执行出效果。
  3. 不会的问chatgpt然后记录下来，不要中途沮丧做其他事情。
  4. 总是碰到不懂的焦虑，你要定时活动，这个问题解决不了你就久坐这肯定是不行的。
  5. 实在糊涂的就先跳过，因为这不是你的问题。
  6. 不要换来换去，你就一条道走到黑。
avatar
文章
阅读
粉丝