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

namespace cg = cooperative_groups;

// Kernel for sparse matrix-matrix multiplication in CSR format
__global__ void csrSpMMKernel(int m, int n, int k, const int *rowPtrA, const int *colIdxA, const float *valA, const float *B, float *C) {
    cg::grid_group grid = cg::this_grid();
    int row = blockIdx.x * blockDim.x + threadIdx.x;

    if (row >= m) return;

    // Each row of A is processed by a single block
    float dotProduct = 0.0f;
    int rowStart = rowPtrA[row];
    int rowEnd = rowPtrA[row + 1];

    for (int jj = rowStart; jj < rowEnd; ++jj) {
        int colA = colIdxA[jj];
        float valAjj = valA[jj];
        for (int kk = 0; kk < k; ++kk) {
            dotProduct += valAjj * B[colA * k + kk];
        }
    }

    // Use cooperative groups to ensure all threads in the block have computed their part
    cg::sync(grid);

    // Write the result to the output matrix C
    for (int kk = 0; kk < k; ++kk) {
        C[row * k + kk] = dotProduct;
    }
}

int main() {
    // Example matrices in CSR format
    int m = 4, n = 4, k = 4;
    int rowPtrA[] = {0, 2, 4, 6, 8};
    int colIdxA[] = {0, 1, 0, 2, 1, 3, 2, 3};
    float valA[] = {1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0};
    float B[] = {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};
    float *C = new float[m * k];

    // Device pointers
    int *d_rowPtrA, *d_colIdxA;
    float *d_valA, *d_B, *d_C;

    // Allocate device memory
    cudaMalloc((void**)&d_rowPtrA, (m + 1) * sizeof(int));
    cudaMalloc((void**)&d_colIdxA, rowPtrA[m] * sizeof(int));
    cudaMalloc((void**)&d_valA, rowPtrA[m] * sizeof(float));
    cudaMalloc((void**)&d_B, n * k * sizeof(float));
    cudaMalloc((void**)&d_C, m * k * sizeof(float));

    // Copy data to device
    cudaMemcpy(d_rowPtrA, rowPtrA, (m + 1) * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_colIdxA, colIdxA, rowPtrA[m] * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_valA, valA, rowPtrA[m] * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, B, n * k * sizeof(float), cudaMemcpyHostToDevice);

    // Launch kernel
    int threadsPerBlock = 256;
    int blocksPerGrid = (m + threadsPerBlock - 1) / threadsPerBlock;
    csrSpMMKernel<<<blocksPerGrid, threadsPerBlock>>>(m, n, k, d_rowPtrA, d_colIdxA, d_valA, d_B, d_C);

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

    // Copy result back to host
    cudaMemcpy(C, d_C, m * k * sizeof(float), cudaMemcpyDeviceToHost);

    // Print result
    std::cout << "Result matrix C:" << std::endl;
    for (int i = 0; i < m; ++i) {
        for (int j = 0; j < k; ++j) {
            std::cout << C[i * k + j] << " ";
        }
        std::cout << std::endl;
    }

    // Free device memory
    cudaFree(d_rowPtrA);
    cudaFree(d_colIdxA);
    cudaFree(d_valA);
    cudaFree(d_B);
    cudaFree(d_C);

    // Free host memory
    delete[] C;

    return 0;
}
	.file	"tmpxft_001e4c8e_00000000-6_cuda_code_039847.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB6835:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE6835:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z48__device_stub__Z13csrSpMMKerneliiiPKiS0_PKfS2_PfiiiPKiS0_PKfS2_Pf
	.type	_Z48__device_stub__Z13csrSpMMKerneliiiPKiS0_PKfS2_PfiiiPKiS0_PKfS2_Pf, @function
_Z48__device_stub__Z13csrSpMMKerneliiiPKiS0_PKfS2_PfiiiPKiS0_PKfS2_Pf:
.LFB6857:
	.cfi_startproc
	endbr64
	subq	$216, %rsp
	.cfi_def_cfa_offset 224
	movq	224(%rsp), %rax
	movl	%edi, 60(%rsp)
	leaq	88(%rsp), %rdi
	movl	%esi, 56(%rsp)
	leaq	100(%rsp), %rsi
	movq	%rax, 16(%rsp)
	movq	232(%rsp), %rax
	movl	%edx, 52(%rsp)
	leaq	72(%rsp), %rdx
	movq	%rcx, 40(%rsp)
	leaq	80(%rsp), %rcx
	movq	%r8, 32(%rsp)
	movq	%r9, 24(%rsp)
	movq	%rax, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 200(%rsp)
	xorl	%eax, %eax
	leaq	60(%rsp), %rax
	movl	$1, 96(%rsp)
	movq	%rax, 136(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	52(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 192(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 88(%rsp)
	movq	%rax, 100(%rsp)
	movl	$1, 108(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	80(%rsp)
	.cfi_def_cfa_offset 232
	leaq	_Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf(%rip), %rdi
	pushq	80(%rsp)
	.cfi_def_cfa_offset 240
	movq	116(%rsp), %rcx
	movl	124(%rsp), %r8d
	movq	104(%rsp), %rsi
	movl	112(%rsp), %edx
	leaq	152(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 232
	popq	%rdx
	.cfi_def_cfa_offset 224
.L2:
	movq	200(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$216, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE6857:
	.size	_Z48__device_stub__Z13csrSpMMKerneliiiPKiS0_PKfS2_PfiiiPKiS0_PKfS2_Pf, .-_Z48__device_stub__Z13csrSpMMKerneliiiPKiS0_PKfS2_PfiiiPKiS0_PKfS2_Pf
	.globl	_Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf
	.type	_Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf, @function
_Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf:
.LFB6858:
	.cfi_startproc
	endbr64
	jmp	_Z48__device_stub__Z13csrSpMMKerneliiiPKiS0_PKfS2_PfiiiPKiS0_PKfS2_Pf
	.cfi_endproc
.LFE6858:
	.size	_Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf, .-_Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"CUDA error: "
.LC2:
	.string	"Result matrix C:"
.LC3:
	.string	" "
	.section	.rodata
	.align 32
.LC0:
	.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
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB6832:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	leaq	.LC0(%rip), %rsi
	movl	$16, %ecx
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movabsq	$25769803780, %rbx
	subq	$232, %rsp
	.cfi_def_cfa_offset 288
	movq	%fs:40, %rax
	movq	%rax, 216(%rsp)
	movabsq	$8589934592, %rax
	movq	%rbx, 76(%rsp)
	movl	$1, %ebx
	leaq	152(%rsp), %rdi
	rep movsl
	salq	$32, %rbx
	movl	$64, %edi
	movl	$8, 84(%rsp)
	movq	%rbx, 88(%rsp)
	movq	%rax, 68(%rsp)
	movq	%rax, 96(%rsp)
	movabsq	$12884901889, %rax
	movq	%rax, 104(%rsp)
	incq	%rax
	movq	%rax, 112(%rsp)
	movabsq	$4611686019492741120, %rax
	movq	%rax, 120(%rsp)
	movabsq	$4647714816524288000, %rax
	movq	%rax, 128(%rsp)
	movabsq	$4665729215040061440, %rax
	movq	%rax, 136(%rsp)
	movabsq	$4683743613553737728, %rax
	movq	%rax, 144(%rsp)
	call	_Znam@PLT
	movq	%rsp, %rdi
	movl	$20, %esi
	movq	%rax, %rbx
	call	cudaMalloc@PLT
	movslq	84(%rsp), %rsi
	leaq	8(%rsp), %rdi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	movslq	84(%rsp), %rsi
	leaq	16(%rsp), %rdi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	leaq	24(%rsp), %rdi
	movl	$64, %esi
	call	cudaMalloc@PLT
	leaq	32(%rsp), %rdi
	movl	$64, %esi
	call	cudaMalloc@PLT
	movq	(%rsp), %rdi
	leaq	68(%rsp), %rsi
	movl	$1, %ecx
	movl	$20, %edx
	call	cudaMemcpy@PLT
	movslq	84(%rsp), %rdx
	movq	8(%rsp), %rdi
	leaq	88(%rsp), %rsi
	movl	$1, %ecx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movslq	84(%rsp), %rdx
	movq	16(%rsp), %rdi
	leaq	120(%rsp), %rsi
	movl	$1, %ecx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movq	24(%rsp), %rdi
	movl	$1, %ecx
	leaq	152(%rsp), %rsi
	movl	$64, %edx
	call	cudaMemcpy@PLT
	movl	$16777217, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	salq	$8, %rdx
	movl	$1, %ecx
	movl	$1, %esi
	movabsq	$4294967297, %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L9
	pushq	32(%rsp)
	.cfi_def_cfa_offset 296
	movl	$4, %edx
	movl	$4, %esi
	movl	$4, %edi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 304
	movq	32(%rsp), %r9
	movq	24(%rsp), %r8
	movq	16(%rsp), %rcx
	call	_Z48__device_stub__Z13csrSpMMKerneliiiPKiS0_PKfS2_PfiiiPKiS0_PKfS2_Pf
	popq	%rax
	.cfi_def_cfa_offset 296
	popq	%rdx
	.cfi_def_cfa_offset 288
.L9:
	call	cudaGetLastError@PLT
	movl	%eax, %ebp
	testl	%eax, %eax
	je	.L10
	leaq	.LC1(%rip), %rsi
	leaq	_ZSt4cerr(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebp, %edi
	movq	%rax, %rbx
	call	cudaGetErrorString@PLT
	movq	%rbx, %rdi
	movq	%rax, %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	orl	$-1, %eax
	jmp	.L8
.L10:
	movq	32(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdi
	xorl	%r12d, %r12d
	movl	$64, %edx
	leaq	_ZSt4cout(%rip), %r14
	leaq	.LC3(%rip), %r15
	call	cudaMemcpy@PLT
	leaq	.LC2(%rip), %rsi
	movq	%r14, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
.L12:
	movq	%r12, %r13
	xorl	%ebp, %ebp
	salq	$4, %r13
	addq	%rbx, %r13
.L13:
	movq	%r14, %rdi
	cvtss2sd	0(%r13,%rbp,4), %xmm0
	incq	%rbp
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	%r15, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cmpq	$4, %rbp
	jne	.L13
	movq	%r14, %rdi
	incq	%r12
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	cmpq	$4, %r12
	jne	.L12
	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
	movq	%rbx, %rdi
	call	_ZdaPv@PLT
	xorl	%eax, %eax
.L8:
	movq	216(%rsp), %rdx
	subq	%fs:40, %rdx
	je	.L15
	call	__stack_chk_fail@PLT
.L15:
	addq	$232, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.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
	ret
	.cfi_endproc
.LFE6832:
	.size	main, .-main
	.section	.rodata.str1.1
.LC4:
	.string	"_Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf"
.LC5:
	.string	"_ZN50_INTERNAL_c110cd6b_19_cuda_code_039847_cu_407a74ef4cuda3std3__419piecewise_constructE"
.LC6:
	.string	"_ZN50_INTERNAL_c110cd6b_19_cuda_code_039847_cu_407a74ef4cuda3std6ranges3__45__cpo4swapE"
.LC7:
	.string	"_ZN50_INTERNAL_c110cd6b_19_cuda_code_039847_cu_407a74ef4cuda3std6ranges3__45__cpo9iter_moveE"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB6860:
	.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	.LC4(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf(%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	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC5(%rip), %rdx
	movl	$1, %r9d
	leaq	_ZN4cuda3std3__419piecewise_constructE(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC6(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$1, %r9d
	leaq	_ZN4cuda3std6ranges3__45__cpo4swapE(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC7(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movl	$1, %r9d
	movq	%rdx, %rcx
	leaq	_ZN4cuda3std6ranges3__45__cpo9iter_moveE(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rdi
	.cfi_def_cfa_offset 24
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%r8
	.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
.LFE6860:
	.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
	.weak	_ZN4cuda3std3__419piecewise_constructE
	.section	.rodata._ZN4cuda3std3__419piecewise_constructE,"aG",@progbits,_ZN4cuda3std3__419piecewise_constructE,comdat
	.type	_ZN4cuda3std3__419piecewise_constructE, @gnu_unique_object
	.size	_ZN4cuda3std3__419piecewise_constructE, 1
_ZN4cuda3std3__419piecewise_constructE:
	.zero	1
	.weak	_ZN4cuda3std6ranges3__45__cpo9iter_moveE
	.section	.rodata._ZN4cuda3std6ranges3__45__cpo9iter_moveE,"aG",@progbits,_ZN4cuda3std6ranges3__45__cpo9iter_moveE,comdat
	.type	_ZN4cuda3std6ranges3__45__cpo9iter_moveE, @gnu_unique_object
	.size	_ZN4cuda3std6ranges3__45__cpo9iter_moveE, 1
_ZN4cuda3std6ranges3__45__cpo9iter_moveE:
	.zero	1
	.weak	_ZN4cuda3std6ranges3__45__cpo4swapE
	.section	.rodata._ZN4cuda3std6ranges3__45__cpo4swapE,"aG",@progbits,_ZN4cuda3std6ranges3__45__cpo4swapE,comdat
	.type	_ZN4cuda3std6ranges3__45__cpo4swapE, @gnu_unique_object
	.size	_ZN4cuda3std6ranges3__45__cpo4swapE, 1
_ZN4cuda3std6ranges3__45__cpo4swapE:
	.zero	1
	.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>
#include <hip/hip_cooperative_groups.h>

namespace cg = cooperative_groups;

// Kernel for sparse matrix-matrix multiplication in CSR format
__global__ void csrSpMMKernel(int m, int n, int k, const int *rowPtrA, const int *colIdxA, const float *valA, const float *B, float *C) {
    cg::grid_group grid = cg::this_grid();
    int row = blockIdx.x * blockDim.x + threadIdx.x;

    if (row >= m) return;

    // Each row of A is processed by a single block
    float dotProduct = 0.0f;
    int rowStart = rowPtrA[row];
    int rowEnd = rowPtrA[row + 1];

    for (int jj = rowStart; jj < rowEnd; ++jj) {
        int colA = colIdxA[jj];
        float valAjj = valA[jj];
        for (int kk = 0; kk < k; ++kk) {
            dotProduct += valAjj * B[colA * k + kk];
        }
    }

    // Use cooperative groups to ensure all threads in the block have computed their part
    cg::sync(grid);

    // Write the result to the output matrix C
    for (int kk = 0; kk < k; ++kk) {
        C[row * k + kk] = dotProduct;
    }
}

int main() {
    // Example matrices in CSR format
    int m = 4, n = 4, k = 4;
    int rowPtrA[] = {0, 2, 4, 6, 8};
    int colIdxA[] = {0, 1, 0, 2, 1, 3, 2, 3};
    float valA[] = {1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0};
    float B[] = {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};
    float *C = new float[m * k];

    // Device pointers
    int *d_rowPtrA, *d_colIdxA;
    float *d_valA, *d_B, *d_C;

    // Allocate device memory
    hipMalloc((void**)&d_rowPtrA, (m + 1) * sizeof(int));
    hipMalloc((void**)&d_colIdxA, rowPtrA[m] * sizeof(int));
    hipMalloc((void**)&d_valA, rowPtrA[m] * sizeof(float));
    hipMalloc((void**)&d_B, n * k * sizeof(float));
    hipMalloc((void**)&d_C, m * k * sizeof(float));

    // Copy data to device
    hipMemcpy(d_rowPtrA, rowPtrA, (m + 1) * sizeof(int), hipMemcpyHostToDevice);
    hipMemcpy(d_colIdxA, colIdxA, rowPtrA[m] * sizeof(int), hipMemcpyHostToDevice);
    hipMemcpy(d_valA, valA, rowPtrA[m] * sizeof(float), hipMemcpyHostToDevice);
    hipMemcpy(d_B, B, n * k * sizeof(float), hipMemcpyHostToDevice);

    // Launch kernel
    int threadsPerBlock = 256;
    int blocksPerGrid = (m + threadsPerBlock - 1) / threadsPerBlock;
    csrSpMMKernel<<<blocksPerGrid, threadsPerBlock>>>(m, n, k, d_rowPtrA, d_colIdxA, d_valA, d_B, d_C);

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

    // Copy result back to host
    hipMemcpy(C, d_C, m * k * sizeof(float), hipMemcpyDeviceToHost);

    // Print result
    std::cout << "Result matrix C:" << std::endl;
    for (int i = 0; i < m; ++i) {
        for (int j = 0; j < k; ++j) {
            std::cout << C[i * k + j] << " ";
        }
        std::cout << std::endl;
    }

    // Free device memory
    hipFree(d_rowPtrA);
    hipFree(d_colIdxA);
    hipFree(d_valA);
    hipFree(d_B);
    hipFree(d_C);

    // Free host memory
    delete[] C;

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

                                        # End of file scope inline assembly
	.globl	_Z28__device_stub__csrSpMMKerneliiiPKiS0_PKfS2_Pf # -- Begin function _Z28__device_stub__csrSpMMKerneliiiPKiS0_PKfS2_Pf
	.type	_Z28__device_stub__csrSpMMKerneliiiPKiS0_PKfS2_Pf,@function
_Z28__device_stub__csrSpMMKerneliiiPKiS0_PKfS2_Pf: # @_Z28__device_stub__csrSpMMKerneliiiPKiS0_PKfS2_Pf
	.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	$160, %rsp
	.cfi_def_cfa_offset 208
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	20(%rsp), %rax
	movl	%edi, (%rax)
	leaq	16(%rsp), %rdi
	movl	%esi, (%rdi)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	56(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	48(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	40(%rsp), %r8
	movq	%r9, (%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	208(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	216(%rsp), %rax
	movq	%rax, 56(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	32(%rsp), %r12
	leaq	24(%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	$_Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$176, %rsp
	.cfi_adjust_cfa_offset -176
	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__csrSpMMKerneliiiPKiS0_PKfS2_Pf, .Lfunc_end0-_Z28__device_stub__csrSpMMKerneliiiPKiS0_PKfS2_Pf
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function main
.LCPI1_0:
	.long	0                               # 0x0
	.long	2                               # 0x2
	.long	4                               # 0x4
	.long	6                               # 0x6
.LCPI1_1:
	.long	0                               # 0x0
	.long	1                               # 0x1
	.long	0                               # 0x0
	.long	2                               # 0x2
.LCPI1_2:
	.long	1                               # 0x1
	.long	3                               # 0x3
	.long	2                               # 0x2
	.long	3                               # 0x3
.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
.LCPI1_6:
	.long	0x41500000                      # float 13
	.long	0x41600000                      # float 14
	.long	0x41700000                      # float 15
	.long	0x41800000                      # float 16
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$216, %rsp
	.cfi_def_cfa_offset 272
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movaps	.LCPI1_0(%rip), %xmm0           # xmm0 = [0,2,4,6]
	leaq	112(%rsp), %r15
	movaps	%xmm0, (%r15)
	movl	$8, 16(%r15)
	movaps	.LCPI1_1(%rip), %xmm0           # xmm0 = [0,1,0,2]
	leaq	80(%rsp), %rax
	movaps	%xmm0, (%rax)
	movaps	.LCPI1_2(%rip), %xmm0           # xmm0 = [1,3,2,3]
	movaps	%xmm0, 16(%rax)
	movaps	.LCPI1_3(%rip), %xmm0           # xmm0 = [1.0E+0,2.0E+0,3.0E+0,4.0E+0]
	leaq	48(%rsp), %rax
	movaps	%xmm0, (%rax)
	movaps	.LCPI1_4(%rip), %xmm1           # xmm1 = [5.0E+0,6.0E+0,7.0E+0,8.0E+0]
	movaps	%xmm1, 16(%rax)
	leaq	144(%rsp), %r14
	movaps	%xmm0, (%r14)
	movaps	%xmm1, 16(%r14)
	movaps	.LCPI1_5(%rip), %xmm0           # xmm0 = [9.0E+0,1.0E+1,1.1E+1,1.2E+1]
	movaps	%xmm0, 32(%r14)
	movaps	.LCPI1_6(%rip), %xmm0           # xmm0 = [1.3E+1,1.4E+1,1.5E+1,1.6E+1]
	movaps	%xmm0, 48(%r14)
	movl	$64, %edi
	callq	_Znam
	movq	%rax, 40(%rsp)                  # 8-byte Spill
	leaq	32(%rsp), %r13
	movl	$20, %esi
	movq	%r13, %rdi
	callq	hipMalloc
	leaq	24(%rsp), %r12
	movl	$32, %esi
	movq	%r12, %rdi
	callq	hipMalloc
	leaq	16(%rsp), %rbx
	movl	$32, %esi
	movq	%rbx, %rdi
	callq	hipMalloc
	leaq	8(%rsp), %rbp
	movl	$64, %esi
	movq	%rbp, %rdi
	callq	hipMalloc
	movq	%rsp, %rdi
	movl	$64, %esi
	callq	hipMalloc
	movq	(%r13), %rdi
	movl	$20, %edx
	movq	%r15, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r12), %rdi
	movslq	16(%r15), %rdx
	shlq	$2, %rdx
	leaq	80(%rsp), %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%rbx), %rdi
	movslq	16(%r15), %rdx
	shlq	$2, %rdx
	leaq	48(%rsp), %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%rbp), %rdi
	movl	$64, %edx
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	leaq	255(%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), %rcx
	movq	24(%rsp), %r8
	movq	16(%rsp), %r9
	movl	$4, %edi
	movl	$4, %esi
	movl	$4, %edx
	pushq	(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	16(%rsp)
	.cfi_adjust_cfa_offset 8
	callq	_Z28__device_stub__csrSpMMKerneliiiPKiS0_PKfS2_Pf
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
.LBB1_2:
	callq	hipGetLastError
	testl	%eax, %eax
	je	.LBB1_7
# %bb.3:
	movl	%eax, %ebx
	movl	$_ZSt4cerr, %edi
	movl	$.L.str, %esi
	movl	$12, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	%ebx, %edi
	callq	hipGetErrorString
	testq	%rax, %rax
	je	.LBB1_4
# %bb.5:
	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_6
.LBB1_7:
	movq	(%rsp), %rsi
	movl	$64, %edx
	movq	40(%rsp), %r13                  # 8-byte Reload
	movq	%r13, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$_ZSt4cout, %ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.1, %esi
	movl	$16, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	xorl	%r14d, %r14d
	movq	%r13, %r15
.LBB1_8:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB1_9 Depth 2
	xorl	%r12d, %r12d
.LBB1_9:                                #   Parent Loop BB1_8 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	xorps	%xmm0, %xmm0
	cvtss2sd	(%r15,%r12,4), %xmm0
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movl	$.L.str.2, %esi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	incq	%r12
	cmpq	$4, %r12
	jne	.LBB1_9
# %bb.10:                               #   in Loop: Header=BB1_8 Depth=1
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	incq	%r14
	addq	$16, %r15
	cmpq	$4, %r14
	jne	.LBB1_8
# %bb.11:
	movq	32(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	(%rsp), %rdi
	callq	hipFree
	movq	%r13, %rdi
	callq	_ZdaPv
	xorl	%eax, %eax
	jmp	.LBB1_12
.LBB1_4:
	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_6:                                # %_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
	movl	$-1, %eax
.LBB1_12:
	addq	$216, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	main, .Lfunc_end1-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB2_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB2_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf, %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	_Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf,@object # @_Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf
	.section	.rodata,"a",@progbits
	.globl	_Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf
	.p2align	3, 0x0
_Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf:
	.quad	_Z28__device_stub__csrSpMMKerneliiiPKiS0_PKfS2_Pf
	.size	_Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"CUDA error: "
	.size	.L.str, 13

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"Result matrix C:"
	.size	.L.str.1, 17

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

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf"
	.size	.L__unnamed_1, 35

	.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__csrSpMMKerneliiiPKiS0_PKfS2_Pf
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf
	.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	_Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf ; -- Begin function _Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf
	.globl	_Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf
	.p2align	8
	.type	_Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf,@function
_Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf:     ; @_Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf
; %bb.0:
	s_clause 0x1
	s_load_b32 s4, s[0:1], 0x44
	s_load_b32 s5, s[0:1], 0x0
	v_and_b32_e32 v3, 0x3ff, v0
	s_add_u32 s2, s0, 56
	s_addc_u32 s3, s1, 0
	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[3:4]
	s_mov_b32 s4, exec_lo
	v_cmpx_gt_i32_e64 s5, v1
	s_cbranch_execz .LBB0_21
; %bb.1:
	s_load_b256 s[4:11], s[0:1], 0x10
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[4:5], 2, v[1:2]
	v_mov_b32_e32 v2, 0
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v4, vcc_lo, s4, v4
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v5, vcc_lo, s5, v5, vcc_lo
	s_mov_b32 s5, exec_lo
	global_load_b64 v[4:5], v[4:5], off
	s_clause 0x1
	s_load_b32 s4, s[0:1], 0x8
	s_load_b64 s[0:1], s[0:1], 0x30
	s_waitcnt vmcnt(0)
	v_cmpx_lt_i32_e64 v4, v5
	s_cbranch_execz .LBB0_8
; %bb.2:                                ; %.lr.ph47
	v_mov_b32_e32 v2, 0
	v_mov_b32_e32 v6, v4
	s_waitcnt lgkmcnt(0)
	s_cmp_gt_i32 s4, 0
	s_mov_b32 s13, 0
	s_cselect_b32 s12, -1, 0
.LBB0_3:                                ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_5 Depth 2
	v_ashrrev_i32_e32 v7, 31, v6
	s_and_not1_b32 vcc_lo, exec_lo, s12
	s_cbranch_vccnz .LBB0_6
; %bb.4:                                ; %.lr.ph
                                        ;   in Loop: Header=BB0_3 Depth=1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[7:8], 2, v[6:7]
	s_mov_b32 s14, s4
	v_add_co_u32 v9, vcc_lo, s6, v7
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v10, vcc_lo, s7, v8, vcc_lo
	v_add_co_u32 v7, vcc_lo, s8, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s9, v8, vcc_lo
	global_load_b32 v9, v[9:10], off
	global_load_b32 v4, v[7:8], off
	s_waitcnt vmcnt(1)
	v_mul_lo_u32 v7, v9, s4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v8, 31, v7
	v_lshlrev_b64 v[7:8], 2, v[7:8]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v7, vcc_lo, s10, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s11, v8, vcc_lo
.LBB0_5:                                ;   Parent Loop BB0_3 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b32 v9, v[7:8], off
	v_add_co_u32 v7, vcc_lo, v7, 4
	v_add_co_ci_u32_e32 v8, vcc_lo, 0, v8, vcc_lo
	s_add_i32 s14, s14, -1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_eq_u32 s14, 0
	s_waitcnt vmcnt(0)
	v_fmac_f32_e32 v2, v4, v9
	s_cbranch_scc0 .LBB0_5
.LBB0_6:                                ; %._crit_edge
                                        ;   in Loop: Header=BB0_3 Depth=1
	v_add_nc_u32_e32 v6, 1, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cmp_ge_i32_e32 vcc_lo, v6, v5
	s_or_b32 s13, vcc_lo, s13
	s_and_not1_b32 exec_lo, exec_lo, s13
	s_cbranch_execnz .LBB0_3
; %bb.7:                                ; %Flow91
	s_or_b32 exec_lo, exec_lo, s13
.LBB0_8:                                ; %Flow92
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s5
	v_bfe_u32 v4, v0, 20, 10
	v_bfe_u32 v0, v0, 10, 10
	s_mov_b32 s5, exec_lo
	s_waitcnt lgkmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	s_barrier
	v_or3_b32 v0, v0, v4, v3
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v0
	s_cbranch_execz .LBB0_18
; %bb.9:
	s_load_b64 s[2:3], s[2:3], 0x58
	s_mov_b32 s8, exec_lo
	s_mov_b32 s7, exec_lo
	v_mbcnt_lo_u32_b32 v0, s8, 0
                                        ; implicit-def: $vgpr3
	s_waitcnt lgkmcnt(0)
	s_load_b32 s6, s[2:3], 0x28
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v0
	s_cbranch_execz .LBB0_11
; %bb.10:
	s_bcnt1_i32_b32 s8, s8
	s_delay_alu instid0(SALU_CYCLE_1)
	v_dual_mov_b32 v3, 0 :: v_dual_mov_b32 v4, s8
	global_atomic_add_u32 v3, v3, v4, s[2:3] offset:32 glc
.LBB0_11:
	s_or_b32 exec_lo, exec_lo, s7
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s7, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v0, s7, v0
	s_waitcnt lgkmcnt(0)
	s_add_i32 s7, s6, -1
	v_and_b32_e32 v3, 0xffff, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cmp_ne_u32_e32 vcc_lo, s7, v3
	s_and_saveexec_b32 s7, vcc_lo
	s_xor_b32 s7, exec_lo, s7
	s_cbranch_execz .LBB0_15
; %bb.12:                               ; %.preheader.i.i.i.i.preheader
	v_mov_b32_e32 v3, 0
	s_mov_b32 s8, 0
.LBB0_13:                               ; %.preheader.i.i.i.i
                                        ; =>This Inner Loop Header: Depth=1
	s_sleep 1
	global_load_b32 v4, v3, s[2:3] offset:32 glc
	s_waitcnt vmcnt(0)
	v_xor_b32_e32 v4, v4, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cmp_lt_u32_e32 vcc_lo, 0xffff, v4
	s_or_b32 s8, vcc_lo, s8
	s_and_not1_b32 exec_lo, exec_lo, s8
	s_cbranch_execnz .LBB0_13
; %bb.14:                               ; %Flow84
	s_or_b32 exec_lo, exec_lo, s8
.LBB0_15:                               ; %Flow86
	s_and_not1_saveexec_b32 s7, s7
	s_cbranch_execz .LBB0_18
; %bb.16:
	s_mov_b32 s7, exec_lo
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mbcnt_lo_u32_b32 v0, s7, 0
	v_cmp_eq_u32_e32 vcc_lo, 0, v0
	s_and_b32 s8, exec_lo, vcc_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 exec_lo, s8
	s_cbranch_execz .LBB0_18
; %bb.17:
	s_sub_i32 s6, 0x10000, s6
	s_bcnt1_i32_b32 s7, s7
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_mul_i32 s6, s6, s7
	v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v3, s6
	global_atomic_add_u32 v0, v3, s[2:3] offset:32
.LBB0_18:                               ; %Flow88
	s_or_b32 exec_lo, exec_lo, s5
	s_cmp_lt_i32 s4, 1
	s_waitcnt_vscnt null, 0x0
	s_barrier
	s_cbranch_scc1 .LBB0_21
; %bb.19:                               ; %.lr.ph51
	v_mul_lo_u32 v0, v1, s4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v1, 31, v0
	v_lshlrev_b64 v[0:1], 2, v[0:1]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s0, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
.LBB0_20:                               ; =>This Inner Loop Header: Depth=1
	global_store_b32 v[0:1], v2, off
	v_add_co_u32 v0, vcc_lo, v0, 4
	v_add_co_ci_u32_e32 v1, vcc_lo, 0, v1, vcc_lo
	s_add_i32 s4, s4, -1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_lg_u32 s4, 0
	s_cbranch_scc1 .LBB0_20
.LBB0_21:                               ; %.loopexit
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 312
		.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 2
		.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_end0:
	.size	_Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf, .Lfunc_end0-_Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 828
; 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: 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: 2
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .offset:         4
        .size:           4
        .value_kind:     by_value
      - .offset:         8
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         40
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         48
        .size:           8
        .value_kind:     global_buffer
      - .offset:         56
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         60
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         64
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         68
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         70
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         72
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         74
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         76
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         78
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         104
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         112
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         120
        .size:           2
        .value_kind:     hidden_grid_dims
      - .offset:         144
        .size:           8
        .value_kind:     hidden_multigrid_sync_arg
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 312
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z13csrSpMMKerneliiiPKiS0_PKfS2_Pf.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
