#include "triangle.cuh"
#include <iostream>
#include <fstream>

using namespace std;

// point class
// This constructor will help us create an instance of v3 from binary data found in the STL file.
v3::v3(char* facet)
{
    float xx = *((float*)facet);
    float yy = *((float*)facet + 1);
    float zz = *((float*)facet + 2);

    x = double(xx);
    y = double(yy);
    z = double(zz);
}

v3::v3(double x, double y, double z) : x(x), y(y), z(z) {}

//triangle class
triangle::triangle(v3 p1, v3 p2, v3 p3) : p1(p1), p2(p2), p3(p3) {}

// util
__host__
void read_stl(string fname, vector <triangle>& v) {
    ifstream myFile(
        fname.c_str(), ios::in | ios::binary);

    char header_info[80] = "";
    char nTri[4];
    unsigned nTriLong;
    unsigned count;

    //read 80 byte header
    if (myFile) {
        myFile.read(header_info, 80);
        cout << "header: " << header_info << endl;
    }
    else {
        cout << "error" << endl;
    }

    //read 4-byte ulong
    if (myFile) {
        myFile.read(nTri, 4);
        nTriLong = *((unsigned*)nTri);
        cout << "Number of triangles in file: " << nTriLong << endl;
    }
    else {
        cout << "error" << endl;
    }

    count = 0;

    //now read in all the triangles
    for (int i = 0; i < nTriLong; i++) {
        char facet[50];
        if (myFile) {
            //read one 50-byte triangle
            myFile.read(facet, 50);
            //populate each point of the triangle
            //using v3::v3(char* bin);
            //Ignore triangles that are parallel to some pixel ray
            v3 norm(facet);
            if (norm.z == 0) continue;
            //facet + 12 skips the triangle's unit normal
            v3 p1(facet + 12);
            v3 p2(facet + 24);
            v3 p3(facet + 36);
            //add a new triangle to the array
            v.push_back(triangle(p1, p2, p3));
            count++;
        }
    }
    cout << "Number of triangles added: " << count << endl;

    return;

}
	.file	"tmpxft_0023b7fe_00000000-6_triangle.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4141:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4141:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.align 2
	.type	_ZNSolsEPFRSoS_E.isra.0, @function
_ZNSolsEPFRSoS_E.isra.0:
.LFB4895:
	.cfi_startproc
	jmp	*%rsi
	.cfi_endproc
.LFE4895:
	.size	_ZNSolsEPFRSoS_E.isra.0, .-_ZNSolsEPFRSoS_E.isra.0
	.align 2
	.globl	_ZN2v3C2EPc
	.type	_ZN2v3C2EPc, @function
_ZN2v3C2EPc:
.LFB4130:
	.cfi_startproc
	endbr64
	cvtss2sd	(%rsi), %xmm0
	movsd	%xmm0, (%rdi)
	cvtss2sd	4(%rsi), %xmm0
	movsd	%xmm0, 8(%rdi)
	cvtss2sd	8(%rsi), %xmm0
	movsd	%xmm0, 16(%rdi)
	ret
	.cfi_endproc
.LFE4130:
	.size	_ZN2v3C2EPc, .-_ZN2v3C2EPc
	.globl	_ZN2v3C1EPc
	.set	_ZN2v3C1EPc,_ZN2v3C2EPc
	.align 2
	.globl	_ZN2v3C2Eddd
	.type	_ZN2v3C2Eddd, @function
_ZN2v3C2Eddd:
.LFB4133:
	.cfi_startproc
	endbr64
	movsd	%xmm0, (%rdi)
	movsd	%xmm1, 8(%rdi)
	movsd	%xmm2, 16(%rdi)
	ret
	.cfi_endproc
.LFE4133:
	.size	_ZN2v3C2Eddd, .-_ZN2v3C2Eddd
	.globl	_ZN2v3C1Eddd
	.set	_ZN2v3C1Eddd,_ZN2v3C2Eddd
	.align 2
	.globl	_ZN8triangleC2E2v3S0_S0_
	.type	_ZN8triangleC2E2v3S0_S0_, @function
_ZN8triangleC2E2v3S0_S0_:
.LFB4136:
	.cfi_startproc
	endbr64
	movq	24(%rsp), %rax
	movups	8(%rsp), %xmm0
	movups	32(%rsp), %xmm1
	movups	56(%rsp), %xmm2
	movq	%rax, 16(%rdi)
	movq	48(%rsp), %rax
	movups	%xmm0, (%rdi)
	movq	%rax, 40(%rdi)
	movq	72(%rsp), %rax
	movups	%xmm1, 24(%rdi)
	movq	%rax, 64(%rdi)
	movups	%xmm2, 48(%rdi)
	ret
	.cfi_endproc
.LFE4136:
	.size	_ZN8triangleC2E2v3S0_S0_, .-_ZN8triangleC2E2v3S0_S0_
	.globl	_ZN8triangleC1E2v3S0_S0_
	.set	_ZN8triangleC1E2v3S0_S0_,_ZN8triangleC2E2v3S0_S0_
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4164:
	.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
.LFE4164:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.rodata._ZNSt6vectorI8triangleSaIS0_EE17_M_realloc_insertIJS0_EEEvN9__gnu_cxx17__normal_iteratorIPS0_S2_EEDpOT_.str1.1,"aMS",@progbits,1
.LC0:
	.string	"vector::_M_realloc_insert"
	.section	.text._ZNSt6vectorI8triangleSaIS0_EE17_M_realloc_insertIJS0_EEEvN9__gnu_cxx17__normal_iteratorIPS0_S2_EEDpOT_,"axG",@progbits,_ZNSt6vectorI8triangleSaIS0_EE17_M_realloc_insertIJS0_EEEvN9__gnu_cxx17__normal_iteratorIPS0_S2_EEDpOT_,comdat
	.align 2
	.weak	_ZNSt6vectorI8triangleSaIS0_EE17_M_realloc_insertIJS0_EEEvN9__gnu_cxx17__normal_iteratorIPS0_S2_EEDpOT_
	.type	_ZNSt6vectorI8triangleSaIS0_EE17_M_realloc_insertIJS0_EEEvN9__gnu_cxx17__normal_iteratorIPS0_S2_EEDpOT_, @function
_ZNSt6vectorI8triangleSaIS0_EE17_M_realloc_insertIJS0_EEEvN9__gnu_cxx17__normal_iteratorIPS0_S2_EEDpOT_:
.LFB4757:
	.cfi_startproc
	endbr64
	movabsq	$-8198552921648689607, %rax
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rsi, %rbx
	movq	%rdx, %rsi
	movabsq	$128102389400760775, %rdx
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	movq	8(%rdi), %r15
	movq	(%rdi), %r13
	movq	%r15, %rbp
	subq	%r13, %rbp
	sarq	$3, %rbp
	imulq	%rax, %rbp
	cmpq	%rdx, %rbp
	jne	.L9
	leaq	.LC0(%rip), %rdi
	call	_ZSt20__throw_length_errorPKc@PLT
.L9:
	cmpq	%r13, %r15
	movl	$1, %eax
	movq	%rbx, %rcx
	movq	%rdi, %r12
	cmovne	%rbp, %rax
	addq	%rax, %rbp
	setc	%al
	subq	%r13, %rcx
	movzbl	%al, %eax
	movq	%rcx, (%rsp)
	testq	%rax, %rax
	jne	.L21
	xorl	%r14d, %r14d
	testq	%rbp, %rbp
	je	.L14
	movabsq	$128102389400760775, %rax
	cmpq	%rax, %rbp
	cmova	%rax, %rbp
	jmp	.L13
.L21:
	movq	%rdx, %rbp
.L13:
	imulq	$72, %rbp, %rdi
	movq	%rsi, 8(%rsp)
	call	_Znwm@PLT
	movq	8(%rsp), %rsi
	movq	%rax, %r14
.L14:
	movq	(%rsp), %rdi
	movl	$18, %ecx
	movq	%r14, %rdx
	movq	%r13, %rax
	addq	%r14, %rdi
	rep movsl
.L15:
	addq	$72, %rdx
	cmpq	%rax, %rbx
	je	.L29
	movq	%rax, %rsi
	leaq	-72(%rdx), %rdi
	movl	$18, %ecx
	addq	$72, %rax
	rep movsl
	jmp	.L15
.L29:
	movabsq	$1024819115206086201, %rax
	subq	%r13, %rbx
	shrq	$3, %rbx
	imulq	%rax, %rbx
	movabsq	$2305843009213693951, %rax
	andq	%rax, %rbx
	imulq	$72, %rbx, %r9
	leaq	1(%rbx), %rax
	imulq	$72, %rax, %rax
	addq	%r13, %r9
	addq	%r14, %rax
	movq	%r9, %rdx
	movq	%rax, %r10
.L17:
	cmpq	%r15, %rdx
	je	.L30
	movq	%r10, %rdi
	movq	%rdx, %rsi
	movl	$18, %ecx
	addq	$72, %rdx
	rep movsl
	addq	$72, %r10
	jmp	.L17
.L30:
	movabsq	$1024819115206086201, %rcx
	subq	%r9, %rdx
	shrq	$3, %rdx
	imulq	%rcx, %rdx
	movabsq	$2305843009213693951, %rcx
	andq	%rcx, %rdx
	imulq	$72, %rdx, %rdx
	leaq	(%rax,%rdx), %rbx
	testq	%r13, %r13
	je	.L19
	movq	16(%r12), %rsi
	movq	%r13, %rdi
	subq	%r13, %rsi
	call	_ZdlPvm@PLT
.L19:
	imulq	$72, %rbp, %rbp
	movq	%r14, (%r12)
	movq	%rbx, 8(%r12)
	addq	%rbp, %r14
	movq	%r14, 16(%r12)
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4757:
	.size	_ZNSt6vectorI8triangleSaIS0_EE17_M_realloc_insertIJS0_EEEvN9__gnu_cxx17__normal_iteratorIPS0_S2_EEDpOT_, .-_ZNSt6vectorI8triangleSaIS0_EE17_M_realloc_insertIJS0_EEEvN9__gnu_cxx17__normal_iteratorIPS0_S2_EEDpOT_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"header: "
.LC2:
	.string	"error"
.LC3:
	.string	"Number of triangles in file: "
.LC5:
	.string	"Number of triangles added: "
	.text
	.globl	_Z8read_stlNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEERSt6vectorI8triangleSaIS6_EE
	.type	_Z8read_stlNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEERSt6vectorI8triangleSaIS6_EE, @function
_Z8read_stlNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEERSt6vectorI8triangleSaIS6_EE:
.LFB4138:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4138
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	$12, %edx
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	leaq	_ZSt4cout(%rip), %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rsi, %rbx
	subq	$760, %rsp
	.cfi_def_cfa_offset 816
	movq	(%rdi), %rsi
	movq	%fs:40, %rax
	movq	%rax, 744(%rsp)
	xorl	%eax, %eax
	leaq	88(%rsp), %rax
	movq	%rax, %rdi
	movq	%rax, 8(%rsp)
.LEHB0:
	call	_ZNSt14basic_ifstreamIcSt11char_traitsIcEEC1EPKcSt13_Ios_Openmode@PLT
.LEHE0:
	xorl	%eax, %eax
	movl	$20, %ecx
	leaq	664(%rsp), %rdi
	rep stosl
	testb	$5, 376(%rsp)
	jne	.L32
	leaq	664(%rsp), %rbp
	movq	8(%rsp), %rdi
	movl	$80, %edx
	movq	%rbp, %rsi
.LEHB1:
	call	_ZNSi4readEPcl@PLT
	leaq	.LC1(%rip), %rsi
	movq	%r12, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	movq	%rbp, %rsi
.L57:
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	testb	$5, 376(%rsp)
	je	.L33
	jmp	.L58
.L32:
	leaq	.LC2(%rip), %rsi
	movq	%r12, %rdi
	jmp	.L57
.L33:
	movq	8(%rsp), %rdi
	leaq	610(%rsp), %rsi
	movl	$4, %edx
	call	_ZNSi4readEPcl@PLT
	leaq	.LC3(%rip), %rsi
	movq	%r12, %rdi
	movl	610(%rsp), %r14d
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	movl	%r14d, %esi
	call	_ZNSo9_M_insertImEERSoT_@PLT
	jmp	.L55
.L58:
	leaq	.LC2(%rip), %rsi
	movq	%r12, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
.L55:
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	xorl	%r13d, %r13d
	xorl	%ebp, %ebp
	leaq	614(%rsp), %r15
.L37:
	cmpl	%r14d, %r13d
	je	.L59
	testb	$5, 376(%rsp)
	jne	.L39
	movq	8(%rsp), %rdi
	movl	$50, %edx
	movq	%r15, %rsi
	call	_ZNSi4readEPcl@PLT
	xorps	%xmm0, %xmm0
	ucomiss	622(%rsp), %xmm0
	jp	.L48
	je	.L39
.L48:
	movq	8(%rbx), %rax
	cmpq	16(%rbx), %rax
	cvtss2sd	626(%rsp), %xmm8
	cvtss2sd	630(%rsp), %xmm7
	cvtss2sd	634(%rsp), %xmm6
	movsd	%xmm8, 16(%rsp)
	cvtss2sd	638(%rsp), %xmm5
	cvtss2sd	642(%rsp), %xmm4
	cvtss2sd	646(%rsp), %xmm3
	movsd	%xmm7, 24(%rsp)
	cvtss2sd	650(%rsp), %xmm2
	cvtss2sd	654(%rsp), %xmm1
	cvtss2sd	658(%rsp), %xmm0
	movsd	%xmm6, 32(%rsp)
	leaq	16(%rsp), %rsi
	movsd	%xmm5, 40(%rsp)
	movsd	%xmm4, 48(%rsp)
	movsd	%xmm3, 56(%rsp)
	movsd	%xmm2, 64(%rsp)
	movsd	%xmm1, 72(%rsp)
	movsd	%xmm0, 80(%rsp)
	je	.L41
	movq	%rax, %rdi
	addq	$72, %rax
	movl	$18, %ecx
	rep movsl
	movq	%rax, 8(%rbx)
	jmp	.L42
.L41:
	movq	%rsi, %rdx
	movq	%rbx, %rdi
	movq	%rax, %rsi
	call	_ZNSt6vectorI8triangleSaIS0_EE17_M_realloc_insertIJS0_EEEvN9__gnu_cxx17__normal_iteratorIPS0_S2_EEDpOT_
.L42:
	incl	%ebp
.L39:
	incl	%r13d
	jmp	.L37
.L59:
	leaq	.LC5(%rip), %rsi
	movq	%r12, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	movl	%ebp, %esi
	call	_ZNSo9_M_insertImEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
.LEHE1:
	movq	8(%rsp), %rdi
	call	_ZNSt14basic_ifstreamIcSt11char_traitsIcEED1Ev@PLT
	movq	744(%rsp), %rax
	subq	%fs:40, %rax
	je	.L46
	jmp	.L50
.L47:
	endbr64
	movq	%rax, %rbx
.L44:
	movq	8(%rsp), %rdi
	call	_ZNSt14basic_ifstreamIcSt11char_traitsIcEED1Ev@PLT
	movq	744(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L50
	movq	%rbx, %rdi
.LEHB2:
	call	_Unwind_Resume@PLT
.LEHE2:
.L50:
	call	__stack_chk_fail@PLT
.L46:
	addq	$760, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4138:
	.globl	__gxx_personality_v0
	.section	.gcc_except_table,"a",@progbits
.LLSDA4138:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4138-.LLSDACSB4138
.LLSDACSB4138:
	.uleb128 .LEHB0-.LFB4138
	.uleb128 .LEHE0-.LEHB0
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB1-.LFB4138
	.uleb128 .LEHE1-.LEHB1
	.uleb128 .L47-.LFB4138
	.uleb128 0
	.uleb128 .LEHB2-.LFB4138
	.uleb128 .LEHE2-.LEHB2
	.uleb128 0
	.uleb128 0
.LLSDACSE4138:
	.text
	.size	_Z8read_stlNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEERSt6vectorI8triangleSaIS6_EE, .-_Z8read_stlNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEERSt6vectorI8triangleSaIS6_EE
	.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
	.hidden	DW.ref.__gxx_personality_v0
	.weak	DW.ref.__gxx_personality_v0
	.section	.data.rel.local.DW.ref.__gxx_personality_v0,"awG",@progbits,DW.ref.__gxx_personality_v0,comdat
	.align 8
	.type	DW.ref.__gxx_personality_v0, @object
	.size	DW.ref.__gxx_personality_v0, 8
DW.ref.__gxx_personality_v0:
	.quad	__gxx_personality_v0
	.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:

//source: http://www.sgh1.net/posts/read-stl-file.md
#ifndef TRIANGLE
#define TRIANGLE

#ifdef __HIPCC__
#define CUDA_CALLABLE_MEMBER __host__ __device__
#else
#define CUDA_CALLABLE_MEMBER
#endif 

#include <string>
#include <vector>

using std::string;
using std::vector;

typedef
struct v3
{
public:
    __host__ __device__ v3(char* bin);
    __host__ __device__ v3(double x = 0.0, double y = 0.0, double z = 0.0);
    // ~v3();

    double x, y, z;
} v3;

typedef
struct triangle
{
public:
    __host__ __device__ triangle(v3 p1 = v3(), v3 p2 = v3(), v3 p3 = v3());
    // ~triangle();
    v3 p1, p2, p3;
} triangle;

// utils
__host__ void read_stl(string fname, vector<triangle>& v);

#endif

	.text
	.file	"triangle.hip"
	.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_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
