将 cuda kernel 编译成 ptx 和 rocm的hip asm

2024-05-31 11:44

本文主要是介绍将 cuda kernel 编译成 ptx 和 rocm的hip asm,希望对大家解决编程问题提供一定的参考价值,需要的开发者们随着小编来一起学习吧!

1,cuda 源码编译

cuda_a_one.cu __global__ void NNNNNVVVVV_one(int *A)
{A[333] = 777;
}

编译命令:

%.ptx: %.cu
    nvcc -arch=sm_70 -ptx $< -o $@

生成的结果:

2, hip 源码编译

hip_a_one.hip__global__ void AAAAAMMMMM_one(int *A)
{A[0x333] = 0x777;
}

编译命令:

%.hip.s: %.hip
    $(HIPCC) $< -o $@ -S --offload-device-only

生成的结果:

存储为文本:

	.text.amdgcn_target "amdgcn-amd-amdhsa--gfx906".protected	_Z14AAAAAMMMMM_onePi    ; -- Begin function _Z14AAAAAMMMMM_onePi.globl	_Z14AAAAAMMMMM_onePi.p2align	8.type	_Z14AAAAAMMMMM_onePi,@function
_Z14AAAAAMMMMM_onePi:                   ; @_Z14AAAAAMMMMM_onePi
; %bb.0:s_load_dwordx2 s[0:1], s[4:5], 0x0v_mov_b32_e32 v0, 0v_mov_b32_e32 v1, 0x777s_waitcnt lgkmcnt(0)global_store_dword v0, v1, s[0:1] offset:3276s_endpgm.section	.rodata,#alloc.p2align	6, 0x0.amdhsa_kernel _Z14AAAAAMMMMM_onePi.amdhsa_group_segment_fixed_size 0.amdhsa_private_segment_fixed_size 0.amdhsa_kernarg_size 8.amdhsa_user_sgpr_count 6.amdhsa_user_sgpr_private_segment_buffer 1.amdhsa_user_sgpr_dispatch_ptr 0.amdhsa_user_sgpr_queue_ptr 0.amdhsa_user_sgpr_kernarg_segment_ptr 1.amdhsa_user_sgpr_dispatch_id 0.amdhsa_user_sgpr_flat_scratch_init 0.amdhsa_user_sgpr_private_segment_size 0.amdhsa_uses_dynamic_stack 0.amdhsa_system_sgpr_private_segment_wavefront_offset 0.amdhsa_system_sgpr_workgroup_id_x 1.amdhsa_system_sgpr_workgroup_id_y 0.amdhsa_system_sgpr_workgroup_id_z 0.amdhsa_system_sgpr_workgroup_info 0.amdhsa_system_vgpr_workitem_id 0.amdhsa_next_free_vgpr 2.amdhsa_next_free_sgpr 6.amdhsa_reserve_vcc 0.amdhsa_reserve_flat_scratch 0.amdhsa_reserve_xnack_mask 1.amdhsa_float_round_mode_32 0.amdhsa_float_round_mode_16_64 0.amdhsa_float_denorm_mode_32 3.amdhsa_float_denorm_mode_16_64 3.amdhsa_dx10_clamp 1.amdhsa_ieee_mode 1.amdhsa_fp16_overflow 0.amdhsa_exception_fp_ieee_invalid_op 0.amdhsa_exception_fp_denorm_src 0.amdhsa_exception_fp_ieee_div_zero 0.amdhsa_exception_fp_ieee_overflow 0.amdhsa_exception_fp_ieee_underflow 0.amdhsa_exception_fp_ieee_inexact 0.amdhsa_exception_int_div_zero 0.end_amdhsa_kernel.text
.Lfunc_end0:.size	_Z14AAAAAMMMMM_onePi, .Lfunc_end0-_Z14AAAAAMMMMM_onePi; -- End function.section	.AMDGPU.csdata
; Kernel info:
; codeLenInByte = 36
; NumSgprs: 10
; NumVgprs: 2
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 1
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 10
; NumVGPRsForWavesPerEU: 2
; Occupancy: 8
; WaveLimiterHint : 1
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 6
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0.ident	"AMD clang version 17.0.0 (https://github.com/RadeonOpenCompute/llvm-project roc-6.0.2 24012 af27734ed982b52a9f1be0f035ac91726fc697e4)".section	".note.GNU-stack".addrsig.amdgpu_metadata
---
amdhsa.kernels:- .args:- .address_space:  global.offset:         0.size:           8.value_kind:     global_buffer.group_segment_fixed_size: 0.kernarg_segment_align: 8.kernarg_segment_size: 8.language:       OpenCL C.language_version:- 2- 0.max_flat_workgroup_size: 1024.name:           _Z14AAAAAMMMMM_onePi.private_segment_fixed_size: 0.sgpr_count:     10.sgpr_spill_count: 0.symbol:         _Z14AAAAAMMMMM_onePi.kd.uniform_work_group_size: 1.uses_dynamic_stack: false.vgpr_count:     2.vgpr_spill_count: 0.wavefront_size: 64
amdhsa.target:   amdgcn-amd-amdhsa--gfx906
amdhsa.version:- 1- 2
....end_amdgpu_metadata

3,hipcc 的概述编译流程

首先,hipcc是一个perl脚本:

#!/usr/bin/env perl
# Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved.
#
# Permission is hereby granted, free of charge, to any person obtaining a copy
# of this software and associated documentation files (the "Software"), to deal
# in the Software without restriction, including without limitation the rights
# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
# copies of the Software, and to permit persons to whom the Software is
# furnished to do so, subject to the following conditions:
#
# The above copyright notice and this permission notice shall be included in
# all copies or substantial portions of the Software.
#
# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
# THE SOFTWARE.# Need perl > 5.10 to use logic-defined or
use 5.006; use v5.10.1;use warnings;use File::Basename;
use File::Spec::Functions 'catfile';# TODO: By default select perl script until change incorporated in HIP build script.
my $USE_PERL_SCRIPT = $ENV{'HIP_USE_PERL_SCRIPTS'};
$USE_PERL_SCRIPT //= 1;  # use defined-or assignment operator.  Use env var, but if not defined default to 1.my $isWindows =  ($^O eq 'MSWin32' or $^O eq 'msys');
# escapes args with quotes SWDEV-341955
foreach $arg (@ARGV) {if ($isWindows) {$arg =~ s/[^-a-zA-Z0-9_=+,.:\/\\ ]/\\$&/g;}
}my $SCRIPT_DIR=dirname(__FILE__);
if ($USE_PERL_SCRIPT) {#Invoke hipcc.plmy $HIPCC_PERL=catfile($SCRIPT_DIR, '/hipcc.pl');system($^X, $HIPCC_PERL, @ARGV);
} else {$BIN_NAME="/hipcc.bin";if ($isWindows) {$BIN_NAME="/hipcc.bin.exe";}my $HIPCC_BIN=catfile($SCRIPT_DIR, $BIN_NAME);if ( -e $HIPCC_BIN ) {#Invoke hipcc.binsystem($HIPCC_BIN, @ARGV);} else {print "hipcc.bin not present; install HIPCC binaries before proceeding\n";exit(-1);}
}# Because of this wrapper we need to check
# the output of the system command for perl and bin
# else the failures are ignored and build fails silently
if ($? == -1) {exit($?);
}
elsif ($? & 127) {exit($?);
}
else {$CMD_EXIT_CODE = $? >> 8;
}
exit($CMD_EXIT_CODE);

具体工作流程:

未完待续。。。

这篇关于将 cuda kernel 编译成 ptx 和 rocm的hip asm的文章就介绍到这儿,希望我们推荐的文章对编程师们有所帮助!



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

相关文章

Linux_kernel驱动开发11

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

笔记整理—内核!启动!—kernel部分(2)从汇编阶段到start_kernel

kernel起始与ENTRY(stext),和uboot一样,都是从汇编阶段开始的,因为对于kernel而言,还没进行栈的维护,所以无法使用c语言。_HEAD定义了后面代码属于段名为.head .text的段。         内核起始部分代码被解压代码调用,前面关于uboot的文章中有提到过(eg:zImage)。uboot启动是无条件的,只要代码的位置对,上电就工作,kern

欧拉系统 kernel 升级、降级

系统版本  cat  /etc/os-release  NAME="openEuler"VERSION="22.03 (LTS-SP1)"ID="openEuler"VERSION_ID="22.03"PRETTY_NAME="openEuler 22.03 (LTS-SP1)"ANSI_COLOR="0;31" 系统初始 kernel 版本 5.10.0-136.12.0.

[Linux Kernel Block Layer第一篇] block layer架构设计

目录 1. single queue架构 2. multi-queue架构(blk-mq)  3. 问题 随着SSD快速存储设备的发展,内核社区越发发现,存储的性能瓶颈从硬件存储设备转移到了内核block layer,主要因为当时的内核block layer是single hw queue的架构,导致cpu锁竞争问题严重,本文先提纲挈领的介绍内核block layer的架构演进,然

PyInstaller问题解决 onnxruntime-gpu 使用GPU和CUDA加速模型推理

前言 在模型推理时,需要使用GPU加速,相关的CUDA和CUDNN安装好后,通过onnxruntime-gpu实现。 直接运行python程序是正常使用GPU的,如果使用PyInstaller将.py文件打包为.exe,发现只能使用CPU推理了。 本文分析这个问题和提供解决方案,供大家参考。 问题分析——找不到ONNX Runtime GPU 动态库 首先直接运行python程序

Kernel 中MakeFile 使用if条件编译

有时需要通过if  else来选择编译哪个驱动,单纯的obj-$(CONFIG_)就不是很方便,下面提供两种参考案例: 案例一: 来源:drivers/char/tpm/Makefileifdef CONFIG_ACPItpm-y += tpm_eventlog.o tpm_acpi.oelseifdef CONFIG_TCG_IBMVTPMtpm-y += tpm_eventlog.o

笔记整理—内核!启动!—kernel部分(1)驱动与内核的关系

首先,恭喜完成了uboot部分的内容整理,其次补充一点,uboot第一部分和第二部分的工作不是一定的,在不同的版本中,可能这个初始化早一点,那个的又放在了第二部分,版本不同,造成的工作顺序不同,但终归是要完成基本内容初始化并传参给kernel的。         那么至于驱动与内核的关系,用一张图来说明最适合不过:         驱动位于OS层的中下层与硬件相接。驱动是内

ASM 10G 基于RMAN 迁移

ASM 10G 基于RMAN 迁移 场景 单节点基于10G R2 的数据库,其数据文件及日志文件均存放在ASM 里,现在为业务需求,将此数据库做迁 移,迁移到另个机房,但是两个机房的网络是通畅的,为尽量减少数据的丢失及平稳迁移和经济实惠,迁 移时,数据库需停应用 工具 本次采用RMAN 的duplicate 命令来进行迁移,运用此命令简化复杂度; 一、源库和目标库的

CUDA:用并行计算的方法对图像进行直方图均衡处理

(一)目的 将所学算法运用于图像处理中。 (二)内容 用并行计算的方法对图像进行直方图均衡处理。 要求: 利用直方图均衡算法处理lena_salt图像 版本1:CPU实现 版本2:GPU实现  实验步骤一 软件设计分析: 数据类型: 根据实验要求,本实验的数据类型为一个256*256*8的整型矩阵,其中元素的值为256*256个0-255的灰度值。 存储方式: 图像在内存中