/************************************************************
Known issues:
This program only works on matrices smaller than or equal to 
256x256. 1024x1024 will cause segmentation faults and 512x512
simply causes the program to almost crash and return times of
0 for each kernel call. 
The matrices must be square and all matrices must be the same
size.
*/

#include <stdio.h>
#include <stdlib.h>
#define MATSIZE 128
#define THREADS_PER_BLOCK 32 
 

 
//serial matrix multiplication kernel
__global__ void smultiply(int* g_a, int* g_b, int* g_c)
{
  int x, y, z;

  for (x = 0; x < MATSIZE; ++ x)
  {
    for (y = 0; y < MATSIZE; ++ y)
    {
      for (z = 0; z < MATSIZE; ++ z)
      {
	g_c[(x * MATSIZE) + y] += g_a[(x * MATSIZE) + z] * g_b[(z * MATSIZE) + y];
      }
    }
  }
}

//parallel matrix multiplication kernel
__global__ void pmultiply(double* g_a, double* g_b, double* g_d , int dim)
{
  int z;
  double sum = 0.0;
  int x = blockIdx.x * blockDim.x + threadIdx.x;
  int y = blockIdx.y * blockDim.y + threadIdx.y;

  for (z = 0; z < dim; ++ z)
  {
    sum += g_a[x * dim + z] * g_b[y + z * dim];
  }
    
  g_d[(x * dim) + y] = sum;
}

extern "C" void Cudamultiply(double* a, double* b, double* d, int Dim)
{
//  double *d;
  int i;
  double *g_a, *g_b, *g_d;
  int g_size = Dim * Dim * sizeof(double);
  cudaEvent_t start, stop;
  float time; 

//  d=new double[Dim*Dim];

  cudaEventCreate(&start);
  cudaEventCreate(&stop);
  //used for timing the Cuda run

  cudaMalloc(&g_a, g_size); //allocate memory on Cuda device
  cudaMemcpy(g_a, a, g_size, cudaMemcpyHostToDevice);
  //copy matrix A onto the Cuda device  

  cudaMalloc(&g_b, g_size);
  cudaMemcpy(g_b, b, g_size, cudaMemcpyHostToDevice);
  
//  cudaMalloc(&g_c, g_size);
//  cudaMemcpy(g_c, c, g_size, cudaMemcpyHostToDevice);

  dim3 dimGrid((Dim / THREADS_PER_BLOCK), (Dim / THREADS_PER_BLOCK));
  //create the needed number of grids
  dim3 threads(THREADS_PER_BLOCK, THREADS_PER_BLOCK);
  //create the needed number of threads in each grid

  //serial Cuda kernel call
  //cudaEventRecord(start, 0);
  //smultiply<<<1,1>>>(g_a, g_b, g_c);
  //cudaEventRecord(stop, 0);  

  //cudaEventSynchronize(stop);
  //cudaEventElapsedTime(&time, start, stop);
  //get run time

//  cudaMemcpy(c, g_c, g_size, cudaMemcpyDeviceToHost);
  //copy results back to host device
//  cudaFree(g_c);  
  //free up unused user allocated memory on Cuda device

  printf("Time = %f milliseconds\n", time);
  
  //create a second answer matrix to use
  //This is not done until now so that memory on the
  //Cuda device is not wasted.
  cudaMalloc(&g_d, g_size);
  cudaMemcpy(g_d, d, g_size, cudaMemcpyHostToDevice);

  //parallel Cuda kernel call
  cudaEventRecord(start, 0);
  pmultiply<<<dimGrid,threads>>>(g_a, g_b, g_d,Dim);
  cudaEventRecord(stop, 0);

  cudaEventSynchronize(stop);
  cudaEventElapsedTime(&time, start, stop);
  printf("Time = %f milliseconds\n", time);

  cudaMemcpy(d, g_d, g_size, cudaMemcpyDeviceToHost);

  cudaFree(g_a);
  cudaFree(g_b);
  cudaFree(g_d);
  //free up all unused user allocated memory on Cuda device

  //printf("\n");  

  /*The next 2 for loops print out the values of both
  answer matrices. This can be used to ensure that both
  kernel calls are producing the same results, and that
  the results are correct. This section can be commented 
  out when the user only wants the timing of a run.*/

//  for (i = 1; i <= Dim * Dim; ++ i)
//  {
//      if (i % Dim == 0)
//      {
////        printf("%f ", c[i-1]);
//	printf("\n");
//      }
//   
//      else
////        printf("%f ", c[i-1]);
//  }

 // printf("\n");

//  for (i = 1; i <= Dim * Dim; ++ i)
//  {
//    if ( i % Dim == 0)
//    {
//     printf("%f ", d[i-1]);
//     printf("\n");
//    }
//
//    else
//     printf("%f ", d[i-1]);
//  }
//  delete d;
}

	.file	"tmpxft_00329966_00000000-6_cudamatrix.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	_Z32__device_stub__Z9smultiplyPiS_S_PiS_S_
	.type	_Z32__device_stub__Z9smultiplyPiS_S_PiS_S_, @function
_Z32__device_stub__Z9smultiplyPiS_S_PiS_S_:
.LFB2052:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z9smultiplyPiS_S_(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L2:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2052:
	.size	_Z32__device_stub__Z9smultiplyPiS_S_PiS_S_, .-_Z32__device_stub__Z9smultiplyPiS_S_PiS_S_
	.globl	_Z9smultiplyPiS_S_
	.type	_Z9smultiplyPiS_S_, @function
_Z9smultiplyPiS_S_:
.LFB2053:
	.cfi_startproc
	endbr64
	jmp	_Z32__device_stub__Z9smultiplyPiS_S_PiS_S_
	.cfi_endproc
.LFE2053:
	.size	_Z9smultiplyPiS_S_, .-_Z9smultiplyPiS_S_
	.globl	_Z33__device_stub__Z9pmultiplyPdS_S_iPdS_S_i
	.type	_Z33__device_stub__Z9pmultiplyPdS_S_iPdS_S_i, @function
_Z33__device_stub__Z9pmultiplyPdS_S_iPdS_S_i:
.LFB2054:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 4(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%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)
	leaq	4(%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	.L8
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z9pmultiplyPdS_S_i(%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
.L8:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2054:
	.size	_Z33__device_stub__Z9pmultiplyPdS_S_iPdS_S_i, .-_Z33__device_stub__Z9pmultiplyPdS_S_iPdS_S_i
	.globl	_Z9pmultiplyPdS_S_i
	.type	_Z9pmultiplyPdS_S_i, @function
_Z9pmultiplyPdS_S_i:
.LFB2055:
	.cfi_startproc
	endbr64
	jmp	_Z33__device_stub__Z9pmultiplyPdS_S_iPdS_S_i
	.cfi_endproc
.LFE2055:
	.size	_Z9pmultiplyPdS_S_i, .-_Z9pmultiplyPdS_S_i
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"Time = %f milliseconds\n"
	.text
	.globl	Cudamultiply
	.type	Cudamultiply, @function
Cudamultiply:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movq	%rdi, %r14
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	movq	%rsi, %r13
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	movq	%rdx, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	movl	%ecx, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	movl	%ebp, %ebx
	imull	%ebp, %ebx
	subq	$80, %rsp
	.cfi_def_cfa_offset 128
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	leaq	32(%rsp), %rdi
	sall	$3, %ebx
	call	cudaEventCreate@PLT
	movslq	%ebx, %rbx
	leaq	40(%rsp), %rdi
	call	cudaEventCreate@PLT
	leaq	8(%rsp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	movq	8(%rsp), %rdi
	movq	%rbx, %rdx
	movq	%r14, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	leaq	16(%rsp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	movq	16(%rsp), %rdi
	movq	%r13, %rsi
	movq	%rbx, %rdx
	movl	$1, %ecx
	leaq	.LC0(%rip), %r13
	call	cudaMemcpy@PLT
	movl	%ebp, %eax
	movl	$32, %ecx
	movq	%r13, %rsi
	cltd
	movl	$2, %edi
	cvtss2sd	4(%rsp), %xmm0
	movl	$1, 56(%rsp)
	idivl	%ecx
	movl	$1, 68(%rsp)
	movl	%eax, 48(%rsp)
	movl	%eax, 52(%rsp)
	movabsq	$137438953504, %rax
	movq	%rax, 60(%rsp)
	movb	$1, %al
	call	__printf_chk@PLT
	leaq	24(%rsp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	movq	24(%rsp), %rdi
	movq	%rbx, %rdx
	movq	%r12, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movq	32(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	68(%rsp), %ecx
	movl	56(%rsp), %esi
	xorl	%r9d, %r9d
	movq	60(%rsp), %rdx
	movq	48(%rsp), %rdi
	xorl	%r8d, %r8d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L14
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rsi
	movl	%ebp, %ecx
	movq	8(%rsp), %rdi
	call	_Z33__device_stub__Z9pmultiplyPdS_S_iPdS_S_i
.L14:
	movq	40(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	40(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	40(%rsp), %rdx
	movq	32(%rsp), %rsi
	leaq	4(%rsp), %rdi
	call	cudaEventElapsedTime@PLT
	movq	%r13, %rsi
	movl	$2, %edi
	movb	$1, %al
	cvtss2sd	4(%rsp), %xmm0
	call	__printf_chk@PLT
	movq	24(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	%r12, %rdi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L15
	call	__stack_chk_fail@PLT
.L15:
	addq	$80, %rsp
	.cfi_def_cfa_offset 48
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	Cudamultiply, .-Cudamultiply
	.section	.rodata.str1.1
.LC1:
	.string	"_Z9pmultiplyPdS_S_i"
.LC2:
	.string	"_Z9smultiplyPiS_S_"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2057:
	.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
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC1(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z9pmultiplyPdS_S_i(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC2(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z9smultiplyPiS_S_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	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	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2057:
	.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:

/************************************************************
Known issues:
This program only works on matrices smaller than or equal to 
256x256. 1024x1024 will cause segmentation faults and 512x512
simply causes the program to almost crash and return times of
0 for each kernel call. 
The matrices must be square and all matrices must be the same
size.
*/


#include <hip/hip_runtime.h>
#include <stdio.h>
#include <stdlib.h>
#define MATSIZE 128
#define THREADS_PER_BLOCK 32 
 

 
//serial matrix multiplication kernel
__global__ void smultiply(int* g_a, int* g_b, int* g_c)
{
  int x, y, z;

  for (x = 0; x < MATSIZE; ++ x)
  {
    for (y = 0; y < MATSIZE; ++ y)
    {
      for (z = 0; z < MATSIZE; ++ z)
      {
	g_c[(x * MATSIZE) + y] += g_a[(x * MATSIZE) + z] * g_b[(z * MATSIZE) + y];
      }
    }
  }
}

//parallel matrix multiplication kernel
__global__ void pmultiply(double* g_a, double* g_b, double* g_d , int dim)
{
  int z;
  double sum = 0.0;
  int x = blockIdx.x * blockDim.x + threadIdx.x;
  int y = blockIdx.y * blockDim.y + threadIdx.y;

  for (z = 0; z < dim; ++ z)
  {
    sum += g_a[x * dim + z] * g_b[y + z * dim];
  }
    
  g_d[(x * dim) + y] = sum;
}

extern "C" void Cudamultiply(double* a, double* b, double* d, int Dim)
{
//  double *d;
  int i;
  double *g_a, *g_b, *g_d;
  int g_size = Dim * Dim * sizeof(double);
  hipEvent_t start, stop;
  float time; 

//  d=new double[Dim*Dim];

  hipEventCreate(&start);
  hipEventCreate(&stop);
  //used for timing the Cuda run

  hipMalloc(&g_a, g_size); //allocate memory on Cuda device
  hipMemcpy(g_a, a, g_size, hipMemcpyHostToDevice);
  //copy matrix A onto the Cuda device  

  hipMalloc(&g_b, g_size);
  hipMemcpy(g_b, b, g_size, hipMemcpyHostToDevice);
  
//  cudaMalloc(&g_c, g_size);
//  cudaMemcpy(g_c, c, g_size, cudaMemcpyHostToDevice);

  dim3 dimGrid((Dim / THREADS_PER_BLOCK), (Dim / THREADS_PER_BLOCK));
  //create the needed number of grids
  dim3 threads(THREADS_PER_BLOCK, THREADS_PER_BLOCK);
  //create the needed number of threads in each grid

  //serial Cuda kernel call
  //cudaEventRecord(start, 0);
  //smultiply<<<1,1>>>(g_a, g_b, g_c);
  //cudaEventRecord(stop, 0);  

  //cudaEventSynchronize(stop);
  //cudaEventElapsedTime(&time, start, stop);
  //get run time

//  cudaMemcpy(c, g_c, g_size, cudaMemcpyDeviceToHost);
  //copy results back to host device
//  cudaFree(g_c);  
  //free up unused user allocated memory on Cuda device

  printf("Time = %f milliseconds\n", time);
  
  //create a second answer matrix to use
  //This is not done until now so that memory on the
  //Cuda device is not wasted.
  hipMalloc(&g_d, g_size);
  hipMemcpy(g_d, d, g_size, hipMemcpyHostToDevice);

  //parallel Cuda kernel call
  hipEventRecord(start, 0);
  pmultiply<<<dimGrid,threads>>>(g_a, g_b, g_d,Dim);
  hipEventRecord(stop, 0);

  hipEventSynchronize(stop);
  hipEventElapsedTime(&time, start, stop);
  printf("Time = %f milliseconds\n", time);

  hipMemcpy(d, g_d, g_size, hipMemcpyDeviceToHost);

  hipFree(g_a);
  hipFree(g_b);
  hipFree(g_d);
  //free up all unused user allocated memory on Cuda device

  //printf("\n");  

  /*The next 2 for loops print out the values of both
  answer matrices. This can be used to ensure that both
  kernel calls are producing the same results, and that
  the results are correct. This section can be commented 
  out when the user only wants the timing of a run.*/

//  for (i = 1; i <= Dim * Dim; ++ i)
//  {
//      if (i % Dim == 0)
//      {
////        printf("%f ", c[i-1]);
//	printf("\n");
//      }
//   
//      else
////        printf("%f ", c[i-1]);
//  }

 // printf("\n");

//  for (i = 1; i <= Dim * Dim; ++ i)
//  {
//    if ( i % Dim == 0)
//    {
//     printf("%f ", d[i-1]);
//     printf("\n");
//    }
//
//    else
//     printf("%f ", d[i-1]);
//  }
//  delete d;
}

	.text
	.file	"cudamatrix.hip"
	.globl	_Z24__device_stub__smultiplyPiS_S_ # -- Begin function _Z24__device_stub__smultiplyPiS_S_
	.type	_Z24__device_stub__smultiplyPiS_S_,@function
_Z24__device_stub__smultiplyPiS_S_:     # @_Z24__device_stub__smultiplyPiS_S_
	.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	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%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	$_Z9smultiplyPiS_S_, %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	_Z24__device_stub__smultiplyPiS_S_, .Lfunc_end0-_Z24__device_stub__smultiplyPiS_S_
	.cfi_endproc
                                        # -- End function
	.globl	_Z24__device_stub__pmultiplyPdS_S_i # -- Begin function _Z24__device_stub__pmultiplyPdS_S_i
	.type	_Z24__device_stub__pmultiplyPdS_S_i,@function
_Z24__device_stub__pmultiplyPdS_S_i:    # @_Z24__device_stub__pmultiplyPdS_S_i
	.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	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	4(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 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	$_Z9pmultiplyPdS_S_i, %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_end1:
	.size	_Z24__device_stub__pmultiplyPdS_S_i, .Lfunc_end1-_Z24__device_stub__pmultiplyPdS_S_i
	.cfi_endproc
                                        # -- End function
	.globl	Cudamultiply                    # -- Begin function Cudamultiply
	.type	Cudamultiply,@function
Cudamultiply:                           # @Cudamultiply
	.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	$56, %rsp
	.cfi_def_cfa_offset 112
	.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	%ecx, %r15d
	movq	%rdx, %rbx
	movq	%rsi, %r13
	movq	%rdi, %rbp
	movl	%ecx, %r14d
	imull	%ecx, %r14d
	shll	$3, %r14d
	leaq	24(%rsp), %rdi
	callq	hipEventCreate
	leaq	8(%rsp), %rdi
	callq	hipEventCreate
	movslq	%r14d, %r14
	leaq	40(%rsp), %r12
	movq	%r12, %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	movq	(%r12), %rdi
	movq	%rbp, %rsi
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	32(%rsp), %r12
	movq	%r12, %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	movq	(%r12), %rdi
	movq	%r13, %rsi
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leal	31(%r15), %eax
	testl	%r15d, %r15d
	cmovnsl	%r15d, %eax
	sarl	$5, %eax
	movq	%rax, %r13
	shlq	$32, %r13
	orq	%rax, %r13
	movl	$.L.str, %edi
	movb	$1, %al
	callq	printf
	leaq	16(%rsp), %r12
	movq	%r12, %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	movq	(%r12), %rdi
	movq	%rbx, %rsi
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	24(%rsp), %rax
	movq	(%rax), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movabsq	$137438953504, %rdx             # imm = 0x2000000020
	movq	%r13, %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_2
# %bb.1:
	movq	40(%rsp), %rdi
	movq	32(%rsp), %rsi
	movq	16(%rsp), %rdx
	movl	%r15d, %ecx
	callq	_Z24__device_stub__pmultiplyPdS_S_i
.LBB2_2:
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	8(%rsp), %rdi
	callq	hipEventSynchronize
	movq	24(%rsp), %rsi
	movq	8(%rsp), %rdx
	leaq	52(%rsp), %r15
	movq	%r15, %rdi
	callq	hipEventElapsedTime
	cvtss2sd	(%r15), %xmm0
	movl	$.L.str, %edi
	movb	$1, %al
	callq	printf
	movq	16(%rsp), %rsi
	movq	%rbx, %rdi
	movq	%r14, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	40(%rsp), %rdi
	callq	hipFree
	movq	32(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	addq	$56, %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_end2:
	.size	Cudamultiply, .Lfunc_end2-Cudamultiply
	.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:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB3_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB3_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z9smultiplyPiS_S_, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z9pmultiplyPdS_S_i, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end3:
	.size	__hip_module_ctor, .Lfunc_end3-__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	.LBB4_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
.LBB4_2:
	retq
.Lfunc_end4:
	.size	__hip_module_dtor, .Lfunc_end4-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z9smultiplyPiS_S_,@object      # @_Z9smultiplyPiS_S_
	.section	.rodata,"a",@progbits
	.globl	_Z9smultiplyPiS_S_
	.p2align	3, 0x0
_Z9smultiplyPiS_S_:
	.quad	_Z24__device_stub__smultiplyPiS_S_
	.size	_Z9smultiplyPiS_S_, 8

	.type	_Z9pmultiplyPdS_S_i,@object     # @_Z9pmultiplyPdS_S_i
	.globl	_Z9pmultiplyPdS_S_i
	.p2align	3, 0x0
_Z9pmultiplyPdS_S_i:
	.quad	_Z24__device_stub__pmultiplyPdS_S_i
	.size	_Z9pmultiplyPdS_S_i, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Time = %f milliseconds\n"
	.size	.L.str, 24

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z9smultiplyPiS_S_"
	.size	.L__unnamed_1, 19

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z9pmultiplyPdS_S_i"
	.size	.L__unnamed_2, 20

	.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 _Z24__device_stub__smultiplyPiS_S_
	.addrsig_sym _Z24__device_stub__pmultiplyPdS_S_i
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z9smultiplyPiS_S_
	.addrsig_sym _Z9pmultiplyPdS_S_i
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z9smultiplyPiS_S_      ; -- Begin function _Z9smultiplyPiS_S_
	.globl	_Z9smultiplyPiS_S_
	.p2align	8
	.type	_Z9smultiplyPiS_S_,@function
_Z9smultiplyPiS_S_:                     ; @_Z9smultiplyPiS_S_
; %bb.0:
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[0:1], s[0:1], 0x10
	v_mov_b32_e32 v0, 0
	s_mov_b32 s3, 0
	s_mov_b32 s16, 0
.LBB0_1:                                ; %.preheader19
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_2 Depth 2
                                        ;       Child Loop BB0_3 Depth 3
	s_delay_alu instid0(SALU_CYCLE_1)
	s_lshl_b32 s17, s16, 7
	s_waitcnt lgkmcnt(0)
	s_mov_b64 s[8:9], s[6:7]
	s_mov_b32 s18, s3
.LBB0_2:                                ; %.preheader
                                        ;   Parent Loop BB0_1 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB0_3 Depth 3
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s2, s18, s17
	s_mov_b64 s[12:13], 0
	s_lshl_b64 s[10:11], s[2:3], 2
	s_mov_b64 s[14:15], s[8:9]
	s_add_u32 s10, s0, s10
	s_addc_u32 s11, s1, s11
	global_load_b32 v1, v0, s[10:11]
.LBB0_3:                                ;   Parent Loop BB0_1 Depth=1
                                        ;     Parent Loop BB0_2 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	s_add_u32 s20, s4, s12
	s_addc_u32 s21, s5, s13
	s_add_u32 s12, s12, 4
	s_clause 0x1
	global_load_b32 v2, v0, s[20:21]
	global_load_b32 v3, v0, s[14:15]
	s_addc_u32 s13, s13, 0
	s_add_u32 s14, s14, 0x200
	s_addc_u32 s15, s15, 0
	s_cmpk_lg_i32 s12, 0x200
	s_waitcnt vmcnt(0)
	v_mul_lo_u32 v2, v3, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v1, v1, v2
	global_store_b32 v0, v1, s[10:11]
	s_cbranch_scc1 .LBB0_3
; %bb.4:                                ;   in Loop: Header=BB0_2 Depth=2
	s_add_i32 s18, s18, 1
	s_add_u32 s8, s8, 4
	s_addc_u32 s9, s9, 0
	s_cmpk_lg_i32 s18, 0x80
	s_cbranch_scc1 .LBB0_2
; %bb.5:                                ;   in Loop: Header=BB0_1 Depth=1
	s_add_i32 s16, s16, 1
	s_add_u32 s4, s4, 0x200
	s_addc_u32 s5, s5, 0
	s_cmpk_lg_i32 s16, 0x80
	s_cbranch_scc1 .LBB0_1
; %bb.6:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z9smultiplyPiS_S_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 24
		.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 4
		.amdhsa_next_free_sgpr 22
		.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	_Z9smultiplyPiS_S_, .Lfunc_end0-_Z9smultiplyPiS_S_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 228
; NumSgprs: 22
; NumVgprs: 4
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 22
; NumVGPRsForWavesPerEU: 4
; 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
	.protected	_Z9pmultiplyPdS_S_i     ; -- Begin function _Z9pmultiplyPdS_S_i
	.globl	_Z9pmultiplyPdS_S_i
	.p2align	8
	.type	_Z9pmultiplyPdS_S_i,@function
_Z9pmultiplyPdS_S_i:                    ; @_Z9pmultiplyPdS_S_i
; %bb.0:
	s_clause 0x1
	s_load_b32 s3, s[0:1], 0x2c
	s_load_b32 s2, s[0:1], 0x18
	v_and_b32_e32 v1, 0x3ff, v0
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[0:1], s[0:1], 0x10
	s_waitcnt lgkmcnt(0)
	s_and_b32 s8, s3, 0xffff
	s_lshr_b32 s3, s3, 16
	v_mad_u64_u32 v[2:3], null, s14, s8, v[1:2]
	v_bfe_u32 v3, v0, 10, 10
	s_cmp_lt_i32 s2, 1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mad_u64_u32 v[0:1], null, s15, s3, v[3:4]
	v_mul_lo_u32 v1, v2, s2
	s_cbranch_scc1 .LBB1_3
; %bb.1:                                ; %.lr.ph.preheader
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_ashrrev_i32_e32 v2, 31, v1
	s_mov_b32 s3, s2
	v_mov_b32_e32 v6, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[4:5], 3, v[1:2]
	v_mov_b32_e32 v2, 0
	v_mov_b32_e32 v3, 0
	v_add_co_u32 v4, vcc_lo, s4, v4
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_ci_u32_e32 v5, vcc_lo, s5, v5, vcc_lo
.LBB1_2:                                ; %.lr.ph
                                        ; =>This Inner Loop Header: Depth=1
	v_ashrrev_i32_e32 v7, 31, v6
	s_add_i32 s3, s3, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_cmp_lg_u32 s3, 0
	v_lshlrev_b64 v[7:8], 3, v[6:7]
	v_add_nc_u32_e32 v6, s2, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v7, vcc_lo, s6, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s7, v8, vcc_lo
	global_load_b64 v[9:10], v[4:5], off
	global_load_b64 v[7:8], v[7:8], off
	v_add_co_u32 v4, vcc_lo, v4, 8
	v_add_co_ci_u32_e32 v5, vcc_lo, 0, v5, vcc_lo
	s_waitcnt vmcnt(0)
	v_fma_f64 v[2:3], v[9:10], v[7:8], v[2:3]
	s_cbranch_scc1 .LBB1_2
	s_branch .LBB1_4
.LBB1_3:
	v_mov_b32_e32 v2, 0
	v_mov_b32_e32 v3, 0
.LBB1_4:                                ; %._crit_edge
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v0, v1, v0
	v_ashrrev_i32_e32 v1, 31, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 3, v[0:1]
	v_add_co_u32 v0, vcc_lo, s0, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_store_b64 v[0:1], v[2:3], off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z9pmultiplyPdS_S_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.amdhsa_user_sgpr_count 14
		.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 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 11
		.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_end1:
	.size	_Z9pmultiplyPdS_S_i, .Lfunc_end1-_Z9pmultiplyPdS_S_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 320
; NumSgprs: 18
; NumVgprs: 11
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 11
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.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
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 24
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z9smultiplyPiS_S_
    .private_segment_fixed_size: 0
    .sgpr_count:     22
    .sgpr_spill_count: 0
    .symbol:         _Z9smultiplyPiS_S_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     4
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .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:           _Z9pmultiplyPdS_S_i
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z9pmultiplyPdS_S_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     11
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
