cuda二进制文件中到底有些什么

大家好。今天我们来讨论一下,相比gcc编译器编译的二进制elf文件,包含有 cuda kernel 的源文件编译出来的 elf 文件有什么不同呢?

之前研究过一点 tvm。从 BYOC 的框架中可以得知,前端将模型 partition 成 host 和 accel(accel 表示后端,比如加速卡,NPU或者其他AI加速模块) 两部分,对 accel 部分,会切分成多个 regions,对应到多个子图,这部分每一个 regions 会被封装成一个独立的 function 进行处理,这些 function 都带有 annotation,附带硬件相关的标签信息,能知道由哪个 accel 后端来处理,在 host 侧对这些函数的处理,仅仅是简单的封装成对一个外部函数的调用,而实际的编译是由 accel-specific 的编译器来编译和 codegen 的。而每一个 sub-graph 会编译成一个 sub-module,最终与 host sub-module 一起封装成一个 heterogenous blob。

而 nvcc 应该也是类似的道理。

在 《Cuda Compiler Driver NVCC.pdf》中,有这么一段介绍

Dispatching GPU jobs by the host process is supported by the CUDA Toolkit in the form of remote procedure calling. The GPU code is implemented as a collection of functions in a language that is essentially C, but with some annotations for distinguishing them from the host code, plus annotations for distinguishing different types of data memory that exists on the GPU. Such functions may have parameters, and they can be called using a syntax that is very similar to regular C function calling, but slightly extended for being able to specify the matrix of GPU threads that must execute the called function. During its life time, the host process may dispatch many parallel GPU tasks.

大致意思就是说,CUDA ToolKit 可以支持主机端程序通过 RPC 的方式调度 GPU 任务。GPU 代码与 C 代码类似,但是带有一些额外的 annotations 信息来做区分。而GPU 代码实现的函数的调用也与传统 c 函数调用相同。

直接来编译一个 helloword 程序

/*
*hello_world.cu
*/
#include<stdio.h>
__global__ void hello_world(void)
{
  printf("GPU: Hello world!\n");
}
int main(int argc,char **argv)
{
  printf("CPU: Hello world!\n");
  hello_world<<<1,10>>>();
  cudaDeviceReset();//if no this line ,it can not output hello world from gpu
  return 0;
}

编译

nvcc --cudart shared -o device helloworld.cu --verbose

使用 --cudart shared 而不使用静态链接的方式,是为了不将 libcudart.a 链接到二进制文件中,使得目标程序大小偏大。

objdump -ds device

观察 hello_world 函数
请添加图片描述
可以看到,本质上是一个函数调用,对 _Z30__device_stub__Z11hello_worldvv 函数的一个调用。

推测: 对 device 设备端的函数,也是封装成一个 external function 的函数调用,而该函数实际是通过 device设备端(也就是GPU) 的code gen 来生成的,最终会将其合并成一个二进制文件。而在编写 cuda 代码的过程中,所使用的这些 c++ 扩展,就类似于 annotation 的作用,注明了这属于 device 设备端的代码。

compile process

请添加图片描述

a CUDA executable can exist in two forms:

  • a binary one that can only target specific devices and an intermediate assembly one that can target any device by JIT compilation.
  • a PTX Assembler (ptxas) performs the compilation during execution time, adding a start-up overhead, at least during the first invocation of a kernel.

cuda 可执行文件可以有两种形式,一种是针对特定设备的二进制文件和一种中间表示的汇编形式,可以通过 JIT 的方式运行与任何设备上。JIT 也就是 Just In Time,java 虚拟机,python,v8 都有 JIT 机制。而另一种,就是 PTX 汇编,这种形式是通过 cuda runtime 在运行时加载编译然后执行的,第一次加载编译时会比较耗时。

A CUDA program can still target different devices by embedding multiple cubins into a single file (called a fat binary ). The appropriate cubin is selected at run-time.

请添加图片描述

从上面可以发现,使用nvcc进行编译,将包含有cuda kernel 的 c++ 代码,分成了 device 的代码和 host 的代码,host 代码通过 clang/gcc 以传统 c++ 代码的方式进行编译,而 device 代码以 nvcc cuda 编译的流程进行编译。我们使用 --verbose 的方式来观察一下具体的编译流程。

$ nvcc --cudart shared -o device helloworld.cu --verbose ---keep

请添加图片描述
helloworld.cu 编译后生成了 hellworld.cpp1.ii 和 helloworld.ptx,ptx 也就是 cuda 汇编代码。然后 helloworld.ptx 编译成了 cubin 二进制文件,而 fatbinary 最终会被嵌入到最终的 elf 二进制文件 devicde 中。

elf 二进制文件分析

使用 readelf 观察一下 device 这个文件

readelf -a device

请添加图片描述

在该 elf 文件中,多了两个段 .nv_fatbin 和 .nvFatBinSegment
请添加图片描述
从 Program Headers 中可以发现,这两个段分别位于代码段和数据段中。

第一个 LOAD 属性为 RE,表示可读可执行表示代码段,而第二个 LOAD 属性为 RW,表示可读可写,为数据段。从上往下索引分别是 02 和 03,所以 .nv_fatbin 位于代码段,而 .nvFatBinSegment 位于数据段中。

.nv_fatbin

It is split into an arbitrary number of distinct regions, each of which contains one or more GPU ELF files, PTX code files, and/or cubin files .

该段中保存的通常是 PTX 汇编代码或者 cubin 二进制代码,正好与上面的分析相符,位于代码段中。

.nvFatBinSegment

It contains metadata about the .nv_fatbin section , such as the starting addresses of its regions. Its size is a multiple of six words (24 bytes), where the third word in each group of six is an address inside of the .nv_fatbin section. If we modify the .nv_fatbin, then these addresses need to be changed to match it.

该段保存的是 .nv_fatbin 的一些 metadata。

文件分析

先来看下 device 文件头的信息,可以通过 readelf -h 的方式查看
请添加图片描述
文件头是 64 个字节,program headers 是 56 个字节,共有 9 个 program headers,每一个 section header 是 64 字节。

先看下 elf.h 中 Elf64_Ehdr 文件的数据结构
请添加图片描述

e_ident 就是上面 readelf -h 结果中的 Magic,也就是 elf 格式的魔数。

文件头大小是 64 字节,使用 od 来分析一下。

od -Ax -tx1 -N 64 deviceß

解释一下这里的参数

  • -Ax: 显示地址的时候,用十六进制来表示。如果使用 -Ad,意思就是用十进制来显示地址;
  • -t -x1: 显示字节码内容的时候,使用十六进制(x),每次显示一个字节(1);
  • -N 64:只需要读取 64 个字节;

请添加图片描述
e_type 为 0x0003(小端),e_type 的取值可以在 elf.h 中查看
在这里插入图片描述
3 表示该文件是 shared object file。
而 e_machine 为 0x003e,
在这里插入图片描述
从 elf.h 中可以看出,0x3e 的 10 进制为 62,也就是 ADM x86_64架构。而 cuda 的 e_machine 应该是 190,也就是 0xbe。
请添加图片描述

也就是说如果是 cuda bin,Elf64_Ehdr 中 e_machine 成员的值,应该是 190,16 进制就是 0xbe。

我们在最上面分析 device 文件时,使用 readelf -a 查看,发现 .nv_fatbin 在 Section Header 中的索引是 17,section header 的起始偏移是 0x42f8 = 17144,从 elf 文件中获取 nv_fatbin 这个 section 的信息,计算偏移为 0x42f8 + 17 * 64,64 就是 e_shentsize 的大小,为 0x40,即每一个 section header item 的大小为 64 字节

请添加图片描述
而 section header 的数据结构为
在这里插入图片描述
该 section 的大小和偏移与在 readelf -S 中看到的一致,看下内容 .nvfatbin 段的内容
请添加图片描述

nvcc 编译时,–keep 将临时文件保存下来,device_dlink.fatbin 与上面的内容一致
请添加图片描述
fatbin 是 device-only 的代码

上面这个 fatbin 文件其实是一个包裹着 elf 文件的二进制文件,
请添加图片描述

文件 e_machine 为 0xbe,就是 cuda elf 格式的文件。

总结

cuda 二进制文件,分成两部分,一个是 host 部分的代码,一个是 device 段的代码。device 段的代码,作为一个 section 的方式,以 fatbin 的方式或者 ptx 汇编代码的方式嵌入到了最终的 elf 文件中。这部分代码,有 cuda runtime 来负责编译运行。

reference

  1. PCI BARs and other means of accessing the GPU
  2. https://www.ofweek.com/ai/2021-05/ART-201721-11000-30500304_3.html

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

如若内容造成侵权/违法违规/事实不符,请联系我们进行投诉反馈qq邮箱809451989@qq.com,一经查实,立即删除!

相关文章

日期处理第四篇(终)- Java日期时间处理大总结

文章目录 日期时间概念通用标准日期字段解析国际化的日期格式 日期的实战第一个问题&#xff1a;日期常用时间操作第二个问题&#xff1a;时区的问题时区概念时区的处理ZoneID的使用 ZoneOffset的使用让人恼火的夏令时 第三个问题&#xff1a;MySQL存储时间用什么类型&#xff…

HCIA——23DNS层次域名空间、域名服务器、域名解析的原理的选择、解答

学习目标&#xff1a; 计算机网络 1.掌握计算机网络的基本概念、基本原理和基本方法。 2.掌握计算机网络的体系结构和典型网络协议&#xff0c;了解典型网络设备的组成和特点&#xff0c;理解典型网络设备的工作原理。 3.能够运用计算机网络的基本概念、基本原理和基本方法进行…

【笔记】Blender4.0建模入门-3物体的基本操作

Blender入门 ——邵发 3.1 物体的移动 演示&#xff1a; 1、选中一个物体 2、选中移动工具 3、移动 - 沿坐标轴移动 - 在坐标平面内移动 - 自由移动&#xff08;不好控制&#xff09; 选中物体&#xff1a;右上的大纲窗口&#xff0c;点击物体名称&#xff0c;物体的轮…

LabVIEW电火花线切割放电点位置

介绍了一个电火花线切割放电点位置分布评价系统&#xff0c;特别是在系统组成、硬件选择和LabVIEW软件应用方面。 本系统由两个主要部分组成&#xff1a;硬件和软件。硬件部分包括电流传感器、高速数据采集卡、开关电源、电阻和导线。软件部分则由LabVIEW编程环境构成&#xf…

2024如何入局云计算?亚麻云助力您成为云专家!一次不过,免费再考

福利派送 2024年伊始&#xff0c;一波又一波的裁员潮又开始啦&#xff01;还在做传统行业&#xff1f;还在做传统程序员&#xff1f;亦或是对未来依旧迷茫&#xff1f;赶紧趁着行业东风&#xff0c;开始了解入门云计算吧&#xff01;亚马逊云科技&#xff0c;认证福利免费送&a…

C++——结构体

1&#xff0c;结构体基本概念 结构体属于用户自定义的数据类型&#xff0c;允许用户存储不同的数据类型。像int&#xff08;整型&#xff09;&#xff0c;浮点型&#xff0c;bool型&#xff0c;字符串型等都是属于系统内置的数据类型。而今天要学习的结构体则是属于我们自定义…

Java前端——HTTP协议中get和post的区别

get 和 post是 HTTP 请求的两种方法 应用场景&#xff1a; 一般 get 请求用于对服务器资源不会产生影响的场景&#xff0c;如请求一个网页的资源 post 请求一般用于对服务器资源会产生影响的情景&#xff0c;如注册用户这一类的操作。 get请求一般用于向服务器请求数据&…

Elastic Stack 8.12:通过对 ES|QL 等的改进增强了向量搜索

作者&#xff1a;来自 Elastic Tyler Perkins, Shani Sagiv, Gilad Gal, Ninoslav Miskovic Elastic Stack 8.12 构建于 Apache Lucene 9.9&#xff08;有史以来最快的 Lucene 版本&#xff09;之上&#xff0c;基于我们对标量量化和搜索并发性的贡献&#xff0c;为文本、向量和…

Bert详解

Bert框架 基本架构Embeddingpre-trainingMLM&#xff08;Mask Language Model&#xff09;NSP&#xff08;Next Sentence Prediction&#xff09; fine-tuning优缺点 基本架构 由Transformer的Encoder层堆叠而来 每个部分组成如下&#xff1a; Embedding Embedding由三种E…

最小生成树(Java实现)

一、Prim算法 Prim算法基本思想为&#xff1a;从联通网络 N{V,E}中某一顶点 v0 出发&#xff0c;此后就从一个顶点在 S 集中&#xff0c; 另一个顶点不在 S 集中的所有顶点中选择出权值最小的边&#xff0c;把对应顶点加入到 S 集 中&#xff0c; 直到所有的顶点都加入到 S 集中…

D-Tale SSRF漏洞复现(CVE-2024-21642)

0x01 产品简介 D-tale 是一个在 2020 年 2 月推出的库, 是 Pandas 数据结构的可视化工具。它具有许多功能,对于探索性数据分析非常方便、支持交互式绘图、3d 绘图、热图、特征之间的相关性、构建自定义列等等。 0x02 漏洞概述 D-Tale 是 Pandas 数据结构的可视化工具。3.9…

swift基础语法

swift学习笔记 参考教程 https://www.runoob.com/swift/swift-data-types.html swift代码规范 https://juejin.cn/post/7129465308376465422 1 环境搭建 必须要有苹果电脑且安装Xcode 2 基本语法 Swift是类型安全的语言&#xff0c;编译时会进行类型检查 import Cocoa var m…

Git学习笔记(第7章):IDEA实现Git操作(VSCode)

目录 7.1 配置忽略文件 7.2 初始化本地库 7.3 添加暂存区、提交本地库 7.4 修改文件 补充&#xff1a;工具栏简介 7.1 配置忽略文件 问题引入 在版本控制系统中&#xff0c;有些文件或目录是不需要纳入版本管理的&#xff0c;比如编译产生的临时文件、日志文件、缓存文件等…

基于springboot+vue的网上购物商城(前后端分离)

博主主页&#xff1a;猫头鹰源码 博主简介&#xff1a;Java领域优质创作者、CSDN博客专家、公司架构师、全网粉丝5万、专注Java技术领域和毕业设计项目实战 主要内容&#xff1a;毕业设计(Javaweb项目|小程序等)、简历模板、学习资料、面试题库、技术咨询 文末联系获取 项目背景…

hugo的常规使用操作

hugo的常规使用操作&#xff08;不断完善中&#xff09; 找到theme主题中config.toml 一般都会通过theme中复制到自己项目的config.toml中做修改和补充&#xff0c;来完善不同的业务需求 Hugo静态资源载入逻辑 原理 将图片信息放到static中&#xff0c;但是在文章中写的时…

电脑存储位置不够怎么办

电脑内存不够怎么办&#xff01;&#xff01;&#xff01; 我前段时间经常因为电脑D盘内存不够而苦恼&#xff08;毕竟电脑内存就那么丁点&#xff0c;C盘作为系统盘不能随便下东西的情况下&#xff0c;就只能选择其他盘进 方法一&#xff1a;检查电脑硬盘的分区情况&#xf…

Unity下实现跨平台的RTMP推流|轻量级RTSP服务|RTMP播放|RTSP播放低延迟解决方案

2018年&#xff0c;我们开始在原生RTSP|RTMP直播播放器的基础上&#xff0c;对接了Unity环境下的低延迟播放&#xff0c;毫秒级延迟&#xff0c;发布后&#xff0c;就得到了业内一致的认可。然后我们覆盖了Windows、Android、iOS、Linux的RTMP推送、轻量级RTSP服务和RTSP|RTMP播…

《WebKit 技术内幕》学习之五(3): HTML解释器和DOM 模型

3 DOM的事件机制 基于 WebKit 的浏览器事件处理过程&#xff1a;首先检测事件发生处的元素有无监听者&#xff0c;如果网页的相关节点注册了事件的监听者则浏览器会将事件派发给 WebKit 内核来处理。另外浏览器可能也需要处理这样的事件&#xff08;浏览器对于有些事件必须响应…

BGP Local-preferenct 、AS-Path、 Origin 综合选路实验

Local-preference&#xff1a; 本地优先级&#xff0c;公认任意&#xff0c;仅能在 AS 内使用&#xff08;IBGP内传递&#xff09;&#xff0c;不能在EBGP传递&#xff0c;默认值 100&#xff0c;越大越优。用于离开本 AS &#xff0c;在 IBGP 的入、出方向都可使用&#xff0c…

双端Diff算法

双端Diff算法 双端Diff算法指的是&#xff0c;在新旧两组子节点的四个端点之间分别进行比较&#xff0c;并试图找到可复用的节点。相比简单Diff算法&#xff0c;双端Diff算法的优势在于&#xff0c;对于同样的更新场景&#xff0c;执行的DOM移动操作次数更少。 简单 Diff 算法…