#include <stdio.h>
#include <cuda.h>
#include <stdint.h>
#include <math.h>

__device__ bool deviceFiFlag = false;
__device__ long long deviceFiInstCount;
__device__ long long deviceFiThreadIndex;
__device__ double deviceSeedFactor; // Should be from [0,1)
__device__ int deviceFiBit;
__device__ long deviceBambooIndex;

/////////////////////////////////////////////////
__device__ long long currentCycle = 0;

extern "C" __device__ void injectFault(char* buf, int size, long bambooIndex)
{
	int blockId = blockIdx.x 
		+ blockIdx.y * gridDim.x 
		+ gridDim.x * gridDim.y * blockIdx.z; 
	long long index = blockId * (blockDim.x * blockDim.y * blockDim.z)
		+ (threadIdx.z * (blockDim.x * blockDim.y))
		+ (threadIdx.y * blockDim.x)
		+ threadIdx.x;

	//////////////////////////////////////////////////////////////////////////////

	if(deviceFiFlag == true && index == deviceFiThreadIndex){
		//printf("Cycle: %lld\n", currentCycle);
		// There is not instCount=0, instCount means Nth inst executed.
		currentCycle++;
	}

	if(deviceFiFlag == true && deviceFiInstCount == currentCycle && index == deviceFiThreadIndex){
		deviceFiFlag = false;

		//printf("\n\\\\\\\\\\\\\\\\\\/////////////////\nInjecting fault ... \n\n");

		unsigned char oldbuf;
		unsigned fiBit, fiBytepos, fiBitpos;
		int randint = deviceSeedFactor*size;
		fiBit = (unsigned) randint;
		fiBytepos = fiBit / 8;
		fiBitpos = fiBit % 8;
		memcpy(&oldbuf, &buf[fiBytepos], 1);

		printf("	|** fiBit: %d, fiBambooIndex: %ld\n", fiBit, bambooIndex);
		printf("	|** FI Config Runtime ** deviceFiThreadIndex: %lld, deviceFiInstCount: %lld, fiBit: %u(size %d)\n", deviceFiThreadIndex, deviceFiInstCount, fiBit, size);
		printf("	|** Original Value: %#010x **\n", oldbuf);
		buf[fiBytepos] ^= 0x1 << fiBitpos;
		printf("	|** Corrupted Value: %#010x **\n", buf[fiBytepos]);

		deviceFiBit = fiBit;
		deviceBambooIndex = bambooIndex;

	}

}


	.file	"tmpxft_0023ba81_00000000-6_injection_runtime.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2030:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2030:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	injectFault
	.type	injectFault, @function
injectFault:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2027:
	.size	injectFault, .-injectFault
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"deviceFiFlag"
.LC1:
	.string	"deviceFiInstCount"
.LC2:
	.string	"deviceFiThreadIndex"
.LC3:
	.string	"deviceSeedFactor"
.LC4:
	.string	"deviceFiBit"
.LC5:
	.string	"deviceBambooIndex"
.LC6:
	.string	"currentCycle"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2053:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	movl	$1, %r9d
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC0(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	movq	%rdx, %rcx
	leaq	_ZL12deviceFiFlag(%rip), %rsi
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC1(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8, %r9d
	leaq	_ZL17deviceFiInstCount(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC2(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8, %r9d
	leaq	_ZL19deviceFiThreadIndex(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rdi
	.cfi_def_cfa_offset 24
	popq	%r8
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC3(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8, %r9d
	leaq	_ZL16deviceSeedFactor(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%r9
	.cfi_def_cfa_offset 24
	popq	%r10
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC4(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$4, %r9d
	leaq	_ZL11deviceFiBit(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%r11
	.cfi_def_cfa_offset 24
	popq	%rax
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC5(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8, %r9d
	leaq	_ZL17deviceBambooIndex(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC6(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8, %r9d
	leaq	_ZL12currentCycle(%rip), %rsi
	call	__cudaRegisterVar@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2053:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.local	_ZL12currentCycle
	.comm	_ZL12currentCycle,8,8
	.local	_ZL17deviceBambooIndex
	.comm	_ZL17deviceBambooIndex,8,8
	.local	_ZL11deviceFiBit
	.comm	_ZL11deviceFiBit,4,4
	.local	_ZL16deviceSeedFactor
	.comm	_ZL16deviceSeedFactor,8,8
	.local	_ZL19deviceFiThreadIndex
	.comm	_ZL19deviceFiThreadIndex,8,8
	.local	_ZL17deviceFiInstCount
	.comm	_ZL17deviceFiInstCount,8,8
	.local	_ZL12deviceFiFlag
	.comm	_ZL12deviceFiFlag,1,1
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

#include <stdio.h>
#include <hip/hip_runtime.h>
#include <stdint.h>
#include <math.h>

__device__ bool deviceFiFlag = false;
__device__ long long deviceFiInstCount;
__device__ long long deviceFiThreadIndex;
__device__ double deviceSeedFactor; // Should be from [0,1)
__device__ int deviceFiBit;
__device__ long deviceBambooIndex;

/////////////////////////////////////////////////
__device__ long long currentCycle = 0;

extern "C" __device__ void injectFault(char* buf, int size, long bambooIndex)
{
	int blockId = blockIdx.x 
		+ blockIdx.y * gridDim.x 
		+ gridDim.x * gridDim.y * blockIdx.z; 
	long long index = blockId * (blockDim.x * blockDim.y * blockDim.z)
		+ (threadIdx.z * (blockDim.x * blockDim.y))
		+ (threadIdx.y * blockDim.x)
		+ threadIdx.x;

	//////////////////////////////////////////////////////////////////////////////

	if(deviceFiFlag == true && index == deviceFiThreadIndex){
		//printf("Cycle: %lld\n", currentCycle);
		// There is not instCount=0, instCount means Nth inst executed.
		currentCycle++;
	}

	if(deviceFiFlag == true && deviceFiInstCount == currentCycle && index == deviceFiThreadIndex){
		deviceFiFlag = false;

		//printf("\n\\\\\\\\\\\\\\\\\\/////////////////\nInjecting fault ... \n\n");

		unsigned char oldbuf;
		unsigned fiBit, fiBytepos, fiBitpos;
		int randint = deviceSeedFactor*size;
		fiBit = (unsigned) randint;
		fiBytepos = fiBit / 8;
		fiBitpos = fiBit % 8;
		memcpy(&oldbuf, &buf[fiBytepos], 1);

		printf("	|** fiBit: %d, fiBambooIndex: %ld\n", fiBit, bambooIndex);
		printf("	|** FI Config Runtime ** deviceFiThreadIndex: %lld, deviceFiInstCount: %lld, fiBit: %u(size %d)\n", deviceFiThreadIndex, deviceFiInstCount, fiBit, size);
		printf("	|** Original Value: %#010x **\n", oldbuf);
		buf[fiBytepos] ^= 0x1 << fiBitpos;
		printf("	|** Corrupted Value: %#010x **\n", buf[fiBytepos]);

		deviceFiBit = fiBit;
		deviceBambooIndex = bambooIndex;

	}

}


	.text
	.file	"injection_runtime.hip"
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB0_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB0_2:
	movl	$deviceFiFlag, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$1, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$deviceFiInstCount, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movl	$8, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$deviceFiThreadIndex, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movl	$8, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$deviceSeedFactor, %esi
	movl	$.L__unnamed_4, %edx
	movl	$.L__unnamed_4, %ecx
	movl	$8, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$deviceFiBit, %esi
	movl	$.L__unnamed_5, %edx
	movl	$.L__unnamed_5, %ecx
	movl	$4, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$deviceBambooIndex, %esi
	movl	$.L__unnamed_6, %edx
	movl	$.L__unnamed_6, %ecx
	movl	$8, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$currentCycle, %esi
	movl	$.L__unnamed_7, %edx
	movl	$.L__unnamed_7, %ecx
	movl	$8, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$__hip_module_dtor, %edi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end0:
	.size	__hip_module_ctor, .Lfunc_end0-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB1_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB1_2:
	retq
.Lfunc_end1:
	.size	__hip_module_dtor, .Lfunc_end1-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	deviceFiFlag,@object            # @deviceFiFlag
	.local	deviceFiFlag
	.comm	deviceFiFlag,1,1
	.type	deviceFiInstCount,@object       # @deviceFiInstCount
	.local	deviceFiInstCount
	.comm	deviceFiInstCount,8,8
	.type	deviceFiThreadIndex,@object     # @deviceFiThreadIndex
	.local	deviceFiThreadIndex
	.comm	deviceFiThreadIndex,8,8
	.type	deviceSeedFactor,@object        # @deviceSeedFactor
	.local	deviceSeedFactor
	.comm	deviceSeedFactor,8,8
	.type	deviceFiBit,@object             # @deviceFiBit
	.local	deviceFiBit
	.comm	deviceFiBit,4,4
	.type	deviceBambooIndex,@object       # @deviceBambooIndex
	.local	deviceBambooIndex
	.comm	deviceBambooIndex,8,8
	.type	currentCycle,@object            # @currentCycle
	.local	currentCycle
	.comm	currentCycle,8,8
	.type	.L__unnamed_1,@object           # @0
	.section	.rodata.str1.1,"aMS",@progbits,1
.L__unnamed_1:
	.asciz	"deviceFiFlag"
	.size	.L__unnamed_1, 13

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"deviceFiInstCount"
	.size	.L__unnamed_2, 18

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"deviceFiThreadIndex"
	.size	.L__unnamed_3, 20

	.type	.L__unnamed_4,@object           # @3
.L__unnamed_4:
	.asciz	"deviceSeedFactor"
	.size	.L__unnamed_4, 17

	.type	.L__unnamed_5,@object           # @4
.L__unnamed_5:
	.asciz	"deviceFiBit"
	.size	.L__unnamed_5, 12

	.type	.L__unnamed_6,@object           # @5
.L__unnamed_6:
	.asciz	"deviceBambooIndex"
	.size	.L__unnamed_6, 18

	.type	.L__unnamed_7,@object           # @6
.L__unnamed_7:
	.asciz	"currentCycle"
	.size	.L__unnamed_7, 13

	.type	__hip_fatbin_wrapper,@object    # @__hip_fatbin_wrapper
	.section	.hipFatBinSegment,"a",@progbits
	.p2align	3, 0x0
__hip_fatbin_wrapper:
	.long	1212764230                      # 0x48495046
	.long	1                               # 0x1
	.quad	__hip_fatbin
	.quad	0
	.size	__hip_fatbin_wrapper, 24

	.type	__hip_gpubin_handle,@object     # @__hip_gpubin_handle
	.local	__hip_gpubin_handle
	.comm	__hip_gpubin_handle,8,8
	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	__hip_module_ctor
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym deviceFiFlag
	.addrsig_sym deviceFiInstCount
	.addrsig_sym deviceFiThreadIndex
	.addrsig_sym deviceSeedFactor
	.addrsig_sym deviceFiBit
	.addrsig_sym deviceBambooIndex
	.addrsig_sym currentCycle
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.protected	deviceFiFlag            ; @deviceFiFlag
	.type	deviceFiFlag,@object
	.section	.bss,"aw",@nobits
	.globl	deviceFiFlag
deviceFiFlag:
	.byte	0                               ; 0x0
	.size	deviceFiFlag, 1

	.protected	deviceFiInstCount       ; @deviceFiInstCount
	.type	deviceFiInstCount,@object
	.globl	deviceFiInstCount
	.p2align	3, 0x0
deviceFiInstCount:
	.quad	0                               ; 0x0
	.size	deviceFiInstCount, 8

	.protected	deviceFiThreadIndex     ; @deviceFiThreadIndex
	.type	deviceFiThreadIndex,@object
	.globl	deviceFiThreadIndex
	.p2align	3, 0x0
deviceFiThreadIndex:
	.quad	0                               ; 0x0
	.size	deviceFiThreadIndex, 8

	.protected	deviceSeedFactor        ; @deviceSeedFactor
	.type	deviceSeedFactor,@object
	.globl	deviceSeedFactor
	.p2align	3, 0x0
deviceSeedFactor:
	.quad	0x0000000000000000              ; double 0
	.size	deviceSeedFactor, 8

	.protected	deviceFiBit             ; @deviceFiBit
	.type	deviceFiBit,@object
	.globl	deviceFiBit
	.p2align	2, 0x0
deviceFiBit:
	.long	0                               ; 0x0
	.size	deviceFiBit, 4

	.protected	deviceBambooIndex       ; @deviceBambooIndex
	.type	deviceBambooIndex,@object
	.globl	deviceBambooIndex
	.p2align	3, 0x0
deviceBambooIndex:
	.quad	0                               ; 0x0
	.size	deviceBambooIndex, 8

	.protected	currentCycle            ; @currentCycle
	.type	currentCycle,@object
	.globl	currentCycle
	.p2align	3, 0x0
currentCycle:
	.quad	0                               ; 0x0
	.size	currentCycle, 8

	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym deviceFiFlag
	.addrsig_sym deviceFiInstCount
	.addrsig_sym deviceFiThreadIndex
	.addrsig_sym deviceSeedFactor
	.addrsig_sym deviceFiBit
	.addrsig_sym deviceBambooIndex
	.addrsig_sym currentCycle
	.addrsig_sym __hip_cuid_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
