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

#define DATASIZE 32768
#define BATCH 16384

/********************/
/* CUDA ERROR CHECK */
/********************/
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

/********/
/* MAIN */
/********/
int main ()
{
    // --- Host side input data allocation and initialization
    cufftReal *hostInputData = (cufftReal*)malloc(DATASIZE*BATCH*sizeof(cufftReal));
    for (int i=0; i<BATCH; i++)
        for (int j=0; j<DATASIZE; j++) hostInputData[i*DATASIZE + j] = (cufftReal)(i + 1);

    // --- Device side input data allocation and initialization
    cufftReal *deviceInputData; gpuErrchk(cudaMalloc((void**)&deviceInputData, DATASIZE * BATCH * sizeof(cufftReal)));
    cudaMemcpy(deviceInputData, hostInputData, DATASIZE * BATCH * sizeof(cufftReal), cudaMemcpyHostToDevice);

    // --- Host side output data allocation
    cufftComplex *hostOutputData = (cufftComplex*)malloc((DATASIZE / 2 + 1) * BATCH * sizeof(cufftComplex));

    // --- Device side output data allocation
    cufftComplex *deviceOutputData; gpuErrchk(cudaMalloc((void**)&deviceOutputData, (DATASIZE / 2 + 1) * BATCH * sizeof(cufftComplex)));

    // --- Batched 1D FFTs
    cufftHandle handle;
    int rank = 1;                           // --- 1D FFTs
    int n[] = { DATASIZE };                 // --- Size of the Fourier transform
    int istride = 1, ostride = 1;           // --- Distance between two successive input/output elements
    int idist = DATASIZE, odist = (DATASIZE / 2 + 1); // --- Distance between batches
    int inembed[] = { 0 };                  // --- Input size with pitch (ignored for 1D transforms)
    int onembed[] = { 0 };                  // --- Output size with pitch (ignored for 1D transforms)
    int batch = BATCH;                      // --- Number of batched executions
    cufftPlanMany(&handle, rank, n, 
                  inembed, istride, idist,
                  onembed, ostride, odist, CUFFT_R2C, batch);

    //cufftPlan1d(&handle, DATASIZE, CUFFT_R2C, BATCH);
    cufftExecR2C(handle,  deviceInputData, deviceOutputData);

    // --- Device->Host copy of the results
    gpuErrchk(cudaMemcpy(hostOutputData, deviceOutputData, (DATASIZE / 2 + 1) * BATCH * sizeof(cufftComplex), cudaMemcpyDeviceToHost));

    /*for (int i=0; i<BATCH; i++)
        for (int j=0; j<(DATASIZE / 2 + 1); j++)
            printf("%i %i %f %f\n", i, j, hostOutputData[i*(DATASIZE / 2 + 1) + j].x, hostOutputData[i*(DATASIZE / 2 + 1) + j].y);
    */
    printf("Done!\n");        
    cufftDestroy(handle);
    gpuErrchk(cudaFree(deviceOutputData));
    gpuErrchk(cudaFree(deviceInputData));

}
	.file	"tmpxft_00294dbe_00000000-6_fft_example.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2054:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2054:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.rodata._Z9gpuAssert9cudaErrorPKcib.str1.1,"aMS",@progbits,1
.LC0:
	.string	"GPUassert: %s %s %d\n"
	.section	.text._Z9gpuAssert9cudaErrorPKcib,"axG",@progbits,_Z9gpuAssert9cudaErrorPKcib,comdat
	.weak	_Z9gpuAssert9cudaErrorPKcib
	.type	_Z9gpuAssert9cudaErrorPKcib, @function
_Z9gpuAssert9cudaErrorPKcib:
.LFB2050:
	.cfi_startproc
	endbr64
	testl	%edi, %edi
	je	.L8
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movl	%ecx, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movl	%edi, %ebx
	subq	$24, %rsp
	.cfi_def_cfa_offset 48
	movl	%edx, 12(%rsp)
	movq	%rsi, (%rsp)
	call	cudaGetErrorString@PLT
	movl	12(%rsp), %r9d
	movq	(%rsp), %r8
	leaq	.LC0(%rip), %rdx
	movq	stderr(%rip), %rdi
	movq	%rax, %rcx
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	testb	%bpl, %bpl
	je	.L2
	movl	%ebx, %edi
	call	exit@PLT
.L2:
	addq	$24, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
.L8:
	.cfi_restore 3
	.cfi_restore 6
	ret
	.cfi_endproc
.LFE2050:
	.size	_Z9gpuAssert9cudaErrorPKcib, .-_Z9gpuAssert9cudaErrorPKcib
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"/home/ubuntu/Datasets/stackv2/train-structured/ekonyagin/wave_propagation/master/cpp/fft_example.cu"
.LC2:
	.string	"Done!\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2051:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movl	$2147483648, %edi
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	subq	$48, %rsp
	.cfi_def_cfa_offset 80
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	call	malloc@PLT
	movq	%rax, %rbp
	movq	%rax, %rcx
	xorl	%eax, %eax
.L13:
	incl	%eax
	xorl	%edx, %edx
	cvtsi2ssl	%eax, %xmm0
.L14:
	movss	%xmm0, (%rcx,%rdx,4)
	incq	%rdx
	cmpq	$32768, %rdx
	jne	.L14
	addq	$131072, %rcx
	cmpl	$16384, %eax
	jne	.L13
	movl	$2147483648, %r12d
	leaq	8(%rsp), %rdi
	leaq	.LC1(%rip), %rbx
	movq	%r12, %rsi
	call	cudaMalloc@PLT
	movl	$1, %ecx
	movl	$33, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z9gpuAssert9cudaErrorPKcib
	movq	8(%rsp), %rdi
	movq	%r12, %rdx
	movq	%rbp, %rsi
	movl	$1, %ecx
	addq	$131072, %r12
	call	cudaMemcpy@PLT
	movq	%r12, %rdi
	call	malloc@PLT
	leaq	16(%rsp), %rdi
	movq	%r12, %rsi
	movq	%rax, %rbp
	call	cudaMalloc@PLT
	movq	%rbx, %rsi
	movl	$1, %ecx
	movl	$40, %edx
	movl	%eax, %edi
	call	_Z9gpuAssert9cudaErrorPKcib
	xorl	%eax, %eax
	leaq	32(%rsp), %rcx
	movl	$32768, 28(%rsp)
	movl	%eax, 32(%rsp)
	leaq	28(%rsp), %rdx
	leaq	4(%rsp), %rdi
	movl	$32768, %r9d
	movl	%eax, 36(%rsp)
	movl	$1, %r8d
	pushq	%rsi
	.cfi_def_cfa_offset 88
	movl	$1, %esi
	pushq	$16384
	.cfi_def_cfa_offset 96
	pushq	$42
	.cfi_def_cfa_offset 104
	pushq	$16385
	.cfi_def_cfa_offset 112
	pushq	$1
	.cfi_def_cfa_offset 120
	leaq	76(%rsp), %rax
	pushq	%rax
	.cfi_def_cfa_offset 128
	call	cufftPlanMany@PLT
	movq	64(%rsp), %rdx
	movq	56(%rsp), %rsi
	addq	$48, %rsp
	.cfi_def_cfa_offset 80
	movl	4(%rsp), %edi
	call	cufftExecR2C@PLT
	movq	16(%rsp), %rsi
	movq	%r12, %rdx
	movq	%rbp, %rdi
	movl	$2, %ecx
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$59, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z9gpuAssert9cudaErrorPKcib
	leaq	.LC2(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	4(%rsp), %edi
	call	cufftDestroy@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movl	$1, %ecx
	movl	$67, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z9gpuAssert9cudaErrorPKcib
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movl	$1, %ecx
	movl	$68, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z9gpuAssert9cudaErrorPKcib
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L16
	call	__stack_chk_fail@PLT
.L16:
	addq	$48, %rsp
	.cfi_def_cfa_offset 32
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2051:
	.size	main, .-main
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2077:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	movq	%rax, %rdi
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2077:
	.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
	.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 <hip/hip_runtime.h>
#include <hipfft/hipfft.h>
#include <stdio.h>
#include <math.h>

#define DATASIZE 32768
#define BATCH 16384

/********************/
/* CUDA ERROR CHECK */
/********************/
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(hipError_t code, const char *file, int line, bool abort=true)
{
   if (code != hipSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", hipGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

/********/
/* MAIN */
/********/
int main ()
{
    // --- Host side input data allocation and initialization
    hipfftReal *hostInputData = (hipfftReal*)malloc(DATASIZE*BATCH*sizeof(hipfftReal));
    for (int i=0; i<BATCH; i++)
        for (int j=0; j<DATASIZE; j++) hostInputData[i*DATASIZE + j] = (hipfftReal)(i + 1);

    // --- Device side input data allocation and initialization
    hipfftReal *deviceInputData; gpuErrchk(hipMalloc((void**)&deviceInputData, DATASIZE * BATCH * sizeof(hipfftReal)));
    hipMemcpy(deviceInputData, hostInputData, DATASIZE * BATCH * sizeof(hipfftReal), hipMemcpyHostToDevice);

    // --- Host side output data allocation
    hipfftComplex *hostOutputData = (hipfftComplex*)malloc((DATASIZE / 2 + 1) * BATCH * sizeof(hipfftComplex));

    // --- Device side output data allocation
    hipfftComplex *deviceOutputData; gpuErrchk(hipMalloc((void**)&deviceOutputData, (DATASIZE / 2 + 1) * BATCH * sizeof(hipfftComplex)));

    // --- Batched 1D FFTs
    hipfftHandle handle;
    int rank = 1;                           // --- 1D FFTs
    int n[] = { DATASIZE };                 // --- Size of the Fourier transform
    int istride = 1, ostride = 1;           // --- Distance between two successive input/output elements
    int idist = DATASIZE, odist = (DATASIZE / 2 + 1); // --- Distance between batches
    int inembed[] = { 0 };                  // --- Input size with pitch (ignored for 1D transforms)
    int onembed[] = { 0 };                  // --- Output size with pitch (ignored for 1D transforms)
    int batch = BATCH;                      // --- Number of batched executions
    hipfftPlanMany(&handle, rank, n, 
                  inembed, istride, idist,
                  onembed, ostride, odist, HIPFFT_R2C, batch);

    //cufftPlan1d(&handle, DATASIZE, CUFFT_R2C, BATCH);
    hipfftExecR2C(handle,  deviceInputData, deviceOutputData);

    // --- Device->Host copy of the results
    gpuErrchk(hipMemcpy(hostOutputData, deviceOutputData, (DATASIZE / 2 + 1) * BATCH * sizeof(hipfftComplex), hipMemcpyDeviceToHost));

    /*for (int i=0; i<BATCH; i++)
        for (int j=0; j<(DATASIZE / 2 + 1); j++)
            printf("%i %i %f %f\n", i, j, hostOutputData[i*(DATASIZE / 2 + 1) + j].x, hostOutputData[i*(DATASIZE / 2 + 1) + j].y);
    */
    printf("Done!\n");        
    hipfftDestroy(handle);
    gpuErrchk(hipFree(deviceOutputData));
    gpuErrchk(hipFree(deviceInputData));

}
	.text
	.file	"fft_example.hip"
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	subq	$40, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	$2147483648, %ebx               # imm = 0x80000000
	movl	$2147483648, %edi               # imm = 0x80000000
	callq	malloc
	movq	%rax, %r14
	xorl	%eax, %eax
	movq	%r14, %rcx
.LBB0_1:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB0_2 Depth 2
	incq	%rax
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	xorl	%edx, %edx
.LBB0_2:                                #   Parent Loop BB0_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movss	%xmm0, (%rcx,%rdx,4)
	incq	%rdx
	cmpq	$32768, %rdx                    # imm = 0x8000
	jne	.LBB0_2
# %bb.3:                                #   in Loop: Header=BB0_1 Depth=1
	addq	$131072, %rcx                   # imm = 0x20000
	cmpq	$16384, %rax                    # imm = 0x4000
	jne	.LBB0_1
# %bb.4:
	leaq	8(%rsp), %rdi
	movl	$2147483648, %esi               # imm = 0x80000000
	callq	hipMalloc
	testl	%eax, %eax
	jne	.LBB0_5
# %bb.7:                                # %_Z9gpuAssert10hipError_tPKcib.exit
	movq	8(%rsp), %rdi
	movl	$2147483648, %edx               # imm = 0x80000000
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	addq	$131072, %rbx                   # imm = 0x20000
	movq	%rbx, %rdi
	callq	malloc
	movq	%rax, %r14
	movq	%rsp, %rdi
	movq	%rbx, %rsi
	callq	hipMalloc
	testl	%eax, %eax
	jne	.LBB0_8
# %bb.9:                                # %_Z9gpuAssert10hipError_tPKcib.exit18
	leaq	36(%rsp), %rdx
	movl	$32768, (%rdx)                  # imm = 0x8000
	xorl	%eax, %eax
	leaq	32(%rsp), %rcx
	movl	%eax, (%rcx)
	leaq	28(%rsp), %r10
	movl	%eax, (%r10)
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	leaq	24(%rsp), %r15
	movq	%r15, %rdi
	movl	$1, %esi
	movl	$1, %r8d
	movl	$32768, %r9d                    # imm = 0x8000
	pushq	$16384                          # imm = 0x4000
	.cfi_adjust_cfa_offset 8
	pushq	$42
	.cfi_adjust_cfa_offset 8
	pushq	$16385                          # imm = 0x4001
	.cfi_adjust_cfa_offset 8
	pushq	$1
	.cfi_adjust_cfa_offset 8
	pushq	%r10
	.cfi_adjust_cfa_offset 8
	callq	hipfftPlanMany
	addq	$48, %rsp
	.cfi_adjust_cfa_offset -48
	movq	(%r15), %rdi
	movq	8(%rsp), %rsi
	movq	(%rsp), %rdx
	callq	hipfftExecR2C
	movq	(%rsp), %rsi
	movq	%r14, %rdi
	movq	%rbx, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	jne	.LBB0_10
# %bb.12:                               # %_Z9gpuAssert10hipError_tPKcib.exit20
	movl	$.Lstr, %edi
	callq	puts@PLT
	movq	16(%rsp), %rdi
	callq	hipfftDestroy
	movq	(%rsp), %rdi
	callq	hipFree
	testl	%eax, %eax
	jne	.LBB0_13
# %bb.14:                               # %_Z9gpuAssert10hipError_tPKcib.exit22
	movq	8(%rsp), %rdi
	callq	hipFree
	testl	%eax, %eax
	jne	.LBB0_15
# %bb.16:                               # %_Z9gpuAssert10hipError_tPKcib.exit24
	xorl	%eax, %eax
	addq	$40, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.LBB0_5:
	.cfi_def_cfa_offset 80
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.2, %esi
	movl	$.L.str, %ecx
	movq	%rbx, %rdi
	movq	%rax, %rdx
	movl	$33, %r8d
	jmp	.LBB0_6
.LBB0_8:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.2, %esi
	movl	$.L.str, %ecx
	movq	%rbx, %rdi
	movq	%rax, %rdx
	movl	$40, %r8d
.LBB0_6:
	xorl	%eax, %eax
	callq	fprintf
	movl	%ebp, %edi
	callq	exit
.LBB0_10:
	movl	%eax, %ebx
	movq	stderr(%rip), %r14
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.2, %esi
	movl	$.L.str, %ecx
	movq	%r14, %rdi
	movq	%rax, %rdx
	movl	$59, %r8d
	jmp	.LBB0_11
.LBB0_13:
	movl	%eax, %ebx
	movq	stderr(%rip), %r14
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.2, %esi
	movl	$.L.str, %ecx
	movq	%r14, %rdi
	movq	%rax, %rdx
	movl	$67, %r8d
	jmp	.LBB0_11
.LBB0_15:
	movl	%eax, %ebx
	movq	stderr(%rip), %r14
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.2, %esi
	movl	$.L.str, %ecx
	movq	%r14, %rdi
	movq	%rax, %rdx
	movl	$68, %r8d
.LBB0_11:
	xorl	%eax, %eax
	callq	fprintf
	movl	%ebx, %edi
	callq	exit
.Lfunc_end0:
	.size	main, .Lfunc_end0-main
	.cfi_endproc
                                        # -- End function
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"/home/ubuntu/Datasets/stackv2/train-structured-repos-hip-sm89/ekonyagin/wave_propagation/master/cpp/fft_example.hip"
	.size	.L.str, 116

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"GPUassert: %s %s %d\n"
	.size	.L.str.2, 21

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

	.type	.Lstr,@object                   # @str
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr:
	.asciz	"Done!"
	.size	.Lstr, 6

	.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_cuid_

	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.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_cuid_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
