#include<stdio.h>
//gpu code
__global__ void square(float * d_out,float *d_in) // arguments are pointers to input and the return is a value to the pointer in the argument list
{
    int idx=threadIdx.x;
    float f= d_in[idx];
    d_out[idx]=f*f;
}
//cpu code
int main()
{
    //generate input array
    const int ARRAY_SIZE=64;
    const int ARRAY_BYTES=ARRAY_SIZE * sizeof(float);
    float h_in[ARRAY_SIZE];
    for (int i=0;i<64;i++)
    {
        h_in[i]=float(i);
    }
    float h_out[ARRAY_SIZE];
    //declare gpu pointers
    float * d_in;
    float * d_out;
    //allocate gpu memory
    //The cudamalloc function returns an integer(instead of pointer) as error code to the memory block
    //All CUDA API function follows the convention of returning an integer error code 
    //when cudamalloc is called, a local variable named d_array is created and assigned with the value of the first function argument.
    //There is no way we can retrieve the value in that local variable outside the function's scope. That why we need to a pointer to a pointer here.
    cudaMalloc((void **) &d_in,ARRAY_BYTES);
    cudaMalloc((void **) &d_out,ARRAY_BYTES);
    //Copy the input array from cpu to gpu
    cudaMemcpy(d_in,h_in,ARRAY_BYTES,cudaMemcpyHostToDevice);
    //Launch The Kernal
    square<<<1,ARRAY_SIZE>>>(d_out,d_in); // launching on one Block with 64 elements 
    //copy back the result array to the CPU
    cudaMemcpy(h_out,d_out,ARRAY_BYTES,cudaMemcpyDeviceToHost);
    //print
    for (int i=0;i<64;i++)
    {
        printf("%f",h_out[i]);
        printf(((i%4)!=3)?"\t":"\n");
    }
    //free the gpu memory 
    cudaFree(d_in);
    cudaFree(d_out);

    return 0;


}
// output console
/*
256.000000      289.000000      324.000000      361.000000
400.000000      441.000000      484.000000      529.000000
576.000000      625.000000      676.000000      729.000000
784.000000      841.000000      900.000000      961.000000
1024.000000     1089.000000     1156.000000     1225.000000
1296.000000     1369.000000     1444.000000     1521.000000
1600.000000     1681.000000     1764.000000     1849.000000
1936.000000     2025.000000     2116.000000     2209.000000
2304.000000     2401.000000     2500.000000     2601.000000
2704.000000     2809.000000     2916.000000     3025.000000
3136.000000     3249.000000     3364.000000     3481.000000
3600.000000     3721.000000     3844.000000     3969.000000
*/
	.file	"tmpxft_002505a7_00000000-6_square.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	_Z27__device_stub__Z6squarePfS_PfS_
	.type	_Z27__device_stub__Z6squarePfS_PfS_, @function
_Z27__device_stub__Z6squarePfS_PfS_:
.LFB2052:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z6squarePfS_(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L2:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2052:
	.size	_Z27__device_stub__Z6squarePfS_PfS_, .-_Z27__device_stub__Z6squarePfS_PfS_
	.globl	_Z6squarePfS_
	.type	_Z6squarePfS_, @function
_Z6squarePfS_:
.LFB2053:
	.cfi_startproc
	endbr64
	jmp	_Z27__device_stub__Z6squarePfS_PfS_
	.cfi_endproc
.LFE2053:
	.size	_Z6squarePfS_, .-_Z6squarePfS_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"\t"
.LC1:
	.string	"\n"
.LC2:
	.string	"%f"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	subq	$560, %rsp
	.cfi_def_cfa_offset 592
	movq	%fs:40, %rax
	movq	%rax, 552(%rsp)
	xorl	%eax, %eax
.L9:
	cvtsi2ssl	%eax, %xmm0
	leaq	40(%rsp), %rbx
	movss	%xmm0, (%rbx,%rax,4)
	incq	%rax
	cmpq	$64, %rax
	jne	.L9
	movq	%rsp, %rdi
	movl	$256, %esi
	call	cudaMalloc@PLT
	leaq	8(%rsp), %rdi
	movl	$256, %esi
	call	cudaMalloc@PLT
	movq	(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rsi
	movl	$256, %edx
	call	cudaMemcpy@PLT
	movl	$67108865, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	salq	$6, %rdx
	movl	$1, %ecx
	movabsq	$4294967297, %rdi
	movl	$1, %esi
	movq	%rdx, 28(%rsp)
	movl	$1, 36(%rsp)
	movq	%rdi, 16(%rsp)
	movl	$1, 24(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L10
	movq	(%rsp), %rsi
	movq	8(%rsp), %rdi
	call	_Z27__device_stub__Z6squarePfS_PfS_
.L10:
	movq	8(%rsp), %rsi
	leaq	296(%rsp), %rbp
	movl	$2, %ecx
	xorl	%ebx, %ebx
	movl	$256, %edx
	movq	%rbp, %rdi
	leaq	.LC2(%rip), %r12
	call	cudaMemcpy@PLT
.L12:
	movq	%r12, %rsi
	movl	$2, %edi
	movb	$1, %al
	cvtss2sd	0(%rbp,%rbx,4), %xmm0
	call	__printf_chk@PLT
	movl	%ebx, %eax
	movl	$2, %edi
	leaq	.LC1(%rip), %rsi
	notl	%eax
	testb	$3, %al
	leaq	.LC0(%rip), %rax
	cmovne	%rax, %rsi
	xorl	%eax, %eax
	incq	%rbx
	call	__printf_chk@PLT
	cmpq	$64, %rbx
	jne	.L12
	movq	(%rsp), %rdi
	call	cudaFree@PLT
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	552(%rsp), %rax
	subq	%fs:40, %rax
	je	.L13
	call	__stack_chk_fail@PLT
.L13:
	addq	$560, %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
.LFE2027:
	.size	main, .-main
	.section	.rodata.str1.1
.LC3:
	.string	"_Z6squarePfS_"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2055:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC3(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z6squarePfS_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2055:
	.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<stdio.h>
//gpu code
__global__ void square(float * d_out,float *d_in) // arguments are pointers to input and the return is a value to the pointer in the argument list
{
    int idx=threadIdx.x;
    float f= d_in[idx];
    d_out[idx]=f*f;
}
//cpu code
int main()
{
    //generate input array
    const int ARRAY_SIZE=64;
    const int ARRAY_BYTES=ARRAY_SIZE * sizeof(float);
    float h_in[ARRAY_SIZE];
    for (int i=0;i<64;i++)
    {
        h_in[i]=float(i);
    }
    float h_out[ARRAY_SIZE];
    //declare gpu pointers
    float * d_in;
    float * d_out;
    //allocate gpu memory
    //The cudamalloc function returns an integer(instead of pointer) as error code to the memory block
    //All CUDA API function follows the convention of returning an integer error code 
    //when cudamalloc is called, a local variable named d_array is created and assigned with the value of the first function argument.
    //There is no way we can retrieve the value in that local variable outside the function's scope. That why we need to a pointer to a pointer here.
    hipMalloc((void **) &d_in,ARRAY_BYTES);
    hipMalloc((void **) &d_out,ARRAY_BYTES);
    //Copy the input array from cpu to gpu
    hipMemcpy(d_in,h_in,ARRAY_BYTES,hipMemcpyHostToDevice);
    //Launch The Kernal
    square<<<1,ARRAY_SIZE>>>(d_out,d_in); // launching on one Block with 64 elements 
    //copy back the result array to the CPU
    hipMemcpy(h_out,d_out,ARRAY_BYTES,hipMemcpyDeviceToHost);
    //print
    for (int i=0;i<64;i++)
    {
        printf("%f",h_out[i]);
        printf(((i%4)!=3)?"\t":"\n");
    }
    //free the gpu memory 
    hipFree(d_in);
    hipFree(d_out);

    return 0;


}
// output console
/*
256.000000      289.000000      324.000000      361.000000
400.000000      441.000000      484.000000      529.000000
576.000000      625.000000      676.000000      729.000000
784.000000      841.000000      900.000000      961.000000
1024.000000     1089.000000     1156.000000     1225.000000
1296.000000     1369.000000     1444.000000     1521.000000
1600.000000     1681.000000     1764.000000     1849.000000
1936.000000     2025.000000     2116.000000     2209.000000
2304.000000     2401.000000     2500.000000     2601.000000
2704.000000     2809.000000     2916.000000     3025.000000
3136.000000     3249.000000     3364.000000     3481.000000
3600.000000     3721.000000     3844.000000     3969.000000
*/
	.text
	.file	"square.hip"
	.globl	_Z21__device_stub__squarePfS_   # -- Begin function _Z21__device_stub__squarePfS_
	.type	_Z21__device_stub__squarePfS_,@function
_Z21__device_stub__squarePfS_:          # @_Z21__device_stub__squarePfS_
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$80, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	16(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%rsp, %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z6squarePfS_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$96, %rsp
	.cfi_adjust_cfa_offset -96
	popq	%rbx
	.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
	retq
.Lfunc_end0:
	.size	_Z21__device_stub__squarePfS_, .Lfunc_end0-_Z21__device_stub__squarePfS_
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	subq	$536, %rsp                      # imm = 0x218
	.cfi_def_cfa_offset 560
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	xorl	%eax, %eax
.LBB1_1:                                # =>This Inner Loop Header: Depth=1
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	movss	%xmm0, 272(%rsp,%rax,4)
	incq	%rax
	cmpq	$64, %rax
	jne	.LBB1_1
# %bb.2:
	leaq	8(%rsp), %rbx
	movl	$256, %esi                      # imm = 0x100
	movq	%rbx, %rdi
	callq	hipMalloc
	movq	%rsp, %rdi
	movl	$256, %esi                      # imm = 0x100
	callq	hipMalloc
	movq	(%rbx), %rdi
	leaq	272(%rsp), %rsi
	movl	$256, %edx                      # imm = 0x100
	movl	$1, %ecx
	callq	hipMemcpy
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	leaq	63(%rdi), %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_4
# %bb.3:
	movq	(%rsp), %rdi
	movq	8(%rsp), %rsi
	callq	_Z21__device_stub__squarePfS_
.LBB1_4:
	movq	(%rsp), %rsi
	leaq	16(%rsp), %rdi
	movl	$256, %edx                      # imm = 0x100
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$.L.str.2, %ebx
	xorl	%r14d, %r14d
.LBB1_5:                                # =>This Inner Loop Header: Depth=1
	xorps	%xmm0, %xmm0
	cvtss2sd	16(%rsp,%r14,4), %xmm0
	movl	$.L.str, %edi
	movb	$1, %al
	callq	printf
	movl	%r14d, %eax
	notl	%eax
	testb	$3, %al
	movl	$.L.str.1, %edi
	cmoveq	%rbx, %rdi
	xorl	%eax, %eax
	callq	printf
	incq	%r14
	cmpq	$64, %r14
	jne	.LBB1_5
# %bb.6:
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	(%rsp), %rdi
	callq	hipFree
	xorl	%eax, %eax
	addq	$536, %rsp                      # imm = 0x218
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	main, .Lfunc_end1-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB2_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB2_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z6squarePfS_, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end2:
	.size	__hip_module_ctor, .Lfunc_end2-__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	.LBB3_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
.LBB3_2:
	retq
.Lfunc_end3:
	.size	__hip_module_dtor, .Lfunc_end3-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z6squarePfS_,@object           # @_Z6squarePfS_
	.section	.rodata,"a",@progbits
	.globl	_Z6squarePfS_
	.p2align	3, 0x0
_Z6squarePfS_:
	.quad	_Z21__device_stub__squarePfS_
	.size	_Z6squarePfS_, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"%f"
	.size	.L.str, 3

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"\t"
	.size	.L.str.1, 2

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"\n"
	.size	.L.str.2, 2

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z6squarePfS_"
	.size	.L__unnamed_1, 14

	.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 _Z21__device_stub__squarePfS_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z6squarePfS_
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z6squarePfS_           ; -- Begin function _Z6squarePfS_
	.globl	_Z6squarePfS_
	.p2align	8
	.type	_Z6squarePfS_,@function
_Z6squarePfS_:                          ; @_Z6squarePfS_
; %bb.0:
	s_load_b128 s[0:3], s[0:1], 0x0
	v_lshlrev_b32_e32 v0, 2, v0
	s_waitcnt lgkmcnt(0)
	global_load_b32 v1, v0, s[2:3]
	s_waitcnt vmcnt(0)
	v_mul_f32_e32 v1, v1, v1
	global_store_b32 v0, v1, s[0:1]
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z6squarePfS_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 16
		.amdhsa_user_sgpr_count 15
		.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_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 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 4
		.amdhsa_reserve_vcc 0
		.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_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 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	_Z6squarePfS_, .Lfunc_end0-_Z6squarePfS_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 52
; NumSgprs: 4
; NumVgprs: 2
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 0
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 4
; NumVGPRsForWavesPerEU: 2
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; 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
	.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_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 16
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z6squarePfS_
    .private_segment_fixed_size: 0
    .sgpr_count:     4
    .sgpr_spill_count: 0
    .symbol:         _Z6squarePfS_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     2
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
