hipcc 编译 amd gpu kernel 和 打包与解包的流程实验

2024-06-17 21:12

本文主要是介绍hipcc 编译 amd gpu kernel 和 打包与解包的流程实验,希望对大家解决编程问题提供一定的参考价值,需要的开发者们随着小编来一起学习吧!

1, hip cuda kernel 编译概观

编译的文件流:

.hip kernel    --(clang++)-->                    .o

.o                    --(lld)-->                           .out

.out      --(clang-offload-bundler)-->     .hipfb

2,示例 hipcc -###

代码:

__global__ void WWWWW()
{((int*)0x8888888)[3] = 0x77777;
}

操作过程:


$ hipcc -### param_00.hip --cuda-device-only --offload-arch=gfx906 
AMD clang version 17.0.0 (https://github.com/RadeonOpenCompute/llvm-project roc-6.0.2 24012 af27734ed982b52a9f1be0f035ac91726fc697e4)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /opt/rocm-6.0.2/llvm/bin
Configuration file: /opt/rocm-6.0.2/lib/llvm/bin/clang++.cfg
clang: warning: argument unused during compilation: '--rtlib=compiler-rt' [-Wunused-command-line-argument]
clang: warning: argument unused during compilation: '-unwindlib=libgcc' [-Wunused-command-line-argument]"/opt/rocm-6.0.2/lib/llvm/bin/clang-17" \"-cc1" "-triple" "amdgcn-amd-amdhsa" "-aux-triple" \"x86_64-unknown-linux-gnu" "-emit-obj" "-disable-free" \"-clear-ast-before-backend" "-disable-llvm-verifier" \"-discard-value-names" "-main-file-name" "param_00.hip" \"-mrelocation-model" "pic" "-pic-level" "2" "-fhalf-no-semantic-interposition" \"-mframe-pointer=none" "-fno-rounding-math" "-mconstructor-aliases" "-aux-target-cpu" \"x86-64" "-fcuda-is-device" "-mllvm" "-amdgpu-internalize-symbols" "-fcuda-allow-variadic-functions" \"-fvisibility=hidden" "-fapply-global-visibility-to-externs" \"-mlink-builtin-bitcode" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/lib/amdgcn/bitcode/hip.bc" \"-mlink-builtin-bitcode" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/lib/amdgcn/bitcode/ocml.bc" \"-mlink-builtin-bitcode" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/lib/amdgcn/bitcode/ockl.bc" \"-mlink-builtin-bitcode" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/lib/amdgcn/bitcode/oclc_daz_opt_off.bc" \"-mlink-builtin-bitcode" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/lib/amdgcn/bitcode/oclc_unsafe_math_off.bc" \"-mlink-builtin-bitcode" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/lib/amdgcn/bitcode/oclc_finite_only_off.bc" \"-mlink-builtin-bitcode" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/lib/amdgcn/bitcode/oclc_correctly_rounded_sqrt_on.bc" \"-mlink-builtin-bitcode" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/lib/amdgcn/bitcode/oclc_wavefrontsize64_on.bc" \"-mlink-builtin-bitcode" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/lib/amdgcn/bitcode/oclc_isa_version_906.bc" \"-mlink-builtin-bitcode" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/lib/amdgcn/bitcode/oclc_abi_version_500.bc" \"-target-cpu" "gfx906" \"-debugger-tuning=gdb" "-resource-dir" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0" \"-internal-isystem" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/include/cuda_wrappers" \"-idirafter" "/opt/rocm-6.0.2/lib/llvm/bin/../../../include" "-include" \"__clang_hip_runtime_wrapper.h" "-c-isystem" "/opt/rocm-6.0.2/llvm/include/gpu-none-llvm" \"-isystem" "/opt/rocm-6.0.2/include" "-internal-isystem" "/usr/lib/gcc/x86_64-linux-gnu/12/../../../../include/c++/12" \"-internal-isystem" "/usr/lib/gcc/x86_64-linux-gnu/12/../../../../include/x86_64-linux-gnu/c++/12" \"-internal-isystem" "/usr/lib/gcc/x86_64-linux-gnu/12/../../../../include/c++/12/backward" \"-internal-isystem" "/usr/lib/gcc/x86_64-linux-gnu/12/../../../../include/c++/12" \"-internal-isystem" "/usr/lib/gcc/x86_64-linux-gnu/12/../../../../include/x86_64-linux-gnu/c++/12" \"-internal-isystem" "/usr/lib/gcc/x86_64-linux-gnu/12/../../../../include/c++/12/backward" \"-internal-isystem" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/include" \"-internal-isystem" "/usr/local/include" \"-internal-isystem" "/usr/lib/gcc/x86_64-linux-gnu/12/../../../../x86_64-linux-gnu/include" \"-internal-externc-isystem" "/usr/include/x86_64-linux-gnu" "-internal-externc-isystem" "/include" \"-internal-externc-isystem" "/usr/include" "-internal-isystem" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/include" \"-internal-isystem" "/usr/local/include" "-internal-isystem" "/usr/lib/gcc/x86_64-linux-gnu/12/../../../../x86_64-linux-gnu/include" \"-internal-externc-isystem" "/usr/include/x86_64-linux-gnu" "-internal-externc-isystem" "/include" \"-internal-externc-isystem" "/usr/include" "-O3" "-fdeprecated-macro" "-fno-autolink" \"-fdebug-compilation-dir=/home/hipper/ex_amd_gpu_compiler/ex/parameters_hip/param_00" \"-ferror-limit" "19" "-fhip-new-launch-api" "-fgnuc-version=4.2.1" "-fcxx-exceptions" "-fexceptions" \"-fcolor-diagnostics" "-vectorize-loops" "-vectorize-slp" "-mllvm" "-amdgpu-early-inline-all=true" "-mllvm" \"-amdgpu-function-calls=false" "-cuid=3e1885b9958b336f" "-fcuda-allow-variadic-functions" "-faddrsig" \"-D__GCC_HAVE_DWARF2_CFI_ASM=1" \"-o" "/tmp/param_00-gfx906-a7e858.o" \"-x" "hip" "param_00.hip""/opt/rocm-6.0.2/llvm/bin/lld" \"-flavor" "gnu" "-m" "elf64_amdgpu" \"--no-undefined" "-shared" \"-plugin-opt=-amdgpu-internalize-symbols" \"-plugin-opt=mcpu=gfx906" \"-plugin-opt=O3" "--lto-CGO3" \"-plugin-opt=-amdgpu-early-inline-all=true" \"-plugin-opt=-amdgpu-function-calls=false" \"--whole-archive" \"-o" "/tmp/param_00-gfx906-65b179.out" \"/tmp/param_00-gfx906-a7e858.o" \"--no-whole-archive""/opt/rocm-6.0.2/llvm/bin/clang-offload-bundler" \"-type=o" "-bundle-align=4096" \"-targets=host-x86_64-unknown-linux,hipv4-amdgcn-amd-amdhsa--gfx906" \"-input=/dev/null" \"-input=/tmp/param_00-gfx906-65b179.out" \"-output=param_00.hip-hip-amdgcn-amd-amdhsa.hipfb"

涉及到了三个命令:

clang++                                   -o xxx.o

lld                                             -o  xxx.out

clang-offload-bundler           -output=xxx.hipfb

3,分析中间文件

 

3.1 clang++ 编译生成的 .o

.o 是一个elf文件

这个.o 是使用自己编译出来的clang++ 编译的,

其中,将 cp /opt/rocm/bin/clang++.cfg  local_amdgpu/llvm/bin/

稍作路径调整,编译生成 .o  :

$ "../../../../local_amdgpu/bin/clang-19" \"-cc1" "-triple" "amdgcn-amd-amdhsa" "-aux-triple" \"x86_64-unknown-linux-gnu" "-emit-obj" "-disable-free" \"-clear-ast-before-backend" "-disable-llvm-verifier" \"-discard-value-names" "-main-file-name" "param_00.hip" \"-mrelocation-model" "pic" "-pic-level" "2" "-fhalf-no-semantic-interposition" \"-mframe-pointer=none" "-fno-rounding-math" "-mconstructor-aliases" "-aux-target-cpu" \"x86-64" "-fcuda-is-device" "-mllvm" "-amdgpu-internalize-symbols" "-fcuda-allow-variadic-functions" \"-fvisibility=hidden" "-fapply-global-visibility-to-externs" \"-mlink-builtin-bitcode" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/lib/amdgcn/bitcode/hip.bc" \"-mlink-builtin-bitcode" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/lib/amdgcn/bitcode/ocml.bc" \"-mlink-builtin-bitcode" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/lib/amdgcn/bitcode/ockl.bc" \"-mlink-builtin-bitcode" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/lib/amdgcn/bitcode/oclc_daz_opt_off.bc" \"-mlink-builtin-bitcode" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/lib/amdgcn/bitcode/oclc_unsafe_math_off.bc" \"-mlink-builtin-bitcode" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/lib/amdgcn/bitcode/oclc_finite_only_off.bc" \"-mlink-builtin-bitcode" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/lib/amdgcn/bitcode/oclc_correctly_rounded_sqrt_on.bc" \"-mlink-builtin-bitcode" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/lib/amdgcn/bitcode/oclc_wavefrontsize64_on.bc" \"-mlink-builtin-bitcode" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/lib/amdgcn/bitcode/oclc_isa_version_906.bc" \"-mlink-builtin-bitcode" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/lib/amdgcn/bitcode/oclc_abi_version_500.bc" \"-target-cpu" "gfx906" \"-debugger-tuning=gdb" "-resource-dir" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0" \"-internal-isystem" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/include/cuda_wrappers" \"-idirafter" "/opt/rocm-6.0.2/lib/llvm/bin/../../../include" "-include" \"__clang_hip_runtime_wrapper.h" "-c-isystem" "/opt/rocm-6.0.2/llvm/include/gpu-none-llvm" \"-isystem" "/opt/rocm-6.0.2/include" "-internal-isystem" "/usr/lib/gcc/x86_64-linux-gnu/12/../../../../include/c++/12" \"-internal-isystem" "/usr/lib/gcc/x86_64-linux-gnu/12/../../../../include/x86_64-linux-gnu/c++/12" \"-internal-isystem" "/usr/lib/gcc/x86_64-linux-gnu/12/../../../../include/c++/12/backward" \"-internal-isystem" "/usr/lib/gcc/x86_64-linux-gnu/12/../../../../include/c++/12" \"-internal-isystem" "/usr/lib/gcc/x86_64-linux-gnu/12/../../../../include/x86_64-linux-gnu/c++/12" \"-internal-isystem" "/usr/lib/gcc/x86_64-linux-gnu/12/../../../../include/c++/12/backward" \"-internal-isystem" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/include" \"-internal-isystem" "/usr/local/include" \"-internal-isystem" "/usr/lib/gcc/x86_64-linux-gnu/12/../../../../x86_64-linux-gnu/include" \"-internal-externc-isystem" "/usr/include/x86_64-linux-gnu" "-internal-externc-isystem" "/include" \"-internal-externc-isystem" "/usr/include" "-internal-isystem" "/opt/rocm-6.0.2/lib/llvm/lib/clang/17.0.0/include" \"-internal-isystem" "/usr/local/include" "-internal-isystem" "/usr/lib/gcc/x86_64-linux-gnu/12/../../../../x86_64-linux-gnu/include" \"-internal-externc-isystem" "/usr/include/x86_64-linux-gnu" "-internal-externc-isystem" "/include" \"-internal-externc-isystem" "/usr/include" "-O3" "-fdeprecated-macro" "-fno-autolink" \"-fdebug-compilation-dir=/home/hipper/ex_amd_gpu_compiler/ex/parameters_hip/param_00" \"-ferror-limit" "19" "-fhip-new-launch-api" "-fgnuc-version=4.2.1" "-fcxx-exceptions" "-fexceptions" \"-fcolor-diagnostics" "-vectorize-loops" "-vectorize-slp" "-mllvm" "-amdgpu-early-inline-all=true" "-mllvm" \"-amdgpu-function-calls=false" "-cuid=3e1885b9958b336f" "-fcuda-allow-variadic-functions" "-faddrsig" \"-D__GCC_HAVE_DWARF2_CFI_ASM=1" \"-o" "./param_00-gfx906-a7e858.o" \"-x" "hip" "param_00.hip"

3.2 lld 链接得到的.out 文件

也属于 elf 文件,但是是DYN (Shared object file) 类型,不再是 relocationable 类型。

其生成命令也是稍作了路径改动:

"../../../../local_amdgpu/bin/lld" \"-flavor" "gnu" "-m" "elf64_amdgpu" \"--no-undefined" "-shared" \"-plugin-opt=-amdgpu-internalize-symbols" \"-plugin-opt=mcpu=gfx906" \"-plugin-opt=O3" "--lto-CGO3" \"-plugin-opt=-amdgpu-early-inline-all=true" \"-plugin-opt=-amdgpu-function-calls=false" \"--whole-archive" \"-o" "./param_00-gfx906-65b179.out" \"./param_00-gfx906-a7e858.o" \"--no-whole-archive"

3.3 hipfb 文件

这个类型的文件 是由 clang-offload-bundler 打包而成,这里没有什么新意,对 clang-offload-bundler做一个介绍:

clang-offload-bundler 是一个工具,它是 Clang/LLVM 编译器工具链的一部分,用于支持在多种设备上进行异构计算。这个工具主要用于处理和打包不同目标设备(如 CPU、GPU、FPGA 等)的代码,以便在一个单一的程序中支持多种计算设备。这种技术通常用于加速应用程序,特别是在需要大量并行处理的场景中。

3.3.1  生成 hipfb 文件的方法

"../../../../local_amdgpu/bin/clang-offload-bundler" \"-type=o" "-bundle-align=4096" \"-targets=host-x86_64-unknown-linux,hipv4-amdgcn-amd-amdhsa--gfx906" \"-input=/dev/null" \"-input=./param_00-gfx906-65b179.out" \"-output=param_00.hip-hip-amdgcn-amd-amdhsa.hipfb"

3.3.2  主要功能和作用

1. 打包和解包代码对象:

  • clang-offload-bundler 能够将针对不同设备的代码对象(如 CPU 和 GPU 的代码)打包到一个单一的文件中。这使得管理和分发针对异构计算平台的应用程序变得更加简单。
  • 同样,它也可以从这种打包的文件中提取特定目标设备的代码对象,以便在适当的设备上执行。

2. 支持异构编程:

  • 在异构编程中,开发者可能需要将程序的不同部分优化并编译到不同的硬件平台上。clang-offload-bundler 通过管理这些不同的代码段,简化了构建和部署过程。

3. 简化编译和链接流程:

  • 在使用 OpenMP 或 CUDA 等并行编程模型时,clang-offload-bundler 能够处理主机代码和加速器代码之间的交互,包括数据传输和执行控制。这样,开发者可以更专注于代码的并行部分,而不是底层的数据管理和设备控制。

4. 提高性能和可移植性:

  • 通过允许代码针对特定硬件进行优化,clang-offload-bundler 帮助提高应用程序的性能。同时,它也支持代码的可移植性,因为同一个应用程序可以针对多种硬件平台进行编译和打包。

3.3.3  使用场景

  • 并行计算应用:在需要大量计算资源的应用中,如科学计算、图像处理、机器学习等,clang-offload-bundler 可以帮助开发者有效地利用多种计算资源。
  • 开发跨平台应用:对于需要在多种硬件设备上运行的软件,如桌面和移动设备,或者 CPU 和 GPU,clang-offload-bundler 提供了一种统一的方式来处理不同平台的代码。

意义:clang-offload-bundler 还是挺强大的,用于支持和简化异构计算环境中的编程和部署过程。它通过管理针对不同硬件平台的代码,使得开发高性能并行应用程序变得更加高效和可行。

3.3.4 解析 hipfb的方法

clang-offload-bundler -type=o -targets=hip-amdgcn-amd-amdhsa--gfx906 -input=param_00.hip-hip-amdgcn-amd-amdhsa.hipfb -output=device_output.o -unbundle

于是又得到了 bundle前的out 文件:

甚至连文件大小都一样:

这篇关于hipcc 编译 amd gpu kernel 和 打包与解包的流程实验的文章就介绍到这儿,希望我们推荐的文章对编程师们有所帮助!



http://www.chinasem.cn/article/1070495

相关文章

Security OAuth2 单点登录流程

单点登录(英语:Single sign-on,缩写为 SSO),又译为单一签入,一种对于许多相互关连,但是又是各自独立的软件系统,提供访问控制的属性。当拥有这项属性时,当用户登录时,就可以获取所有系统的访问权限,不用对每个单一系统都逐一登录。这项功能通常是以轻型目录访问协议(LDAP)来实现,在服务器上会将用户信息存储到LDAP数据库中。相同的,单一注销(single sign-off)就是指

Spring Security基于数据库验证流程详解

Spring Security 校验流程图 相关解释说明(认真看哦) AbstractAuthenticationProcessingFilter 抽象类 /*** 调用 #requiresAuthentication(HttpServletRequest, HttpServletResponse) 决定是否需要进行验证操作。* 如果需要验证,则会调用 #attemptAuthentica

springboot3打包成war包,用tomcat8启动

1、在pom中,将打包类型改为war <packaging>war</packaging> 2、pom中排除SpringBoot内置的Tomcat容器并添加Tomcat依赖,用于编译和测试,         *依赖时一定设置 scope 为 provided (相当于 tomcat 依赖只在本地运行和测试的时候有效,         打包的时候会排除这个依赖)<scope>provided

Linux_kernel驱动开发11

一、改回nfs方式挂载根文件系统         在产品将要上线之前,需要制作不同类型格式的根文件系统         在产品研发阶段,我们还是需要使用nfs的方式挂载根文件系统         优点:可以直接在上位机中修改文件系统内容,延长EMMC的寿命         【1】重启上位机nfs服务         sudo service nfs-kernel-server resta

AI Toolkit + H100 GPU,一小时内微调最新热门文生图模型 FLUX

上个月,FLUX 席卷了互联网,这并非没有原因。他们声称优于 DALLE 3、Ideogram 和 Stable Diffusion 3 等模型,而这一点已被证明是有依据的。随着越来越多的流行图像生成工具(如 Stable Diffusion Web UI Forge 和 ComyUI)开始支持这些模型,FLUX 在 Stable Diffusion 领域的扩展将会持续下去。 自 FLU

如何用GPU算力卡P100玩黑神话悟空?

精力有限,只记录关键信息,希望未来能够有助于其他人。 文章目录 综述背景评估游戏性能需求显卡需求CPU和内存系统需求主机需求显式需求 实操硬件安装安装操作系统Win11安装驱动修改注册表选择程序使用什么GPU 安装黑神话悟空其他 综述 用P100 + PCIe Gen3.0 + Dell720服务器(32C64G),运行黑神话悟空画质中等流畅运行。 背景 假设有一张P100-

maven 编译构建可以执行的jar包

💝💝💝欢迎莅临我的博客,很高兴能够在这里和您见面!希望您在这里可以感受到一份轻松愉快的氛围,不仅可以获得有趣的内容和知识,也可以畅所欲言、分享您的想法和见解。 推荐:「stormsha的主页」👈,「stormsha的知识库」👈持续学习,不断总结,共同进步,为了踏实,做好当下事儿~ 专栏导航 Python系列: Python面试题合集,剑指大厂Git系列: Git操作技巧GO

kubelet组件的启动流程源码分析

概述 摘要: 本文将总结kubelet的作用以及原理,在有一定基础认识的前提下,通过阅读kubelet源码,对kubelet组件的启动流程进行分析。 正文 kubelet的作用 这里对kubelet的作用做一个简单总结。 节点管理 节点的注册 节点状态更新 容器管理(pod生命周期管理) 监听apiserver的容器事件 容器的创建、删除(CRI) 容器的网络的创建与删除

STM32(十一):ADC数模转换器实验

AD单通道: 1.RCC开启GPIO和ADC时钟。配置ADCCLK分频器。 2.配置GPIO,把GPIO配置成模拟输入的模式。 3.配置多路开关,把左面通道接入到右面规则组列表里。 4.配置ADC转换器, 包括AD转换器和AD数据寄存器。单次转换,连续转换;扫描、非扫描;有几个通道,触发源是什么,数据对齐是左对齐还是右对齐。 5.ADC_CMD 开启ADC。 void RCC_AD

火语言RPA流程组件介绍--浏览网页

🚩【组件功能】:浏览器打开指定网址或本地html文件 配置预览 配置说明 网址URL 支持T或# 默认FLOW输入项 输入需要打开的网址URL 超时时间 支持T或# 打开网页超时时间 执行后后等待时间(ms) 支持T或# 当前组件执行完成后继续等待的时间 UserAgent 支持T或# User Agent中文名为用户代理,简称 UA,它是一个特殊字符串头,使得服务器