/*
    Copyright 2018 - The OPRECOMP Project Consortium, Alma Mater Studiorum
    Università di Bologna. All rights reserved.

    Licensed under the Apache License, Version 2.0 (the "License");
    you may not use this file except in compliance with the License.
    You may obtain a copy of the License at

        http://www.apache.org/licenses/LICENSE-2.0

    Unless required by applicable law or agreed to in writing, software
    distributed under the License is distributed on an "AS IS" BASIS,
    WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
    See the License for the specific language governing permissions and
    limitations under the License.
*/

#include <stdio.h>
#include <cuda_fp16.h>

#define SIZE 1000000

#define nTPB 256

#ifndef FLOAT
#define FLOAT double
#endif

__global__ void saxpy(int n, FLOAT a, FLOAT *x, FLOAT *y)
{
  int idx = threadIdx.x+blockDim.x*blockIdx.x;
  if (idx < n)
  {
    y[idx] = a * x[idx] + y[idx];
  }
}

int main()
{

  FLOAT *hin, *hout, *din, *dout;
  hin  = (FLOAT *)malloc(SIZE*sizeof(FLOAT));
  hout = (FLOAT *)malloc(SIZE*sizeof(FLOAT));
  for (int i = 0; i < SIZE; i++) hin[i] = i%15;
  for (int i = 0; i < SIZE; i++) hout[i] = i%15;
  cudaMalloc(&din,  SIZE*sizeof(FLOAT));
  cudaMalloc(&dout, SIZE*sizeof(FLOAT));
  cudaMemcpy(din, hin, SIZE*sizeof(FLOAT), cudaMemcpyHostToDevice);
  cudaMemcpy(dout, hout, SIZE*sizeof(FLOAT), cudaMemcpyHostToDevice);

  int k;
  for(k=0; k<5; ++k)
    saxpy<<<(SIZE+nTPB-1)/nTPB,nTPB>>>(SIZE, 0.5124353, din, dout);
  cudaMemcpy(hout, dout, SIZE*sizeof(FLOAT), cudaMemcpyDeviceToHost);
  for (int i = 0; i < SIZE; i++)
    printf("%f,", hout[i]);
  return 0;
}


	.file	"tmpxft_002e4a01_00000000-6_saxpy_ref.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2404:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2404:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z28__device_stub__Z5saxpyidPdS_idPdS_
	.type	_Z28__device_stub__Z5saxpyidPdS_idPdS_, @function
_Z28__device_stub__Z5saxpyidPdS_idPdS_:
.LFB2426:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movl	%edi, 28(%rsp)
	leaq	48(%rsp), %rcx
	leaq	56(%rsp), %rdi
	movq	%rsi, 8(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, (%rsp)
	leaq	40(%rsp), %rdx
	movsd	%xmm0, 16(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	%rsp, %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z5saxpyidPdS_(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 176
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L2:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2426:
	.size	_Z28__device_stub__Z5saxpyidPdS_idPdS_, .-_Z28__device_stub__Z5saxpyidPdS_idPdS_
	.globl	_Z5saxpyidPdS_
	.type	_Z5saxpyidPdS_, @function
_Z5saxpyidPdS_:
.LFB2427:
	.cfi_startproc
	endbr64
	jmp	_Z28__device_stub__Z5saxpyidPdS_idPdS_
	.cfi_endproc
.LFE2427:
	.size	_Z5saxpyidPdS_, .-_Z5saxpyidPdS_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"%f,"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2401:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movl	$8000000, %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
	movl	$8000000, %edi
	movq	%rax, %rbp
	call	malloc@PLT
	xorl	%ecx, %ecx
	movl	$15, %esi
	movq	%rax, %rbx
.L9:
	movl	%ecx, %eax
	cltd
	idivl	%esi
	cvtsi2sdl	%edx, %xmm0
	movsd	%xmm0, 0(%rbp,%rcx,8)
	incq	%rcx
	cmpq	$1000000, %rcx
	jne	.L9
	xorl	%ecx, %ecx
	movl	$15, %esi
.L10:
	movl	%ecx, %eax
	cltd
	idivl	%esi
	cvtsi2sdl	%edx, %xmm0
	movsd	%xmm0, (%rbx,%rcx,8)
	incq	%rcx
	cmpq	$1000000, %rcx
	jne	.L10
	movq	%rsp, %rdi
	movl	$8000000, %esi
	movabsq	$4294967552, %r12
	call	cudaMalloc@PLT
	leaq	8(%rsp), %rdi
	movl	$8000000, %esi
	call	cudaMalloc@PLT
	movq	(%rsp), %rdi
	movq	%rbp, %rsi
	movl	$1, %ecx
	movl	$8000000, %edx
	movl	$5, %ebp
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rsi
	movl	$8000000, %edx
	call	cudaMemcpy@PLT
.L12:
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%r12, %rdx
	movl	$1, %ecx
	movabsq	$4294971203, %rdi
	movl	$1, %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L11
	movq	8(%rsp), %rdx
	movq	(%rsp), %rsi
	movl	$1000000, %edi
	movsd	.LC0(%rip), %xmm0
	call	_Z28__device_stub__Z5saxpyidPdS_idPdS_
.L11:
	decl	%ebp
	jne	.L12
	movq	8(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdi
	xorl	%ebp, %ebp
	movl	$8000000, %edx
	leaq	.LC1(%rip), %r12
	call	cudaMemcpy@PLT
.L13:
	movsd	(%rbx,%rbp,8), %xmm0
	movq	%r12, %rsi
	movl	$2, %edi
	movb	$1, %al
	incq	%rbp
	call	__printf_chk@PLT
	cmpq	$1000000, %rbp
	jne	.L13
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L14
	call	__stack_chk_fail@PLT
.L14:
	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
.LFE2401:
	.size	main, .-main
	.section	.rodata.str1.1
.LC2:
	.string	"_Z5saxpyidPdS_"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2429:
	.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	.LC2(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z5saxpyidPdS_(%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
.LFE2429:
	.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
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC0:
	.long	-1227219903
	.long	1071670750
	.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:

/*
    Copyright 2018 - The OPRECOMP Project Consortium, Alma Mater Studiorum
    Università di Bologna. All rights reserved.

    Licensed under the Apache License, Version 2.0 (the "License");
    you may not use this file except in compliance with the License.
    You may obtain a copy of the License at

        http://www.apache.org/licenses/LICENSE-2.0

    Unless required by applicable law or agreed to in writing, software
    distributed under the License is distributed on an "AS IS" BASIS,
    WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
    See the License for the specific language governing permissions and
    limitations under the License.
*/


#include <hip/hip_runtime.h>
#include <stdio.h>
#include <hip/hip_fp16.h>

#define SIZE 1000000

#define nTPB 256

#ifndef FLOAT
#define FLOAT double
#endif

__global__ void saxpy(int n, FLOAT a, FLOAT *x, FLOAT *y)
{
  int idx = threadIdx.x+blockDim.x*blockIdx.x;
  if (idx < n)
  {
    y[idx] = a * x[idx] + y[idx];
  }
}

int main()
{

  FLOAT *hin, *hout, *din, *dout;
  hin  = (FLOAT *)malloc(SIZE*sizeof(FLOAT));
  hout = (FLOAT *)malloc(SIZE*sizeof(FLOAT));
  for (int i = 0; i < SIZE; i++) hin[i] = i%15;
  for (int i = 0; i < SIZE; i++) hout[i] = i%15;
  hipMalloc(&din,  SIZE*sizeof(FLOAT));
  hipMalloc(&dout, SIZE*sizeof(FLOAT));
  hipMemcpy(din, hin, SIZE*sizeof(FLOAT), hipMemcpyHostToDevice);
  hipMemcpy(dout, hout, SIZE*sizeof(FLOAT), hipMemcpyHostToDevice);

  int k;
  for(k=0; k<5; ++k)
    saxpy<<<(SIZE+nTPB-1)/nTPB,nTPB>>>(SIZE, 0.5124353, din, dout);
  hipMemcpy(hout, dout, SIZE*sizeof(FLOAT), hipMemcpyDeviceToHost);
  for (int i = 0; i < SIZE; i++)
    printf("%f,", hout[i]);
  return 0;
}


	.text
	.file	"saxpy_ref.hip"
	.globl	_Z20__device_stub__saxpyidPdS_  # -- Begin function _Z20__device_stub__saxpyidPdS_
	.type	_Z20__device_stub__saxpyidPdS_,@function
_Z20__device_stub__saxpyidPdS_:         # @_Z20__device_stub__saxpyidPdS_
	.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	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	4(%rsp), %rax
	movl	%edi, (%rax)
	leaq	40(%rsp), %rcx
	movsd	%xmm0, (%rcx)
	leaq	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rdi, 16(%rbx)
	movq	%rsi, 24(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%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	$_Z5saxpyidPdS_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	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	_Z20__device_stub__saxpyidPdS_, .Lfunc_end0-_Z20__device_stub__saxpyidPdS_
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI1_0:
	.quad	0x3fe065deb6da1c41              # double 0.51243530000000004
	.text
	.globl	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	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	$8000000, %edi                  # imm = 0x7A1200
	callq	malloc
	movq	%rax, %r14
	movl	$8000000, %edi                  # imm = 0x7A1200
	callq	malloc
	movq	%rax, %rbx
	xorl	%eax, %eax
	movl	$2290649225, %ecx               # imm = 0x88888889
.LBB1_1:                                # =>This Inner Loop Header: Depth=1
	movl	%eax, %edx
	imulq	%rcx, %rdx
	shrq	$35, %rdx
	leal	(%rdx,%rdx,4), %edx
	leal	(%rdx,%rdx,2), %edx
	movl	%eax, %esi
	subl	%edx, %esi
	xorps	%xmm0, %xmm0
	cvtsi2sd	%esi, %xmm0
	movsd	%xmm0, (%r14,%rax,8)
	incq	%rax
	cmpq	$1000000, %rax                  # imm = 0xF4240
	jne	.LBB1_1
# %bb.2:                                # %.preheader.preheader
	xorl	%eax, %eax
	movl	$2290649225, %ecx               # imm = 0x88888889
.LBB1_3:                                # %.preheader
                                        # =>This Inner Loop Header: Depth=1
	movl	%eax, %edx
	imulq	%rcx, %rdx
	shrq	$35, %rdx
	leal	(%rdx,%rdx,4), %edx
	leal	(%rdx,%rdx,2), %edx
	movl	%eax, %esi
	subl	%edx, %esi
	xorps	%xmm0, %xmm0
	cvtsi2sd	%esi, %xmm0
	movsd	%xmm0, (%rbx,%rax,8)
	incq	%rax
	cmpq	$1000000, %rax                  # imm = 0xF4240
	jne	.LBB1_3
# %bb.4:
	movabsq	$4294967552, %r15               # imm = 0x100000100
	leaq	16(%rsp), %r12
	movl	$8000000, %esi                  # imm = 0x7A1200
	movq	%r12, %rdi
	callq	hipMalloc
	leaq	8(%rsp), %r13
	movl	$8000000, %esi                  # imm = 0x7A1200
	movq	%r13, %rdi
	callq	hipMalloc
	movq	(%r12), %rdi
	movl	$8000000, %edx                  # imm = 0x7A1200
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r13), %rdi
	movl	$8000000, %edx                  # imm = 0x7A1200
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$5, %ebp
	leaq	3651(%r15), %r14
.LBB1_5:                                # =>This Inner Loop Header: Depth=1
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_7
# %bb.6:                                #   in Loop: Header=BB1_5 Depth=1
	movq	16(%rsp), %rsi
	movq	8(%rsp), %rdx
	movl	$1000000, %edi                  # imm = 0xF4240
	movsd	.LCPI1_0(%rip), %xmm0           # xmm0 = mem[0],zero
	callq	_Z20__device_stub__saxpyidPdS_
.LBB1_7:                                #   in Loop: Header=BB1_5 Depth=1
	decl	%ebp
	jne	.LBB1_5
# %bb.8:
	movq	8(%rsp), %rsi
	movl	$8000000, %edx                  # imm = 0x7A1200
	movq	%rbx, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	xorl	%r14d, %r14d
.LBB1_9:                                # =>This Inner Loop Header: Depth=1
	movsd	(%rbx,%r14,8), %xmm0            # xmm0 = mem[0],zero
	movl	$.L.str, %edi
	movb	$1, %al
	callq	printf
	incq	%r14
	cmpq	$1000000, %r14                  # imm = 0xF4240
	jne	.LBB1_9
# %bb.10:
	xorl	%eax, %eax
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.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
.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	$_Z5saxpyidPdS_, %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	_Z5saxpyidPdS_,@object          # @_Z5saxpyidPdS_
	.section	.rodata,"a",@progbits
	.globl	_Z5saxpyidPdS_
	.p2align	3, 0x0
_Z5saxpyidPdS_:
	.quad	_Z20__device_stub__saxpyidPdS_
	.size	_Z5saxpyidPdS_, 8

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

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z5saxpyidPdS_"
	.size	.L__unnamed_1, 15

	.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 _Z20__device_stub__saxpyidPdS_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z5saxpyidPdS_
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z5saxpyidPdS_          ; -- Begin function _Z5saxpyidPdS_
	.globl	_Z5saxpyidPdS_
	.p2align	8
	.type	_Z5saxpyidPdS_,@function
_Z5saxpyidPdS_:                         ; @_Z5saxpyidPdS_
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x2c
	s_load_b32 s3, s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_i32_e64 s3, v1
	s_cbranch_execz .LBB0_2
; %bb.1:
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x8
	s_load_b64 s[0:1], s[0:1], 0x18
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 3, v[1:2]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s6, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v3, vcc_lo, s7, v1, vcc_lo
	v_add_co_u32 v0, vcc_lo, s0, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_load_b64 v[2:3], v[2:3], off
	global_load_b64 v[4:5], v[0:1], off
	s_waitcnt vmcnt(0)
	v_fma_f64 v[2:3], v[2:3], s[4:5], v[4:5]
	global_store_b64 v[0:1], v[2:3], off
.LBB0_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z5saxpyidPdS_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.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 6
		.amdhsa_next_free_sgpr 16
		.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	_Z5saxpyidPdS_, .Lfunc_end0-_Z5saxpyidPdS_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 176
; NumSgprs: 18
; NumVgprs: 6
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 6
; 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:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .offset:         8
        .size:           8
        .value_kind:     by_value
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z5saxpyidPdS_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z5saxpyidPdS_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     6
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
