CUDA编程学习 (1)——CUDA C介绍

news/2025/3/17 4:15:07/文章来源:https://www.cnblogs.com/Astron-fjh/p/18503824

1. 内存分配和数据移动 API 函数

CUDA编程模型是一个异构模型,需要CPU和GPU协同工作。在CUDA中,hostdevic e 是两个重要的概念,我们用host指代CPU及其内存,而用device指代GPU及其内存。

image-20240921213601913

CUDA程序中既包含host程序,又包含device程序,它们分别在CPU和GPU上运行。同时,host与device之间可以进行通信,这样它们之间可以进行数据拷贝。典型的CUDA程序的执行流程如下:

  1. 分配host内存,并进行数据初始化;
  2. 分配device内存,并从host将数据拷贝到device上;
  3. 调用CUDA的核函数在device上完成指定的运算;
  4. 将device上的运算结果拷贝到host上;
  5. 释放device和host上分配的内存。
#include <cuda.h>
void vecAdd(float *h_A, float *h_B, float *h_C, int n){int size = n* sizeof(float);float *d_A, *d_B, *d_C;// Part 1// 为 A,B,C 分配 device 内存// copy A 和 B 到 device 内存// Part 2// 内核启动代码 - 设备执行实际的矢量加法// Part 3// 从 device 内存 copy C
}
image-20240921221415897

Device 代码可以:

  • 写/读 per-thread register
  • 写/读 all-shared global memory

Host 代码可以:

  • 将数据传输到每个 grid global memory / 从每个 grid global memory 传输数据

1.1 Device Memory 管理 API

image-20240921223550002

cudaMalloc()

  • 在 device global memory 中分配一个对象

  • 两个参数

    • 分配对象指针的地址
    • 分配对象的大小(以字节为单位)

cudaFree()

  • 从 device global memory 中释放对象
  • 一个参数
    • 释放对象的指针

1.2 Host-Device 数据传输 API

image-20240921225324676

cudaMemCpy()

  • 内存数据传输

  • 4个参数

    • 目的地指针
    • 源指针
    • 复制的字节数
    • 传输类型/方向
  • 与 host 同步向 device 传输数据

1.3 矢量加法-显式内存管理

image-20240921230232238

1.4 统一内存(Unified Memory)

image-20240921231200588

cudaMallocManaged(void** ptr, size_t size)

  • 统一内存是 CPU 和 GPU 共享的相同内存空间。这意味着,CPU 和 GPU 使用相同的地址空间,无需像传统的 CUDA 编程那样手动管理内存传输(如cudaMemcpy())。

    • 它维持数据的单一副本,即不再需要独立的主机内存和设备内存副本。
  • CUDA-managed 数据

    • 其中 ptr 是指向分配内存的指针,size 是所需内存的大小。
    • 按需页面迁移:在GPU执行代码时,所需的数据会按需从主机(Host)迁移到设备(Device)。
  • cudaMalloc()cudaFree() 兼容

  • 可优化

    • cudaMemAdvise(): 给CUDA提供建议,提示内存的使用模式(例如,数据主要在哪个处理器上使用)。
    • cudaMemPrefetchAsync(): 允许程序员预取数据到设备或主机,减少内存访问延迟。
    • cudaMemcpyAsync(): 在异步传输数据时,可以优化性能。

1.5 矢量加法-统一内存

image-20240921234243047

1.6 检查 host 代码中的 API 错误

cudaError_t err = cudaMalloc((void **) &d_A, size);
if(err != cudaSuccess){printf("%s in %s at line %d\n", cudaGetErrorString(err), __FILE__, __LINE__);exit(EXIT_FAILURE);
}

2. 线程和内核函数

2.1 Kernel 线程层次结构

首先GPU上很多并行化的轻量级线程,kernel 在 device 上执行时实际上是启动很多线程。

一个 kernel 所启动的所有线程称为一个网格(grid),同一个网格上的线程共享相同的 global memory,grid 是线程结构的第一层次,而网格又可以分为很多 线程块(block),一个线程块里面包含很多线程,这是第二个层次。

img

2.2 数据并行 - 向量加法示例

image-20240922003424983

2.3 CUDA 执行模型

  • 异构 host(CPU)+ device(GPU)应用 C 程序
    • host C 代码中的串行部件
    • device SPMD 内核代码中的并行部分
image-20240922003839725

2.4 并行 thread 阵列

  • CUDA kernel 由 threads 组成的一个 grid(array)执行
    • grid 中的所有 thread 运行相同的 kernel 代码(单程序多数据)
    • 每个 thread 都有索引,用于计算内存地址和做出控制决定
image-20240922124715191

2.4.1 线程块:可扩展的协作

image-20240922125350352
  • 将 thread array 划分为多个 block
    • block 内的 thread 通过 共享内存(shared memory)原子操作(atomic operations)障碍同步(barrier synchronization)进行协作
    • 不同 block 中的 thread 不会相互影响

2.4.2 block 索引和 thread 索引

image-20240922125856835
  • 每个 thread 使用索引来决定处理哪些数据

    • blockIdx:1D, 2D or 3D (CUDA 4.0)
    • threadIdx:1D, 2D or 3D
  • 处理多维数据时简化内存寻址

    • 图像处理
    • 解决体积上的 PDEs

对于一个 2-dim 的 $block(D_x,D_y)$,线程 $(x,y)$ 的 ID 值为 $(x+y∗D_x)$ ,如果是 3-dim 的 $block(D_x,D_y,D_z)$,线程 $(x,y,z)$ 的ID值为 $(x+y∗D_x+z∗D_x∗D_y)$ 。另外线程还有内置变量 gridDim,用于获得网格块各个维度的大小。

3. CUDA 工具包简介

3.1 NVCC Compiler

  • 英伟达™(NVIDIA®)提供 CUDA-C 编译器,简单的编译命令为:
nvcc vectorAdd.cu -o vectorAdd

这将使用默认参数编译 vectorAdd.cu,并输出一个叫 vectorAdd 的可执行文件。

  • nvcc 编译设备代码,然后将代码转发给 host 编译器(如 g++)。

  • 可用于编译和链接 host 应用程序

两种源文件:

对于 CUDA 程序,我们通常有:

  1. .cu 文件:包含 CUDA C/C++ 代码,包括主程序代码和 CUDA kernels,需要 nvcc 编译。
  2. .cpp/.c 文件:只包含主程序代码,不含 CUDA kernels,可以用 gcc/g++ 编译。

编译 CUDA 程序的一般步骤是:

  1. nvcc 编译 .cu 文件,生成 .o 文件。
  2. g++ 编译 .cpp 文件,也生成 .o 文件。
  3. 再用 g++/nvcc 链接所有 .o 文件生成最终可执行文件。
nvcc -c demo.cu -o demo.o
g++ -c main.cpp -o main.o
nvcc demo.o main.o -o demo -lcudart

3.2 调试工具(Debugger)

image-20240922172325240
  • NVIDIA 提供的调试工具
    • Nsight:这是一个用于 GPU 编程调试和优化的工具,专门设计用来分析 GPU 代码的性能问题,并进行代码调试。支持对 CUDA 和 OpenGL 等代码进行深入的分析。
    • Nsight Systems:一个面向系统级性能分析的工具,可以帮助开发者查找瓶颈,了解程序在 CPU 和 GPU 之间的负载分布。
    • CUDA-GDB:这是一个基于 GDB 的调试器,专门用于 CUDA 编程的调试工作,支持查看 GPU 上的内存状态、线程执行情况等。
    • CUDA MEMCHECK:用于检测 CUDA 代码中的内存错误,包括非法内存访问、未初始化内存使用等,类似于 CPU 上的 valgrind 工具。
  • 第三方调试工具
    • ARM Forge:支持跨平台的高性能计算(HPC)调试工具,适用于调试大规模并行程序,包括CUDA应用。
    • TotalView:另一个广泛使用的高性能并行调试器,适用于调试复杂的CUDA程序,提供对多核、多线程代码的深入支持。

3.3 性能分析工具(Profilers)

image-20240922173202824
  • NVIDIA 提供的性能分析工具

    • NSIGHT:与调试工具一样,Nsight 还可用于性能分析,帮助开发者分析 GPU 代码的瓶颈、内存使用、指令执行效率等。它支持 CUDA、OpenGL 和 Vulkan 等多种 API。
    • NVVP(NVIDIA Visual Profiler):专门为 CUDA 程序设计的图形化性能分析工具,提供可视化的性能瓶颈分析,帮助开发者定位代码中导致性能下降的具体部分。
    • NVPROF:是一个命令行性能分析工具,专注于 CUDA 程序的性能数据收集。它可以通过收集时间信息、内存传输、计算效率等指标,提供详细的性能分析结果。
  • 第三方性能分析工具

    • TAU:一款适用于并行程序的全面性能分析工具,支持多种硬件架构和 API,帮助用户从全局上评估程序的性能表现。
    • VampirTrace:一个可视化的性能分析工具,用于追踪并行程序的执行情况,帮助开发者分析程序的时间线、通信延迟等。

4. Nsight 计算和 Nsight 系统

image-20240922174403339
  • Phasing Out(逐步淘汰的工具)

    • NVPROF
    • NVVP(NVIDIA Visual Profiler)
  • Current(当前主流的工具)

    • Nsight Compute
    • Nsight Systems

参考文献

[1] CUDA编程入门极简教程 - 知乎 (zhihu.com)

[2] nvcc 编译 .cu 源代码的基本用法_多个文件 nvcc怎么编译-CSDN博客

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处:http://www.hqwc.cn/news/822279.html

如若内容造成侵权/违法违规/事实不符,请联系编程知识网进行投诉反馈email:809451989@qq.com,一经查实,立即删除!

相关文章

博图SCL均值计算

这一篇学习笔记在新浪博客记录过,这里再记录一次。 工作中有时候会需要做一些均值计算,比如计算某个测量值近一分钟的均值,近一小时均值,近一天的均值。今天在家休息,试着做一下分钟均值,按照每秒一个数据,比如现在时刻10:07:10,那么计算从10:06:11到10:07:10这个时间段…

CUDA编程学习 (2)——CUDA并行性模型

这篇文章深入探讨了CUDA编程中的并行性模型,重点介绍了基于kernel的单指令多数据(SPMD)编程。首先,通过向量加法示例,展示了CUDA内核函数的编写与调用,并解释了__global__、__device__和__host__的不同用法。接着,文章扩展到多维网格配置,演示如何处理图像数据,如RGB转灰…

SLF4J 中的适配器模式

什么是适配器模式 适配器模式中,适配器包装不兼容指定接口的对象,来实现不同兼容指定接口。 SLF4J 中的适配器模式 SLF4J 是一个日志门面系统,其中提供了统一的 Logger 等接口,许多框架都会面向 SLF4J 打印日志,这样就不会和具体的日志框架耦合在一起,框架使用者也就能够…

Pbootcms留言“提交成功”的提示语怎么修改

要在 PbootCMS 中修改留言“提交成功”的提示语,可以按照以下步骤操作:定位文件:打开 apps/home/controller/MessageController.php 文件。查找代码段:在文件中找到大约第 103 行的代码段,该段代码如下:if ($this->model->addMessage($data)) {session(lastsub, ti…

Protues中51单片机按键无法复位(已解决)

前言 昨晚用 Protues 搭建了 51 的最小系统电路,在实物中好用的复位电路,到仿真里不能正常复位了。 51 单片机是高电平复位,所以在运行时 RST 引脚应该是低电平,但在仿真中 RST 引脚一直保持高电平,导致按下按键也不能复位单片机。解决方法 我在网上搜索的解决方法一共有两…

南昌航空大学-22207316-涂高杰-JAVA第一次blog作业

一.前言 本学期新增JAVA的面向程序设计课程,为增加学生编写能力开始了本学期的PTA作业,以及接下来我将根据我的实际情况总结前三次PTA题目集中最后一题并讲诉自己对Java的学习心得。从这三次PTA作业中学习到的了对ArrayLis、Vector等自动增长的数组的使用方法,学习到了许多Ja…

安装网站出现404 not found如何解决?

遇到404 Not Found错误时,可以尝试以下几个步骤来解决问题:检查URL:确认输入的网址是否正确,包括大小写和拼写。 检查是否有遗漏或多余的字符。清除浏览器缓存:有时候旧的缓存数据会导致页面加载错误,清除缓存后重新尝试访问。刷新页面:使用F5键或点击浏览器的刷新按钮重…

pbootcms模板上线推广百度竞价后打不开网站出现404错误

PbootCMS V3.2.5 版本中为了增强安全性或优化URL结构,加入了对URL参数的严格判断。当URL中包含?但不符合特定条件(如/?tag=、/?page=、/?ext_)时,系统会自动返回404错误页面。这种做法虽然有助于防止一些非法请求,但也可能导致合法的请求被误判为无效,特别是对于那些…

分布式数据库TDSQL搭建

TDSQL介绍TDSQL是腾讯基于MySQL/Mariadb社区版本打造的一款金融级分布式数据库集群方案,目前腾讯主推TDSQL MySQL版。TDSQL MySQL版具备强一致高可用、全球部署架构、分布式水平扩展、高性能、企业级安全等特性,同时提供智能 DBA、自动化运营、监控告警等配套设施,为客户提供…

BUUCTF_BUU SQL COURSE 1

BUUCTF_BUU SQL COURSE 1 打开实例发现“热点”及“登录”两个选项 根据题目提示,sql注入,尝试在登录界面寻找注入点,无果 接着进入热点界面,发现三篇新闻,依次点击发现url变化,burp抓包发现是通过对id值的控制访问不同界面,由此发现注入点 1.判断注入类型,输入1,2-1 ?…