#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <time.h>
#include <sys/time.h>

void initialData( float *ip, int size )
{
    // generate different seed for random number
    time_t t;
    srand( (unsigned int) time (&t) );

    for (int i=0; i<size; i++) {
        ip[i] = (float)( rand() & 0xFF ) / 10.0f;
    }
}

void sumMatrixOnHost_1( float *A, float *B, float *C, const int nx,
        const int ny )
{
    float *ia = A;
    float *ib = B;
    float *ic = C;
    for (int iy=0; iy<ny; iy++) {
        for (int ix=0; ix<nx; ix++) {
            ic[ix] = ia[ix] + ib[ix];
        }
        ia += nx;
        ib += nx;
        ic += nx;
    }
}

void sumMatrixOnHost_2( float *A, float *B, float *C, const int nx,
        const int ny )
{
    for (int iy=0; iy<ny; iy++) {
        for (int ix=0; ix<nx; ix++) {
            C[ix + iy * nx] = A[ix + iy * nx] + B[ix + iy * nx];
        }
    }
}

void sumMatrixOnHost_3( float *A, float *B, float *C, const int nx,
        const int ny )
{
    int nxy = nx * ny;
    for (int i=0; i<nxy; i++) {
            C[i] = A[i] + B[i];
    }
}

__global__ void sumMatrixOnDevice_1( float *A, float *B, float *C,
        const int nx, const int ny )
{
    size_t ix = threadIdx.x + blockIdx.x * blockDim.x;
    size_t iy = threadIdx.y + blockIdx.y * blockDim.y;

    if (ix < nx && iy < ny) {
        C[ix + iy * nx] = A[ix + iy * nx] + B[ix + iy * nx];
    }
    return;
}

double cpuSecond()
{
    struct timeval tp;
    gettimeofday( &tp, NULL );
    return ( (double)tp.tv_sec + (double)tp.tv_usec*1e-6 );
}

int main( int argc, char **argv )
{
    // timing...
    double startTime, cpuElapsed_1, cpuElapsed_2, cpuElapsed_3, gpuElapsed;

    // data
    int nx = 1<<14;
    int ny = 1<<14;
    int nxy = nx * ny;
    size_t nBytes = nxy * sizeof(float);
    printf("matrix dimensions  = (%i, %i)\n", nx, ny);
    printf("number of elements = %i\n", nxy);
    printf("number of bytes in each matrix = %lu\n", nBytes);

    float *h_A, *h_B, *h_C;
    h_A = (float *)malloc(nBytes);
    h_B = (float *)malloc(nBytes);
    h_C = (float *)malloc(nBytes);

    initialData( h_A, nxy );
    initialData( h_B, nxy );

    // set up device
    int dev = 0;
    cudaSetDevice(dev);

    float *d_A, *d_B, *d_C, *gpuRes;
    cudaMalloc( (float **)&d_A, nBytes );
    cudaMalloc( (float **)&d_B, nBytes );
    cudaMalloc( (float **)&d_C, nBytes );
    gpuRes = (float *)malloc(nBytes);

    cudaMemcpy( d_A, h_A, nBytes, cudaMemcpyHostToDevice );
    cudaMemcpy( d_B, h_B, nBytes, cudaMemcpyHostToDevice );

    memset( h_C, 0, nBytes );
    memset( gpuRes, 0, nBytes );

    startTime = cpuSecond();
    sumMatrixOnHost_1( h_A, h_B, h_C, nx, ny );
    cpuElapsed_1 = cpuSecond() - startTime;

    startTime = cpuSecond();
    sumMatrixOnHost_2( h_A, h_B, h_C, nx, ny );
    cpuElapsed_2 = cpuSecond() - startTime;

    startTime = cpuSecond();
    sumMatrixOnHost_3( h_A, h_B, h_C, nx, ny );
    cpuElapsed_3 = cpuSecond() - startTime;

    dim3 block(32, 32);
    dim3 grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y);

    printf("block = (%i, %i, %i)\n", block.x, block.y, block.z);
    printf("grid  = (%i, %i, %i)\n", grid.x, grid.y, grid.z);
    printf("number of threads / block = %i\n", block.x * block.y * block.z);
    if (block.x * block.y * block.z > 1024) {
        printf("Can only launch 1024 threads per block, not %i\n",
                block.x * block.y * block.z);
        exit(1);
    }
    printf("total launched = %i\n", block.x * grid.x * block.y * grid.y);
    printf("total needed   = %i\n", nxy);

    startTime = cpuSecond();
    sumMatrixOnDevice_1<<< grid, block >>>( d_A, d_B, d_C, nx, ny );
    cudaDeviceSynchronize();
    printf("Launched with blockDim = (%i, %i, %i)\n", block.x, block.y, block.z);
    printf("Launched with  gridDim = (%i, %i, %i)\n", grid.x, grid.y, grid.z);
    gpuElapsed = cpuSecond() - startTime;

    cudaError_t error = cudaMemcpy( gpuRes, d_C, nBytes,
            cudaMemcpyDeviceToHost );

    if (error != cudaSuccess) {
        printf("Error: %s:%d, ", __FILE__, __LINE__);
        printf("code:%d, reason: %s\n", error, cudaGetErrorString(error));
        exit(1);
    }

    double err = 0.0;
    for (int i=0; i<nxy; i++) {
        err += abs(h_C[i] - gpuRes[i]);
        // printf("%3i:: %7.3f %7.3f %7.3f %7.3f\n", i, h_C[i], gpuRes[i], h_A[i], h_B[i]);
    }

    printf("Total error is %f\n", err);
    printf("Time on CPU_v1 is %f\n", cpuElapsed_1);
    printf("Time on CPU_v2 is %f\n", cpuElapsed_2);
    printf("Time on CPU_v3 is %f\n", cpuElapsed_3);
    printf("Time on GPU    is %f\n", gpuElapsed);
    printf("GPU speed-up over CPU is %.2f x\n",
            (cpuElapsed_1 < cpuElapsed_2 ? 
               (cpuElapsed_1 < cpuElapsed_3 ? cpuElapsed_1 : cpuElapsed_3) 
             : (cpuElapsed_2 < cpuElapsed_3 ? cpuElapsed_2 : cpuElapsed_3)) / gpuElapsed);

    free(h_A);
    free(h_B);
    free(h_C);

    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    free(gpuRes);

    return 0;
}

int comp(const void * elem1, const void * elem2) 
{
    int f = *((int*)elem1);
    int s = *((int*)elem2);
    if (f > s) return  1;
    if (f < s) return -1;
    return 0;
}

void checkSol()
{
    int indices[] = { 96, 97 , 98 , 99 , 100, 101, 102, 103, 104, 105, 106, 107, 108, 109, 110, 111, 112, 113, 114, 115,
                     116, 117, 118, 119, 120, 121, 122, 123, 124, 125, 126, 127, 32 , 33 , 34 , 35 , 36 , 37 , 38 , 39 ,
                     40 , 41 , 42 , 43 , 44 , 45 , 46 , 47 , 48 , 49 , 50 , 51 , 52 , 53 , 54 , 55 , 56 , 57 , 58 , 59 ,
                     60 , 61 , 62 , 63 , 128, 129, 130, 131, 132, 133, 134, 135, 136, 137, 138, 139, 140, 141, 142, 143,
                     144, 145, 146, 147, 148, 149, 150, 151, 152, 153, 154, 155, 156, 157, 158, 159, 160, 161, 162, 163,
                     164, 165, 166, 167, 168, 169, 170, 171, 172, 173, 174, 175, 176, 177, 178, 179, 180, 181, 182, 183,
                     184, 185, 186, 187, 188, 189, 190, 191, 64 , 65 , 66 , 67 , 68 , 69 , 70 , 71 , 72 , 73 , 74 , 75 ,
                     76 , 77 , 78 , 79 , 80 , 81 , 82 , 83 , 84 , 85 , 86 , 87 , 88 , 89 , 90 , 91 , 92 , 93 , 94 , 95 ,
                     0  , 1  , 2  , 3  , 4  , 5  , 6  , 7  , 8  , 9  , 10 , 11 , 12 , 13 , 14 , 15 , 16 , 17 , 18 , 19 ,
                     20 , 21 , 22 , 23 , 24 , 25 , 26 , 27 , 28 , 29 , 30 , 31 , 192, 193, 194, 195, 196, 197, 198, 199,
                     200, 201, 202, 203, 204, 205, 206, 207, 208, 209, 210, 211, 212, 213, 214, 215, 216, 217, 218, 219,
                     220, 221, 222, 223, 224, 225, 226, 227, 228, 229, 230, 231, 232, 233, 234, 235, 236, 237, 238, 239,
                     240, 241, 242, 243, 244, 245, 246, 247, 248, 249, 250, 251, 252, 253, 254, 255 };

    qsort(indices, sizeof(indices)/sizeof(*indices), sizeof(*indices), comp);

    for (int i = 0 ; i < sizeof(indices)/sizeof(*indices) ; i++) {
        // printf ("%d ", indices[i]);
        if (i != indices[i]) {
            printf("failed:: %d:%d\n", i, indices[i]);
            break;
        }
    }
    printf("success:: all matched\n");
    return;
}

	.file	"tmpxft_002de5aa_00000000-6_sumMatrices.cudafe1.cpp"
	.text
#APP

#NO_APP
	.globl	_Z4compPKvS0_
	.type	_Z4compPKvS0_, @function
_Z4compPKvS0_:
.LFB2033:
	.cfi_startproc
	endbr64
	movl	(%rsi), %edx
	movl	$1, %eax
	cmpl	%edx, (%rdi)
	jg	.L1
	setl	%al
	movzbl	%al, %eax
	negl	%eax
.L1:
	ret
	.cfi_endproc
.LFE2033:
	.size	_Z4compPKvS0_, .-_Z4compPKvS0_
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2037:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2037:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z11initialDataPfi
	.type	_Z11initialDataPfi, @function
_Z11initialDataPfi:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movq	%rdi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movl	%esi, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	xorl	%ebx, %ebx
	subq	$16, %rsp
	.cfi_def_cfa_offset 48
	movq	%fs:40, %rax
	movq	%rax, 8(%rsp)
	xorl	%eax, %eax
	movq	%rsp, %rdi
	call	time@PLT
	movl	%eax, %edi
	call	srand@PLT
.L7:
	cmpl	%ebx, %ebp
	jle	.L11
	call	rand@PLT
	movzbl	%al, %eax
	cvtsi2ssl	%eax, %xmm0
	divss	.LC1(%rip), %xmm0
	movss	%xmm0, (%r12,%rbx,4)
	incq	%rbx
	jmp	.L7
.L11:
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	je	.L9
	call	__stack_chk_fail@PLT
.L9:
	addq	$16, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z11initialDataPfi, .-_Z11initialDataPfi
	.globl	_Z17sumMatrixOnHost_1PfS_S_ii
	.type	_Z17sumMatrixOnHost_1PfS_S_ii, @function
_Z17sumMatrixOnHost_1PfS_S_ii:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movslq	%ecx, %r11
	movq	%rsi, %r9
	movq	%rdx, %r10
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	salq	$2, %r11
	xorl	%edx, %edx
	xorl	%esi, %esi
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
.L13:
	leaq	(%r10,%rdx), %r12
	leaq	(%r9,%rdx), %rbp
	leaq	(%rdi,%rdx), %rbx
	cmpl	%r8d, %esi
	jge	.L18
	xorl	%eax, %eax
.L15:
	cmpl	%eax, %ecx
	jle	.L19
	movss	(%rbx,%rax,4), %xmm0
	addss	0(%rbp,%rax,4), %xmm0
	movss	%xmm0, (%r12,%rax,4)
	incq	%rax
	jmp	.L15
.L19:
	incl	%esi
	addq	%r11, %rdx
	jmp	.L13
.L18:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2028:
	.size	_Z17sumMatrixOnHost_1PfS_S_ii, .-_Z17sumMatrixOnHost_1PfS_S_ii
	.globl	_Z17sumMatrixOnHost_2PfS_S_ii
	.type	_Z17sumMatrixOnHost_2PfS_S_ii, @function
_Z17sumMatrixOnHost_2PfS_S_ii:
.LFB2029:
	.cfi_startproc
	endbr64
	movq	%rdi, %r10
	movq	%rsi, %r11
	xorl	%edi, %edi
	xorl	%r9d, %r9d
.L21:
	cmpl	%r8d, %r9d
	jge	.L20
	movslq	%edi, %rax
	xorl	%esi, %esi
	salq	$2, %rax
.L24:
	cmpl	%ecx, %esi
	jge	.L25
	movss	(%r10,%rax), %xmm0
	addss	(%r11,%rax), %xmm0
	incl	%esi
	movss	%xmm0, (%rdx,%rax)
	addq	$4, %rax
	jmp	.L24
.L25:
	incl	%r9d
	addl	%ecx, %edi
	jmp	.L21
.L20:
	ret
	.cfi_endproc
.LFE2029:
	.size	_Z17sumMatrixOnHost_2PfS_S_ii, .-_Z17sumMatrixOnHost_2PfS_S_ii
	.globl	_Z17sumMatrixOnHost_3PfS_S_ii
	.type	_Z17sumMatrixOnHost_3PfS_S_ii, @function
_Z17sumMatrixOnHost_3PfS_S_ii:
.LFB2030:
	.cfi_startproc
	endbr64
	imull	%r8d, %ecx
	xorl	%eax, %eax
.L27:
	cmpl	%eax, %ecx
	jle	.L29
	movss	(%rdi,%rax,4), %xmm0
	addss	(%rsi,%rax,4), %xmm0
	movss	%xmm0, (%rdx,%rax,4)
	incq	%rax
	jmp	.L27
.L29:
	ret
	.cfi_endproc
.LFE2030:
	.size	_Z17sumMatrixOnHost_3PfS_S_ii, .-_Z17sumMatrixOnHost_3PfS_S_ii
	.globl	_Z9cpuSecondv
	.type	_Z9cpuSecondv, @function
_Z9cpuSecondv:
.LFB2031:
	.cfi_startproc
	endbr64
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	xorl	%esi, %esi
	movq	%fs:40, %rax
	movq	%rax, 24(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rdi
	call	gettimeofday@PLT
	cvtsi2sdq	16(%rsp), %xmm0
	mulsd	.LC2(%rip), %xmm0
	cvtsi2sdq	8(%rsp), %xmm1
	addsd	%xmm1, %xmm0
	movq	24(%rsp), %rax
	subq	%fs:40, %rax
	je	.L31
	call	__stack_chk_fail@PLT
.L31:
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2031:
	.size	_Z9cpuSecondv, .-_Z9cpuSecondv
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC3:
	.string	"failed:: %d:%d\n"
.LC4:
	.string	"success:: all matched\n"
	.section	.rodata
	.align 32
.LC0:
	.long	96
	.long	97
	.long	98
	.long	99
	.long	100
	.long	101
	.long	102
	.long	103
	.long	104
	.long	105
	.long	106
	.long	107
	.long	108
	.long	109
	.long	110
	.long	111
	.long	112
	.long	113
	.long	114
	.long	115
	.long	116
	.long	117
	.long	118
	.long	119
	.long	120
	.long	121
	.long	122
	.long	123
	.long	124
	.long	125
	.long	126
	.long	127
	.long	32
	.long	33
	.long	34
	.long	35
	.long	36
	.long	37
	.long	38
	.long	39
	.long	40
	.long	41
	.long	42
	.long	43
	.long	44
	.long	45
	.long	46
	.long	47
	.long	48
	.long	49
	.long	50
	.long	51
	.long	52
	.long	53
	.long	54
	.long	55
	.long	56
	.long	57
	.long	58
	.long	59
	.long	60
	.long	61
	.long	62
	.long	63
	.long	128
	.long	129
	.long	130
	.long	131
	.long	132
	.long	133
	.long	134
	.long	135
	.long	136
	.long	137
	.long	138
	.long	139
	.long	140
	.long	141
	.long	142
	.long	143
	.long	144
	.long	145
	.long	146
	.long	147
	.long	148
	.long	149
	.long	150
	.long	151
	.long	152
	.long	153
	.long	154
	.long	155
	.long	156
	.long	157
	.long	158
	.long	159
	.long	160
	.long	161
	.long	162
	.long	163
	.long	164
	.long	165
	.long	166
	.long	167
	.long	168
	.long	169
	.long	170
	.long	171
	.long	172
	.long	173
	.long	174
	.long	175
	.long	176
	.long	177
	.long	178
	.long	179
	.long	180
	.long	181
	.long	182
	.long	183
	.long	184
	.long	185
	.long	186
	.long	187
	.long	188
	.long	189
	.long	190
	.long	191
	.long	64
	.long	65
	.long	66
	.long	67
	.long	68
	.long	69
	.long	70
	.long	71
	.long	72
	.long	73
	.long	74
	.long	75
	.long	76
	.long	77
	.long	78
	.long	79
	.long	80
	.long	81
	.long	82
	.long	83
	.long	84
	.long	85
	.long	86
	.long	87
	.long	88
	.long	89
	.long	90
	.long	91
	.long	92
	.long	93
	.long	94
	.long	95
	.long	0
	.long	1
	.long	2
	.long	3
	.long	4
	.long	5
	.long	6
	.long	7
	.long	8
	.long	9
	.long	10
	.long	11
	.long	12
	.long	13
	.long	14
	.long	15
	.long	16
	.long	17
	.long	18
	.long	19
	.long	20
	.long	21
	.long	22
	.long	23
	.long	24
	.long	25
	.long	26
	.long	27
	.long	28
	.long	29
	.long	30
	.long	31
	.long	192
	.long	193
	.long	194
	.long	195
	.long	196
	.long	197
	.long	198
	.long	199
	.long	200
	.long	201
	.long	202
	.long	203
	.long	204
	.long	205
	.long	206
	.long	207
	.long	208
	.long	209
	.long	210
	.long	211
	.long	212
	.long	213
	.long	214
	.long	215
	.long	216
	.long	217
	.long	218
	.long	219
	.long	220
	.long	221
	.long	222
	.long	223
	.long	224
	.long	225
	.long	226
	.long	227
	.long	228
	.long	229
	.long	230
	.long	231
	.long	232
	.long	233
	.long	234
	.long	235
	.long	236
	.long	237
	.long	238
	.long	239
	.long	240
	.long	241
	.long	242
	.long	243
	.long	244
	.long	245
	.long	246
	.long	247
	.long	248
	.long	249
	.long	250
	.long	251
	.long	252
	.long	253
	.long	254
	.long	255
	.text
	.globl	_Z8checkSolv
	.type	_Z8checkSolv, @function
_Z8checkSolv:
.LFB2034:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movl	$4, %edx
	movl	$256, %ecx
	leaq	.LC0(%rip), %rsi
	subq	$1040, %rsp
	.cfi_def_cfa_offset 1056
	movq	%fs:40, %rax
	movq	%rax, 1032(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rbx
	leaq	8(%rsp), %rdi
	rep movsl
	leaq	_Z4compPKvS0_(%rip), %rcx
	movl	$256, %esi
	movq	%rbx, %rdi
	call	qsort@PLT
	xorl	%edx, %edx
.L36:
	movl	(%rbx,%rdx,4), %ecx
	cmpl	%edx, %ecx
	je	.L34
	leaq	.LC3(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	jmp	.L35
.L34:
	incq	%rdx
	cmpq	$256, %rdx
	jne	.L36
.L35:
	xorl	%eax, %eax
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	1032(%rsp), %rax
	subq	%fs:40, %rax
	je	.L37
	call	__stack_chk_fail@PLT
.L37:
	addq	$1040, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2034:
	.size	_Z8checkSolv, .-_Z8checkSolv
	.globl	_Z45__device_stub__Z19sumMatrixOnDevice_1PfS_S_iiPfS_S_ii
	.type	_Z45__device_stub__Z19sumMatrixOnDevice_1PfS_S_iiPfS_S_ii, @function
_Z45__device_stub__Z19sumMatrixOnDevice_1PfS_S_iiPfS_S_ii:
.LFB2059:
	.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	.L40
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z19sumMatrixOnDevice_1PfS_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
.L40:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L42
	call	__stack_chk_fail@PLT
.L42:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2059:
	.size	_Z45__device_stub__Z19sumMatrixOnDevice_1PfS_S_iiPfS_S_ii, .-_Z45__device_stub__Z19sumMatrixOnDevice_1PfS_S_iiPfS_S_ii
	.globl	_Z19sumMatrixOnDevice_1PfS_S_ii
	.type	_Z19sumMatrixOnDevice_1PfS_S_ii, @function
_Z19sumMatrixOnDevice_1PfS_S_ii:
.LFB2060:
	.cfi_startproc
	endbr64
	jmp	_Z45__device_stub__Z19sumMatrixOnDevice_1PfS_S_iiPfS_S_ii
	.cfi_endproc
.LFE2060:
	.size	_Z19sumMatrixOnDevice_1PfS_S_ii, .-_Z19sumMatrixOnDevice_1PfS_S_ii
	.section	.rodata.str1.1
.LC6:
	.string	"matrix dimensions  = (%i, %i)\n"
.LC7:
	.string	"number of elements = %i\n"
.LC8:
	.string	"number of bytes in each matrix = %lu\n"
.LC9:
	.string	"block = (%i, %i, %i)\n"
.LC10:
	.string	"grid  = (%i, %i, %i)\n"
.LC11:
	.string	"number of threads / block = %i\n"
.LC12:
	.string	"total launched = %i\n"
.LC13:
	.string	"total needed   = %i\n"
.LC14:
	.string	"Launched with blockDim = (%i, %i, %i)\n"
.LC15:
	.string	"Launched with  gridDim = (%i, %i, %i)\n"
.LC16:
	.string	"/home/ubuntu/Datasets/stackv2/train-structured/jamieborder/cuda_c_programming/master/chapter2/sumMatrices.cu"
.LC17:
	.string	"Error: %s:%d, "
.LC18:
	.string	"code:%d, reason: %s\n"
.LC20:
	.string	"Total error is %f\n"
.LC21:
	.string	"Time on CPU_v1 is %f\n"
.LC22:
	.string	"Time on CPU_v2 is %f\n"
.LC23:
	.string	"Time on CPU_v3 is %f\n"
.LC24:
	.string	"Time on GPU    is %f\n"
.LC25:
	.string	"GPU speed-up over CPU is %.2f x\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2032:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	$16384, %ecx
	movl	$16384, %edx
	leaq	.LC6(%rip), %rsi
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movl	$2, %edi
	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
	subq	$104, %rsp
	.cfi_def_cfa_offset 160
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$268435456, %edx
	leaq	.LC7(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	$1073741824, %edx
	leaq	.LC8(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	$1073741824, %edi
	call	malloc@PLT
	movl	$1073741824, %edi
	movq	%rax, %r12
	call	malloc@PLT
	movl	$1073741824, %edi
	movq	%rax, %rbp
	call	malloc@PLT
	movl	$268435456, %esi
	movq	%r12, %rdi
	movq	%rax, %rbx
	call	_Z11initialDataPfi
	movl	$268435456, %esi
	movq	%rbp, %rdi
	call	_Z11initialDataPfi
	xorl	%edi, %edi
	call	cudaSetDevice@PLT
	leaq	40(%rsp), %rdi
	movl	$1073741824, %esi
	call	cudaMalloc@PLT
	leaq	48(%rsp), %rdi
	movl	$1073741824, %esi
	call	cudaMalloc@PLT
	movl	$1073741824, %esi
	leaq	56(%rsp), %rdi
	call	cudaMalloc@PLT
	movl	$1073741824, %edi
	call	malloc@PLT
	movq	40(%rsp), %rdi
	movl	$1, %ecx
	movq	%r12, %rsi
	movl	$1073741824, %edx
	movq	%rax, %r13
	call	cudaMemcpy@PLT
	movq	48(%rsp), %rdi
	movl	$1073741824, %edx
	movq	%rbp, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	xorl	%eax, %eax
	movl	$268435456, %ecx
	movq	%rbx, %rdi
	rep stosl
	movl	$268435456, %ecx
	movq	%r13, %rdi
	rep stosl
	call	_Z9cpuSecondv
	movl	$16384, %ecx
	movq	%rbx, %rdx
	movq	%rbp, %rsi
	movl	$16384, %r8d
	movq	%r12, %rdi
	movsd	%xmm0, 8(%rsp)
	call	_Z17sumMatrixOnHost_1PfS_S_ii
	call	_Z9cpuSecondv
	subsd	8(%rsp), %xmm0
	movsd	%xmm0, 16(%rsp)
	call	_Z9cpuSecondv
	movl	$16384, %ecx
	movq	%rbx, %rdx
	movq	%rbp, %rsi
	movl	$16384, %r8d
	movq	%r12, %rdi
	movsd	%xmm0, 8(%rsp)
	call	_Z17sumMatrixOnHost_2PfS_S_ii
	call	_Z9cpuSecondv
	subsd	8(%rsp), %xmm0
	movq	%xmm0, %r15
	call	_Z9cpuSecondv
	movl	$16384, %ecx
	movq	%rbx, %rdx
	movq	%rbp, %rsi
	movl	$16384, %r8d
	movq	%r12, %rdi
	movsd	%xmm0, 8(%rsp)
	call	_Z17sumMatrixOnHost_3PfS_S_ii
	call	_Z9cpuSecondv
	subsd	8(%rsp), %xmm0
	movl	$1, %r8d
	movl	$32, %ecx
	movl	$32, %edx
	leaq	.LC9(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movsd	%xmm0, 8(%rsp)
	call	__printf_chk@PLT
	xorl	%eax, %eax
	movl	$1, %r8d
	movl	$512, %ecx
	movl	$512, %edx
	leaq	.LC10(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	$1024, %edx
	leaq	.LC11(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	$268435456, %edx
	leaq	.LC12(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	$268435456, %edx
	leaq	.LC13(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	call	_Z9cpuSecondv
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movabsq	$137438953504, %rdx
	movl	$1, %esi
	movabsq	$2199023256064, %rdi
	movsd	%xmm0, 24(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L46
	movq	56(%rsp), %rdx
	movq	48(%rsp), %rsi
	movl	$16384, %ecx
	movl	$16384, %r8d
	movq	40(%rsp), %rdi
	call	_Z45__device_stub__Z19sumMatrixOnDevice_1PfS_S_iiPfS_S_ii
.L46:
	call	cudaDeviceSynchronize@PLT
	movl	$1, %r8d
	xorl	%eax, %eax
	movl	$32, %ecx
	movl	$32, %edx
	leaq	.LC14(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	$1, %r8d
	xorl	%eax, %eax
	movl	$512, %ecx
	movl	$512, %edx
	leaq	.LC15(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	call	_Z9cpuSecondv
	subsd	24(%rsp), %xmm0
	movq	56(%rsp), %rsi
	movq	%r13, %rdi
	movl	$2, %ecx
	movl	$1073741824, %edx
	movsd	%xmm0, 24(%rsp)
	call	cudaMemcpy@PLT
	xorps	%xmm1, %xmm1
	movss	.LC19(%rip), %xmm2
	movl	%eax, %r14d
	xorl	%eax, %eax
	testl	%r14d, %r14d
	je	.L47
	movl	$147, %ecx
	leaq	.LC16(%rip), %rdx
	leaq	.LC17(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	%r14d, %edi
	call	cudaGetErrorString@PLT
	movl	$2, %edi
	movl	%r14d, %edx
	leaq	.LC18(%rip), %rsi
	movq	%rax, %rcx
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$1, %edi
	call	exit@PLT
.L47:
	movss	(%rbx,%rax), %xmm0
	subss	0(%r13,%rax), %xmm0
	addq	$4, %rax
	andps	%xmm2, %xmm0
	cvtss2sd	%xmm0, %xmm0
	addsd	%xmm0, %xmm1
	cmpq	$1073741824, %rax
	jne	.L47
	movaps	%xmm1, %xmm0
	leaq	.LC20(%rip), %rsi
	movl	$2, %edi
	movb	$1, %al
	call	__printf_chk@PLT
	movsd	16(%rsp), %xmm0
	leaq	.LC21(%rip), %rsi
	movb	$1, %al
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	%r15, %xmm0
	leaq	.LC22(%rip), %rsi
	movb	$1, %al
	movl	$2, %edi
	call	__printf_chk@PLT
	movsd	8(%rsp), %xmm0
	leaq	.LC23(%rip), %rsi
	movb	$1, %al
	movl	$2, %edi
	call	__printf_chk@PLT
	movsd	24(%rsp), %xmm0
	leaq	.LC24(%rip), %rsi
	movb	$1, %al
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	%r15, %xmm5
	comisd	16(%rsp), %xmm5
	jbe	.L59
	movsd	16(%rsp), %xmm3
	minsd	8(%rsp), %xmm3
	movsd	%xmm3, 8(%rsp)
	jmp	.L50
.L59:
	movq	%r15, %xmm4
	minsd	8(%rsp), %xmm4
	movsd	%xmm4, 8(%rsp)
.L50:
	leaq	.LC25(%rip), %rsi
	movl	$2, %edi
	movb	$1, %al
	movsd	8(%rsp), %xmm0
	divsd	24(%rsp), %xmm0
	call	__printf_chk@PLT
	movq	%r12, %rdi
	call	free@PLT
	movq	%rbp, %rdi
	call	free@PLT
	movq	%rbx, %rdi
	call	free@PLT
	movq	40(%rsp), %rdi
	call	cudaFree@PLT
	movq	48(%rsp), %rdi
	call	cudaFree@PLT
	movq	56(%rsp), %rdi
	call	cudaFree@PLT
	movq	%r13, %rdi
	call	free@PLT
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L51
	call	__stack_chk_fail@PLT
.L51:
	addq	$104, %rsp
	.cfi_def_cfa_offset 56
	xorl	%eax, %eax
	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
.LFE2032:
	.size	main, .-main
	.section	.rodata.str1.1
.LC26:
	.string	"_Z19sumMatrixOnDevice_1PfS_S_ii"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2062:
	.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	.LC26(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z19sumMatrixOnDevice_1PfS_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
.LFE2062:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC1:
	.long	1092616192
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC2:
	.long	-1598689907
	.long	1051772663
	.section	.rodata.cst16,"aM",@progbits,16
	.align 16
.LC19:
	.long	2147483647
	.long	0
	.long	0
	.long	0
	.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 <hip/hip_runtime.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <time.h>
#include <sys/time.h>

void initialData( float *ip, int size )
{
    // generate different seed for random number
    time_t t;
    srand( (unsigned int) time (&t) );

    for (int i=0; i<size; i++) {
        ip[i] = (float)( rand() & 0xFF ) / 10.0f;
    }
}

void sumMatrixOnHost_1( float *A, float *B, float *C, const int nx,
        const int ny )
{
    float *ia = A;
    float *ib = B;
    float *ic = C;
    for (int iy=0; iy<ny; iy++) {
        for (int ix=0; ix<nx; ix++) {
            ic[ix] = ia[ix] + ib[ix];
        }
        ia += nx;
        ib += nx;
        ic += nx;
    }
}

void sumMatrixOnHost_2( float *A, float *B, float *C, const int nx,
        const int ny )
{
    for (int iy=0; iy<ny; iy++) {
        for (int ix=0; ix<nx; ix++) {
            C[ix + iy * nx] = A[ix + iy * nx] + B[ix + iy * nx];
        }
    }
}

void sumMatrixOnHost_3( float *A, float *B, float *C, const int nx,
        const int ny )
{
    int nxy = nx * ny;
    for (int i=0; i<nxy; i++) {
            C[i] = A[i] + B[i];
    }
}

__global__ void sumMatrixOnDevice_1( float *A, float *B, float *C,
        const int nx, const int ny )
{
    size_t ix = threadIdx.x + blockIdx.x * blockDim.x;
    size_t iy = threadIdx.y + blockIdx.y * blockDim.y;

    if (ix < nx && iy < ny) {
        C[ix + iy * nx] = A[ix + iy * nx] + B[ix + iy * nx];
    }
    return;
}

double cpuSecond()
{
    struct timeval tp;
    gettimeofday( &tp, NULL );
    return ( (double)tp.tv_sec + (double)tp.tv_usec*1e-6 );
}

int main( int argc, char **argv )
{
    // timing...
    double startTime, cpuElapsed_1, cpuElapsed_2, cpuElapsed_3, gpuElapsed;

    // data
    int nx = 1<<14;
    int ny = 1<<14;
    int nxy = nx * ny;
    size_t nBytes = nxy * sizeof(float);
    printf("matrix dimensions  = (%i, %i)\n", nx, ny);
    printf("number of elements = %i\n", nxy);
    printf("number of bytes in each matrix = %lu\n", nBytes);

    float *h_A, *h_B, *h_C;
    h_A = (float *)malloc(nBytes);
    h_B = (float *)malloc(nBytes);
    h_C = (float *)malloc(nBytes);

    initialData( h_A, nxy );
    initialData( h_B, nxy );

    // set up device
    int dev = 0;
    hipSetDevice(dev);

    float *d_A, *d_B, *d_C, *gpuRes;
    hipMalloc( (float **)&d_A, nBytes );
    hipMalloc( (float **)&d_B, nBytes );
    hipMalloc( (float **)&d_C, nBytes );
    gpuRes = (float *)malloc(nBytes);

    hipMemcpy( d_A, h_A, nBytes, hipMemcpyHostToDevice );
    hipMemcpy( d_B, h_B, nBytes, hipMemcpyHostToDevice );

    memset( h_C, 0, nBytes );
    memset( gpuRes, 0, nBytes );

    startTime = cpuSecond();
    sumMatrixOnHost_1( h_A, h_B, h_C, nx, ny );
    cpuElapsed_1 = cpuSecond() - startTime;

    startTime = cpuSecond();
    sumMatrixOnHost_2( h_A, h_B, h_C, nx, ny );
    cpuElapsed_2 = cpuSecond() - startTime;

    startTime = cpuSecond();
    sumMatrixOnHost_3( h_A, h_B, h_C, nx, ny );
    cpuElapsed_3 = cpuSecond() - startTime;

    dim3 block(32, 32);
    dim3 grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y);

    printf("block = (%i, %i, %i)\n", block.x, block.y, block.z);
    printf("grid  = (%i, %i, %i)\n", grid.x, grid.y, grid.z);
    printf("number of threads / block = %i\n", block.x * block.y * block.z);
    if (block.x * block.y * block.z > 1024) {
        printf("Can only launch 1024 threads per block, not %i\n",
                block.x * block.y * block.z);
        exit(1);
    }
    printf("total launched = %i\n", block.x * grid.x * block.y * grid.y);
    printf("total needed   = %i\n", nxy);

    startTime = cpuSecond();
    sumMatrixOnDevice_1<<< grid, block >>>( d_A, d_B, d_C, nx, ny );
    hipDeviceSynchronize();
    printf("Launched with blockDim = (%i, %i, %i)\n", block.x, block.y, block.z);
    printf("Launched with  gridDim = (%i, %i, %i)\n", grid.x, grid.y, grid.z);
    gpuElapsed = cpuSecond() - startTime;

    hipError_t error = hipMemcpy( gpuRes, d_C, nBytes,
            hipMemcpyDeviceToHost );

    if (error != hipSuccess) {
        printf("Error: %s:%d, ", __FILE__, __LINE__);
        printf("code:%d, reason: %s\n", error, hipGetErrorString(error));
        exit(1);
    }

    double err = 0.0;
    for (int i=0; i<nxy; i++) {
        err += abs(h_C[i] - gpuRes[i]);
        // printf("%3i:: %7.3f %7.3f %7.3f %7.3f\n", i, h_C[i], gpuRes[i], h_A[i], h_B[i]);
    }

    printf("Total error is %f\n", err);
    printf("Time on CPU_v1 is %f\n", cpuElapsed_1);
    printf("Time on CPU_v2 is %f\n", cpuElapsed_2);
    printf("Time on CPU_v3 is %f\n", cpuElapsed_3);
    printf("Time on GPU    is %f\n", gpuElapsed);
    printf("GPU speed-up over CPU is %.2f x\n",
            (cpuElapsed_1 < cpuElapsed_2 ? 
               (cpuElapsed_1 < cpuElapsed_3 ? cpuElapsed_1 : cpuElapsed_3) 
             : (cpuElapsed_2 < cpuElapsed_3 ? cpuElapsed_2 : cpuElapsed_3)) / gpuElapsed);

    free(h_A);
    free(h_B);
    free(h_C);

    hipFree(d_A);
    hipFree(d_B);
    hipFree(d_C);
    free(gpuRes);

    return 0;
}

int comp(const void * elem1, const void * elem2) 
{
    int f = *((int*)elem1);
    int s = *((int*)elem2);
    if (f > s) return  1;
    if (f < s) return -1;
    return 0;
}

void checkSol()
{
    int indices[] = { 96, 97 , 98 , 99 , 100, 101, 102, 103, 104, 105, 106, 107, 108, 109, 110, 111, 112, 113, 114, 115,
                     116, 117, 118, 119, 120, 121, 122, 123, 124, 125, 126, 127, 32 , 33 , 34 , 35 , 36 , 37 , 38 , 39 ,
                     40 , 41 , 42 , 43 , 44 , 45 , 46 , 47 , 48 , 49 , 50 , 51 , 52 , 53 , 54 , 55 , 56 , 57 , 58 , 59 ,
                     60 , 61 , 62 , 63 , 128, 129, 130, 131, 132, 133, 134, 135, 136, 137, 138, 139, 140, 141, 142, 143,
                     144, 145, 146, 147, 148, 149, 150, 151, 152, 153, 154, 155, 156, 157, 158, 159, 160, 161, 162, 163,
                     164, 165, 166, 167, 168, 169, 170, 171, 172, 173, 174, 175, 176, 177, 178, 179, 180, 181, 182, 183,
                     184, 185, 186, 187, 188, 189, 190, 191, 64 , 65 , 66 , 67 , 68 , 69 , 70 , 71 , 72 , 73 , 74 , 75 ,
                     76 , 77 , 78 , 79 , 80 , 81 , 82 , 83 , 84 , 85 , 86 , 87 , 88 , 89 , 90 , 91 , 92 , 93 , 94 , 95 ,
                     0  , 1  , 2  , 3  , 4  , 5  , 6  , 7  , 8  , 9  , 10 , 11 , 12 , 13 , 14 , 15 , 16 , 17 , 18 , 19 ,
                     20 , 21 , 22 , 23 , 24 , 25 , 26 , 27 , 28 , 29 , 30 , 31 , 192, 193, 194, 195, 196, 197, 198, 199,
                     200, 201, 202, 203, 204, 205, 206, 207, 208, 209, 210, 211, 212, 213, 214, 215, 216, 217, 218, 219,
                     220, 221, 222, 223, 224, 225, 226, 227, 228, 229, 230, 231, 232, 233, 234, 235, 236, 237, 238, 239,
                     240, 241, 242, 243, 244, 245, 246, 247, 248, 249, 250, 251, 252, 253, 254, 255 };

    qsort(indices, sizeof(indices)/sizeof(*indices), sizeof(*indices), comp);

    for (int i = 0 ; i < sizeof(indices)/sizeof(*indices) ; i++) {
        // printf ("%d ", indices[i]);
        if (i != indices[i]) {
            printf("failed:: %d:%d\n", i, indices[i]);
            break;
        }
    }
    printf("success:: all matched\n");
    return;
}

	.text
	.file	"sumMatrices.hip"
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _Z11initialDataPfi
.LCPI0_0:
	.long	0x41200000                      # float 10
	.text
	.globl	_Z11initialDataPfi
	.type	_Z11initialDataPfi,@function
_Z11initialDataPfi:                     # @_Z11initialDataPfi
	.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	%rbx
	.cfi_def_cfa_offset 40
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%esi, %ebp
	movq	%rdi, %rbx
	movq	%rsp, %rdi
	callq	time
	movl	%eax, %edi
	callq	srand
	testl	%ebp, %ebp
	jle	.LBB0_3
# %bb.1:                                # %.lr.ph.preheader
	movl	%ebp, %r14d
	xorl	%r15d, %r15d
.LBB0_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	callq	rand
	movzbl	%al, %eax
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	divss	.LCPI0_0(%rip), %xmm0
	movss	%xmm0, (%rbx,%r15,4)
	incq	%r15
	cmpq	%r15, %r14
	jne	.LBB0_2
.LBB0_3:                                # %._crit_edge
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.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_end0:
	.size	_Z11initialDataPfi, .Lfunc_end0-_Z11initialDataPfi
	.cfi_endproc
                                        # -- End function
	.globl	_Z17sumMatrixOnHost_1PfS_S_ii   # -- Begin function _Z17sumMatrixOnHost_1PfS_S_ii
	.type	_Z17sumMatrixOnHost_1PfS_S_ii,@function
_Z17sumMatrixOnHost_1PfS_S_ii:          # @_Z17sumMatrixOnHost_1PfS_S_ii
	.cfi_startproc
# %bb.0:
	testl	%r8d, %r8d
	jle	.LBB1_6
# %bb.1:                                # %.preheader.lr.ph
	movslq	%ecx, %rax
	movl	%eax, %r9d
	shlq	$2, %rax
	xorl	%r10d, %r10d
.LBB1_2:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB1_4 Depth 2
	testl	%ecx, %ecx
	jle	.LBB1_5
# %bb.3:                                # %.lr.ph.preheader
                                        #   in Loop: Header=BB1_2 Depth=1
	xorl	%r11d, %r11d
.LBB1_4:                                # %.lr.ph
                                        #   Parent Loop BB1_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movss	(%rdi,%r11,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
	addss	(%rsi,%r11,4), %xmm0
	movss	%xmm0, (%rdx,%r11,4)
	incq	%r11
	cmpq	%r11, %r9
	jne	.LBB1_4
.LBB1_5:                                # %._crit_edge
                                        #   in Loop: Header=BB1_2 Depth=1
	incl	%r10d
	addq	%rax, %rdx
	addq	%rax, %rsi
	addq	%rax, %rdi
	cmpl	%r8d, %r10d
	jne	.LBB1_2
.LBB1_6:                                # %._crit_edge30
	retq
.Lfunc_end1:
	.size	_Z17sumMatrixOnHost_1PfS_S_ii, .Lfunc_end1-_Z17sumMatrixOnHost_1PfS_S_ii
	.cfi_endproc
                                        # -- End function
	.globl	_Z17sumMatrixOnHost_2PfS_S_ii   # -- Begin function _Z17sumMatrixOnHost_2PfS_S_ii
	.type	_Z17sumMatrixOnHost_2PfS_S_ii,@function
_Z17sumMatrixOnHost_2PfS_S_ii:          # @_Z17sumMatrixOnHost_2PfS_S_ii
	.cfi_startproc
# %bb.0:
	testl	%r8d, %r8d
	jle	.LBB2_7
# %bb.1:                                # %.preheader.lr.ph
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movl	%r8d, %eax
	movl	%ecx, %r8d
	xorl	%r9d, %r9d
	xorl	%r10d, %r10d
.LBB2_2:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_4 Depth 2
	testl	%ecx, %ecx
	jle	.LBB2_5
# %bb.3:                                # %.lr.ph
                                        #   in Loop: Header=BB2_2 Depth=1
	movl	%r9d, %r14d
	leaq	(%rdx,%r14,4), %r11
	leaq	(%rsi,%r14,4), %rbx
	leaq	(%rdi,%r14,4), %r14
	xorl	%r15d, %r15d
.LBB2_4:                                #   Parent Loop BB2_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movss	(%r14,%r15,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
	addss	(%rbx,%r15,4), %xmm0
	movss	%xmm0, (%r11,%r15,4)
	incq	%r15
	cmpq	%r15, %r8
	jne	.LBB2_4
.LBB2_5:                                # %._crit_edge
                                        #   in Loop: Header=BB2_2 Depth=1
	incq	%r10
	addl	%ecx, %r9d
	cmpq	%rax, %r10
	jne	.LBB2_2
# %bb.6:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	.cfi_restore %rbx
	.cfi_restore %r14
	.cfi_restore %r15
.LBB2_7:                                # %._crit_edge21
	retq
.Lfunc_end2:
	.size	_Z17sumMatrixOnHost_2PfS_S_ii, .Lfunc_end2-_Z17sumMatrixOnHost_2PfS_S_ii
	.cfi_endproc
                                        # -- End function
	.globl	_Z17sumMatrixOnHost_3PfS_S_ii   # -- Begin function _Z17sumMatrixOnHost_3PfS_S_ii
	.type	_Z17sumMatrixOnHost_3PfS_S_ii,@function
_Z17sumMatrixOnHost_3PfS_S_ii:          # @_Z17sumMatrixOnHost_3PfS_S_ii
	.cfi_startproc
# %bb.0:
	imull	%r8d, %ecx
	testl	%ecx, %ecx
	jle	.LBB3_3
# %bb.1:                                # %.lr.ph.preheader
	movl	%ecx, %eax
	xorl	%ecx, %ecx
.LBB3_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movss	(%rdi,%rcx,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
	addss	(%rsi,%rcx,4), %xmm0
	movss	%xmm0, (%rdx,%rcx,4)
	incq	%rcx
	cmpq	%rcx, %rax
	jne	.LBB3_2
.LBB3_3:                                # %._crit_edge
	retq
.Lfunc_end3:
	.size	_Z17sumMatrixOnHost_3PfS_S_ii, .Lfunc_end3-_Z17sumMatrixOnHost_3PfS_S_ii
	.cfi_endproc
                                        # -- End function
	.globl	_Z34__device_stub__sumMatrixOnDevice_1PfS_S_ii # -- Begin function _Z34__device_stub__sumMatrixOnDevice_1PfS_S_ii
	.type	_Z34__device_stub__sumMatrixOnDevice_1PfS_S_ii,@function
_Z34__device_stub__sumMatrixOnDevice_1PfS_S_ii: # @_Z34__device_stub__sumMatrixOnDevice_1PfS_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	$_Z19sumMatrixOnDevice_1PfS_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_end4:
	.size	_Z34__device_stub__sumMatrixOnDevice_1PfS_S_ii, .Lfunc_end4-_Z34__device_stub__sumMatrixOnDevice_1PfS_S_ii
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _Z9cpuSecondv
.LCPI5_0:
	.quad	0x3eb0c6f7a0b5ed8d              # double 9.9999999999999995E-7
	.text
	.globl	_Z9cpuSecondv
	.type	_Z9cpuSecondv,@function
_Z9cpuSecondv:                          # @_Z9cpuSecondv
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$16, %rsp
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -16
	movq	%rsp, %rbx
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	cvtsi2sdq	(%rbx), %xmm1
	cvtsi2sdq	8(%rbx), %xmm0
	mulsd	.LCPI5_0(%rip), %xmm0
	addsd	%xmm1, %xmm0
	addq	$16, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end5:
	.size	_Z9cpuSecondv, .Lfunc_end5-_Z9cpuSecondv
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function main
.LCPI6_0:
	.long	0x7fffffff                      # float NaN
	.long	0x7fffffff                      # float NaN
	.long	0x7fffffff                      # float NaN
	.long	0x7fffffff                      # float NaN
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0
.LCPI6_1:
	.quad	0x3eb0c6f7a0b5ed8d              # double 9.9999999999999995E-7
	.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	$168, %rsp
	.cfi_def_cfa_offset 224
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	xorl	%ebp, %ebp
	movl	$.L.str, %edi
	movl	$16384, %esi                    # imm = 0x4000
	movl	$16384, %edx                    # imm = 0x4000
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.1, %edi
	movl	$268435456, %esi                # imm = 0x10000000
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.2, %edi
	movl	$1073741824, %esi               # imm = 0x40000000
	xorl	%eax, %eax
	callq	printf
	movl	$1073741824, %edi               # imm = 0x40000000
	callq	malloc
	movq	%rax, %rbx
	movl	$1073741824, %edi               # imm = 0x40000000
	callq	malloc
	movq	%rax, %r14
	movl	$1, %edi
	movl	$1073741824, %esi               # imm = 0x40000000
	callq	calloc@PLT
	movq	%rax, %r15
	movq	%rbx, %rdi
	movl	$268435456, %esi                # imm = 0x10000000
	callq	_Z11initialDataPfi
	movq	%r14, %rdi
	movl	$268435456, %esi                # imm = 0x10000000
	callq	_Z11initialDataPfi
	xorl	%edi, %edi
	callq	hipSetDevice
	leaq	72(%rsp), %r13
	movl	$1073741824, %esi               # imm = 0x40000000
	movq	%r13, %rdi
	callq	hipMalloc
	leaq	48(%rsp), %rdi
	movl	$1073741824, %esi               # imm = 0x40000000
	callq	hipMalloc
	leaq	40(%rsp), %rdi
	movl	$1073741824, %esi               # imm = 0x40000000
	callq	hipMalloc
	movl	$1, %edi
	movl	$1073741824, %esi               # imm = 0x40000000
	callq	calloc@PLT
	movq	%rax, %r12
	movq	(%r13), %rdi
	movl	$1073741824, %edx               # imm = 0x40000000
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	48(%rsp), %rax
	movq	(%rax), %rdi
	movl	$1073741824, %edx               # imm = 0x40000000
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	56(%rsp), %r13
	movq	%r13, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movq	(%r13), %rax
	movq	%rax, 32(%rsp)                  # 8-byte Spill
	movq	8(%r13), %rax
	movq	%rax, 24(%rsp)                  # 8-byte Spill
	movl	$65536, %eax                    # imm = 0x10000
	movq	%r15, %rcx
	movq	%r14, %rdx
	movq	%rbx, %rsi
.LBB6_1:                                # %.preheader.i
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB6_2 Depth 2
	xorl	%edi, %edi
.LBB6_2:                                # %.lr.ph.i
                                        #   Parent Loop BB6_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movss	(%rsi,%rdi,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
	addss	(%rdx,%rdi,4), %xmm0
	movss	%xmm0, (%rcx,%rdi,4)
	incq	%rdi
	cmpq	$16384, %rdi                    # imm = 0x4000
	jne	.LBB6_2
# %bb.3:                                # %._crit_edge.i
                                        #   in Loop: Header=BB6_1 Depth=1
	addq	%rax, %rsi
	addq	%rax, %rdx
	addq	%rax, %rcx
	incl	%ebp
	cmpl	$16384, %ebp                    # imm = 0x4000
	jne	.LBB6_1
# %bb.4:                                # %_Z17sumMatrixOnHost_1PfS_S_ii.exit
	xorl	%ebp, %ebp
	leaq	56(%rsp), %r13
	movq	%r13, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movq	(%r13), %rax
	movq	%rax, 8(%rsp)                   # 8-byte Spill
	movq	8(%r13), %rax
	movq	%rax, 16(%rsp)                  # 8-byte Spill
	leaq	56(%rsp), %r13
	movq	%r13, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movq	(%r13), %rax
	movq	%rax, 160(%rsp)                 # 8-byte Spill
	movq	8(%r13), %rax
	movq	%rax, 152(%rsp)                 # 8-byte Spill
	movl	$65536, %eax                    # imm = 0x10000
	movq	%rbx, %rcx
	movq	%r14, %rdx
	movq	%r15, %rsi
.LBB6_5:                                # %.preheader.i95
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB6_6 Depth 2
	xorl	%edi, %edi
.LBB6_6:                                #   Parent Loop BB6_5 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movss	(%rcx,%rdi,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
	addss	(%rdx,%rdi,4), %xmm0
	movss	%xmm0, (%rsi,%rdi,4)
	incq	%rdi
	cmpq	$16384, %rdi                    # imm = 0x4000
	jne	.LBB6_6
# %bb.7:                                # %._crit_edge.i100
                                        #   in Loop: Header=BB6_5 Depth=1
	incq	%rbp
	addq	%rax, %rsi
	addq	%rax, %rdx
	addq	%rax, %rcx
	cmpq	$16384, %rbp                    # imm = 0x4000
	jne	.LBB6_5
# %bb.8:                                # %_Z17sumMatrixOnHost_2PfS_S_ii.exit
	xorl	%ebp, %ebp
	leaq	56(%rsp), %r13
	movq	%r13, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movq	(%r13), %rax
	movq	%rax, 144(%rsp)                 # 8-byte Spill
	movq	8(%r13), %rax
	movq	%rax, 136(%rsp)                 # 8-byte Spill
	leaq	56(%rsp), %r13
	movq	%r13, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movq	(%r13), %rax
	movq	%rax, 128(%rsp)                 # 8-byte Spill
	movq	8(%r13), %rax
	movq	%rax, 120(%rsp)                 # 8-byte Spill
.LBB6_9:                                # %.lr.ph.i101
                                        # =>This Inner Loop Header: Depth=1
	movss	(%rbx,%rbp,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
	addss	(%r14,%rbp,4), %xmm0
	movss	%xmm0, (%r15,%rbp,4)
	incq	%rbp
	cmpq	$268435456, %rbp                # imm = 0x10000000
	jne	.LBB6_9
# %bb.10:
	leaq	56(%rsp), %r13
	movq	%r13, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movq	(%r13), %rax
	movq	%rax, 112(%rsp)                 # 8-byte Spill
	movq	8(%r13), %rax
	movq	%rax, 104(%rsp)                 # 8-byte Spill
	movl	$.L.str.3, %edi
	movl	$32, %esi
	movl	$32, %edx
	movl	$1, %ecx
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.4, %edi
	movl	$512, %esi                      # imm = 0x200
	movl	$512, %edx                      # imm = 0x200
	movl	$1, %ecx
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.5, %edi
	movl	$1024, %esi                     # imm = 0x400
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.7, %edi
	movl	$268435456, %esi                # imm = 0x10000000
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.8, %edi
	movl	$268435456, %esi                # imm = 0x10000000
	xorl	%eax, %eax
	callq	printf
	leaq	56(%rsp), %r13
	movq	%r13, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movq	(%r13), %rax
	movq	%rax, 96(%rsp)                  # 8-byte Spill
	movq	8(%r13), %rax
	movq	%rax, 88(%rsp)                  # 8-byte Spill
	movabsq	$2199023256064, %rdi            # imm = 0x20000000200
	movabsq	$137438953504, %rdx             # imm = 0x2000000020
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB6_12
# %bb.11:
	movq	72(%rsp), %rdi
	movq	48(%rsp), %rsi
	movq	40(%rsp), %rdx
	movl	$16384, %ecx                    # imm = 0x4000
	movl	$16384, %r8d                    # imm = 0x4000
	callq	_Z34__device_stub__sumMatrixOnDevice_1PfS_S_ii
.LBB6_12:
	callq	hipDeviceSynchronize
	movl	$.L.str.9, %edi
	movl	$32, %esi
	movl	$32, %edx
	movl	$1, %ecx
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.10, %edi
	movl	$512, %esi                      # imm = 0x200
	movl	$512, %edx                      # imm = 0x200
	movl	$1, %ecx
	xorl	%eax, %eax
	callq	printf
	leaq	56(%rsp), %r13
	movq	%r13, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movq	(%r13), %rax
	movq	%rax, 80(%rsp)                  # 8-byte Spill
	movq	8(%r13), %r13
	movq	40(%rsp), %rsi
	movl	$1073741824, %edx               # imm = 0x40000000
	movq	%r12, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	jne	.LBB6_16
# %bb.13:                               # %.preheader.preheader
	xorps	%xmm0, %xmm0
	xorl	%eax, %eax
	movaps	.LCPI6_0(%rip), %xmm1           # xmm1 = [NaN,NaN,NaN,NaN]
.LBB6_14:                               # %.preheader
                                        # =>This Inner Loop Header: Depth=1
	movss	(%r15,%rax,4), %xmm2            # xmm2 = mem[0],zero,zero,zero
	subss	(%r12,%rax,4), %xmm2
	andps	%xmm1, %xmm2
	cvtss2sd	%xmm2, %xmm2
	addsd	%xmm2, %xmm0
	incq	%rax
	cmpq	$268435456, %rax                # imm = 0x10000000
	jne	.LBB6_14
# %bb.15:
	xorps	%xmm1, %xmm1
	cvtsi2sd	%r13, %xmm1
	cvtsi2sdq	80(%rsp), %xmm13        # 8-byte Folded Reload
	xorps	%xmm2, %xmm2
	cvtsi2sdq	88(%rsp), %xmm2         # 8-byte Folded Reload
	cvtsi2sdq	96(%rsp), %xmm3         # 8-byte Folded Reload
	cvtsi2sdq	104(%rsp), %xmm4        # 8-byte Folded Reload
	cvtsi2sdq	112(%rsp), %xmm14       # 8-byte Folded Reload
	cvtsi2sdq	120(%rsp), %xmm5        # 8-byte Folded Reload
	cvtsi2sdq	128(%rsp), %xmm6        # 8-byte Folded Reload
	cvtsi2sdq	136(%rsp), %xmm7        # 8-byte Folded Reload
	cvtsi2sdq	144(%rsp), %xmm15       # 8-byte Folded Reload
	cvtsi2sdq	152(%rsp), %xmm8        # 8-byte Folded Reload
	cvtsi2sdq	160(%rsp), %xmm9        # 8-byte Folded Reload
	cvtsi2sdq	16(%rsp), %xmm10        # 8-byte Folded Reload
	cvtsi2sdq	8(%rsp), %xmm11         # 8-byte Folded Reload
	movsd	%xmm11, 8(%rsp)                 # 8-byte Spill
	xorps	%xmm11, %xmm11
	cvtsi2sdq	24(%rsp), %xmm11        # 8-byte Folded Reload
	movsd	.LCPI6_1(%rip), %xmm12          # xmm12 = mem[0],zero
	mulsd	%xmm12, %xmm1
	addsd	%xmm1, %xmm13
	xorps	%xmm1, %xmm1
	cvtsi2sdq	32(%rsp), %xmm1         # 8-byte Folded Reload
	mulsd	%xmm12, %xmm2
	addsd	%xmm2, %xmm3
	subsd	%xmm3, %xmm13
	movsd	%xmm13, 32(%rsp)                # 8-byte Spill
	mulsd	%xmm12, %xmm4
	addsd	%xmm4, %xmm14
	mulsd	%xmm12, %xmm5
	addsd	%xmm5, %xmm6
	subsd	%xmm6, %xmm14
	movsd	%xmm14, 24(%rsp)                # 8-byte Spill
	mulsd	%xmm12, %xmm7
	addsd	%xmm7, %xmm15
	mulsd	%xmm12, %xmm8
	addsd	%xmm8, %xmm9
	subsd	%xmm9, %xmm15
	movsd	%xmm15, 16(%rsp)                # 8-byte Spill
	mulsd	%xmm12, %xmm10
	movsd	8(%rsp), %xmm2                  # 8-byte Reload
                                        # xmm2 = mem[0],zero
	addsd	%xmm10, %xmm2
	mulsd	%xmm12, %xmm11
	addsd	%xmm11, %xmm1
	subsd	%xmm1, %xmm2
	movsd	%xmm2, 8(%rsp)                  # 8-byte Spill
	movl	$.L.str.14, %edi
	movb	$1, %al
	callq	printf
	movl	$.L.str.15, %edi
	movsd	8(%rsp), %xmm0                  # 8-byte Reload
                                        # xmm0 = mem[0],zero
	movb	$1, %al
	callq	printf
	movl	$.L.str.16, %edi
	movsd	16(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	movb	$1, %al
	callq	printf
	movl	$.L.str.17, %edi
	movsd	24(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	movb	$1, %al
	callq	printf
	movl	$.L.str.18, %edi
	movsd	32(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	movb	$1, %al
	callq	printf
	movsd	8(%rsp), %xmm0                  # 8-byte Reload
                                        # xmm0 = mem[0],zero
	minsd	16(%rsp), %xmm0                 # 8-byte Folded Reload
	minsd	24(%rsp), %xmm0                 # 8-byte Folded Reload
	divsd	32(%rsp), %xmm0                 # 8-byte Folded Reload
	movl	$.L.str.19, %edi
	movb	$1, %al
	callq	printf
	movq	%rbx, %rdi
	callq	free
	movq	%r14, %rdi
	callq	free
	movq	%r15, %rdi
	callq	free
	movq	72(%rsp), %rdi
	callq	hipFree
	movq	48(%rsp), %rdi
	callq	hipFree
	movq	40(%rsp), %rdi
	callq	hipFree
	movq	%r12, %rdi
	callq	free
	xorl	%eax, %eax
	addq	$168, %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
.LBB6_16:
	.cfi_def_cfa_offset 224
	movl	%eax, %ebp
	movl	$.L.str.11, %edi
	movl	$.L.str.12, %esi
	movl	$149, %edx
	xorl	%eax, %eax
	callq	printf
	movl	%ebp, %edi
	callq	hipGetErrorString
	movl	$.L.str.13, %edi
	movl	%ebp, %esi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	printf
	movl	$1, %edi
	callq	exit
.Lfunc_end6:
	.size	main, .Lfunc_end6-main
	.cfi_endproc
                                        # -- End function
	.globl	_Z4compPKvS0_                   # -- Begin function _Z4compPKvS0_
	.type	_Z4compPKvS0_,@function
_Z4compPKvS0_:                          # @_Z4compPKvS0_
	.cfi_startproc
# %bb.0:
	movl	(%rdi), %eax
	movl	(%rsi), %ecx
	xorl	%edx, %edx
	cmpl	%ecx, %eax
	setl	%dl
	negl	%edx
	cmpl	%ecx, %eax
	movl	$1, %eax
	cmovlel	%edx, %eax
	retq
.Lfunc_end7:
	.size	_Z4compPKvS0_, .Lfunc_end7-_Z4compPKvS0_
	.cfi_endproc
                                        # -- End function
	.globl	_Z8checkSolv                    # -- Begin function _Z8checkSolv
	.type	_Z8checkSolv,@function
_Z8checkSolv:                           # @_Z8checkSolv
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$1024, %rsp                     # imm = 0x400
	.cfi_def_cfa_offset 1040
	.cfi_offset %rbx, -16
	movq	%rsp, %rbx
	movl	$.L__const._Z8checkSolv.indices, %esi
	movl	$1024, %edx                     # imm = 0x400
	movq	%rbx, %rdi
	callq	memcpy@PLT
	movl	$256, %esi                      # imm = 0x100
	movl	$4, %edx
	movl	$_Z4compPKvS0_, %ecx
	movq	%rbx, %rdi
	callq	qsort
	xorl	%esi, %esi
.LBB8_1:                                # =>This Inner Loop Header: Depth=1
	movl	(%rsp,%rsi,4), %edx
	cmpq	%rdx, %rsi
	jne	.LBB8_2
# %bb.3:                                #   in Loop: Header=BB8_1 Depth=1
	incq	%rsi
	cmpq	$256, %rsi                      # imm = 0x100
	jne	.LBB8_1
	jmp	.LBB8_4
.LBB8_2:
	movl	$.L.str.20, %edi
                                        # kill: def $esi killed $esi killed $rsi
                                        # kill: def $edx killed $edx killed $rdx
	xorl	%eax, %eax
	callq	printf
.LBB8_4:                                # %.loopexit
	movl	$.Lstr, %edi
	callq	puts@PLT
	addq	$1024, %rsp                     # imm = 0x400
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end8:
	.size	_Z8checkSolv, .Lfunc_end8-_Z8checkSolv
	.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	.LBB9_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB9_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z19sumMatrixOnDevice_1PfS_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_end9:
	.size	__hip_module_ctor, .Lfunc_end9-__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	.LBB10_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
.LBB10_2:
	retq
.Lfunc_end10:
	.size	__hip_module_dtor, .Lfunc_end10-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z19sumMatrixOnDevice_1PfS_S_ii,@object # @_Z19sumMatrixOnDevice_1PfS_S_ii
	.section	.rodata,"a",@progbits
	.globl	_Z19sumMatrixOnDevice_1PfS_S_ii
	.p2align	3, 0x0
_Z19sumMatrixOnDevice_1PfS_S_ii:
	.quad	_Z34__device_stub__sumMatrixOnDevice_1PfS_S_ii
	.size	_Z19sumMatrixOnDevice_1PfS_S_ii, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"matrix dimensions  = (%i, %i)\n"
	.size	.L.str, 31

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"number of elements = %i\n"
	.size	.L.str.1, 25

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"number of bytes in each matrix = %lu\n"
	.size	.L.str.2, 38

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"block = (%i, %i, %i)\n"
	.size	.L.str.3, 22

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"grid  = (%i, %i, %i)\n"
	.size	.L.str.4, 22

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"number of threads / block = %i\n"
	.size	.L.str.5, 32

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"total launched = %i\n"
	.size	.L.str.7, 21

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"total needed   = %i\n"
	.size	.L.str.8, 21

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"Launched with blockDim = (%i, %i, %i)\n"
	.size	.L.str.9, 39

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"Launched with  gridDim = (%i, %i, %i)\n"
	.size	.L.str.10, 39

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"Error: %s:%d, "
	.size	.L.str.11, 15

	.type	.L.str.12,@object               # @.str.12
.L.str.12:
	.asciz	"/home/ubuntu/Datasets/stackv2/train-structured-repos-hip-sm89/jamieborder/cuda_c_programming/master/chapter2/sumMatrices.hip"
	.size	.L.str.12, 125

	.type	.L.str.13,@object               # @.str.13
.L.str.13:
	.asciz	"code:%d, reason: %s\n"
	.size	.L.str.13, 21

	.type	.L.str.14,@object               # @.str.14
.L.str.14:
	.asciz	"Total error is %f\n"
	.size	.L.str.14, 19

	.type	.L.str.15,@object               # @.str.15
.L.str.15:
	.asciz	"Time on CPU_v1 is %f\n"
	.size	.L.str.15, 22

	.type	.L.str.16,@object               # @.str.16
.L.str.16:
	.asciz	"Time on CPU_v2 is %f\n"
	.size	.L.str.16, 22

	.type	.L.str.17,@object               # @.str.17
.L.str.17:
	.asciz	"Time on CPU_v3 is %f\n"
	.size	.L.str.17, 22

	.type	.L.str.18,@object               # @.str.18
.L.str.18:
	.asciz	"Time on GPU    is %f\n"
	.size	.L.str.18, 22

	.type	.L.str.19,@object               # @.str.19
.L.str.19:
	.asciz	"GPU speed-up over CPU is %.2f x\n"
	.size	.L.str.19, 33

	.type	.L__const._Z8checkSolv.indices,@object # @__const._Z8checkSolv.indices
	.section	.rodata,"a",@progbits
	.p2align	4, 0x0
.L__const._Z8checkSolv.indices:
	.long	96                              # 0x60
	.long	97                              # 0x61
	.long	98                              # 0x62
	.long	99                              # 0x63
	.long	100                             # 0x64
	.long	101                             # 0x65
	.long	102                             # 0x66
	.long	103                             # 0x67
	.long	104                             # 0x68
	.long	105                             # 0x69
	.long	106                             # 0x6a
	.long	107                             # 0x6b
	.long	108                             # 0x6c
	.long	109                             # 0x6d
	.long	110                             # 0x6e
	.long	111                             # 0x6f
	.long	112                             # 0x70
	.long	113                             # 0x71
	.long	114                             # 0x72
	.long	115                             # 0x73
	.long	116                             # 0x74
	.long	117                             # 0x75
	.long	118                             # 0x76
	.long	119                             # 0x77
	.long	120                             # 0x78
	.long	121                             # 0x79
	.long	122                             # 0x7a
	.long	123                             # 0x7b
	.long	124                             # 0x7c
	.long	125                             # 0x7d
	.long	126                             # 0x7e
	.long	127                             # 0x7f
	.long	32                              # 0x20
	.long	33                              # 0x21
	.long	34                              # 0x22
	.long	35                              # 0x23
	.long	36                              # 0x24
	.long	37                              # 0x25
	.long	38                              # 0x26
	.long	39                              # 0x27
	.long	40                              # 0x28
	.long	41                              # 0x29
	.long	42                              # 0x2a
	.long	43                              # 0x2b
	.long	44                              # 0x2c
	.long	45                              # 0x2d
	.long	46                              # 0x2e
	.long	47                              # 0x2f
	.long	48                              # 0x30
	.long	49                              # 0x31
	.long	50                              # 0x32
	.long	51                              # 0x33
	.long	52                              # 0x34
	.long	53                              # 0x35
	.long	54                              # 0x36
	.long	55                              # 0x37
	.long	56                              # 0x38
	.long	57                              # 0x39
	.long	58                              # 0x3a
	.long	59                              # 0x3b
	.long	60                              # 0x3c
	.long	61                              # 0x3d
	.long	62                              # 0x3e
	.long	63                              # 0x3f
	.long	128                             # 0x80
	.long	129                             # 0x81
	.long	130                             # 0x82
	.long	131                             # 0x83
	.long	132                             # 0x84
	.long	133                             # 0x85
	.long	134                             # 0x86
	.long	135                             # 0x87
	.long	136                             # 0x88
	.long	137                             # 0x89
	.long	138                             # 0x8a
	.long	139                             # 0x8b
	.long	140                             # 0x8c
	.long	141                             # 0x8d
	.long	142                             # 0x8e
	.long	143                             # 0x8f
	.long	144                             # 0x90
	.long	145                             # 0x91
	.long	146                             # 0x92
	.long	147                             # 0x93
	.long	148                             # 0x94
	.long	149                             # 0x95
	.long	150                             # 0x96
	.long	151                             # 0x97
	.long	152                             # 0x98
	.long	153                             # 0x99
	.long	154                             # 0x9a
	.long	155                             # 0x9b
	.long	156                             # 0x9c
	.long	157                             # 0x9d
	.long	158                             # 0x9e
	.long	159                             # 0x9f
	.long	160                             # 0xa0
	.long	161                             # 0xa1
	.long	162                             # 0xa2
	.long	163                             # 0xa3
	.long	164                             # 0xa4
	.long	165                             # 0xa5
	.long	166                             # 0xa6
	.long	167                             # 0xa7
	.long	168                             # 0xa8
	.long	169                             # 0xa9
	.long	170                             # 0xaa
	.long	171                             # 0xab
	.long	172                             # 0xac
	.long	173                             # 0xad
	.long	174                             # 0xae
	.long	175                             # 0xaf
	.long	176                             # 0xb0
	.long	177                             # 0xb1
	.long	178                             # 0xb2
	.long	179                             # 0xb3
	.long	180                             # 0xb4
	.long	181                             # 0xb5
	.long	182                             # 0xb6
	.long	183                             # 0xb7
	.long	184                             # 0xb8
	.long	185                             # 0xb9
	.long	186                             # 0xba
	.long	187                             # 0xbb
	.long	188                             # 0xbc
	.long	189                             # 0xbd
	.long	190                             # 0xbe
	.long	191                             # 0xbf
	.long	64                              # 0x40
	.long	65                              # 0x41
	.long	66                              # 0x42
	.long	67                              # 0x43
	.long	68                              # 0x44
	.long	69                              # 0x45
	.long	70                              # 0x46
	.long	71                              # 0x47
	.long	72                              # 0x48
	.long	73                              # 0x49
	.long	74                              # 0x4a
	.long	75                              # 0x4b
	.long	76                              # 0x4c
	.long	77                              # 0x4d
	.long	78                              # 0x4e
	.long	79                              # 0x4f
	.long	80                              # 0x50
	.long	81                              # 0x51
	.long	82                              # 0x52
	.long	83                              # 0x53
	.long	84                              # 0x54
	.long	85                              # 0x55
	.long	86                              # 0x56
	.long	87                              # 0x57
	.long	88                              # 0x58
	.long	89                              # 0x59
	.long	90                              # 0x5a
	.long	91                              # 0x5b
	.long	92                              # 0x5c
	.long	93                              # 0x5d
	.long	94                              # 0x5e
	.long	95                              # 0x5f
	.long	0                               # 0x0
	.long	1                               # 0x1
	.long	2                               # 0x2
	.long	3                               # 0x3
	.long	4                               # 0x4
	.long	5                               # 0x5
	.long	6                               # 0x6
	.long	7                               # 0x7
	.long	8                               # 0x8
	.long	9                               # 0x9
	.long	10                              # 0xa
	.long	11                              # 0xb
	.long	12                              # 0xc
	.long	13                              # 0xd
	.long	14                              # 0xe
	.long	15                              # 0xf
	.long	16                              # 0x10
	.long	17                              # 0x11
	.long	18                              # 0x12
	.long	19                              # 0x13
	.long	20                              # 0x14
	.long	21                              # 0x15
	.long	22                              # 0x16
	.long	23                              # 0x17
	.long	24                              # 0x18
	.long	25                              # 0x19
	.long	26                              # 0x1a
	.long	27                              # 0x1b
	.long	28                              # 0x1c
	.long	29                              # 0x1d
	.long	30                              # 0x1e
	.long	31                              # 0x1f
	.long	192                             # 0xc0
	.long	193                             # 0xc1
	.long	194                             # 0xc2
	.long	195                             # 0xc3
	.long	196                             # 0xc4
	.long	197                             # 0xc5
	.long	198                             # 0xc6
	.long	199                             # 0xc7
	.long	200                             # 0xc8
	.long	201                             # 0xc9
	.long	202                             # 0xca
	.long	203                             # 0xcb
	.long	204                             # 0xcc
	.long	205                             # 0xcd
	.long	206                             # 0xce
	.long	207                             # 0xcf
	.long	208                             # 0xd0
	.long	209                             # 0xd1
	.long	210                             # 0xd2
	.long	211                             # 0xd3
	.long	212                             # 0xd4
	.long	213                             # 0xd5
	.long	214                             # 0xd6
	.long	215                             # 0xd7
	.long	216                             # 0xd8
	.long	217                             # 0xd9
	.long	218                             # 0xda
	.long	219                             # 0xdb
	.long	220                             # 0xdc
	.long	221                             # 0xdd
	.long	222                             # 0xde
	.long	223                             # 0xdf
	.long	224                             # 0xe0
	.long	225                             # 0xe1
	.long	226                             # 0xe2
	.long	227                             # 0xe3
	.long	228                             # 0xe4
	.long	229                             # 0xe5
	.long	230                             # 0xe6
	.long	231                             # 0xe7
	.long	232                             # 0xe8
	.long	233                             # 0xe9
	.long	234                             # 0xea
	.long	235                             # 0xeb
	.long	236                             # 0xec
	.long	237                             # 0xed
	.long	238                             # 0xee
	.long	239                             # 0xef
	.long	240                             # 0xf0
	.long	241                             # 0xf1
	.long	242                             # 0xf2
	.long	243                             # 0xf3
	.long	244                             # 0xf4
	.long	245                             # 0xf5
	.long	246                             # 0xf6
	.long	247                             # 0xf7
	.long	248                             # 0xf8
	.long	249                             # 0xf9
	.long	250                             # 0xfa
	.long	251                             # 0xfb
	.long	252                             # 0xfc
	.long	253                             # 0xfd
	.long	254                             # 0xfe
	.long	255                             # 0xff
	.size	.L__const._Z8checkSolv.indices, 1024

	.type	.L.str.20,@object               # @.str.20
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.20:
	.asciz	"failed:: %d:%d\n"
	.size	.L.str.20, 16

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z19sumMatrixOnDevice_1PfS_S_ii"
	.size	.L__unnamed_1, 32

	.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

	.type	.Lstr,@object                   # @str
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr:
	.asciz	"success:: all matched"
	.size	.Lstr, 22

	.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 _Z34__device_stub__sumMatrixOnDevice_1PfS_S_ii
	.addrsig_sym _Z4compPKvS0_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z19sumMatrixOnDevice_1PfS_S_ii
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z19sumMatrixOnDevice_1PfS_S_ii ; -- Begin function _Z19sumMatrixOnDevice_1PfS_S_ii
	.globl	_Z19sumMatrixOnDevice_1PfS_S_ii
	.p2align	8
	.type	_Z19sumMatrixOnDevice_1PfS_S_ii,@function
_Z19sumMatrixOnDevice_1PfS_S_ii:        ; @_Z19sumMatrixOnDevice_1PfS_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
	v_and_b32_e32 v2, 0x3ff, v0
	v_bfe_u32 v4, v0, 10, 10
	s_waitcnt lgkmcnt(0)
	s_and_b32 s5, s4, 0xffff
	s_lshr_b32 s4, s4, 16
	v_mad_u64_u32 v[0:1], null, s14, s5, v[2:3]
	v_mov_b32_e32 v1, 0
	v_mad_u64_u32 v[2:3], null, s15, s4, v[4:5]
	s_ashr_i32 s5, s2, 31
	s_mov_b32 s4, s2
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_mov_b32_e32 v3, v1
	s_ashr_i32 s7, s3, 31
	s_mov_b32 s6, s3
	v_cmp_gt_u64_e32 vcc_lo, s[4:5], v[0:1]
	v_cmp_gt_u64_e64 s2, s[6:7], v[2:3]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s2, vcc_lo, s2
	s_and_saveexec_b32 s3, s2
	s_cbranch_execz .LBB0_2
; %bb.1:
	v_mad_u64_u32 v[3:4], null, v2, s4, v[0:1]
	s_clause 0x1
	s_load_b128 s[8:11], s[0:1], 0x0
	s_load_b64 s[0:1], s[0:1], 0x10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mov_b32_e32 v0, v4
	v_mad_u64_u32 v[4:5], null, v2, s5, v[0:1]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[3:4]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s8, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v3, vcc_lo, s9, v1, vcc_lo
	v_add_co_u32 v4, vcc_lo, s10, v0
	v_add_co_ci_u32_e32 v5, vcc_lo, s11, v1, vcc_lo
	v_add_co_u32 v0, vcc_lo, s0, v0
	global_load_b32 v2, v[2:3], off
	global_load_b32 v3, v[4:5], off
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	s_waitcnt vmcnt(0)
	v_add_f32_e32 v2, v2, v3
	global_store_b32 v[0:1], v2, off
.LBB0_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z19sumMatrixOnDevice_1PfS_S_ii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.amdhsa_user_sgpr_count 14
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 6
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z19sumMatrixOnDevice_1PfS_S_ii, .Lfunc_end0-_Z19sumMatrixOnDevice_1PfS_S_ii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 268
; NumSgprs: 18
; NumVgprs: 6
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 6
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .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:           _Z19sumMatrixOnDevice_1PfS_S_ii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z19sumMatrixOnDevice_1PfS_S_ii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     6
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
