#include <stdio.h>
#include <iostream>
#include "cuda_runtime.h"
#include <device_launch_parameters.h>
#include <device_functions.h>
#include <time.h>
#include <cuda_runtime_api.h>
/*Realization of the multiplication between a vector and a matrix, on the GPU.

It realizes multiplication between a vector of 4 element and a matrix (4x4),
so as result is expected a vector of 4 element.


We initialized the input array on the host (CPU) and trasferred to device (GPU).
Then, CPU launches the kernel which is elaborated on GPU.
Finally, the results is trasferred from GPU to CPU and printed out*/



__global__ void vector_matrix_mult(float *d_vec, float *d_mat, float *d_out, const int N, const int M)
{

int tid = threadIdx.x + blockIdx.x*blockDim.x;
float sum = 0;
if (tid < M) {
for (int i = 0; i < N; i++)
sum += d_vec[i] * d_mat[(i*M) + tid];
d_out[tid] = sum;
}
}

int main(int argc, char ** argv) {



	float elapsed = 0;
	cudaEvent_t start, stop;
	cudaEventCreate(&start);
	cudaEventCreate(&stop);

	const int N = 4;
	const int M = N;
	const int ARRAY_SIZE = N;
	const int ARRAY_BYTES = ARRAY_SIZE * sizeof(float);
	const int MATRIX_SIZE = ARRAY_SIZE*ARRAY_SIZE;
	const int MATRIX_BYTES = MATRIX_SIZE * sizeof(float);

	// generate the input vector and input matrix on the host
	float h_vec[ARRAY_SIZE] = { 2, 1, 1, 1 };
	float h_mat[ARRAY_SIZE][ARRAY_SIZE] = { 2, 2, 3, 4, 1, 2, 3, 4, 1, 2, 3, 4, 1, 2, 3, 4 };

	
	float h_out[ARRAY_SIZE];

	//declare GPU memory pointers
	float * d_vec;
	float * d_mat;
	float * d_out;

	//allocate GPU memory
	cudaMalloc((void**)&d_vec, ARRAY_BYTES);
	cudaMalloc((void**)&d_mat, MATRIX_BYTES);
	cudaMalloc((void**)&d_out, ARRAY_BYTES);

	//transfer the input from CPU to GPU

	cudaMemcpy(d_vec, h_vec, ARRAY_BYTES, cudaMemcpyHostToDevice);
	cudaMemcpy(d_mat, h_mat, MATRIX_BYTES, cudaMemcpyHostToDevice);



	cudaEventRecord(start, 0);
	//launch the kernel
	vector_matrix_mult <<<1, ARRAY_SIZE>>> (d_vec, d_mat, d_out, N, M);
	

	cudaEventRecord(stop, 0);

	

	//trasfer the results from GPU to CPU
	cudaMemcpy(h_out, d_out, ARRAY_BYTES, cudaMemcpyDeviceToHost);

	cudaEventSynchronize(stop);
	cudaEventElapsedTime(&elapsed, start, stop);

	//print out the resulting array
	for (int i = 0; i < ARRAY_SIZE; i++) {
		printf("%f", h_out[i]);
		printf(((i % 4) != 3) ? "\t" : "\n");
	}
	printf("The GPU time elapsed is %.6f ms \"", elapsed);

	//free GPU location memory
	cudaFree(d_vec);
	cudaFree(d_mat);
	cudaFree(d_out);

	cudaEventDestroy(start);
	cudaEventDestroy(stop);
	return 0;
}

/*Realization though CPU*/

/*
#include <time.h>
int main(int argc, char ** argv) {

clock_t cpu_startTime, cpu_stopTime;

double cpu_elapsedTime = 0;
cpu_startTime = clock();

const int ARRAY_SIZE = 4;
const int ARRAY_BYTES = ARRAY_SIZE * sizeof(float);

float h_vec[ARRAY_SIZE] = { 1, 1, 1, 1 };
float h_mat[ARRAY_SIZE][ARRAY_SIZE] = { 1, 2, 3, 4, 1, 2, 3, 4, 1, 2, 3, 4, 1, 2, 3, 4 };
float risultato[ARRAY_SIZE];
float h_out[ARRAY_SIZE];
float tot = 0;



int i, j;

for (j = 0; j < ARRAY_SIZE; j++){
for (i = 0; i < ARRAY_SIZE; i++) {
risultato[i] = h_vec[i] * h_mat[i][j];
tot += risultato[i];

}
h_out[j] = tot;
tot = 0;

printf("%f", h_out[j]);
printf(((i % 4) != 3) ? "\t" : "\n");
}

cpu_stopTime = clock();
cpu_elapsedTime = ((cpu_startTime - cpu_stopTime) / CLOCKS_PER_SEC);
printf("The CPU elapsed time is %.6f ms \"", cpu_elapsedTime);
return 0;
}
*/
	.file	"tmpxft_0038cb3b_00000000-6_kernel.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3638:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3638:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z44__device_stub__Z18vector_matrix_multPfS_S_iiPfS_S_ii
	.type	_Z44__device_stub__Z18vector_matrix_multPfS_S_iiPfS_S_ii, @function
_Z44__device_stub__Z18vector_matrix_multPfS_S_iiPfS_S_ii:
.LFB3660:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movl	%ecx, 4(%rsp)
	leaq	40(%rsp), %rcx
	movl	%r8d, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%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)
	leaq	4(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	%rsp, %rax
	movq	%rax, 128(%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 168
	leaq	_Z18vector_matrix_multPfS_S_ii(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 176
	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 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
.LFE3660:
	.size	_Z44__device_stub__Z18vector_matrix_multPfS_S_iiPfS_S_ii, .-_Z44__device_stub__Z18vector_matrix_multPfS_S_iiPfS_S_ii
	.globl	_Z18vector_matrix_multPfS_S_ii
	.type	_Z18vector_matrix_multPfS_S_ii, @function
_Z18vector_matrix_multPfS_S_ii:
.LFB3661:
	.cfi_startproc
	endbr64
	jmp	_Z44__device_stub__Z18vector_matrix_multPfS_S_iiPfS_S_ii
	.cfi_endproc
.LFE3661:
	.size	_Z18vector_matrix_multPfS_S_ii, .-_Z18vector_matrix_multPfS_S_ii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"\t"
.LC2:
	.string	"\n"
.LC4:
	.string	"%f"
.LC5:
	.string	"The GPU time elapsed is %.6f ms \""
	.section	.rodata
	.align 32
.LC0:
	.long	1073741824
	.long	1073741824
	.long	1077936128
	.long	1082130432
	.long	1065353216
	.long	1073741824
	.long	1077936128
	.long	1082130432
	.long	1065353216
	.long	1073741824
	.long	1077936128
	.long	1082130432
	.long	1065353216
	.long	1073741824
	.long	1077936128
	.long	1082130432
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3635:
	.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	$192, %rsp
	.cfi_def_cfa_offset 224
	movq	%fs:40, %rax
	movq	%rax, 184(%rsp)
	xorl	%eax, %eax
	leaq	16(%rsp), %rdi
	movl	$0x00000000, 12(%rsp)
	call	cudaEventCreate@PLT
	leaq	24(%rsp), %rdi
	call	cudaEventCreate@PLT
	movl	$4261412865, %eax
	leaq	112(%rsp), %rdi
	leaq	.LC0(%rip), %rsi
	salq	$30, %rax
	movl	$16, %ecx
	rep movsl
	movq	%rax, 80(%rsp)
	leaq	32(%rsp), %rdi
	subq	$8388608, %rax
	movl	$16, %esi
	movq	%rax, 88(%rsp)
	call	cudaMalloc@PLT
	leaq	40(%rsp), %rdi
	movl	$64, %esi
	call	cudaMalloc@PLT
	leaq	48(%rsp), %rdi
	movl	$16, %esi
	call	cudaMalloc@PLT
	movq	32(%rsp), %rdi
	leaq	80(%rsp), %rsi
	movl	$1, %ecx
	movl	$16, %edx
	call	cudaMemcpy@PLT
	movq	40(%rsp), %rdi
	movl	$1, %ecx
	movl	$64, %edx
	leaq	112(%rsp), %rsi
	call	cudaMemcpy@PLT
	movq	16(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	$1073741825, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	salq	$2, %rdx
	movl	$1, %ecx
	movabsq	$4294967297, %rdi
	movl	$1, %esi
	movq	%rdx, 68(%rsp)
	movl	$1, 76(%rsp)
	movq	%rdi, 56(%rsp)
	movl	$1, 64(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L9
	movq	48(%rsp), %rdx
	movq	40(%rsp), %rsi
	movl	$4, %ecx
	movl	$4, %r8d
	movq	32(%rsp), %rdi
	call	_Z44__device_stub__Z18vector_matrix_multPfS_S_iiPfS_S_ii
.L9:
	movq	24(%rsp), %rdi
	xorl	%esi, %esi
	leaq	96(%rsp), %rbp
	xorl	%ebx, %ebx
	leaq	.LC4(%rip), %r12
	call	cudaEventRecord@PLT
	movq	48(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbp, %rdi
	movl	$16, %edx
	call	cudaMemcpy@PLT
	movq	24(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rsi
	leaq	12(%rsp), %rdi
	call	cudaEventElapsedTime@PLT
.L11:
	movq	%r12, %rsi
	movl	$2, %edi
	movb	$1, %al
	cvtss2sd	0(%rbp,%rbx,4), %xmm0
	call	__printf_chk@PLT
	cmpq	$3, %rbx
	movl	$2, %edi
	leaq	.LC1(%rip), %rax
	leaq	.LC2(%rip), %rsi
	cmovne	%rax, %rsi
	xorl	%eax, %eax
	incq	%rbx
	call	__printf_chk@PLT
	cmpq	$4, %rbx
	jne	.L11
	leaq	.LC5(%rip), %rsi
	movl	$2, %edi
	movb	$1, %al
	cvtss2sd	12(%rsp), %xmm0
	call	__printf_chk@PLT
	movq	32(%rsp), %rdi
	call	cudaFree@PLT
	movq	40(%rsp), %rdi
	call	cudaFree@PLT
	movq	48(%rsp), %rdi
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movq	24(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movq	184(%rsp), %rax
	subq	%fs:40, %rax
	je	.L12
	call	__stack_chk_fail@PLT
.L12:
	addq	$192, %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
.LFE3635:
	.size	main, .-main
	.section	.rodata.str1.1
.LC7:
	.string	"_Z18vector_matrix_multPfS_S_ii"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3663:
	.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	.LC7(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z18vector_matrix_multPfS_S_ii(%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
.LFE3663:
	.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 <stdio.h>
#include <iostream>
#include "hip/hip_runtime.h"

#include <hip/device_functions.h>
#include <time.h>
#include <hip/hip_runtime_api.h>
/*Realization of the multiplication between a vector and a matrix, on the GPU.

It realizes multiplication between a vector of 4 element and a matrix (4x4),
so as result is expected a vector of 4 element.


We initialized the input array on the host (CPU) and trasferred to device (GPU).
Then, CPU launches the kernel which is elaborated on GPU.
Finally, the results is trasferred from GPU to CPU and printed out*/



__global__ void vector_matrix_mult(float *d_vec, float *d_mat, float *d_out, const int N, const int M)
{

int tid = threadIdx.x + blockIdx.x*blockDim.x;
float sum = 0;
if (tid < M) {
for (int i = 0; i < N; i++)
sum += d_vec[i] * d_mat[(i*M) + tid];
d_out[tid] = sum;
}
}

int main(int argc, char ** argv) {



	float elapsed = 0;
	hipEvent_t start, stop;
	hipEventCreate(&start);
	hipEventCreate(&stop);

	const int N = 4;
	const int M = N;
	const int ARRAY_SIZE = N;
	const int ARRAY_BYTES = ARRAY_SIZE * sizeof(float);
	const int MATRIX_SIZE = ARRAY_SIZE*ARRAY_SIZE;
	const int MATRIX_BYTES = MATRIX_SIZE * sizeof(float);

	// generate the input vector and input matrix on the host
	float h_vec[ARRAY_SIZE] = { 2, 1, 1, 1 };
	float h_mat[ARRAY_SIZE][ARRAY_SIZE] = { 2, 2, 3, 4, 1, 2, 3, 4, 1, 2, 3, 4, 1, 2, 3, 4 };

	
	float h_out[ARRAY_SIZE];

	//declare GPU memory pointers
	float * d_vec;
	float * d_mat;
	float * d_out;

	//allocate GPU memory
	hipMalloc((void**)&d_vec, ARRAY_BYTES);
	hipMalloc((void**)&d_mat, MATRIX_BYTES);
	hipMalloc((void**)&d_out, ARRAY_BYTES);

	//transfer the input from CPU to GPU

	hipMemcpy(d_vec, h_vec, ARRAY_BYTES, hipMemcpyHostToDevice);
	hipMemcpy(d_mat, h_mat, MATRIX_BYTES, hipMemcpyHostToDevice);



	hipEventRecord(start, 0);
	//launch the kernel
	vector_matrix_mult <<<1, ARRAY_SIZE>>> (d_vec, d_mat, d_out, N, M);
	

	hipEventRecord(stop, 0);

	

	//trasfer the results from GPU to CPU
	hipMemcpy(h_out, d_out, ARRAY_BYTES, hipMemcpyDeviceToHost);

	hipEventSynchronize(stop);
	hipEventElapsedTime(&elapsed, start, stop);

	//print out the resulting array
	for (int i = 0; i < ARRAY_SIZE; i++) {
		printf("%f", h_out[i]);
		printf(((i % 4) != 3) ? "\t" : "\n");
	}
	printf("The GPU time elapsed is %.6f ms \"", elapsed);

	//free GPU location memory
	hipFree(d_vec);
	hipFree(d_mat);
	hipFree(d_out);

	hipEventDestroy(start);
	hipEventDestroy(stop);
	return 0;
}

/*Realization though CPU*/

/*
#include <time.h>
int main(int argc, char ** argv) {

clock_t cpu_startTime, cpu_stopTime;

double cpu_elapsedTime = 0;
cpu_startTime = clock();

const int ARRAY_SIZE = 4;
const int ARRAY_BYTES = ARRAY_SIZE * sizeof(float);

float h_vec[ARRAY_SIZE] = { 1, 1, 1, 1 };
float h_mat[ARRAY_SIZE][ARRAY_SIZE] = { 1, 2, 3, 4, 1, 2, 3, 4, 1, 2, 3, 4, 1, 2, 3, 4 };
float risultato[ARRAY_SIZE];
float h_out[ARRAY_SIZE];
float tot = 0;



int i, j;

for (j = 0; j < ARRAY_SIZE; j++){
for (i = 0; i < ARRAY_SIZE; i++) {
risultato[i] = h_vec[i] * h_mat[i][j];
tot += risultato[i];

}
h_out[j] = tot;
tot = 0;

printf("%f", h_out[j]);
printf(((i % 4) != 3) ? "\t" : "\n");
}

cpu_stopTime = clock();
cpu_elapsedTime = ((cpu_startTime - cpu_stopTime) / CLOCKS_PER_SEC);
printf("The CPU elapsed time is %.6f ms \"", cpu_elapsedTime);
return 0;
}
*/
	.text
	.file	"kernel.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.globl	_Z33__device_stub__vector_matrix_multPfS_S_ii # -- Begin function _Z33__device_stub__vector_matrix_multPfS_S_ii
	.type	_Z33__device_stub__vector_matrix_multPfS_S_ii,@function
_Z33__device_stub__vector_matrix_multPfS_S_ii: # @_Z33__device_stub__vector_matrix_multPfS_S_ii
	.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	$128, %rsp
	.cfi_def_cfa_offset 176
	.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)
	movq	%rsp, %rcx
	movl	%r8d, (%rcx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%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	$_Z18vector_matrix_multPfS_S_ii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$144, %rsp
	.cfi_adjust_cfa_offset -144
	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	_Z33__device_stub__vector_matrix_multPfS_S_ii, .Lfunc_end0-_Z33__device_stub__vector_matrix_multPfS_S_ii
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function main
.LCPI1_0:
	.long	0x40000000                      # float 2
	.long	0x3f800000                      # float 1
	.long	0x3f800000                      # float 1
	.long	0x3f800000                      # float 1
.LCPI1_1:
	.long	0x40000000                      # float 2
	.long	0x40000000                      # float 2
	.long	0x40400000                      # float 3
	.long	0x40800000                      # float 4
.LCPI1_2:
	.long	0x3f800000                      # float 1
	.long	0x40000000                      # float 2
	.long	0x40400000                      # float 3
	.long	0x40800000                      # float 4
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.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	$144, %rsp
	.cfi_def_cfa_offset 192
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movl	$0, 4(%rsp)
	leaq	40(%rsp), %rbx
	movq	%rbx, %rdi
	callq	hipEventCreate
	leaq	8(%rsp), %rdi
	callq	hipEventCreate
	movaps	.LCPI1_0(%rip), %xmm0           # xmm0 = [2.0E+0,1.0E+0,1.0E+0,1.0E+0]
	leaq	64(%rsp), %r15
	movaps	%xmm0, (%r15)
	movaps	.LCPI1_1(%rip), %xmm0           # xmm0 = [2.0E+0,2.0E+0,3.0E+0,4.0E+0]
	leaq	80(%rsp), %r14
	movaps	%xmm0, (%r14)
	movaps	.LCPI1_2(%rip), %xmm0           # xmm0 = [1.0E+0,2.0E+0,3.0E+0,4.0E+0]
	movaps	%xmm0, 16(%r14)
	movaps	%xmm0, 32(%r14)
	movaps	%xmm0, 48(%r14)
	leaq	32(%rsp), %r12
	movl	$16, %esi
	movq	%r12, %rdi
	callq	hipMalloc
	leaq	24(%rsp), %r13
	movl	$64, %esi
	movq	%r13, %rdi
	callq	hipMalloc
	leaq	16(%rsp), %rdi
	movl	$16, %esi
	callq	hipMalloc
	movq	(%r12), %rdi
	movl	$16, %edx
	movq	%r15, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r13), %rdi
	movl	$64, %edx
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%rbx), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	leaq	3(%rdi), %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_2
# %bb.1:
	movq	32(%rsp), %rdi
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdx
	movl	$4, %ecx
	movl	$4, %r8d
	callq	_Z33__device_stub__vector_matrix_multPfS_S_ii
.LBB1_2:
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	16(%rsp), %rsi
	leaq	48(%rsp), %rdi
	movl	$16, %edx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	8(%rsp), %rdi
	callq	hipEventSynchronize
	movq	40(%rsp), %rsi
	movq	8(%rsp), %rdx
	leaq	4(%rsp), %rdi
	callq	hipEventElapsedTime
	movl	$.L.str.2, %ebx
	xorl	%r14d, %r14d
.LBB1_3:                                # =>This Inner Loop Header: Depth=1
	xorps	%xmm0, %xmm0
	cvtss2sd	48(%rsp,%r14,4), %xmm0
	movl	$.L.str, %edi
	movb	$1, %al
	callq	printf
	cmpq	$3, %r14
	movl	$.L.str.1, %edi
	cmoveq	%rbx, %rdi
	xorl	%eax, %eax
	callq	printf
	incq	%r14
	cmpq	$4, %r14
	jne	.LBB1_3
# %bb.4:
	xorps	%xmm0, %xmm0
	cvtss2sd	4(%rsp), %xmm0
	movl	$.L.str.3, %edi
	movb	$1, %al
	callq	printf
	movq	32(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	40(%rsp), %rdi
	callq	hipEventDestroy
	movq	8(%rsp), %rdi
	callq	hipEventDestroy
	xorl	%eax, %eax
	addq	$144, %rsp
	.cfi_def_cfa_offset 48
	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	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	$_Z18vector_matrix_multPfS_S_ii, %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	_Z18vector_matrix_multPfS_S_ii,@object # @_Z18vector_matrix_multPfS_S_ii
	.section	.rodata,"a",@progbits
	.globl	_Z18vector_matrix_multPfS_S_ii
	.p2align	3, 0x0
_Z18vector_matrix_multPfS_S_ii:
	.quad	_Z33__device_stub__vector_matrix_multPfS_S_ii
	.size	_Z18vector_matrix_multPfS_S_ii, 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.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"The GPU time elapsed is %.6f ms \""
	.size	.L.str.3, 34

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z18vector_matrix_multPfS_S_ii"
	.size	.L__unnamed_1, 31

	.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 _Z33__device_stub__vector_matrix_multPfS_S_ii
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z18vector_matrix_multPfS_S_ii
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z18vector_matrix_multPfS_S_ii ; -- Begin function _Z18vector_matrix_multPfS_S_ii
	.globl	_Z18vector_matrix_multPfS_S_ii
	.p2align	8
	.type	_Z18vector_matrix_multPfS_S_ii,@function
_Z18vector_matrix_multPfS_S_ii:         ; @_Z18vector_matrix_multPfS_S_ii
; %bb.0:
	s_clause 0x1
	s_load_b32 s4, s[0:1], 0x2c
	s_load_b64 s[2:3], s[0:1], 0x18
	s_waitcnt lgkmcnt(0)
	s_and_b32 s4, s4, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s4, v[0:1]
	s_mov_b32 s4, exec_lo
	v_cmpx_gt_i32_e64 s3, v1
	s_cbranch_execz .LBB0_5
; %bb.1:                                ; %.preheader
	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_cmp_lt_i32 s2, 1
	s_cbranch_scc1 .LBB0_4
; %bb.2:                                ; %.lr.ph.preheader
	v_mov_b32_e32 v2, v1
.LBB0_3:                                ; %.lr.ph
                                        ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	s_waitcnt lgkmcnt(0)
	s_load_b32 s8, s[4:5], 0x0
	s_add_i32 s2, s2, -1
	s_add_u32 s4, s4, 4
	s_addc_u32 s5, s5, 0
	v_lshlrev_b64 v[3:4], 2, v[2:3]
	v_add_nc_u32_e32 v2, s3, v2
	s_cmp_eq_u32 s2, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v3, vcc_lo, s6, v3
	v_add_co_ci_u32_e32 v4, vcc_lo, s7, v4, vcc_lo
	global_load_b32 v3, v[3:4], off
	s_waitcnt vmcnt(0) lgkmcnt(0)
	v_fmac_f32_e32 v0, s8, v3
	s_cbranch_scc0 .LBB0_3
.LBB0_4:                                ; %Flow39
	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[1:2], 2, v[1:2]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v1, vcc_lo, s0, v1
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v2, vcc_lo, s1, v2, vcc_lo
	global_store_b32 v[1:2], v0, off
.LBB0_5:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z18vector_matrix_multPfS_S_ii
		.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 5
		.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	_Z18vector_matrix_multPfS_S_ii, .Lfunc_end0-_Z18vector_matrix_multPfS_S_ii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 236
; NumSgprs: 18
; NumVgprs: 5
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 5
; 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
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         28
        .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:           _Z18vector_matrix_multPfS_S_ii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z18vector_matrix_multPfS_S_ii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     5
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
