#include <cuda_runtime.h>
#include <cuda_runtime.h>
#include <stdio.h>
#include <time.h>
#include <stdlib.h>  //srand()
#include <iostream>  //cout
#include <string.h>  //memset()
extern "C" void gpuTestAll(float *MatA, float *MatB, float *MatC, int nx, int ny);
// grid 1D block 1D
// grid 2D block 2D
// grid 2D block 1D
// grid 2D block 2D

#define CHECK(status);									\
{														\
	if (status != 0)									\
	{													\
		std::cout << "Cuda failure: " << status;		\
		abort();										\
	}													\
}


__global__ void sumMatrixOnGPU2D(float *MatA, float *MatB, float *MatC, int nx,
                                 int ny)
{
    unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;
    unsigned int iy = threadIdx.y + blockIdx.y * blockDim.y;
    unsigned int idx = iy * nx + ix;

    if (ix < nx && iy < ny)
        MatC[idx] = MatA[idx] + MatB[idx];
}

// grid 1D block 1D
__global__ void sumMatrixOnGPU1D(float *MatA, float *MatB, float *MatC, int nx,
                                 int ny)
{
    unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;

    if (ix < nx )
        for (int iy = 0; iy < ny; iy++)
        {
            int idx = iy * nx + ix;
            MatC[idx] = MatA[idx] + MatB[idx];
        }


}

// grid 2D block 1D
__global__ void sumMatrixOnGPUMix(float *MatA, float *MatB, float *MatC, int nx,
                                  int ny)
{
    unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;
    unsigned int iy = blockIdx.y;
    unsigned int idx = iy * nx + ix;

    if (ix < nx && iy < ny)
        MatC[idx] = MatA[idx] + MatB[idx];
}

////=================================================
////=================================================




void gpuTestAll(float *d_MatA, float *d_MatB, float *d_MatC, int nx, int ny)
{
    clock_t t1, t2, t3, t4, t5, t6, t7, t8, t9, t10, t11, t12, t13, t14, t15, t16, t17, t18, t19, t20, t21, t22, t23, t24, t25, t26, t27, t28, t29, t30, t31, t32, t33, t34;
    int dimx = 32;
    int dimy = 32;
    dim3 block(dimx, dimy);
    dim3 grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y);

t5=clock();
    sumMatrixOnGPU2D<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(cudaDeviceSynchronize());
t6=clock();
    printf("sumMatrixOnGPU2D  <<<  (512,512),  (32,32)  >>> elapsed %f sec\n", (double)(t6-t5)/(CLOCKS_PER_SEC));

    // adjust block size
    block.x = 16;
    grid.x  = (nx + block.x - 1) / block.x;
    grid.y  = (ny + block.y - 1) / block.y;

t7=clock();
    sumMatrixOnGPU2D<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(cudaDeviceSynchronize());
t8=clock();
    printf("sumMatrixOnGPU2D  <<<  (1024,512), (16,32)  >>> elapsed %f sec\n", (double)(t8-t7)/(CLOCKS_PER_SEC));

    // adjust block size
    block.y = 16;
    block.x = 32;
    grid.x  = (nx + block.x - 1) / block.x;
    grid.y  = (ny + block.y - 1) / block.y;

t9=clock();
    sumMatrixOnGPU2D<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(cudaDeviceSynchronize());
t10=clock();
    printf("sumMatrixOnGPU2D  <<<  (512,1024), (32,16)  >>> elapsed %f sec\n", (double)(t10-t9)/(CLOCKS_PER_SEC));

    block.y = 16;
    block.x = 16;
    grid.x  = (nx + block.x - 1) / block.x;
    grid.y  = (ny + block.y - 1) / block.y;

t11=clock();
    sumMatrixOnGPU2D<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(cudaDeviceSynchronize());
t12=clock();
    printf("sumMatrixOnGPU2D  <<<  (1024,1024),(16,16)  >>> elapsed %f sec\n", (double)(t12-t11)/(CLOCKS_PER_SEC));

    block.y = 16;
    block.x = 64;
    grid.x  = (nx + block.x - 1) / block.x;
    grid.y  = (ny + block.y - 1) / block.y;

t13=clock();
    sumMatrixOnGPU2D<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(cudaDeviceSynchronize());
t14=clock();
    printf("sumMatrixOnGPU2D  <<<  (256,1024), (64,16)  >>> elapsed %f sec\n", (double)(t14-t13)/(CLOCKS_PER_SEC));

    block.y = 64;
    block.x = 16;
    grid.x  = (nx + block.x - 1) / block.x;
    grid.y  = (ny + block.y - 1) / block.y;

t15=clock();
    sumMatrixOnGPU2D<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(cudaDeviceSynchronize());
t16=clock();
    printf("sumMatrixOnGPU2D  <<<  (1024,256), (16,64)  >>> elapsed %f sec\n", (double)(t16-t15)/(CLOCKS_PER_SEC));
printf("\n");

    block.x = 32;
    grid.x  = (nx + block.x - 1) / block.x;
    block.y = 1;
    grid.y  = 1;

t17=clock();
    sumMatrixOnGPU1D<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(cudaDeviceSynchronize());
t18=clock();
    printf("sumMatrixOnGPU1D  <<<  (512,1)   , (32,1)   >>> elapsed %f sec\n", (double)(t18-t17)/(CLOCKS_PER_SEC));

    block.x = 64;
    grid.x  = (nx + block.x - 1) / block.x;
    block.y = 1;
    grid.y  = 1;

t19=clock();
    sumMatrixOnGPU1D<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(cudaDeviceSynchronize());
t20=clock();
    printf("sumMatrixOnGPU1D  <<<  (256,1)   , (64,1)   >>> elapsed %f sec\n", (double)(t20-t19)/(CLOCKS_PER_SEC));

    block.x = 128;
    grid.x  = (nx + block.x - 1) / block.x;
    block.y = 1;
    grid.y  = 1;

t21=clock();
    sumMatrixOnGPU1D<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(cudaDeviceSynchronize());
t22=clock();;
    printf("sumMatrixOnGPU1D  <<<  (128,1)   , (128,1)  >>> elapsed %f sec\n", (double)(t22-t21)/(CLOCKS_PER_SEC));

printf("\n");
    // grid 2D and block 1D
    block.x = 32;
    grid.x  = (nx + block.x - 1) / block.x;
    block.y = 1;
    grid.y  = ny;

t23=clock();
    sumMatrixOnGPUMix<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(cudaDeviceSynchronize());
t24=clock();
    printf("sumMatrixOnGPUMix <<<  (512,16384),(32,1)   >>> elapsed %f sec\n",(double)(t24-t23)/(CLOCKS_PER_SEC));

    block.x = 64;
    grid.x  = (nx + block.x - 1) / block.x;
    block.y = 1;
    grid.y  = ny;

t25=clock();
    sumMatrixOnGPUMix<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(cudaDeviceSynchronize());
t26=clock();
    printf("sumMatrixOnGPUMix <<<  (256,16384),(64,1)   >>> elapsed %f sec\n",(double)(t26-t25)/(CLOCKS_PER_SEC));

    block.x = 128;
    grid.x  = (nx + block.x - 1) / block.x;
    block.y = 1;
    grid.y  = ny;

t27=clock();
    sumMatrixOnGPUMix<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(cudaDeviceSynchronize());
t28=clock();
    printf("sumMatrixOnGPUMix <<<  (128,16384),(128,1)  >>> elapsed %f sec\n",(double)(t28-t27)/(CLOCKS_PER_SEC));

    block.x = 256;
    grid.x  = (nx + block.x - 1) / block.x;
    block.y = 1;
    grid.y  = ny;

t29=clock();
    sumMatrixOnGPUMix<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(cudaDeviceSynchronize());
t30=clock();
    printf("sumMatrixOnGPUMix <<<  (64,16384), (256,1)  >>> elapsed %f sec\n",(double)(t30-t29)/(CLOCKS_PER_SEC));

    block.x = 512;
    grid.x  = (nx + block.x - 1) / block.x;
    block.y = 1;
    grid.y  = ny;

t31=clock();
    sumMatrixOnGPUMix<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(cudaDeviceSynchronize());
t32=clock();
    printf("sumMatrixOnGPUMix <<<  (32,16384), (512,1)  >>> elapsed %f sec\n",(double)(t32-t31)/(CLOCKS_PER_SEC));




}














	.file	"tmpxft_002048b7_00000000-6_test6_cu.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3638:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3638:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z42__device_stub__Z16sumMatrixOnGPU2DPfS_S_iiPfS_S_ii
	.type	_Z42__device_stub__Z16sumMatrixOnGPU2DPfS_S_iiPfS_S_ii, @function
_Z42__device_stub__Z16sumMatrixOnGPU2DPfS_S_iiPfS_S_ii:
.LFB3660:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movl	%ecx, 4(%rsp)
	leaq	40(%rsp), %rcx
	movl	%r8d, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	%rsp, %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z16sumMatrixOnGPU2DPfS_S_ii(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 176
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L2:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3660:
	.size	_Z42__device_stub__Z16sumMatrixOnGPU2DPfS_S_iiPfS_S_ii, .-_Z42__device_stub__Z16sumMatrixOnGPU2DPfS_S_iiPfS_S_ii
	.globl	_Z16sumMatrixOnGPU2DPfS_S_ii
	.type	_Z16sumMatrixOnGPU2DPfS_S_ii, @function
_Z16sumMatrixOnGPU2DPfS_S_ii:
.LFB3661:
	.cfi_startproc
	endbr64
	jmp	_Z42__device_stub__Z16sumMatrixOnGPU2DPfS_S_iiPfS_S_ii
	.cfi_endproc
.LFE3661:
	.size	_Z16sumMatrixOnGPU2DPfS_S_ii, .-_Z16sumMatrixOnGPU2DPfS_S_ii
	.globl	_Z42__device_stub__Z16sumMatrixOnGPU1DPfS_S_iiPfS_S_ii
	.type	_Z42__device_stub__Z16sumMatrixOnGPU1DPfS_S_iiPfS_S_ii, @function
_Z42__device_stub__Z16sumMatrixOnGPU1DPfS_S_iiPfS_S_ii:
.LFB3662:
	.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	.L8
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z16sumMatrixOnGPU1DPfS_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
.L8:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3662:
	.size	_Z42__device_stub__Z16sumMatrixOnGPU1DPfS_S_iiPfS_S_ii, .-_Z42__device_stub__Z16sumMatrixOnGPU1DPfS_S_iiPfS_S_ii
	.globl	_Z16sumMatrixOnGPU1DPfS_S_ii
	.type	_Z16sumMatrixOnGPU1DPfS_S_ii, @function
_Z16sumMatrixOnGPU1DPfS_S_ii:
.LFB3663:
	.cfi_startproc
	endbr64
	jmp	_Z42__device_stub__Z16sumMatrixOnGPU1DPfS_S_iiPfS_S_ii
	.cfi_endproc
.LFE3663:
	.size	_Z16sumMatrixOnGPU1DPfS_S_ii, .-_Z16sumMatrixOnGPU1DPfS_S_ii
	.globl	_Z43__device_stub__Z17sumMatrixOnGPUMixPfS_S_iiPfS_S_ii
	.type	_Z43__device_stub__Z17sumMatrixOnGPUMixPfS_S_iiPfS_S_ii, @function
_Z43__device_stub__Z17sumMatrixOnGPUMixPfS_S_iiPfS_S_ii:
.LFB3664:
	.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	.L13
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z17sumMatrixOnGPUMixPfS_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
.L13:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L15
	call	__stack_chk_fail@PLT
.L15:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3664:
	.size	_Z43__device_stub__Z17sumMatrixOnGPUMixPfS_S_iiPfS_S_ii, .-_Z43__device_stub__Z17sumMatrixOnGPUMixPfS_S_iiPfS_S_ii
	.globl	_Z17sumMatrixOnGPUMixPfS_S_ii
	.type	_Z17sumMatrixOnGPUMixPfS_S_ii, @function
_Z17sumMatrixOnGPUMixPfS_S_ii:
.LFB3665:
	.cfi_startproc
	endbr64
	jmp	_Z43__device_stub__Z17sumMatrixOnGPUMixPfS_S_iiPfS_S_ii
	.cfi_endproc
.LFE3665:
	.size	_Z17sumMatrixOnGPUMixPfS_S_ii, .-_Z17sumMatrixOnGPUMixPfS_S_ii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"Cuda failure: "
.LC2:
	.string	"sumMatrixOnGPU2D  <<<  (512,512),  (32,32)  >>> elapsed %f sec\n"
.LC3:
	.string	"sumMatrixOnGPU2D  <<<  (1024,512), (16,32)  >>> elapsed %f sec\n"
.LC4:
	.string	"sumMatrixOnGPU2D  <<<  (512,1024), (32,16)  >>> elapsed %f sec\n"
.LC5:
	.string	"sumMatrixOnGPU2D  <<<  (1024,1024),(16,16)  >>> elapsed %f sec\n"
.LC6:
	.string	"sumMatrixOnGPU2D  <<<  (256,1024), (64,16)  >>> elapsed %f sec\n"
.LC7:
	.string	"sumMatrixOnGPU2D  <<<  (1024,256), (16,64)  >>> elapsed %f sec\n"
.LC8:
	.string	"\n"
.LC9:
	.string	"sumMatrixOnGPU1D  <<<  (512,1)   , (32,1)   >>> elapsed %f sec\n"
.LC10:
	.string	"sumMatrixOnGPU1D  <<<  (256,1)   , (64,1)   >>> elapsed %f sec\n"
.LC11:
	.string	"sumMatrixOnGPU1D  <<<  (128,1)   , (128,1)  >>> elapsed %f sec\n"
.LC12:
	.string	"sumMatrixOnGPUMix <<<  (512,16384),(32,1)   >>> elapsed %f sec\n"
.LC13:
	.string	"sumMatrixOnGPUMix <<<  (256,16384),(64,1)   >>> elapsed %f sec\n"
.LC14:
	.string	"sumMatrixOnGPUMix <<<  (128,16384),(128,1)  >>> elapsed %f sec\n"
.LC15:
	.string	"sumMatrixOnGPUMix <<<  (64,16384), (256,1)  >>> elapsed %f sec\n"
.LC16:
	.string	"sumMatrixOnGPUMix <<<  (32,16384), (512,1)  >>> elapsed %f sec\n"
	.text
	.globl	gpuTestAll
	.type	gpuTestAll, @function
gpuTestAll:
.LFB3635:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	leal	31(%r8), %eax
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	shrl	$5, %eax
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	leal	31(%rcx), %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	shrl	$5, %r12d
	movl	%r8d, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%ecx, %ebx
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	movq	%rsi, 16(%rsp)
	movq	%rdx, 24(%rsp)
	movl	%r12d, 52(%rsp)
	movl	%eax, 56(%rsp)
	movl	$1, 60(%rsp)
	call	clock@PLT
	movq	52(%rsp), %rdi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	60(%rsp), %esi
	movl	$1, %ecx
	movq	%rax, %r13
	movabsq	$137438953504, %rdx
	movq	%rdx, 40(%rsp)
	movl	$1, 48(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L19
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rsi
	movl	%ebp, %r8d
	movl	%ebx, %ecx
	movq	8(%rsp), %rdi
	call	_Z42__device_stub__Z16sumMatrixOnGPU2DPfS_S_iiPfS_S_ii
.L19:
	call	cudaDeviceSynchronize@PLT
	testl	%eax, %eax
	jne	.L49
	call	clock@PLT
	leal	15(%rbx), %r14d
	movl	$2, %edi
	leaq	.LC2(%rip), %rsi
	subq	%r13, %rax
	shrl	$4, %r14d
	cvtsi2sdq	%rax, %xmm0
	movb	$1, %al
	divsd	.LC1(%rip), %xmm0
	call	__printf_chk@PLT
	movl	%r14d, 52(%rsp)
	call	clock@PLT
	movl	48(%rsp), %ecx
	movl	60(%rsp), %esi
	xorl	%r9d, %r9d
	movl	$16, 40(%rsp)
	movq	52(%rsp), %rdi
	xorl	%r8d, %r8d
	movq	%rax, %r13
	movq	40(%rsp), %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L21
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rsi
	movl	%ebp, %r8d
	movl	%ebx, %ecx
	movq	8(%rsp), %rdi
	call	_Z42__device_stub__Z16sumMatrixOnGPU2DPfS_S_iiPfS_S_ii
.L21:
	call	cudaDeviceSynchronize@PLT
	testl	%eax, %eax
	jne	.L49
	call	clock@PLT
	leaq	.LC3(%rip), %rsi
	movl	$2, %edi
	subq	%r13, %rax
	cvtsi2sdq	%rax, %xmm0
	movb	$1, %al
	divsd	.LC1(%rip), %xmm0
	call	__printf_chk@PLT
	leal	15(%rbp), %eax
	movl	%r12d, 52(%rsp)
	shrl	$4, %eax
	movl	%eax, 56(%rsp)
	call	clock@PLT
	movl	$2147483649, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	48(%rsp), %ecx
	movq	52(%rsp), %rdi
	salq	$5, %rdx
	movq	%rax, %r13
	movl	60(%rsp), %esi
	movq	%rdx, 40(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L23
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rsi
	movl	%ebp, %r8d
	movl	%ebx, %ecx
	movq	8(%rsp), %rdi
	call	_Z42__device_stub__Z16sumMatrixOnGPU2DPfS_S_iiPfS_S_ii
.L23:
	call	cudaDeviceSynchronize@PLT
	testl	%eax, %eax
	jne	.L49
	call	clock@PLT
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	subq	%r13, %rax
	cvtsi2sdq	%rax, %xmm0
	movb	$1, %al
	divsd	.LC1(%rip), %xmm0
	call	__printf_chk@PLT
	movl	%r14d, 52(%rsp)
	call	clock@PLT
	movl	48(%rsp), %ecx
	movl	60(%rsp), %esi
	xorl	%r9d, %r9d
	movl	$16, 40(%rsp)
	movq	52(%rsp), %rdi
	xorl	%r8d, %r8d
	movq	%rax, %r13
	movq	40(%rsp), %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L25
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rsi
	movl	%ebp, %r8d
	movl	%ebx, %ecx
	movq	8(%rsp), %rdi
	call	_Z42__device_stub__Z16sumMatrixOnGPU2DPfS_S_iiPfS_S_ii
.L25:
	call	cudaDeviceSynchronize@PLT
	testl	%eax, %eax
	jne	.L49
	call	clock@PLT
	leaq	.LC5(%rip), %rsi
	movl	$2, %edi
	subq	%r13, %rax
	leal	63(%rbx), %r13d
	cvtsi2sdq	%rax, %xmm0
	shrl	$6, %r13d
	movb	$1, %al
	divsd	.LC1(%rip), %xmm0
	call	__printf_chk@PLT
	movl	%r13d, 52(%rsp)
	call	clock@PLT
	movl	48(%rsp), %ecx
	movl	60(%rsp), %esi
	xorl	%r9d, %r9d
	movl	$64, 40(%rsp)
	movq	52(%rsp), %rdi
	xorl	%r8d, %r8d
	movq	%rax, %r15
	movq	40(%rsp), %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L27
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rsi
	movl	%ebp, %r8d
	movl	%ebx, %ecx
	movq	8(%rsp), %rdi
	call	_Z42__device_stub__Z16sumMatrixOnGPU2DPfS_S_iiPfS_S_ii
.L27:
	call	cudaDeviceSynchronize@PLT
	testl	%eax, %eax
	jne	.L49
	call	clock@PLT
	leaq	.LC6(%rip), %rsi
	movl	$2, %edi
	subq	%r15, %rax
	cvtsi2sdq	%rax, %xmm0
	movb	$1, %al
	divsd	.LC1(%rip), %xmm0
	call	__printf_chk@PLT
	leal	63(%rbp), %eax
	movl	%r14d, 52(%rsp)
	shrl	$6, %eax
	movl	%eax, 56(%rsp)
	call	clock@PLT
	movl	48(%rsp), %ecx
	movl	60(%rsp), %esi
	xorl	%r9d, %r9d
	movq	52(%rsp), %rdi
	xorl	%r8d, %r8d
	movq	%rax, %r14
	movabsq	$274877906960, %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L29
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rsi
	movl	%ebp, %r8d
	movl	%ebx, %ecx
	movq	8(%rsp), %rdi
	call	_Z42__device_stub__Z16sumMatrixOnGPU2DPfS_S_iiPfS_S_ii
.L29:
	call	cudaDeviceSynchronize@PLT
	testl	%eax, %eax
	jne	.L49
	call	clock@PLT
	leaq	.LC7(%rip), %rsi
	movl	$2, %edi
	subq	%r14, %rax
	cvtsi2sdq	%rax, %xmm0
	movb	$1, %al
	divsd	.LC1(%rip), %xmm0
	call	__printf_chk@PLT
	leaq	.LC8(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	%r12d, 52(%rsp)
	movl	$1, 56(%rsp)
	call	clock@PLT
	movl	$134217729, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	48(%rsp), %ecx
	movq	52(%rsp), %rdi
	salq	$5, %rdx
	movq	%rax, %r14
	movl	60(%rsp), %esi
	movq	%rdx, 40(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L31
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rsi
	movl	%ebp, %r8d
	movl	%ebx, %ecx
	movq	8(%rsp), %rdi
	call	_Z42__device_stub__Z16sumMatrixOnGPU1DPfS_S_iiPfS_S_ii
.L31:
	call	cudaDeviceSynchronize@PLT
	testl	%eax, %eax
	jne	.L49
	call	clock@PLT
	leaq	.LC9(%rip), %rsi
	movl	$2, %edi
	subq	%r14, %rax
	cvtsi2sdq	%rax, %xmm0
	movb	$1, %al
	divsd	.LC1(%rip), %xmm0
	call	__printf_chk@PLT
	movl	%r13d, 52(%rsp)
	call	clock@PLT
	movl	48(%rsp), %ecx
	movl	60(%rsp), %esi
	xorl	%r9d, %r9d
	movl	$64, 40(%rsp)
	movq	52(%rsp), %rdi
	xorl	%r8d, %r8d
	movq	%rax, %r14
	movq	40(%rsp), %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L33
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rsi
	movl	%ebp, %r8d
	movl	%ebx, %ecx
	movq	8(%rsp), %rdi
	call	_Z42__device_stub__Z16sumMatrixOnGPU1DPfS_S_iiPfS_S_ii
.L33:
	call	cudaDeviceSynchronize@PLT
	testl	%eax, %eax
	jne	.L49
	call	clock@PLT
	leaq	.LC10(%rip), %rsi
	movl	$2, %edi
	subq	%r14, %rax
	leal	127(%rbx), %r14d
	cvtsi2sdq	%rax, %xmm0
	shrl	$7, %r14d
	movb	$1, %al
	divsd	.LC1(%rip), %xmm0
	call	__printf_chk@PLT
	movl	%r14d, 52(%rsp)
	call	clock@PLT
	movl	48(%rsp), %ecx
	movl	60(%rsp), %esi
	xorl	%r9d, %r9d
	movl	$128, 40(%rsp)
	movq	52(%rsp), %rdi
	xorl	%r8d, %r8d
	movq	%rax, %r15
	movq	40(%rsp), %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L35
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rsi
	movl	%ebp, %r8d
	movl	%ebx, %ecx
	movq	8(%rsp), %rdi
	call	_Z42__device_stub__Z16sumMatrixOnGPU1DPfS_S_iiPfS_S_ii
.L35:
	call	cudaDeviceSynchronize@PLT
	testl	%eax, %eax
	jne	.L49
	call	clock@PLT
	leaq	.LC11(%rip), %rsi
	movl	$2, %edi
	subq	%r15, %rax
	cvtsi2sdq	%rax, %xmm0
	movb	$1, %al
	divsd	.LC1(%rip), %xmm0
	call	__printf_chk@PLT
	leaq	.LC8(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	%r12d, 52(%rsp)
	movl	%ebp, 56(%rsp)
	call	clock@PLT
	movl	48(%rsp), %ecx
	movl	60(%rsp), %esi
	xorl	%r9d, %r9d
	movl	$32, 40(%rsp)
	movq	52(%rsp), %rdi
	xorl	%r8d, %r8d
	movq	%rax, %r12
	movq	40(%rsp), %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L37
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rsi
	movl	%ebp, %r8d
	movl	%ebx, %ecx
	movq	8(%rsp), %rdi
	call	_Z43__device_stub__Z17sumMatrixOnGPUMixPfS_S_iiPfS_S_ii
.L37:
	call	cudaDeviceSynchronize@PLT
	testl	%eax, %eax
	jne	.L49
	call	clock@PLT
	leaq	.LC12(%rip), %rsi
	movl	$2, %edi
	subq	%r12, %rax
	cvtsi2sdq	%rax, %xmm0
	movb	$1, %al
	divsd	.LC1(%rip), %xmm0
	call	__printf_chk@PLT
	movl	%r13d, 52(%rsp)
	call	clock@PLT
	movl	48(%rsp), %ecx
	movl	60(%rsp), %esi
	xorl	%r9d, %r9d
	movl	$64, 40(%rsp)
	movq	52(%rsp), %rdi
	xorl	%r8d, %r8d
	movq	%rax, %r12
	movq	40(%rsp), %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L39
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rsi
	movl	%ebp, %r8d
	movl	%ebx, %ecx
	movq	8(%rsp), %rdi
	call	_Z43__device_stub__Z17sumMatrixOnGPUMixPfS_S_iiPfS_S_ii
.L39:
	call	cudaDeviceSynchronize@PLT
	testl	%eax, %eax
	jne	.L49
	call	clock@PLT
	leaq	.LC13(%rip), %rsi
	movl	$2, %edi
	subq	%r12, %rax
	cvtsi2sdq	%rax, %xmm0
	movb	$1, %al
	divsd	.LC1(%rip), %xmm0
	call	__printf_chk@PLT
	movl	%r14d, 52(%rsp)
	call	clock@PLT
	movl	48(%rsp), %ecx
	movl	60(%rsp), %esi
	xorl	%r9d, %r9d
	movl	$128, 40(%rsp)
	movq	52(%rsp), %rdi
	xorl	%r8d, %r8d
	movq	%rax, %r12
	movq	40(%rsp), %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rsi
	movl	%ebp, %r8d
	movl	%ebx, %ecx
	movq	8(%rsp), %rdi
	call	_Z43__device_stub__Z17sumMatrixOnGPUMixPfS_S_iiPfS_S_ii
.L41:
	call	cudaDeviceSynchronize@PLT
	testl	%eax, %eax
	jne	.L49
	call	clock@PLT
	leaq	.LC14(%rip), %rsi
	movl	$2, %edi
	subq	%r12, %rax
	cvtsi2sdq	%rax, %xmm0
	movb	$1, %al
	divsd	.LC1(%rip), %xmm0
	call	__printf_chk@PLT
	leal	255(%rbx), %eax
	shrl	$8, %eax
	movl	%eax, 52(%rsp)
	call	clock@PLT
	movl	48(%rsp), %ecx
	movl	60(%rsp), %esi
	xorl	%r9d, %r9d
	movl	$256, 40(%rsp)
	movq	52(%rsp), %rdi
	xorl	%r8d, %r8d
	movq	%rax, %r12
	movq	40(%rsp), %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L43
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rsi
	movl	%ebp, %r8d
	movl	%ebx, %ecx
	movq	8(%rsp), %rdi
	call	_Z43__device_stub__Z17sumMatrixOnGPUMixPfS_S_iiPfS_S_ii
.L43:
	call	cudaDeviceSynchronize@PLT
	testl	%eax, %eax
	jne	.L49
	call	clock@PLT
	leaq	.LC15(%rip), %rsi
	movl	$2, %edi
	subq	%r12, %rax
	cvtsi2sdq	%rax, %xmm0
	movb	$1, %al
	divsd	.LC1(%rip), %xmm0
	call	__printf_chk@PLT
	leal	511(%rbx), %eax
	shrl	$9, %eax
	movl	%eax, 52(%rsp)
	call	clock@PLT
	movl	48(%rsp), %ecx
	movl	60(%rsp), %esi
	xorl	%r9d, %r9d
	movl	$512, 40(%rsp)
	movq	52(%rsp), %rdi
	xorl	%r8d, %r8d
	movq	%rax, %r12
	movq	40(%rsp), %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L45
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rsi
	movl	%ebp, %r8d
	movl	%ebx, %ecx
	movq	8(%rsp), %rdi
	call	_Z43__device_stub__Z17sumMatrixOnGPUMixPfS_S_iiPfS_S_ii
.L45:
	call	cudaDeviceSynchronize@PLT
	testl	%eax, %eax
	je	.L46
.L49:
	leaq	.LC0(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rbx
	call	cudaDeviceSynchronize@PLT
	movq	%rbx, %rdi
	movl	%eax, %esi
	call	_ZNSolsEi@PLT
	call	abort@PLT
.L46:
	call	clock@PLT
	leaq	.LC16(%rip), %rsi
	movl	$2, %edi
	subq	%r12, %rax
	cvtsi2sdq	%rax, %xmm0
	divsd	.LC1(%rip), %xmm0
	movb	$1, %al
	addq	$72, %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
	jmp	__printf_chk@PLT
	.cfi_endproc
.LFE3635:
	.size	gpuTestAll, .-gpuTestAll
	.section	.rodata.str1.1
.LC17:
	.string	"_Z17sumMatrixOnGPUMixPfS_S_ii"
.LC18:
	.string	"_Z16sumMatrixOnGPU1DPfS_S_ii"
.LC19:
	.string	"_Z16sumMatrixOnGPU2DPfS_S_ii"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3667:
	.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	.LC17(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z17sumMatrixOnGPUMixPfS_S_ii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC18(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z16sumMatrixOnGPU1DPfS_S_ii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC19(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z16sumMatrixOnGPU2DPfS_S_ii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE3667:
	.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.cst8,"aM",@progbits,8
	.align 8
.LC1:
	.long	0
	.long	1093567616
	.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 <time.h>
#include <stdlib.h>  //srand()
#include <iostream>  //cout
#include <string.h>  //memset()
extern "C" void gpuTestAll(float *MatA, float *MatB, float *MatC, int nx, int ny);
// grid 1D block 1D
// grid 2D block 2D
// grid 2D block 1D
// grid 2D block 2D

#define CHECK(status);									\
{														\
	if (status != 0)									\
	{													\
		std::cout << "Cuda failure: " << status;		\
		abort();										\
	}													\
}


__global__ void sumMatrixOnGPU2D(float *MatA, float *MatB, float *MatC, int nx,
                                 int ny)
{
    unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;
    unsigned int iy = threadIdx.y + blockIdx.y * blockDim.y;
    unsigned int idx = iy * nx + ix;

    if (ix < nx && iy < ny)
        MatC[idx] = MatA[idx] + MatB[idx];
}

// grid 1D block 1D
__global__ void sumMatrixOnGPU1D(float *MatA, float *MatB, float *MatC, int nx,
                                 int ny)
{
    unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;

    if (ix < nx )
        for (int iy = 0; iy < ny; iy++)
        {
            int idx = iy * nx + ix;
            MatC[idx] = MatA[idx] + MatB[idx];
        }


}

// grid 2D block 1D
__global__ void sumMatrixOnGPUMix(float *MatA, float *MatB, float *MatC, int nx,
                                  int ny)
{
    unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;
    unsigned int iy = blockIdx.y;
    unsigned int idx = iy * nx + ix;

    if (ix < nx && iy < ny)
        MatC[idx] = MatA[idx] + MatB[idx];
}

////=================================================
////=================================================




void gpuTestAll(float *d_MatA, float *d_MatB, float *d_MatC, int nx, int ny)
{
    clock_t t1, t2, t3, t4, t5, t6, t7, t8, t9, t10, t11, t12, t13, t14, t15, t16, t17, t18, t19, t20, t21, t22, t23, t24, t25, t26, t27, t28, t29, t30, t31, t32, t33, t34;
    int dimx = 32;
    int dimy = 32;
    dim3 block(dimx, dimy);
    dim3 grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y);

t5=clock();
    sumMatrixOnGPU2D<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(hipDeviceSynchronize());
t6=clock();
    printf("sumMatrixOnGPU2D  <<<  (512,512),  (32,32)  >>> elapsed %f sec\n", (double)(t6-t5)/(CLOCKS_PER_SEC));

    // adjust block size
    block.x = 16;
    grid.x  = (nx + block.x - 1) / block.x;
    grid.y  = (ny + block.y - 1) / block.y;

t7=clock();
    sumMatrixOnGPU2D<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(hipDeviceSynchronize());
t8=clock();
    printf("sumMatrixOnGPU2D  <<<  (1024,512), (16,32)  >>> elapsed %f sec\n", (double)(t8-t7)/(CLOCKS_PER_SEC));

    // adjust block size
    block.y = 16;
    block.x = 32;
    grid.x  = (nx + block.x - 1) / block.x;
    grid.y  = (ny + block.y - 1) / block.y;

t9=clock();
    sumMatrixOnGPU2D<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(hipDeviceSynchronize());
t10=clock();
    printf("sumMatrixOnGPU2D  <<<  (512,1024), (32,16)  >>> elapsed %f sec\n", (double)(t10-t9)/(CLOCKS_PER_SEC));

    block.y = 16;
    block.x = 16;
    grid.x  = (nx + block.x - 1) / block.x;
    grid.y  = (ny + block.y - 1) / block.y;

t11=clock();
    sumMatrixOnGPU2D<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(hipDeviceSynchronize());
t12=clock();
    printf("sumMatrixOnGPU2D  <<<  (1024,1024),(16,16)  >>> elapsed %f sec\n", (double)(t12-t11)/(CLOCKS_PER_SEC));

    block.y = 16;
    block.x = 64;
    grid.x  = (nx + block.x - 1) / block.x;
    grid.y  = (ny + block.y - 1) / block.y;

t13=clock();
    sumMatrixOnGPU2D<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(hipDeviceSynchronize());
t14=clock();
    printf("sumMatrixOnGPU2D  <<<  (256,1024), (64,16)  >>> elapsed %f sec\n", (double)(t14-t13)/(CLOCKS_PER_SEC));

    block.y = 64;
    block.x = 16;
    grid.x  = (nx + block.x - 1) / block.x;
    grid.y  = (ny + block.y - 1) / block.y;

t15=clock();
    sumMatrixOnGPU2D<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(hipDeviceSynchronize());
t16=clock();
    printf("sumMatrixOnGPU2D  <<<  (1024,256), (16,64)  >>> elapsed %f sec\n", (double)(t16-t15)/(CLOCKS_PER_SEC));
printf("\n");

    block.x = 32;
    grid.x  = (nx + block.x - 1) / block.x;
    block.y = 1;
    grid.y  = 1;

t17=clock();
    sumMatrixOnGPU1D<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(hipDeviceSynchronize());
t18=clock();
    printf("sumMatrixOnGPU1D  <<<  (512,1)   , (32,1)   >>> elapsed %f sec\n", (double)(t18-t17)/(CLOCKS_PER_SEC));

    block.x = 64;
    grid.x  = (nx + block.x - 1) / block.x;
    block.y = 1;
    grid.y  = 1;

t19=clock();
    sumMatrixOnGPU1D<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(hipDeviceSynchronize());
t20=clock();
    printf("sumMatrixOnGPU1D  <<<  (256,1)   , (64,1)   >>> elapsed %f sec\n", (double)(t20-t19)/(CLOCKS_PER_SEC));

    block.x = 128;
    grid.x  = (nx + block.x - 1) / block.x;
    block.y = 1;
    grid.y  = 1;

t21=clock();
    sumMatrixOnGPU1D<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(hipDeviceSynchronize());
t22=clock();;
    printf("sumMatrixOnGPU1D  <<<  (128,1)   , (128,1)  >>> elapsed %f sec\n", (double)(t22-t21)/(CLOCKS_PER_SEC));

printf("\n");
    // grid 2D and block 1D
    block.x = 32;
    grid.x  = (nx + block.x - 1) / block.x;
    block.y = 1;
    grid.y  = ny;

t23=clock();
    sumMatrixOnGPUMix<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(hipDeviceSynchronize());
t24=clock();
    printf("sumMatrixOnGPUMix <<<  (512,16384),(32,1)   >>> elapsed %f sec\n",(double)(t24-t23)/(CLOCKS_PER_SEC));

    block.x = 64;
    grid.x  = (nx + block.x - 1) / block.x;
    block.y = 1;
    grid.y  = ny;

t25=clock();
    sumMatrixOnGPUMix<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(hipDeviceSynchronize());
t26=clock();
    printf("sumMatrixOnGPUMix <<<  (256,16384),(64,1)   >>> elapsed %f sec\n",(double)(t26-t25)/(CLOCKS_PER_SEC));

    block.x = 128;
    grid.x  = (nx + block.x - 1) / block.x;
    block.y = 1;
    grid.y  = ny;

t27=clock();
    sumMatrixOnGPUMix<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(hipDeviceSynchronize());
t28=clock();
    printf("sumMatrixOnGPUMix <<<  (128,16384),(128,1)  >>> elapsed %f sec\n",(double)(t28-t27)/(CLOCKS_PER_SEC));

    block.x = 256;
    grid.x  = (nx + block.x - 1) / block.x;
    block.y = 1;
    grid.y  = ny;

t29=clock();
    sumMatrixOnGPUMix<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(hipDeviceSynchronize());
t30=clock();
    printf("sumMatrixOnGPUMix <<<  (64,16384), (256,1)  >>> elapsed %f sec\n",(double)(t30-t29)/(CLOCKS_PER_SEC));

    block.x = 512;
    grid.x  = (nx + block.x - 1) / block.x;
    block.y = 1;
    grid.y  = ny;

t31=clock();
    sumMatrixOnGPUMix<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(hipDeviceSynchronize());
t32=clock();
    printf("sumMatrixOnGPUMix <<<  (32,16384), (512,1)  >>> elapsed %f sec\n",(double)(t32-t31)/(CLOCKS_PER_SEC));




}














	.text
	.file	"test6_cu.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.globl	_Z31__device_stub__sumMatrixOnGPU2DPfS_S_ii # -- Begin function _Z31__device_stub__sumMatrixOnGPU2DPfS_S_ii
	.type	_Z31__device_stub__sumMatrixOnGPU2DPfS_S_ii,@function
_Z31__device_stub__sumMatrixOnGPU2DPfS_S_ii: # @_Z31__device_stub__sumMatrixOnGPU2DPfS_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	$_Z16sumMatrixOnGPU2DPfS_S_ii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$144, %rsp
	.cfi_adjust_cfa_offset -144
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z31__device_stub__sumMatrixOnGPU2DPfS_S_ii, .Lfunc_end0-_Z31__device_stub__sumMatrixOnGPU2DPfS_S_ii
	.cfi_endproc
                                        # -- End function
	.globl	_Z31__device_stub__sumMatrixOnGPU1DPfS_S_ii # -- Begin function _Z31__device_stub__sumMatrixOnGPU1DPfS_S_ii
	.type	_Z31__device_stub__sumMatrixOnGPU1DPfS_S_ii,@function
_Z31__device_stub__sumMatrixOnGPU1DPfS_S_ii: # @_Z31__device_stub__sumMatrixOnGPU1DPfS_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	$_Z16sumMatrixOnGPU1DPfS_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_end1:
	.size	_Z31__device_stub__sumMatrixOnGPU1DPfS_S_ii, .Lfunc_end1-_Z31__device_stub__sumMatrixOnGPU1DPfS_S_ii
	.cfi_endproc
                                        # -- End function
	.globl	_Z32__device_stub__sumMatrixOnGPUMixPfS_S_ii # -- Begin function _Z32__device_stub__sumMatrixOnGPUMixPfS_S_ii
	.type	_Z32__device_stub__sumMatrixOnGPUMixPfS_S_ii,@function
_Z32__device_stub__sumMatrixOnGPUMixPfS_S_ii: # @_Z32__device_stub__sumMatrixOnGPUMixPfS_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	$_Z17sumMatrixOnGPUMixPfS_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_end2:
	.size	_Z32__device_stub__sumMatrixOnGPUMixPfS_S_ii, .Lfunc_end2-_Z32__device_stub__sumMatrixOnGPUMixPfS_S_ii
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function gpuTestAll
.LCPI3_0:
	.quad	0x412e848000000000              # double 1.0E+6
	.text
	.globl	gpuTestAll
	.type	gpuTestAll,@function
gpuTestAll:                             # @gpuTestAll
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$56, %rsp
	.cfi_def_cfa_offset 112
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
                                        # kill: def $r8d killed $r8d def $r8
	movl	%ecx, %r13d
	movq	%rdx, 24(%rsp)                  # 8-byte Spill
	movq	%rsi, 16(%rsp)                  # 8-byte Spill
	movq	%rdi, 8(%rsp)                   # 8-byte Spill
	movabsq	$137438953488, %rbx             # imm = 0x2000000010
	leal	31(%r13), %ebp
	shrl	$5, %ebp
	movq	%r8, (%rsp)                     # 8-byte Spill
	leal	31(%r8), %r14d
	shrl	$5, %r14d
	shlq	$32, %r14
	leaq	(%r14,%rbp), %r12
	callq	clock
	movq	%rax, %r15
	leaq	16(%rbx), %rdx
	movq	%r12, %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_2
# %bb.1:
	movq	8(%rsp), %rdi                   # 8-byte Reload
	movq	16(%rsp), %rsi                  # 8-byte Reload
	movq	24(%rsp), %rdx                  # 8-byte Reload
	movl	%r13d, %ecx
	movq	(%rsp), %r8                     # 8-byte Reload
                                        # kill: def $r8d killed $r8d killed $r8
	callq	_Z31__device_stub__sumMatrixOnGPU2DPfS_S_ii
.LBB3_2:
	callq	hipDeviceSynchronize
	testl	%eax, %eax
	jne	.LBB3_43
# %bb.3:
	callq	clock
	subq	%r15, %rax
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI3_0(%rip), %xmm0
	movl	$.L.str.1, %edi
	movb	$1, %al
	callq	printf
	movq	%r13, 32(%rsp)                  # 8-byte Spill
	addl	$15, %r13d
	shrl	$4, %r13d
	orq	%r13, %r14
	callq	clock
	movq	%rax, %r15
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_5
# %bb.4:
	movq	8(%rsp), %rdi                   # 8-byte Reload
	movq	16(%rsp), %rsi                  # 8-byte Reload
	movq	24(%rsp), %rdx                  # 8-byte Reload
	movq	32(%rsp), %rcx                  # 8-byte Reload
                                        # kill: def $ecx killed $ecx killed $rcx
	movq	(%rsp), %r8                     # 8-byte Reload
                                        # kill: def $r8d killed $r8d killed $r8
	callq	_Z31__device_stub__sumMatrixOnGPU2DPfS_S_ii
.LBB3_5:
	callq	hipDeviceSynchronize
	testl	%eax, %eax
	jne	.LBB3_43
# %bb.6:
	movabsq	$68719476752, %rbx              # imm = 0x1000000010
	callq	clock
	subq	%r15, %rax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI3_0(%rip), %xmm0
	movl	$.L.str.2, %edi
	movb	$1, %al
	callq	printf
	movq	(%rsp), %rax                    # 8-byte Reload
	leal	15(%rax), %r14d
	shrl	$4, %r14d
	shlq	$32, %r14
	leaq	(%r14,%rbp), %r12
	callq	clock
	movq	%rax, %r15
	leaq	16(%rbx), %rdx
	movq	%r12, %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_8
# %bb.7:
	movq	8(%rsp), %rdi                   # 8-byte Reload
	movq	16(%rsp), %rsi                  # 8-byte Reload
	movq	24(%rsp), %rdx                  # 8-byte Reload
	movq	32(%rsp), %rcx                  # 8-byte Reload
                                        # kill: def $ecx killed $ecx killed $rcx
	movq	(%rsp), %r8                     # 8-byte Reload
                                        # kill: def $r8d killed $r8d killed $r8
	callq	_Z31__device_stub__sumMatrixOnGPU2DPfS_S_ii
.LBB3_8:
	callq	hipDeviceSynchronize
	testl	%eax, %eax
	jne	.LBB3_43
# %bb.9:
	callq	clock
	subq	%r15, %rax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI3_0(%rip), %xmm0
	movl	$.L.str.3, %edi
	movb	$1, %al
	callq	printf
	movq	%r14, %r12
	orq	%r13, %r12
	callq	clock
	movq	%rax, %r15
	movq	%r12, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_11
# %bb.10:
	movq	8(%rsp), %rdi                   # 8-byte Reload
	movq	16(%rsp), %rsi                  # 8-byte Reload
	movq	24(%rsp), %rdx                  # 8-byte Reload
	movq	32(%rsp), %rcx                  # 8-byte Reload
                                        # kill: def $ecx killed $ecx killed $rcx
	movq	(%rsp), %r8                     # 8-byte Reload
                                        # kill: def $r8d killed $r8d killed $r8
	callq	_Z31__device_stub__sumMatrixOnGPU2DPfS_S_ii
.LBB3_11:
	callq	hipDeviceSynchronize
	testl	%eax, %eax
	jne	.LBB3_43
# %bb.12:
	callq	clock
	subq	%r15, %rax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI3_0(%rip), %xmm0
	movl	$.L.str.4, %edi
	movb	$1, %al
	callq	printf
	movq	32(%rsp), %rax                  # 8-byte Reload
	leal	63(%rax), %r12d
	shrl	$6, %r12d
	orq	%r12, %r14
	callq	clock
	movq	%rax, %r15
	addq	$48, %rbx
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_14
# %bb.13:
	movq	8(%rsp), %rdi                   # 8-byte Reload
	movq	16(%rsp), %rsi                  # 8-byte Reload
	movq	24(%rsp), %rdx                  # 8-byte Reload
	movq	32(%rsp), %rcx                  # 8-byte Reload
                                        # kill: def $ecx killed $ecx killed $rcx
	movq	(%rsp), %r8                     # 8-byte Reload
                                        # kill: def $r8d killed $r8d killed $r8
	callq	_Z31__device_stub__sumMatrixOnGPU2DPfS_S_ii
.LBB3_14:
	callq	hipDeviceSynchronize
	testl	%eax, %eax
	jne	.LBB3_43
# %bb.15:
	callq	clock
	subq	%r15, %rax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI3_0(%rip), %xmm0
	movl	$.L.str.5, %edi
	movb	$1, %al
	callq	printf
	movq	(%rsp), %rax                    # 8-byte Reload
	leal	63(%rax), %r14d
	shrl	$6, %r14d
	shlq	$32, %r14
	orq	%r13, %r14
	callq	clock
	movq	%rax, %rbx
	movabsq	$274877906960, %rdx             # imm = 0x4000000010
	movq	%r14, %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	movq	32(%rsp), %r13                  # 8-byte Reload
	jne	.LBB3_17
# %bb.16:
	movq	8(%rsp), %rdi                   # 8-byte Reload
	movq	16(%rsp), %rsi                  # 8-byte Reload
	movq	24(%rsp), %rdx                  # 8-byte Reload
	movl	%r13d, %ecx
	movq	(%rsp), %r8                     # 8-byte Reload
                                        # kill: def $r8d killed $r8d killed $r8
	callq	_Z31__device_stub__sumMatrixOnGPU2DPfS_S_ii
.LBB3_17:
	callq	hipDeviceSynchronize
	testl	%eax, %eax
	jne	.LBB3_43
# %bb.18:
	movabsq	$4294967296, %r15               # imm = 0x100000000
	callq	clock
	subq	%rbx, %rax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI3_0(%rip), %xmm0
	movl	$.L.str.6, %edi
	movb	$1, %al
	callq	printf
	movl	$10, %edi
	callq	putchar@PLT
	movq	%rbp, %r14
	btsq	$32, %r14
	callq	clock
	movq	%rax, %rbx
	leaq	32(%r15), %rdx
	movq	%r14, %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_20
# %bb.19:
	movq	8(%rsp), %rdi                   # 8-byte Reload
	movq	16(%rsp), %rsi                  # 8-byte Reload
	movq	24(%rsp), %rdx                  # 8-byte Reload
	movl	%r13d, %ecx
	movq	(%rsp), %r8                     # 8-byte Reload
                                        # kill: def $r8d killed $r8d killed $r8
	callq	_Z31__device_stub__sumMatrixOnGPU1DPfS_S_ii
.LBB3_20:
	callq	hipDeviceSynchronize
	testl	%eax, %eax
	jne	.LBB3_43
# %bb.21:
	callq	clock
	subq	%rbx, %rax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI3_0(%rip), %xmm0
	movl	$.L.str.8, %edi
	movb	$1, %al
	callq	printf
	movq	%r12, %r14
	orq	%r15, %r14
	callq	clock
	movq	%rax, %rbx
	leaq	64(%r15), %rdx
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%rdx, 48(%rsp)                  # 8-byte Spill
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_23
# %bb.22:
	movq	8(%rsp), %rdi                   # 8-byte Reload
	movq	16(%rsp), %rsi                  # 8-byte Reload
	movq	24(%rsp), %rdx                  # 8-byte Reload
	movl	%r13d, %ecx
	movq	(%rsp), %r8                     # 8-byte Reload
                                        # kill: def $r8d killed $r8d killed $r8
	callq	_Z31__device_stub__sumMatrixOnGPU1DPfS_S_ii
.LBB3_23:
	callq	hipDeviceSynchronize
	testl	%eax, %eax
	jne	.LBB3_43
# %bb.24:
	callq	clock
	subq	%rbx, %rax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI3_0(%rip), %xmm0
	movl	$.L.str.9, %edi
	movb	$1, %al
	callq	printf
	addl	$127, %r13d
	shrl	$7, %r13d
	movq	%r13, %r14
	orq	%r15, %r14
	callq	clock
	movq	%rax, %rbx
	leaq	128(%r15), %rdx
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%rdx, 40(%rsp)                  # 8-byte Spill
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_26
# %bb.25:
	movq	8(%rsp), %rdi                   # 8-byte Reload
	movq	16(%rsp), %rsi                  # 8-byte Reload
	movq	24(%rsp), %rdx                  # 8-byte Reload
	movq	32(%rsp), %rcx                  # 8-byte Reload
                                        # kill: def $ecx killed $ecx killed $rcx
	movq	(%rsp), %r8                     # 8-byte Reload
                                        # kill: def $r8d killed $r8d killed $r8
	callq	_Z31__device_stub__sumMatrixOnGPU1DPfS_S_ii
.LBB3_26:
	callq	hipDeviceSynchronize
	testl	%eax, %eax
	jne	.LBB3_43
# %bb.27:
	callq	clock
	subq	%rbx, %rax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI3_0(%rip), %xmm0
	movl	$.L.str.10, %edi
	movb	$1, %al
	callq	printf
	movl	$10, %edi
	callq	putchar@PLT
	movq	(%rsp), %rbx                    # 8-byte Reload
	shlq	$32, %rbx
	orq	%rbx, %rbp
	callq	clock
	movq	%rax, %r14
	leaq	32(%r15), %rdx
	movq	%rbp, %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_29
# %bb.28:
	movq	8(%rsp), %rdi                   # 8-byte Reload
	movq	16(%rsp), %rsi                  # 8-byte Reload
	movq	24(%rsp), %rdx                  # 8-byte Reload
	movq	32(%rsp), %rcx                  # 8-byte Reload
                                        # kill: def $ecx killed $ecx killed $rcx
	movq	(%rsp), %r8                     # 8-byte Reload
                                        # kill: def $r8d killed $r8d killed $r8
	callq	_Z32__device_stub__sumMatrixOnGPUMixPfS_S_ii
.LBB3_29:
	callq	hipDeviceSynchronize
	testl	%eax, %eax
	jne	.LBB3_43
# %bb.30:
	callq	clock
	subq	%r14, %rax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI3_0(%rip), %xmm0
	movl	$.L.str.11, %edi
	movb	$1, %al
	callq	printf
	orq	%rbx, %r12
	callq	clock
	movq	%rax, %r14
	movq	%r12, %rdi
	movl	$1, %esi
	movq	48(%rsp), %rdx                  # 8-byte Reload
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_32
# %bb.31:
	movq	8(%rsp), %rdi                   # 8-byte Reload
	movq	16(%rsp), %rsi                  # 8-byte Reload
	movq	24(%rsp), %rdx                  # 8-byte Reload
	movq	32(%rsp), %rcx                  # 8-byte Reload
                                        # kill: def $ecx killed $ecx killed $rcx
	movq	(%rsp), %r8                     # 8-byte Reload
                                        # kill: def $r8d killed $r8d killed $r8
	callq	_Z32__device_stub__sumMatrixOnGPUMixPfS_S_ii
.LBB3_32:
	callq	hipDeviceSynchronize
	testl	%eax, %eax
	jne	.LBB3_43
# %bb.33:
	callq	clock
	subq	%r14, %rax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI3_0(%rip), %xmm0
	movl	$.L.str.12, %edi
	movb	$1, %al
	callq	printf
	orq	%rbx, %r13
	callq	clock
	movq	%rax, %r14
	movq	%r13, %rdi
	movl	$1, %esi
	movq	40(%rsp), %rdx                  # 8-byte Reload
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	movq	32(%rsp), %r13                  # 8-byte Reload
	jne	.LBB3_35
# %bb.34:
	movq	8(%rsp), %rdi                   # 8-byte Reload
	movq	16(%rsp), %rsi                  # 8-byte Reload
	movq	24(%rsp), %rdx                  # 8-byte Reload
	movl	%r13d, %ecx
	movq	(%rsp), %r8                     # 8-byte Reload
                                        # kill: def $r8d killed $r8d killed $r8
	callq	_Z32__device_stub__sumMatrixOnGPUMixPfS_S_ii
.LBB3_35:
	callq	hipDeviceSynchronize
	testl	%eax, %eax
	jne	.LBB3_43
# %bb.36:
	callq	clock
	subq	%r14, %rax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI3_0(%rip), %xmm0
	movl	$.L.str.13, %edi
	movb	$1, %al
	callq	printf
	leal	255(%r13), %r12d
	shrl	$8, %r12d
	orq	%rbx, %r12
	callq	clock
	movq	%rax, %r14
	leaq	256(%r15), %rdx
	movq	%r12, %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_38
# %bb.37:
	movq	8(%rsp), %rdi                   # 8-byte Reload
	movq	16(%rsp), %rsi                  # 8-byte Reload
	movq	24(%rsp), %rdx                  # 8-byte Reload
	movl	%r13d, %ecx
	movq	(%rsp), %r8                     # 8-byte Reload
                                        # kill: def $r8d killed $r8d killed $r8
	callq	_Z32__device_stub__sumMatrixOnGPUMixPfS_S_ii
.LBB3_38:
	callq	hipDeviceSynchronize
	testl	%eax, %eax
	jne	.LBB3_43
# %bb.39:
	callq	clock
	subq	%r14, %rax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI3_0(%rip), %xmm0
	movl	$.L.str.14, %edi
	movb	$1, %al
	callq	printf
	leal	511(%r13), %eax
	shrl	$9, %eax
	orq	%rax, %rbx
	callq	clock
	movq	%rax, %r14
	addq	$512, %r15                      # imm = 0x200
	movq	%rbx, %rdi
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_41
# %bb.40:
	movq	8(%rsp), %rdi                   # 8-byte Reload
	movq	16(%rsp), %rsi                  # 8-byte Reload
	movq	24(%rsp), %rdx                  # 8-byte Reload
	movl	%r13d, %ecx
	movq	(%rsp), %r8                     # 8-byte Reload
                                        # kill: def $r8d killed $r8d killed $r8
	callq	_Z32__device_stub__sumMatrixOnGPUMixPfS_S_ii
.LBB3_41:
	callq	hipDeviceSynchronize
	testl	%eax, %eax
	jne	.LBB3_43
# %bb.42:
	callq	clock
	subq	%r14, %rax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI3_0(%rip), %xmm0
	movl	$.L.str.15, %edi
	movb	$1, %al
	addq	$56, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	printf                          # TAILCALL
.LBB3_43:
	.cfi_def_cfa_offset 112
	movl	$_ZSt4cout, %edi
	movl	$.L.str, %esi
	callq	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc
	movq	%rax, %rbx
	callq	hipDeviceSynchronize
	movq	%rbx, %rdi
	movl	%eax, %esi
	callq	_ZNSolsEi
	callq	abort
.Lfunc_end3:
	.size	gpuTestAll, .Lfunc_end3-gpuTestAll
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB4_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB4_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z16sumMatrixOnGPU2DPfS_S_ii, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z16sumMatrixOnGPU1DPfS_S_ii, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z17sumMatrixOnGPUMixPfS_S_ii, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end4:
	.size	__hip_module_ctor, .Lfunc_end4-__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	.LBB5_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
.LBB5_2:
	retq
.Lfunc_end5:
	.size	__hip_module_dtor, .Lfunc_end5-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z16sumMatrixOnGPU2DPfS_S_ii,@object # @_Z16sumMatrixOnGPU2DPfS_S_ii
	.section	.rodata,"a",@progbits
	.globl	_Z16sumMatrixOnGPU2DPfS_S_ii
	.p2align	3, 0x0
_Z16sumMatrixOnGPU2DPfS_S_ii:
	.quad	_Z31__device_stub__sumMatrixOnGPU2DPfS_S_ii
	.size	_Z16sumMatrixOnGPU2DPfS_S_ii, 8

	.type	_Z16sumMatrixOnGPU1DPfS_S_ii,@object # @_Z16sumMatrixOnGPU1DPfS_S_ii
	.globl	_Z16sumMatrixOnGPU1DPfS_S_ii
	.p2align	3, 0x0
_Z16sumMatrixOnGPU1DPfS_S_ii:
	.quad	_Z31__device_stub__sumMatrixOnGPU1DPfS_S_ii
	.size	_Z16sumMatrixOnGPU1DPfS_S_ii, 8

	.type	_Z17sumMatrixOnGPUMixPfS_S_ii,@object # @_Z17sumMatrixOnGPUMixPfS_S_ii
	.globl	_Z17sumMatrixOnGPUMixPfS_S_ii
	.p2align	3, 0x0
_Z17sumMatrixOnGPUMixPfS_S_ii:
	.quad	_Z32__device_stub__sumMatrixOnGPUMixPfS_S_ii
	.size	_Z17sumMatrixOnGPUMixPfS_S_ii, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Cuda failure: "
	.size	.L.str, 15

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"sumMatrixOnGPU2D  <<<  (512,512),  (32,32)  >>> elapsed %f sec\n"
	.size	.L.str.1, 64

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"sumMatrixOnGPU2D  <<<  (1024,512), (16,32)  >>> elapsed %f sec\n"
	.size	.L.str.2, 64

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"sumMatrixOnGPU2D  <<<  (512,1024), (32,16)  >>> elapsed %f sec\n"
	.size	.L.str.3, 64

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"sumMatrixOnGPU2D  <<<  (1024,1024),(16,16)  >>> elapsed %f sec\n"
	.size	.L.str.4, 64

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"sumMatrixOnGPU2D  <<<  (256,1024), (64,16)  >>> elapsed %f sec\n"
	.size	.L.str.5, 64

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"sumMatrixOnGPU2D  <<<  (1024,256), (16,64)  >>> elapsed %f sec\n"
	.size	.L.str.6, 64

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"sumMatrixOnGPU1D  <<<  (512,1)   , (32,1)   >>> elapsed %f sec\n"
	.size	.L.str.8, 64

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"sumMatrixOnGPU1D  <<<  (256,1)   , (64,1)   >>> elapsed %f sec\n"
	.size	.L.str.9, 64

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"sumMatrixOnGPU1D  <<<  (128,1)   , (128,1)  >>> elapsed %f sec\n"
	.size	.L.str.10, 64

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"sumMatrixOnGPUMix <<<  (512,16384),(32,1)   >>> elapsed %f sec\n"
	.size	.L.str.11, 64

	.type	.L.str.12,@object               # @.str.12
.L.str.12:
	.asciz	"sumMatrixOnGPUMix <<<  (256,16384),(64,1)   >>> elapsed %f sec\n"
	.size	.L.str.12, 64

	.type	.L.str.13,@object               # @.str.13
.L.str.13:
	.asciz	"sumMatrixOnGPUMix <<<  (128,16384),(128,1)  >>> elapsed %f sec\n"
	.size	.L.str.13, 64

	.type	.L.str.14,@object               # @.str.14
.L.str.14:
	.asciz	"sumMatrixOnGPUMix <<<  (64,16384), (256,1)  >>> elapsed %f sec\n"
	.size	.L.str.14, 64

	.type	.L.str.15,@object               # @.str.15
.L.str.15:
	.asciz	"sumMatrixOnGPUMix <<<  (32,16384), (512,1)  >>> elapsed %f sec\n"
	.size	.L.str.15, 64

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z16sumMatrixOnGPU2DPfS_S_ii"
	.size	.L__unnamed_1, 29

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z16sumMatrixOnGPU1DPfS_S_ii"
	.size	.L__unnamed_2, 29

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"_Z17sumMatrixOnGPUMixPfS_S_ii"
	.size	.L__unnamed_3, 30

	.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 _Z31__device_stub__sumMatrixOnGPU2DPfS_S_ii
	.addrsig_sym _Z31__device_stub__sumMatrixOnGPU1DPfS_S_ii
	.addrsig_sym _Z32__device_stub__sumMatrixOnGPUMixPfS_S_ii
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z16sumMatrixOnGPU2DPfS_S_ii
	.addrsig_sym _Z16sumMatrixOnGPU1DPfS_S_ii
	.addrsig_sym _Z17sumMatrixOnGPUMixPfS_S_ii
	.addrsig_sym _ZSt4cout
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z16sumMatrixOnGPU2DPfS_S_ii ; -- Begin function _Z16sumMatrixOnGPU2DPfS_S_ii
	.globl	_Z16sumMatrixOnGPU2DPfS_S_ii
	.p2align	8
	.type	_Z16sumMatrixOnGPU2DPfS_S_ii,@function
_Z16sumMatrixOnGPU2DPfS_S_ii:           ; @_Z16sumMatrixOnGPU2DPfS_S_ii
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x2c
	s_load_b64 s[4:5], s[0:1], 0x18
	v_and_b32_e32 v2, 0x3ff, v0
	v_bfe_u32 v3, v0, 10, 10
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s3, s2, 16
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_mad_u64_u32 v[0:1], null, s14, s2, v[2:3]
	v_mad_u64_u32 v[1:2], null, s15, s3, v[3:4]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_gt_u32_e32 vcc_lo, s4, v0
	v_cmp_gt_u32_e64 s2, s5, v1
	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:
	s_load_b128 s[8:11], s[0:1], 0x0
	v_mad_u64_u32 v[2:3], null, v1, s4, v[0:1]
	v_mov_b32_e32 v3, 0
	s_load_b64 s[0:1], s[0:1], 0x10
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[2:3]
	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 _Z16sumMatrixOnGPU2DPfS_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	_Z16sumMatrixOnGPU2DPfS_S_ii, .Lfunc_end0-_Z16sumMatrixOnGPU2DPfS_S_ii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 232
; 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
	.protected	_Z16sumMatrixOnGPU1DPfS_S_ii ; -- Begin function _Z16sumMatrixOnGPU1DPfS_S_ii
	.globl	_Z16sumMatrixOnGPU1DPfS_S_ii
	.p2align	8
	.type	_Z16sumMatrixOnGPU1DPfS_S_ii,@function
_Z16sumMatrixOnGPU1DPfS_S_ii:           ; @_Z16sumMatrixOnGPU1DPfS_S_ii
; %bb.0:
	s_clause 0x1
	s_load_b32 s4, s[0:1], 0x2c
	s_load_b64 s[2:3], s[0:1], 0x18
	s_waitcnt lgkmcnt(0)
	s_and_b32 s4, s4, 0xffff
	s_cmp_gt_i32 s3, 0
	v_mad_u64_u32 v[1:2], null, s15, s4, v[0:1]
	s_cselect_b32 s4, -1, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cmp_gt_u32_e32 vcc_lo, s2, v1
	s_and_b32 s4, vcc_lo, s4
	s_and_saveexec_b32 s5, s4
	s_cbranch_execz .LBB1_3
; %bb.1:                                ; %.lr.ph.preheader
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[0:1], s[0:1], 0x10
.LBB1_2:                                ; %.lr.ph
                                        ; =>This Inner Loop Header: Depth=1
	v_ashrrev_i32_e32 v2, 31, v1
	s_add_i32 s3, s3, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_cmp_lg_u32 s3, 0
	v_lshlrev_b64 v[2:3], 2, v[1:2]
	v_add_nc_u32_e32 v1, s2, v1
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v4, vcc_lo, s4, v2
	v_add_co_ci_u32_e32 v5, vcc_lo, s5, v3, vcc_lo
	v_add_co_u32 v6, vcc_lo, s6, v2
	v_add_co_ci_u32_e32 v7, vcc_lo, s7, v3, vcc_lo
	global_load_b32 v0, v[4:5], off
	global_load_b32 v4, v[6:7], off
	v_add_co_u32 v2, vcc_lo, s0, v2
	v_add_co_ci_u32_e32 v3, vcc_lo, s1, v3, vcc_lo
	s_waitcnt vmcnt(0)
	v_add_f32_e32 v0, v0, v4
	global_store_b32 v[2:3], v0, off
	s_cbranch_scc1 .LBB1_2
.LBB1_3:                                ; %.loopexit
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z16sumMatrixOnGPU1DPfS_S_ii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 8
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end1:
	.size	_Z16sumMatrixOnGPU1DPfS_S_ii, .Lfunc_end1-_Z16sumMatrixOnGPU1DPfS_S_ii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 208
; NumSgprs: 18
; NumVgprs: 8
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 8
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.protected	_Z17sumMatrixOnGPUMixPfS_S_ii ; -- Begin function _Z17sumMatrixOnGPUMixPfS_S_ii
	.globl	_Z17sumMatrixOnGPUMixPfS_S_ii
	.p2align	8
	.type	_Z17sumMatrixOnGPUMixPfS_S_ii,@function
_Z17sumMatrixOnGPUMixPfS_S_ii:          ; @_Z17sumMatrixOnGPUMixPfS_S_ii
; %bb.0:
	s_clause 0x1
	s_load_b32 s4, s[0:1], 0x2c
	s_load_b64 s[2:3], s[0:1], 0x18
	s_waitcnt lgkmcnt(0)
	s_and_b32 s4, s4, 0xffff
	s_cmp_lt_u32 s15, s3
	v_mad_u64_u32 v[1:2], null, s14, s4, v[0:1]
	s_cselect_b32 s3, -1, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cmp_gt_u32_e32 vcc_lo, s2, v1
	s_and_b32 s3, s3, vcc_lo
	s_and_saveexec_b32 s4, s3
	s_cbranch_execz .LBB2_2
; %bb.1:
	s_load_b128 s[4:7], s[0:1], 0x0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[2:3], null, s15, s2, v[1:2]
	v_mov_b32_e32 v3, 0
	s_load_b64 s[0:1], s[0:1], 0x10
	v_lshlrev_b64 v[0:1], 2, v[2:3]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v2, vcc_lo, s4, v0
	v_add_co_ci_u32_e32 v3, vcc_lo, s5, v1, vcc_lo
	v_add_co_u32 v4, vcc_lo, s6, v0
	v_add_co_ci_u32_e32 v5, vcc_lo, s7, 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
.LBB2_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z17sumMatrixOnGPUMixPfS_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 0
		.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_end2:
	.size	_Z17sumMatrixOnGPUMixPfS_S_ii, .Lfunc_end2-_Z17sumMatrixOnGPUMixPfS_S_ii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 196
; 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: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         28
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z16sumMatrixOnGPU2DPfS_S_ii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z16sumMatrixOnGPU2DPfS_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
  - .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:           _Z16sumMatrixOnGPU1DPfS_S_ii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z16sumMatrixOnGPU1DPfS_S_ii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     8
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         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:           _Z17sumMatrixOnGPUMixPfS_S_ii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z17sumMatrixOnGPUMixPfS_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
