#include <iostream>
#include <cuda_runtime.h>

// Define the number of non-zero elements in the sparse matrix
#define NUM_NON_ZERO 32

// Kernel for sparse matrix-vector multiplication using persistent threads
__global__ void spmvCsrKernel(const int* rowPtr, const int* colInd, const float* values, const float* x, float* y, int numRows) {
    // Thread ID in block
    int tid = threadIdx.x;

    // Each thread computes one row of the matrix-vector product
    int row = blockIdx.x * blockDim.x + tid;

    // Ensure thread does not exceed the number of rows
    if (row < numRows) {
        float sum = 0.0f;
        // Loop over the non-zero elements of the row
        for (int j = rowPtr[row]; j < rowPtr[row + 1]; ++j) {
            sum += values[j] * x[colInd[j]];
        }
        // Write the result to the output vector
        y[row] = sum;
    }
}

int main() {
    // Host data
    int h_rowPtr[] = {0, 3, 5, 7, 10, 12, 15, 18, 20, 22, 25, 28, 30, 32};
    int h_colInd[] = {0, 2, 5, 1, 4, 2, 3, 0, 2, 6, 7, 1, 3, 5, 6, 7, 8, 2, 5, 7, 8, 9, 3, 6, 8, 9, 10, 4, 7, 10, 11, 5, 8};
    float h_values[] = {1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0, 17.0, 18.0, 19.0, 20.0, 21.0, 22.0, 23.0, 24.0, 25.0, 26.0, 27.0, 28.0, 29.0, 30.0, 31.0, 32.0};
    float h_x[] = {1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0};
    float h_y[] = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0};

    // Number of rows in the matrix
    int numRows = 10;

    // Device data
    int* d_rowPtr;
    int* d_colInd;
    float* d_values;
    float* d_x;
    float* d_y;

    // Allocate device memory
    cudaError_t err = cudaMalloc((void**)&d_rowPtr, (numRows + 1) * sizeof(int));
    if (err != cudaSuccess) {
        std::cerr << "Error allocating device memory for rowPtr: " << cudaGetErrorString(err) << std::endl;
        return -1;
    }
    err = cudaMalloc((void**)&d_colInd, NUM_NON_ZERO * sizeof(int));
    if (err != cudaSuccess) {
        std::cerr << "Error allocating device memory for colInd: " << cudaGetErrorString(err) << std::endl;
        return -1;
    }
    err = cudaMalloc((void**)&d_values, NUM_NON_ZERO * sizeof(float));
    if (err != cudaSuccess) {
        std::cerr << "Error allocating device memory for values: " << cudaGetErrorString(err) << std::endl;
        return -1;
    }
    err = cudaMalloc((void**)&d_x, 12 * sizeof(float));
    if (err != cudaSuccess) {
        std::cerr << "Error allocating device memory for x: " << cudaGetErrorString(err) << std::endl;
        return -1;
    }
    err = cudaMalloc((void**)&d_y, numRows * sizeof(float));
    if (err != cudaSuccess) {
        std::cerr << "Error allocating device memory for y: " << cudaGetErrorString(err) << std::endl;
        return -1;
    }

    // Copy data from host to device
    err = cudaMemcpy(d_rowPtr, h_rowPtr, (numRows + 1) * sizeof(int), cudaMemcpyHostToDevice);
    if (err != cudaSuccess) {
        std::cerr << "Error copying rowPtr to device: " << cudaGetErrorString(err) << std::endl;
        return -1;
    }
    err = cudaMemcpy(d_colInd, h_colInd, NUM_NON_ZERO * sizeof(int), cudaMemcpyHostToDevice);
    if (err != cudaSuccess) {
        std::cerr << "Error copying colInd to device: " << cudaGetErrorString(err) << std::endl;
        return -1;
    }
    err = cudaMemcpy(d_values, h_values, NUM_NON_ZERO * sizeof(float), cudaMemcpyHostToDevice);
    if (err != cudaSuccess) {
        std::cerr << "Error copying values to device: " << cudaGetErrorString(err) << std::endl;
        return -1;
    }
    err = cudaMemcpy(d_x, h_x, 12 * sizeof(float), cudaMemcpyHostToDevice);
    if (err != cudaSuccess) {
        std::cerr << "Error copying x to device: " << cudaGetErrorString(err) << std::endl;
        return -1;
    }

    // Launch kernel
    int blockSize = 32;
    int numBlocks = (numRows + blockSize - 1) / blockSize;
    spmvCsrKernel<<<numBlocks, blockSize>>>(d_rowPtr, d_colInd, d_values, d_x, d_y, numRows);

    // Check for kernel launch errors
    err = cudaGetLastError();
    if (err != cudaSuccess) {
        std::cerr << "Kernel launch error: " << cudaGetErrorString(err) << std::endl;
        return -1;
    }

    // Copy result back to host
    err = cudaMemcpy(h_y, d_y, numRows * sizeof(float), cudaMemcpyDeviceToHost);
    if (err != cudaSuccess) {
        std::cerr << "Error copying y back to host: " << cudaGetErrorString(err) << std::endl;
        return -1;
    }

    // Print result
    std::cout << "Result vector y:" << std::endl;
    for (int i = 0; i < numRows; ++i) {
        std::cout << h_y[i] << " ";
    }
    std::cout << std::endl;

    // Free device memory
    cudaFree(d_rowPtr);
    cudaFree(d_colInd);
    cudaFree(d_values);
    cudaFree(d_x);
    cudaFree(d_y);

    return 0;
}
	.file	"tmpxft_000bcff6_00000000-6_cuda_code_056401.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
	.align 2
	.type	_ZNSolsEPFRSoS_E.isra.0, @function
_ZNSolsEPFRSoS_E.isra.0:
.LFB4293:
	.cfi_startproc
	jmp	*%rsi
	.cfi_endproc
.LFE4293:
	.size	_ZNSolsEPFRSoS_E.isra.0, .-_ZNSolsEPFRSoS_E.isra.0
	.globl	_Z46__device_stub__Z13spmvCsrKernelPKiS0_PKfS2_PfiPKiS0_PKfS2_Pfi
	.type	_Z46__device_stub__Z13spmvCsrKernelPKiS0_PKfS2_PfiPKiS0_PKfS2_Pfi, @function
_Z46__device_stub__Z13spmvCsrKernelPKiS0_PKfS2_PfiPKiS0_PKfS2_Pfi:
.LFB3660:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movq	%rdi, 40(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	56(%rsp), %rdx
	movq	%rcx, 16(%rsp)
	leaq	64(%rsp), %rcx
	movq	%r8, 8(%rsp)
	movl	%r9d, 4(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 120(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L3
	pushq	64(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z13spmvCsrKernelPKiS0_PKfS2_Pfi(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 208
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	136(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L3:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L5
	call	__stack_chk_fail@PLT
.L5:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3660:
	.size	_Z46__device_stub__Z13spmvCsrKernelPKiS0_PKfS2_PfiPKiS0_PKfS2_Pfi, .-_Z46__device_stub__Z13spmvCsrKernelPKiS0_PKfS2_PfiPKiS0_PKfS2_Pfi
	.globl	_Z13spmvCsrKernelPKiS0_PKfS2_Pfi
	.type	_Z13spmvCsrKernelPKiS0_PKfS2_Pfi, @function
_Z13spmvCsrKernelPKiS0_PKfS2_Pfi:
.LFB3661:
	.cfi_startproc
	endbr64
	jmp	_Z46__device_stub__Z13spmvCsrKernelPKiS0_PKfS2_PfiPKiS0_PKfS2_Pfi
	.cfi_endproc
.LFE3661:
	.size	_Z13spmvCsrKernelPKiS0_PKfS2_Pfi, .-_Z13spmvCsrKernelPKiS0_PKfS2_Pfi
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC4:
	.string	"Error allocating device memory for rowPtr: "
.LC5:
	.string	"Error allocating device memory for colInd: "
.LC6:
	.string	"Error allocating device memory for values: "
.LC7:
	.string	"Error allocating device memory for x: "
.LC8:
	.string	"Error allocating device memory for y: "
.LC9:
	.string	"Error copying rowPtr to device: "
.LC10:
	.string	"Error copying colInd to device: "
.LC11:
	.string	"Error copying values to device: "
.LC12:
	.string	"Error copying x to device: "
.LC13:
	.string	"Kernel launch error: "
.LC14:
	.string	"Error copying y back to host: "
.LC15:
	.string	"Result vector y:"
.LC16:
	.string	" "
	.section	.rodata
	.align 32
.LC0:
	.long	0
	.long	3
	.long	5
	.long	7
	.long	10
	.long	12
	.long	15
	.long	18
	.long	20
	.long	22
	.long	25
	.long	28
	.long	30
	.long	32
	.align 32
.LC1:
	.long	0
	.long	2
	.long	5
	.long	1
	.long	4
	.long	2
	.long	3
	.long	0
	.long	2
	.long	6
	.long	7
	.long	1
	.long	3
	.long	5
	.long	6
	.long	7
	.long	8
	.long	2
	.long	5
	.long	7
	.long	8
	.long	9
	.long	3
	.long	6
	.long	8
	.long	9
	.long	10
	.long	4
	.long	7
	.long	10
	.long	11
	.long	5
	.long	8
	.align 32
.LC2:
	.long	1065353216
	.long	1073741824
	.long	1077936128
	.long	1082130432
	.long	1084227584
	.long	1086324736
	.long	1088421888
	.long	1090519040
	.long	1091567616
	.long	1092616192
	.long	1093664768
	.long	1094713344
	.long	1095761920
	.long	1096810496
	.long	1097859072
	.long	1098907648
	.long	1099431936
	.long	1099956224
	.long	1100480512
	.long	1101004800
	.long	1101529088
	.long	1102053376
	.long	1102577664
	.long	1103101952
	.long	1103626240
	.long	1104150528
	.long	1104674816
	.long	1105199104
	.long	1105723392
	.long	1106247680
	.long	1106771968
	.long	1107296256
	.align 32
.LC3:
	.long	1065353216
	.long	1073741824
	.long	1077936128
	.long	1082130432
	.long	1084227584
	.long	1086324736
	.long	1088421888
	.long	1090519040
	.long	1091567616
	.long	1092616192
	.long	1093664768
	.long	1094713344
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3635:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	leaq	.LC0(%rip), %rsi
	movl	$14, %ecx
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	subq	$488, %rsp
	.cfi_def_cfa_offset 528
	movq	%fs:40, %rax
	movq	%rax, 472(%rsp)
	xorl	%eax, %eax
	leaq	156(%rsp), %rdi
	rep movsl
	leaq	340(%rsp), %rdi
	leaq	.LC1(%rip), %rsi
	movl	$33, %ecx
	rep movsl
	leaq	212(%rsp), %rdi
	leaq	.LC2(%rip), %rsi
	movl	$32, %ecx
	rep movsl
	leaq	108(%rsp), %rdi
	leaq	.LC3(%rip), %rsi
	movl	$12, %ecx
	rep movsl
	leaq	68(%rsp), %rdi
	movl	$10, %ecx
	movl	$44, %esi
	rep stosl
	movq	%rsp, %rdi
	call	cudaMalloc@PLT
	leaq	.LC4(%rip), %rsi
	movl	%eax, %ebx
	testl	%eax, %eax
	jne	.L28
	leaq	8(%rsp), %rdi
	movl	$128, %esi
	call	cudaMalloc@PLT
	movl	%eax, %ebx
	testl	%eax, %eax
	je	.L12
	leaq	.LC5(%rip), %rsi
.L28:
	leaq	_ZSt4cerr(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %edi
	movq	%rax, %rbp
	call	cudaGetErrorString@PLT
	movq	%rbp, %rdi
	movq	%rax, %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	orl	$-1, %eax
	jmp	.L9
.L12:
	movl	$128, %esi
	leaq	16(%rsp), %rdi
	call	cudaMalloc@PLT
	leaq	.LC6(%rip), %rsi
	movl	%eax, %ebx
	testl	%eax, %eax
	jne	.L28
	movl	$48, %esi
	leaq	24(%rsp), %rdi
	call	cudaMalloc@PLT
	leaq	.LC7(%rip), %rsi
	movl	%eax, %ebx
	testl	%eax, %eax
	jne	.L28
	movl	$40, %esi
	leaq	32(%rsp), %rdi
	call	cudaMalloc@PLT
	leaq	.LC8(%rip), %rsi
	movl	%eax, %ebx
	testl	%eax, %eax
	jne	.L28
	movq	(%rsp), %rdi
	movl	$1, %ecx
	movl	$44, %edx
	leaq	156(%rsp), %rsi
	call	cudaMemcpy@PLT
	leaq	.LC9(%rip), %rsi
	movl	%eax, %ebx
	testl	%eax, %eax
	jne	.L28
	movq	8(%rsp), %rdi
	movl	$1, %ecx
	movl	$128, %edx
	leaq	340(%rsp), %rsi
	call	cudaMemcpy@PLT
	leaq	.LC10(%rip), %rsi
	movl	%eax, %ebx
	testl	%eax, %eax
	jne	.L28
	movq	16(%rsp), %rdi
	movl	$1, %ecx
	movl	$128, %edx
	leaq	212(%rsp), %rsi
	call	cudaMemcpy@PLT
	leaq	.LC11(%rip), %rsi
	movl	%eax, %ebx
	testl	%eax, %eax
	jne	.L28
	movq	24(%rsp), %rdi
	leaq	108(%rsp), %rsi
	movl	$1, %ecx
	movl	$48, %edx
	call	cudaMemcpy@PLT
	leaq	.LC12(%rip), %rsi
	movl	%eax, %ebx
	testl	%eax, %eax
	jne	.L28
	movl	$134217729, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	salq	$5, %rdx
	movl	$1, %esi
	movabsq	$4294967297, %rdi
	movl	$1, 64(%rsp)
	movq	%rdx, 56(%rsp)
	leaq	68(%rsp), %rbp
	movq	%rdi, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L21
	movq	32(%rsp), %r8
	movq	24(%rsp), %rcx
	movl	$10, %r9d
	movq	16(%rsp), %rdx
	movq	8(%rsp), %rsi
	movq	(%rsp), %rdi
	call	_Z46__device_stub__Z13spmvCsrKernelPKiS0_PKfS2_PfiPKiS0_PKfS2_Pfi
.L21:
	call	cudaGetLastError@PLT
	leaq	.LC13(%rip), %rsi
	movl	%eax, %ebx
	testl	%eax, %eax
	jne	.L28
	movq	32(%rsp), %rsi
	movl	$2, %ecx
	movl	$40, %edx
	movq	%rbp, %rdi
	call	cudaMemcpy@PLT
	leaq	.LC14(%rip), %rsi
	movl	%eax, %ebx
	testl	%eax, %eax
	jne	.L28
	leaq	_ZSt4cout(%rip), %r12
	leaq	.LC15(%rip), %rsi
	xorl	%ebx, %ebx
	movq	%r12, %rdi
	leaq	.LC16(%rip), %r13
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
.L24:
	movq	%r12, %rdi
	cvtss2sd	0(%rbp,%rbx,4), %xmm0
	incq	%rbx
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	%r13, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cmpq	$10, %rbx
	jne	.L24
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%r12, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	(%rsp), %rdi
	call	cudaFree@PLT
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	32(%rsp), %rdi
	call	cudaFree@PLT
	xorl	%eax, %eax
.L9:
	movq	472(%rsp), %rdx
	subq	%fs:40, %rdx
	je	.L25
	call	__stack_chk_fail@PLT
.L25:
	addq	$488, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3635:
	.size	main, .-main
	.section	.rodata.str1.1
.LC17:
	.string	"_Z13spmvCsrKernelPKiS0_PKfS2_Pfi"
	.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	.LC17(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z13spmvCsrKernelPKiS0_PKfS2_Pfi(%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 <iostream>
#include <hip/hip_runtime.h>

// Define the number of non-zero elements in the sparse matrix
#define NUM_NON_ZERO 32

// Kernel for sparse matrix-vector multiplication using persistent threads
__global__ void spmvCsrKernel(const int* rowPtr, const int* colInd, const float* values, const float* x, float* y, int numRows) {
    // Thread ID in block
    int tid = threadIdx.x;

    // Each thread computes one row of the matrix-vector product
    int row = blockIdx.x * blockDim.x + tid;

    // Ensure thread does not exceed the number of rows
    if (row < numRows) {
        float sum = 0.0f;
        // Loop over the non-zero elements of the row
        for (int j = rowPtr[row]; j < rowPtr[row + 1]; ++j) {
            sum += values[j] * x[colInd[j]];
        }
        // Write the result to the output vector
        y[row] = sum;
    }
}

int main() {
    // Host data
    int h_rowPtr[] = {0, 3, 5, 7, 10, 12, 15, 18, 20, 22, 25, 28, 30, 32};
    int h_colInd[] = {0, 2, 5, 1, 4, 2, 3, 0, 2, 6, 7, 1, 3, 5, 6, 7, 8, 2, 5, 7, 8, 9, 3, 6, 8, 9, 10, 4, 7, 10, 11, 5, 8};
    float h_values[] = {1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0, 17.0, 18.0, 19.0, 20.0, 21.0, 22.0, 23.0, 24.0, 25.0, 26.0, 27.0, 28.0, 29.0, 30.0, 31.0, 32.0};
    float h_x[] = {1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0};
    float h_y[] = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0};

    // Number of rows in the matrix
    int numRows = 10;

    // Device data
    int* d_rowPtr;
    int* d_colInd;
    float* d_values;
    float* d_x;
    float* d_y;

    // Allocate device memory
    hipError_t err = hipMalloc((void**)&d_rowPtr, (numRows + 1) * sizeof(int));
    if (err != hipSuccess) {
        std::cerr << "Error allocating device memory for rowPtr: " << hipGetErrorString(err) << std::endl;
        return -1;
    }
    err = hipMalloc((void**)&d_colInd, NUM_NON_ZERO * sizeof(int));
    if (err != hipSuccess) {
        std::cerr << "Error allocating device memory for colInd: " << hipGetErrorString(err) << std::endl;
        return -1;
    }
    err = hipMalloc((void**)&d_values, NUM_NON_ZERO * sizeof(float));
    if (err != hipSuccess) {
        std::cerr << "Error allocating device memory for values: " << hipGetErrorString(err) << std::endl;
        return -1;
    }
    err = hipMalloc((void**)&d_x, 12 * sizeof(float));
    if (err != hipSuccess) {
        std::cerr << "Error allocating device memory for x: " << hipGetErrorString(err) << std::endl;
        return -1;
    }
    err = hipMalloc((void**)&d_y, numRows * sizeof(float));
    if (err != hipSuccess) {
        std::cerr << "Error allocating device memory for y: " << hipGetErrorString(err) << std::endl;
        return -1;
    }

    // Copy data from host to device
    err = hipMemcpy(d_rowPtr, h_rowPtr, (numRows + 1) * sizeof(int), hipMemcpyHostToDevice);
    if (err != hipSuccess) {
        std::cerr << "Error copying rowPtr to device: " << hipGetErrorString(err) << std::endl;
        return -1;
    }
    err = hipMemcpy(d_colInd, h_colInd, NUM_NON_ZERO * sizeof(int), hipMemcpyHostToDevice);
    if (err != hipSuccess) {
        std::cerr << "Error copying colInd to device: " << hipGetErrorString(err) << std::endl;
        return -1;
    }
    err = hipMemcpy(d_values, h_values, NUM_NON_ZERO * sizeof(float), hipMemcpyHostToDevice);
    if (err != hipSuccess) {
        std::cerr << "Error copying values to device: " << hipGetErrorString(err) << std::endl;
        return -1;
    }
    err = hipMemcpy(d_x, h_x, 12 * sizeof(float), hipMemcpyHostToDevice);
    if (err != hipSuccess) {
        std::cerr << "Error copying x to device: " << hipGetErrorString(err) << std::endl;
        return -1;
    }

    // Launch kernel
    int blockSize = 32;
    int numBlocks = (numRows + blockSize - 1) / blockSize;
    spmvCsrKernel<<<numBlocks, blockSize>>>(d_rowPtr, d_colInd, d_values, d_x, d_y, numRows);

    // Check for kernel launch errors
    err = hipGetLastError();
    if (err != hipSuccess) {
        std::cerr << "Kernel launch error: " << hipGetErrorString(err) << std::endl;
        return -1;
    }

    // Copy result back to host
    err = hipMemcpy(h_y, d_y, numRows * sizeof(float), hipMemcpyDeviceToHost);
    if (err != hipSuccess) {
        std::cerr << "Error copying y back to host: " << hipGetErrorString(err) << std::endl;
        return -1;
    }

    // Print result
    std::cout << "Result vector y:" << std::endl;
    for (int i = 0; i < numRows; ++i) {
        std::cout << h_y[i] << " ";
    }
    std::cout << std::endl;

    // Free device memory
    hipFree(d_rowPtr);
    hipFree(d_colInd);
    hipFree(d_values);
    hipFree(d_x);
    hipFree(d_y);

    return 0;
}
	.text
	.file	"cuda_code_056401.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.globl	_Z28__device_stub__spmvCsrKernelPKiS0_PKfS2_Pfi # -- Begin function _Z28__device_stub__spmvCsrKernelPKiS0_PKfS2_Pfi
	.type	_Z28__device_stub__spmvCsrKernelPKiS0_PKfS2_Pfi,@function
_Z28__device_stub__spmvCsrKernelPKiS0_PKfS2_Pfi: # @_Z28__device_stub__spmvCsrKernelPKiS0_PKfS2_Pfi
	.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
	leaq	56(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	48(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	24(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	4(%rsp), %r8
	movl	%r9d, (%r8)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%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	$_Z13spmvCsrKernelPKiS0_PKfS2_Pfi, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$160, %rsp
	.cfi_adjust_cfa_offset -160
	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	_Z28__device_stub__spmvCsrKernelPKiS0_PKfS2_Pfi, .Lfunc_end0-_Z28__device_stub__spmvCsrKernelPKiS0_PKfS2_Pfi
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function main
.LCPI1_0:
	.long	0                               # 0x0
	.long	3                               # 0x3
	.long	5                               # 0x5
	.long	7                               # 0x7
.LCPI1_1:
	.long	10                              # 0xa
	.long	12                              # 0xc
	.long	15                              # 0xf
	.long	18                              # 0x12
.LCPI1_2:
	.long	20                              # 0x14
	.long	22                              # 0x16
	.long	25                              # 0x19
	.long	28                              # 0x1c
.LCPI1_3:
	.long	0x3f800000                      # float 1
	.long	0x40000000                      # float 2
	.long	0x40400000                      # float 3
	.long	0x40800000                      # float 4
.LCPI1_4:
	.long	0x40a00000                      # float 5
	.long	0x40c00000                      # float 6
	.long	0x40e00000                      # float 7
	.long	0x41000000                      # float 8
.LCPI1_5:
	.long	0x41100000                      # float 9
	.long	0x41200000                      # float 10
	.long	0x41300000                      # float 11
	.long	0x41400000                      # float 12
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$480, %rsp                      # imm = 0x1E0
	.cfi_def_cfa_offset 496
	.cfi_offset %rbx, -16
	movaps	.LCPI1_0(%rip), %xmm0           # xmm0 = [0,3,5,7]
	movaps	%xmm0, 96(%rsp)
	movaps	.LCPI1_1(%rip), %xmm0           # xmm0 = [10,12,15,18]
	movaps	%xmm0, 112(%rsp)
	movaps	.LCPI1_2(%rip), %xmm0           # xmm0 = [20,22,25,28]
	movaps	%xmm0, 128(%rsp)
	movabsq	$137438953502, %rax             # imm = 0x200000001E
	movq	%rax, 144(%rsp)
	leaq	336(%rsp), %rdi
	movl	$.L__const.main.h_colInd, %esi
	movl	$132, %edx
	callq	memcpy@PLT
	leaq	208(%rsp), %rdi
	movl	$16, %ecx
	movl	$.L__const.main.h_values, %esi
	rep;movsq (%rsi), %es:(%rdi)
	movaps	.LCPI1_3(%rip), %xmm0           # xmm0 = [1.0E+0,2.0E+0,3.0E+0,4.0E+0]
	movaps	%xmm0, 160(%rsp)
	movaps	.LCPI1_4(%rip), %xmm0           # xmm0 = [5.0E+0,6.0E+0,7.0E+0,8.0E+0]
	movaps	%xmm0, 176(%rsp)
	movaps	.LCPI1_5(%rip), %xmm0           # xmm0 = [9.0E+0,1.0E+1,1.1E+1,1.2E+1]
	movaps	%xmm0, 192(%rsp)
	xorps	%xmm0, %xmm0
	movaps	%xmm0, 64(%rsp)
	movaps	%xmm0, 48(%rsp)
	movq	$0, 80(%rsp)
	leaq	40(%rsp), %rdi
	movl	$44, %esi
	callq	hipMalloc
	testl	%eax, %eax
	je	.LBB1_2
# %bb.1:
	movl	%eax, %ebx
	movl	$_ZSt4cerr, %edi
	movl	$.L.str, %esi
	jmp	.LBB1_6
.LBB1_2:
	leaq	32(%rsp), %rdi
	movl	$128, %esi
	callq	hipMalloc
	testl	%eax, %eax
	je	.LBB1_4
# %bb.3:
	movl	%eax, %ebx
	movl	$_ZSt4cerr, %edi
	movl	$.L.str.1, %esi
	jmp	.LBB1_6
.LBB1_4:
	leaq	24(%rsp), %rdi
	movl	$128, %esi
	callq	hipMalloc
	testl	%eax, %eax
	je	.LBB1_13
# %bb.5:
	movl	%eax, %ebx
	movl	$_ZSt4cerr, %edi
	movl	$.L.str.2, %esi
.LBB1_6:
	movl	$43, %edx
.LBB1_7:
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	%ebx, %edi
	callq	hipGetErrorString
	testq	%rax, %rax
	je	.LBB1_9
# %bb.8:
	movq	%rax, %rbx
	movq	%rax, %rdi
	callq	strlen
	movl	$_ZSt4cerr, %edi
	movq	%rbx, %rsi
	movq	%rax, %rdx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	jmp	.LBB1_10
.LBB1_9:
	movq	_ZSt4cerr(%rip), %rax
	movq	-24(%rax), %rax
	leaq	_ZSt4cerr(%rax), %rdi
	movl	32(%rdi), %esi
	orl	$1, %esi
	callq	_ZNSt9basic_iosIcSt11char_traitsIcEE5clearESt12_Ios_Iostate
.LBB1_10:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit
	movq	_ZSt4cerr(%rip), %rax
	movl	$_ZSt4cerr, %ecx
	movq	-24(%rax), %rdi
	addq	%rcx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cerr, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.LBB1_11:
	movl	$-1, %eax
.LBB1_12:
	addq	$480, %rsp                      # imm = 0x1E0
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.LBB1_13:
	.cfi_def_cfa_offset 496
	leaq	16(%rsp), %rdi
	movl	$48, %esi
	callq	hipMalloc
	testl	%eax, %eax
	je	.LBB1_15
# %bb.14:
	movl	%eax, %ebx
	movl	$_ZSt4cerr, %edi
	movl	$.L.str.3, %esi
	movl	$38, %edx
	jmp	.LBB1_7
.LBB1_15:
	leaq	8(%rsp), %rdi
	movl	$40, %esi
	callq	hipMalloc
	testl	%eax, %eax
	je	.LBB1_17
# %bb.16:
	movl	%eax, %ebx
	movl	$_ZSt4cerr, %edi
	movl	$.L.str.4, %esi
	movl	$38, %edx
	jmp	.LBB1_24
.LBB1_17:
	movq	40(%rsp), %rdi
	leaq	96(%rsp), %rsi
	movl	$44, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	je	.LBB1_19
# %bb.18:
	movl	%eax, %ebx
	movl	$_ZSt4cerr, %edi
	movl	$.L.str.5, %esi
	jmp	.LBB1_23
.LBB1_19:
	movq	32(%rsp), %rdi
	leaq	336(%rsp), %rsi
	movl	$128, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	je	.LBB1_21
# %bb.20:
	movl	%eax, %ebx
	movl	$_ZSt4cerr, %edi
	movl	$.L.str.6, %esi
	jmp	.LBB1_23
.LBB1_21:
	movq	24(%rsp), %rdi
	leaq	208(%rsp), %rsi
	movl	$128, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	je	.LBB1_25
# %bb.22:
	movl	%eax, %ebx
	movl	$_ZSt4cerr, %edi
	movl	$.L.str.7, %esi
.LBB1_23:
	movl	$32, %edx
.LBB1_24:
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	%ebx, %edi
	callq	hipGetErrorString
	movl	$_ZSt4cerr, %edi
	movq	%rax, %rsi
	callq	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc
	movq	%rax, %rdi
	callq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_
	jmp	.LBB1_11
.LBB1_25:
	movq	16(%rsp), %rdi
	leaq	160(%rsp), %rsi
	movl	$48, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	je	.LBB1_27
# %bb.26:
	movl	%eax, %ebx
	movl	$_ZSt4cerr, %edi
	movl	$.L.str.8, %esi
	movl	$27, %edx
	jmp	.LBB1_24
.LBB1_27:
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	leaq	31(%rdi), %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_29
# %bb.28:
	movq	40(%rsp), %rdi
	movq	32(%rsp), %rsi
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rcx
	movq	8(%rsp), %r8
	movl	$10, %r9d
	callq	_Z28__device_stub__spmvCsrKernelPKiS0_PKfS2_Pfi
.LBB1_29:
	callq	hipGetLastError
	testl	%eax, %eax
	je	.LBB1_31
# %bb.30:
	movl	%eax, %ebx
	movl	$_ZSt4cerr, %edi
	movl	$.L.str.9, %esi
	movl	$21, %edx
	jmp	.LBB1_24
.LBB1_31:
	movq	8(%rsp), %rsi
	leaq	48(%rsp), %rdi
	movl	$40, %edx
	movl	$2, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	je	.LBB1_33
# %bb.32:
	movl	%eax, %ebx
	movl	$_ZSt4cerr, %edi
	movl	$.L.str.10, %esi
	movl	$30, %edx
	jmp	.LBB1_24
.LBB1_33:
	movl	$_ZSt4cout, %edi
	movl	$.L.str.11, %esi
	movl	$16, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	callq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_
	xorl	%ebx, %ebx
.LBB1_34:                               # =>This Inner Loop Header: Depth=1
	xorps	%xmm0, %xmm0
	cvtss2sd	48(%rsp,%rbx,4), %xmm0
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movl	$.L.str.12, %esi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	incq	%rbx
	cmpq	$10, %rbx
	jne	.LBB1_34
# %bb.35:
	movl	$_ZSt4cout, %edi
	callq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_
	movq	40(%rsp), %rdi
	callq	hipFree
	movq	32(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	8(%rsp), %rdi
	callq	hipFree
	xorl	%eax, %eax
	jmp	.LBB1_12
.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	$_Z13spmvCsrKernelPKiS0_PKfS2_Pfi, %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	_Z13spmvCsrKernelPKiS0_PKfS2_Pfi,@object # @_Z13spmvCsrKernelPKiS0_PKfS2_Pfi
	.section	.rodata,"a",@progbits
	.globl	_Z13spmvCsrKernelPKiS0_PKfS2_Pfi
	.p2align	3, 0x0
_Z13spmvCsrKernelPKiS0_PKfS2_Pfi:
	.quad	_Z28__device_stub__spmvCsrKernelPKiS0_PKfS2_Pfi
	.size	_Z13spmvCsrKernelPKiS0_PKfS2_Pfi, 8

	.type	.L__const.main.h_colInd,@object # @__const.main.h_colInd
	.p2align	4, 0x0
.L__const.main.h_colInd:
	.long	0                               # 0x0
	.long	2                               # 0x2
	.long	5                               # 0x5
	.long	1                               # 0x1
	.long	4                               # 0x4
	.long	2                               # 0x2
	.long	3                               # 0x3
	.long	0                               # 0x0
	.long	2                               # 0x2
	.long	6                               # 0x6
	.long	7                               # 0x7
	.long	1                               # 0x1
	.long	3                               # 0x3
	.long	5                               # 0x5
	.long	6                               # 0x6
	.long	7                               # 0x7
	.long	8                               # 0x8
	.long	2                               # 0x2
	.long	5                               # 0x5
	.long	7                               # 0x7
	.long	8                               # 0x8
	.long	9                               # 0x9
	.long	3                               # 0x3
	.long	6                               # 0x6
	.long	8                               # 0x8
	.long	9                               # 0x9
	.long	10                              # 0xa
	.long	4                               # 0x4
	.long	7                               # 0x7
	.long	10                              # 0xa
	.long	11                              # 0xb
	.long	5                               # 0x5
	.long	8                               # 0x8
	.size	.L__const.main.h_colInd, 132

	.type	.L__const.main.h_values,@object # @__const.main.h_values
	.p2align	4, 0x0
.L__const.main.h_values:
	.long	0x3f800000                      # float 1
	.long	0x40000000                      # float 2
	.long	0x40400000                      # float 3
	.long	0x40800000                      # float 4
	.long	0x40a00000                      # float 5
	.long	0x40c00000                      # float 6
	.long	0x40e00000                      # float 7
	.long	0x41000000                      # float 8
	.long	0x41100000                      # float 9
	.long	0x41200000                      # float 10
	.long	0x41300000                      # float 11
	.long	0x41400000                      # float 12
	.long	0x41500000                      # float 13
	.long	0x41600000                      # float 14
	.long	0x41700000                      # float 15
	.long	0x41800000                      # float 16
	.long	0x41880000                      # float 17
	.long	0x41900000                      # float 18
	.long	0x41980000                      # float 19
	.long	0x41a00000                      # float 20
	.long	0x41a80000                      # float 21
	.long	0x41b00000                      # float 22
	.long	0x41b80000                      # float 23
	.long	0x41c00000                      # float 24
	.long	0x41c80000                      # float 25
	.long	0x41d00000                      # float 26
	.long	0x41d80000                      # float 27
	.long	0x41e00000                      # float 28
	.long	0x41e80000                      # float 29
	.long	0x41f00000                      # float 30
	.long	0x41f80000                      # float 31
	.long	0x42000000                      # float 32
	.size	.L__const.main.h_values, 128

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Error allocating device memory for rowPtr: "
	.size	.L.str, 44

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"Error allocating device memory for colInd: "
	.size	.L.str.1, 44

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"Error allocating device memory for values: "
	.size	.L.str.2, 44

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"Error allocating device memory for x: "
	.size	.L.str.3, 39

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"Error allocating device memory for y: "
	.size	.L.str.4, 39

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"Error copying rowPtr to device: "
	.size	.L.str.5, 33

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"Error copying colInd to device: "
	.size	.L.str.6, 33

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"Error copying values to device: "
	.size	.L.str.7, 33

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"Error copying x to device: "
	.size	.L.str.8, 28

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"Kernel launch error: "
	.size	.L.str.9, 22

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"Error copying y back to host: "
	.size	.L.str.10, 31

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"Result vector y:"
	.size	.L.str.11, 17

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

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z13spmvCsrKernelPKiS0_PKfS2_Pfi"
	.size	.L__unnamed_1, 33

	.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 _Z28__device_stub__spmvCsrKernelPKiS0_PKfS2_Pfi
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z13spmvCsrKernelPKiS0_PKfS2_Pfi
	.addrsig_sym _ZSt4cerr
	.addrsig_sym _ZSt4cout
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z13spmvCsrKernelPKiS0_PKfS2_Pfi ; -- Begin function _Z13spmvCsrKernelPKiS0_PKfS2_Pfi
	.globl	_Z13spmvCsrKernelPKiS0_PKfS2_Pfi
	.p2align	8
	.type	_Z13spmvCsrKernelPKiS0_PKfS2_Pfi,@function
_Z13spmvCsrKernelPKiS0_PKfS2_Pfi:       ; @_Z13spmvCsrKernelPKiS0_PKfS2_Pfi
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x3c
	s_load_b32 s3, s[0:1], 0x28
	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_6
; %bb.1:
	s_load_b256 s[4:11], s[0:1], 0x0
	v_ashrrev_i32_e32 v2, 31, v1
	s_load_b64 s[2:3], s[0:1], 0x20
	v_mov_b32_e32 v8, 0
	s_mov_b32 s1, exec_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s4, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v3, vcc_lo, s5, v1, vcc_lo
	global_load_b64 v[2:3], v[2:3], off
	s_waitcnt vmcnt(0)
	v_cmpx_lt_i32_e64 v2, v3
	s_cbranch_execz .LBB0_5
; %bb.2:                                ; %.lr.ph.preheader
	v_ashrrev_i32_e32 v5, 31, v2
	v_mov_b32_e32 v4, v2
	v_mov_b32_e32 v8, 0
	s_mov_b32 s4, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[6:7], 2, v[4:5]
	v_add_co_u32 v4, vcc_lo, s8, v6
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v5, vcc_lo, s9, v7, vcc_lo
	v_add_co_u32 v6, vcc_lo, s6, v6
	v_add_co_ci_u32_e32 v7, vcc_lo, s7, v7, vcc_lo
.LBB0_3:                                ; %.lr.ph
                                        ; =>This Inner Loop Header: Depth=1
	global_load_b32 v9, v[6:7], off
	v_add_nc_u32_e32 v2, 1, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_ge_i32_e64 s0, v2, v3
	s_or_b32 s4, s0, s4
	s_waitcnt vmcnt(0)
	v_ashrrev_i32_e32 v10, 31, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[9:10], 2, v[9:10]
	v_add_co_u32 v9, vcc_lo, s10, v9
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v10, vcc_lo, s11, v10, vcc_lo
	global_load_b32 v11, v[4:5], off
	global_load_b32 v9, v[9:10], off
	v_add_co_u32 v4, vcc_lo, v4, 4
	v_add_co_ci_u32_e32 v5, vcc_lo, 0, v5, vcc_lo
	v_add_co_u32 v6, vcc_lo, v6, 4
	v_add_co_ci_u32_e32 v7, vcc_lo, 0, v7, vcc_lo
	s_waitcnt vmcnt(0)
	v_fmac_f32_e32 v8, v11, v9
	s_and_not1_b32 exec_lo, exec_lo, s4
	s_cbranch_execnz .LBB0_3
; %bb.4:                                ; %Flow
	s_or_b32 exec_lo, exec_lo, s4
.LBB0_5:                                ; %Flow46
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s1
	v_add_co_u32 v0, vcc_lo, s2, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s3, v1, vcc_lo
	global_store_b32 v[0:1], v8, off
.LBB0_6:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z13spmvCsrKernelPKiS0_PKfS2_Pfi
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 304
		.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 12
		.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	_Z13spmvCsrKernelPKiS0_PKfS2_Pfi, .Lfunc_end0-_Z13spmvCsrKernelPKiS0_PKfS2_Pfi
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 364
; NumSgprs: 18
; NumVgprs: 12
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 12
; Occupancy: 16
; WaveLimiterHint : 1
; 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
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .offset:         40
        .size:           4
        .value_kind:     by_value
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         52
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         56
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         60
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         62
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         64
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         66
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         68
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         70
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         104
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         112
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 304
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z13spmvCsrKernelPKiS0_PKfS2_Pfi
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z13spmvCsrKernelPKiS0_PKfS2_Pfi.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     12
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
