#include <cuda.h>
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <cmath>
#include <time.h>
#include <iostream>

#define DTYPE float
__global__ void kernel(float *a, float *x, float* buff,int Xblocks,int size,bool comp,int toreduce)
{
	int i=threadIdx.x+blockIdx.x*blockDim.x;
   	int j=threadIdx.y+blockIdx.y*blockDim.y;
   	__shared__ float xSM[1024];
   	if(comp==true){
		xSM[threadIdx.x+threadIdx.y*blockDim.x]=a[i+j*size]*x[i];
		__syncthreads();
	}
	else{
		xSM[threadIdx.x+threadIdx.y*blockDim.x]=buff[i+j*Xblocks];
	}   	
   	if(i<toreduce){	
        	for(int k=blockDim.x/2;k>0;k/=2){
            		if(threadIdx.x<k&&i+k<toreduce){
                		xSM[threadIdx.x+threadIdx.y*blockDim.x]+=xSM[threadIdx.x+k+threadIdx.y*blockDim.x];
                		__syncthreads();
            		}
        	}
        }
       if(threadIdx.x==0){
		}
       	buff[blockIdx.x+j*Xblocks]=xSM[threadIdx.y*blockDim.x];
}

void fillA(DTYPE *a, int size)
{
   for (int i=0;i<size*size;i++)
      a[i]=1.0;
}

//X mit Werten füllen 
void fillX(DTYPE *x, int size)
{
   for (int i=0;i<size;i++){
	x[i]=1;
//      x[i]= (DTYPE)(i+1);
	}
}
int main(int argc, char**argv)
{   
   int sx=32;
   int sy=32;
   int i=1;
   bool standard=true;
   if (argc>1)   
   {
	standard=false;
	sx=atoi(argv[1]);
      	if (argc>2){ 
      		sy=atoi(argv[2]);
      		if(argc>3){ 
      			i=atoi(argv[3]);      					
      		}
      	}	
   }
   if(standard){
   	std::cout<<"Do experiment with standard settings"<<std::endl;
   	std::cout<<"Sx="<<sx<<"\nSy="<<sy<<"\nSize=1024*"<<i<<std::endl;
   }else 
   {
   	std::cout<<"Do experiment with individual settings"<<std::endl;
   	std::cout<<"Sx="<<sx<<"\nSy="<<sy<<"\nSize=1024*"<<i<<std::endl;
       }
   /*if(sx*sy!=t){
   	std::cout<<"Sx*Sy has to be equal to threads per block"<<std::endl;
   	return -1;
   }*/
   int size=1024*i;
   int xblocks=size/sx;
   //Datenfelder anlegen für Host
   DTYPE *a_host, *buff_host,*x_host;
   //und Device
   DTYPE *a_dev, *buff_dev,*x_dev;
   //Events für die Zeitmessung
   cudaEvent_t start,end;
   //Zeiten: 
   //htd: Host->Device Memcpy von A und x
   //dth: Device->Host Memcpy von y
   //kernelA, kernelAT
   float kernelA_time=0.0;

   //TODO: Host Speicher anlegen und A und x füllen
   a_host = (DTYPE*)malloc(size*size*sizeof(DTYPE));
   x_host = (DTYPE*)malloc(size*sizeof(DTYPE));
   buff_host=(DTYPE*)malloc(xblocks*size*sizeof(DTYPE));

   fillA(a_host,size);
   fillX(x_host,size);
   //TODO: CUDA Events erstellen

   //TODO: CUDA Speicher anlegen für alle Arrays (a_dev,x_dev,y_dev)
   cudaMalloc((void**)&a_dev,size*size*sizeof(DTYPE));
   cudaMalloc((void**)&x_dev,size*sizeof(DTYPE));
   cudaMalloc((void**)&buff_dev,xblocks*size*sizeof(DTYPE));

	cudaMemcpy(a_dev,a_host,size*size*sizeof(DTYPE),cudaMemcpyHostToDevice);
	cudaMemcpy(x_dev,x_host,size*sizeof(DTYPE),cudaMemcpyHostToDevice);
   
   dim3 block(sx,sy);
   dim3 grid(size/block.x,size/block.y);
   
   int toreduce=size;//Anzahl der noch zu reduzierenden Wert, im ersten Kernel=size, da noch nichts reduziert wurde und man somit im ersten Berechnungsschritt size viel Skalare erhällt
   bool doComputation=true;

	//cache Konfiguration
	if(argc>4){
		if(atoi(argv[4])==1){//L1 Prefered
			std::cout<<"16 kB shared, 48kB L1"<<std::endl;
			cudaFuncSetCacheConfig(kernel, cudaFuncCachePreferL1);
		} else if(atoi(argv[4])==2){
			std::cout<<"48kB shared, 16kb L1"<<std::endl;
			cudaFuncSetCacheConfig(kernel, cudaFuncCachePreferShared);
		}else{
			std::cout<<"32kB shared, 32kB L1"<<std::endl;
			cudaFuncSetCacheConfig(kernel, cudaFuncCachePreferNone);
		}
	}
   //Start Zeitmessung, da dies nur der erste Aufruf ist und danach der Kernel nochmal aufgerufen wird, mit einem anderen Argument als du Computation, endet die Zeitmessung erst nach der while Schleife, inder der Kernel noch einmal aufgerufen wird.
   cudaEventCreate(&start);
   cudaEventCreate(&end);
   cudaEventRecord(start,0);
   kernel<<<grid,block>>>(a_dev,x_dev,buff_dev,xblocks,size,doComputation,toreduce);//,y_dev,size);

	//Überprüfung des Zwischenergebnisses
	/*cudaMemcpy(buff_host,buff_dev,xblocks*size*sizeof(DTYPE),cudaMemcpyDeviceToHost);
   for(int lj=0;lj<10;lj++){
   	for(int li=0;li<10;li++)
   		std::cout<<buff_host[li+xblocks*lj]<<" ; ";
   	std::cout<<std::endl;
   }*/
   //Vorbereitung für weitere Durchführung
   doComputation=false; //Matrix Vektor Multiplikation musst nur einmal durchgeführt werden
   if(toreduce%sx==0)
   	toreduce/=sx;
   else
   	toreduce=toreduce/sx+1;
   while(toreduce>1){//toreduce=1 bedeut, dass auf einen Wert reduziert wurde und somit keine weitere Reduzierung notwendig ist
   	int xgrid;
   	int ygrid=size/sy;
   	if(toreduce%sx==0)
   		xgrid=toreduce/sx;
   	else
	   	xgrid=toreduce/sx+1;
   	dim3 itgrid(xgrid,ygrid);
   	kernel<<<itgrid,block>>>(a_dev,x_dev,buff_dev,xblocks,size,doComputation,toreduce);//,y_dev,size);
   	if(toreduce%sx==0)
   	toreduce/=sx;
	else
   	toreduce=toreduce/sx+1;	
   }
   cudaEventRecord(end,0);
   cudaEventSynchronize(end);
   cudaEventElapsedTime(&kernelA_time,start,end);
   if(argc>4){
   	cudaMemcpy(buff_host,buff_dev,xblocks*size*sizeof(DTYPE),cudaMemcpyDeviceToHost);
   	std::cout<<"Das Endergebnis ist"<<std::endl;
   	for(int lj=0;lj<10;lj++){
   		//for(int li=0;li<10;li++)
   		//	std::cout<<buff_host[li+xblocks*lj]<<" ; ";
   //   		std::cout<<std::endl;
   		std::cout<<buff_host[xblocks*lj]<<std::endl;
   	}
   }
   std::cout<<"Computation time: "<<kernelA_time<<std::endl;
   //Bestimmung der Performance
   //float time_s=kernelA_time/1000;
   //float gflops=pow(10,-9)*size*size*2/time_s;
   float gflops=pow(10,-6)*size*size*2/kernelA_time;
   std::cout<<"Computation Performance in GFLOPs: "<<gflops<<std::endl;
   return 0;
}


	.file	"tmpxft_002e9992_00000000-6_onlyshared.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3641:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3641:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.align 2
	.type	_ZNSolsEPFRSoS_E.isra.0, @function
_ZNSolsEPFRSoS_E.isra.0:
.LFB4297:
	.cfi_startproc
	jmp	*%rsi
	.cfi_endproc
.LFE4297:
	.size	_ZNSolsEPFRSoS_E.isra.0, .-_ZNSolsEPFRSoS_E.isra.0
	.globl	_Z5fillAPfi
	.type	_Z5fillAPfi, @function
_Z5fillAPfi:
.LFB3635:
	.cfi_startproc
	endbr64
	imull	%esi, %esi
	movss	.LC0(%rip), %xmm0
	xorl	%eax, %eax
.L4:
	cmpl	%eax, %esi
	jle	.L7
	movss	%xmm0, (%rdi,%rax,4)
	incq	%rax
	jmp	.L4
.L7:
	ret
	.cfi_endproc
.LFE3635:
	.size	_Z5fillAPfi, .-_Z5fillAPfi
	.globl	_Z5fillXPfi
	.type	_Z5fillXPfi, @function
_Z5fillXPfi:
.LFB3636:
	.cfi_startproc
	endbr64
	movss	.LC0(%rip), %xmm0
	xorl	%eax, %eax
.L9:
	cmpl	%eax, %esi
	jle	.L11
	movss	%xmm0, (%rdi,%rax,4)
	incq	%rax
	jmp	.L9
.L11:
	ret
	.cfi_endproc
.LFE3636:
	.size	_Z5fillXPfi, .-_Z5fillXPfi
	.globl	_Z33__device_stub__Z6kernelPfS_S_iibiPfS_S_iibi
	.type	_Z33__device_stub__Z6kernelPfS_S_iibiPfS_S_iibi, @function
_Z33__device_stub__Z6kernelPfS_S_iibiPfS_S_iibi:
.LFB3663:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movq	%rdi, 40(%rsp)
	leaq	64(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	48(%rsp), %rdx
	movl	%ecx, 20(%rsp)
	leaq	56(%rsp), %rcx
	movl	%r8d, 16(%rsp)
	movb	%r9b, 12(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 112(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	192(%rsp), %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L12
	pushq	56(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z6kernelPfS_S_iibi(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 208
	movq	92(%rsp), %rcx
	movl	100(%rsp), %r8d
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	leaq	128(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L12:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L14
	call	__stack_chk_fail@PLT
.L14:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3663:
	.size	_Z33__device_stub__Z6kernelPfS_S_iibiPfS_S_iibi, .-_Z33__device_stub__Z6kernelPfS_S_iibiPfS_S_iibi
	.globl	_Z6kernelPfS_S_iibi
	.type	_Z6kernelPfS_S_iibi, @function
_Z6kernelPfS_S_iibi:
.LFB3664:
	.cfi_startproc
	endbr64
	movzbl	%r9b, %r9d
	jmp	_Z33__device_stub__Z6kernelPfS_S_iibiPfS_S_iibi
	.cfi_endproc
.LFE3664:
	.size	_Z6kernelPfS_S_iibi, .-_Z6kernelPfS_S_iibi
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC2:
	.string	"16 kB shared, 48kB L1"
.LC3:
	.string	"48kB shared, 16kb L1"
.LC4:
	.string	"32kB shared, 32kB L1"
.LC5:
	.string	"Computation time: "
.LC7:
	.string	"Computation Performance in GFLOPs: "
.LC8:
	.string	"Das Endergebnis ist"
.LC9:
	.string	"Do experiment with standard settings"
.LC10:
	.string	"Sx="
.LC11:
	.string	"\nSy="
.LC12:
	.string	"\nSize=1024*"
.LC13:
	.string	"Do experiment with individual settings"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3637:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rsi, %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movl	%edi, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	leaq	_ZSt4cout(%rip), %r13
	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	$168, %rsp
	.cfi_def_cfa_offset 224
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	decl	%edi
	jle	.L18
	movq	8(%rsi), %rdi
	call	atoi@PLT
	movl	%eax, %ebp
	cmpl	$2, %r14d
	je	.L34
	movq	16(%r15), %rdi
	movl	$1, %ebx
	call	atoi@PLT
	movl	%eax, %r12d
	cmpl	$3, %r14d
	je	.L19
	movq	24(%r15), %rdi
	call	atoi@PLT
	movl	%eax, %ebx
	jmp	.L19
.L18:
	leaq	.LC9(%rip), %rsi
	movq	%r13, %rdi
	movl	$1, %ebx
	movl	$32, %r12d
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movl	$32, %ebp
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC10(%rip), %rsi
	movq	%r13, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	$32, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC11(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	$32, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC12(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	$1, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
.L32:
	sall	$10, %ebx
	movl	$0x00000000, 68(%rsp)
	movl	%ebx, %eax
	cltd
	idivl	%ebp
	movslq	%eax, %r13
	movl	%ebx, %eax
	movl	%edx, 60(%rsp)
	imull	%ebx, %eax
	cltq
	salq	$2, %rax
	movq	%rax, %rdi
	movq	%rax, 8(%rsp)
	call	malloc@PLT
	movq	%rax, 40(%rsp)
	movslq	%ebx, %rax
	salq	$2, %rax
	movq	%rax, %rdi
	movq	%rax, 16(%rsp)
	call	malloc@PLT
	movq	%rax, 48(%rsp)
	movl	%ebx, %eax
	imull	%r13d, %eax
	cltq
	salq	$2, %rax
	movq	%rax, %rdi
	movq	%rax, 24(%rsp)
	call	malloc@PLT
	movq	40(%rsp), %rdi
	movl	%ebx, %esi
	movq	%rax, 32(%rsp)
	call	_Z5fillAPfi
	movq	48(%rsp), %rdi
	movl	%ebx, %esi
	call	_Z5fillXPfi
	movq	8(%rsp), %rsi
	leaq	72(%rsp), %rdi
	call	cudaMalloc@PLT
	movq	16(%rsp), %rsi
	leaq	88(%rsp), %rdi
	call	cudaMalloc@PLT
	movq	24(%rsp), %rsi
	leaq	80(%rsp), %rdi
	call	cudaMalloc@PLT
	movq	8(%rsp), %rdx
	movq	40(%rsp), %rsi
	movl	$1, %ecx
	movq	72(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	16(%rsp), %rdx
	movq	48(%rsp), %rsi
	movl	$1, %ecx
	movq	88(%rsp), %rdi
	call	cudaMemcpy@PLT
	movl	%ebx, %eax
	xorl	%edx, %edx
	movl	$1, 124(%rsp)
	divl	%ebp
	movl	$1, 136(%rsp)
	xorl	%edx, %edx
	movl	%eax, 128(%rsp)
	movl	%ebx, %eax
	divl	%r12d
	movl	%eax, 132(%rsp)
	cmpl	$4, %r14d
	jle	.L20
	movq	32(%r15), %rdi
	leaq	_Z6kernelPfS_S_iibi(%rip), %r15
	call	atoi@PLT
	leaq	_ZSt4cout(%rip), %rdi
	cmpl	$1, %eax
	jne	.L21
	leaq	.LC2(%rip), %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movl	$2, %esi
	jmp	.L42
.L21:
	cmpl	$2, %eax
	jne	.L22
	leaq	.LC3(%rip), %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movl	$1, %esi
	jmp	.L42
.L22:
	leaq	.LC4(%rip), %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	xorl	%esi, %esi
.L42:
	movq	%r15, %rdi
	call	cudaFuncSetCacheConfig@PLT
.L20:
	leaq	96(%rsp), %rdi
	call	cudaEventCreate@PLT
	leaq	104(%rsp), %rdi
	call	cudaEventCreate@PLT
	movq	96(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	%ebp, 116(%rsp)
	movl	124(%rsp), %ecx
	xorl	%r9d, %r9d
	movl	%r12d, 120(%rsp)
	movq	128(%rsp), %rdi
	xorl	%r8d, %r8d
	movq	116(%rsp), %rdx
	movl	136(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L23
	pushq	%rsi
	.cfi_def_cfa_offset 232
	movl	%ebx, %r8d
	movl	$1, %r9d
	movl	%r13d, %ecx
	pushq	%rbx
	.cfi_def_cfa_offset 240
	movq	96(%rsp), %rdx
	movq	104(%rsp), %rsi
	movq	88(%rsp), %rdi
	call	_Z33__device_stub__Z6kernelPfS_S_iibiPfS_S_iibi
	popq	%rdi
	.cfi_def_cfa_offset 232
	popq	%r8
	.cfi_def_cfa_offset 224
.L23:
	cmpl	$1, 60(%rsp)
	movl	%r13d, %r10d
	sbbl	$-1, %r10d
.L25:
	cmpl	$1, %r10d
	jle	.L43
	movl	%ebx, %eax
	movl	%ebp, 116(%rsp)
	movl	$1, %esi
	cltd
	movl	%r12d, 120(%rsp)
	idivl	%r12d
	movl	%r10d, 8(%rsp)
	movl	%eax, %ecx
	movl	%r10d, %eax
	cltd
	movl	%ecx, 144(%rsp)
	movl	124(%rsp), %ecx
	idivl	%ebp
	cmpl	$1, %edx
	movq	116(%rsp), %rdx
	sbbl	$-1, %eax
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	%eax, 140(%rsp)
	movq	140(%rsp), %rdi
	movl	%eax, %r15d
	call	__cudaPushCallConfiguration@PLT
	movl	8(%rsp), %r10d
	testl	%eax, %eax
	jne	.L27
	pushq	%rax
	.cfi_def_cfa_offset 232
	movl	%r13d, %ecx
	xorl	%r9d, %r9d
	movl	%ebx, %r8d
	pushq	%r10
	.cfi_def_cfa_offset 240
	movq	96(%rsp), %rdx
	movq	104(%rsp), %rsi
	movq	88(%rsp), %rdi
	call	_Z33__device_stub__Z6kernelPfS_S_iibiPfS_S_iibi
	popq	%rdx
	.cfi_def_cfa_offset 232
	popq	%rcx
	.cfi_def_cfa_offset 224
.L27:
	movl	%r15d, %r10d
	jmp	.L25
.L43:
	movq	104(%rsp), %rdi
	xorl	%esi, %esi
	leaq	_ZSt4cout(%rip), %rbp
	call	cudaEventRecord@PLT
	movq	104(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	104(%rsp), %rdx
	movq	96(%rsp), %rsi
	leaq	68(%rsp), %rdi
	call	cudaEventElapsedTime@PLT
	cmpl	$4, %r14d
	jg	.L29
.L31:
	leaq	.LC5(%rip), %rsi
	movq	%rbp, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cvtss2sd	68(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	cvtsi2sdl	%ebx, %xmm1
	movsd	.LC6(%rip), %xmm0
	movq	%rbp, %rdi
	leaq	.LC7(%rip), %rsi
	mulsd	%xmm1, %xmm0
	mulsd	%xmm1, %xmm0
	cvtss2sd	68(%rsp), %xmm1
	addsd	%xmm0, %xmm0
	divsd	%xmm1, %xmm0
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, 8(%rsp)
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movss	8(%rsp), %xmm0
	movq	%rax, %rdi
	cvtss2sd	%xmm0, %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L33
	call	__stack_chk_fail@PLT
.L29:
	movq	24(%rsp), %rdx
	movq	80(%rsp), %rsi
	movl	$2, %ecx
	salq	$2, %r13
	movq	32(%rsp), %rdi
	movl	$10, %r14d
	call	cudaMemcpy@PLT
	leaq	.LC8(%rip), %rsi
	movq	%rbp, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	32(%rsp), %r12
.L30:
	movq	%rbp, %rdi
	cvtss2sd	(%r12), %xmm0
	addq	%r13, %r12
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	decl	%r14d
	jne	.L30
	jmp	.L31
.L34:
	movl	$1, %ebx
	movl	$32, %r12d
.L19:
	movq	%r13, %rdi
	leaq	.LC13(%rip), %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	%r13, %rdi
	leaq	.LC10(%rip), %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebp, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC11(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%r12d, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC12(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	jmp	.L32
.L33:
	addq	$168, %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
.LFE3637:
	.size	main, .-main
	.section	.rodata.str1.1
.LC14:
	.string	"_Z6kernelPfS_S_iibi"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3666:
	.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	.LC14(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z6kernelPfS_S_iibi(%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
.LFE3666:
	.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
.LC0:
	.long	1065353216
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC6:
	.long	-1598689907
	.long	1051772663
	.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 <math.h>
#include <cmath>
#include <time.h>
#include <iostream>

#define DTYPE float
__global__ void kernel(float *a, float *x, float* buff,int Xblocks,int size,bool comp,int toreduce)
{
	int i=threadIdx.x+blockIdx.x*blockDim.x;
   	int j=threadIdx.y+blockIdx.y*blockDim.y;
   	__shared__ float xSM[1024];
   	if(comp==true){
		xSM[threadIdx.x+threadIdx.y*blockDim.x]=a[i+j*size]*x[i];
		__syncthreads();
	}
	else{
		xSM[threadIdx.x+threadIdx.y*blockDim.x]=buff[i+j*Xblocks];
	}   	
   	if(i<toreduce){	
        	for(int k=blockDim.x/2;k>0;k/=2){
            		if(threadIdx.x<k&&i+k<toreduce){
                		xSM[threadIdx.x+threadIdx.y*blockDim.x]+=xSM[threadIdx.x+k+threadIdx.y*blockDim.x];
                		__syncthreads();
            		}
        	}
        }
       if(threadIdx.x==0){
		}
       	buff[blockIdx.x+j*Xblocks]=xSM[threadIdx.y*blockDim.x];
}

void fillA(DTYPE *a, int size)
{
   for (int i=0;i<size*size;i++)
      a[i]=1.0;
}

//X mit Werten füllen 
void fillX(DTYPE *x, int size)
{
   for (int i=0;i<size;i++){
	x[i]=1;
//      x[i]= (DTYPE)(i+1);
	}
}
int main(int argc, char**argv)
{   
   int sx=32;
   int sy=32;
   int i=1;
   bool standard=true;
   if (argc>1)   
   {
	standard=false;
	sx=atoi(argv[1]);
      	if (argc>2){ 
      		sy=atoi(argv[2]);
      		if(argc>3){ 
      			i=atoi(argv[3]);      					
      		}
      	}	
   }
   if(standard){
   	std::cout<<"Do experiment with standard settings"<<std::endl;
   	std::cout<<"Sx="<<sx<<"\nSy="<<sy<<"\nSize=1024*"<<i<<std::endl;
   }else 
   {
   	std::cout<<"Do experiment with individual settings"<<std::endl;
   	std::cout<<"Sx="<<sx<<"\nSy="<<sy<<"\nSize=1024*"<<i<<std::endl;
       }
   /*if(sx*sy!=t){
   	std::cout<<"Sx*Sy has to be equal to threads per block"<<std::endl;
   	return -1;
   }*/
   int size=1024*i;
   int xblocks=size/sx;
   //Datenfelder anlegen für Host
   DTYPE *a_host, *buff_host,*x_host;
   //und Device
   DTYPE *a_dev, *buff_dev,*x_dev;
   //Events für die Zeitmessung
   hipEvent_t start,end;
   //Zeiten: 
   //htd: Host->Device Memcpy von A und x
   //dth: Device->Host Memcpy von y
   //kernelA, kernelAT
   float kernelA_time=0.0;

   //TODO: Host Speicher anlegen und A und x füllen
   a_host = (DTYPE*)malloc(size*size*sizeof(DTYPE));
   x_host = (DTYPE*)malloc(size*sizeof(DTYPE));
   buff_host=(DTYPE*)malloc(xblocks*size*sizeof(DTYPE));

   fillA(a_host,size);
   fillX(x_host,size);
   //TODO: CUDA Events erstellen

   //TODO: CUDA Speicher anlegen für alle Arrays (a_dev,x_dev,y_dev)
   hipMalloc((void**)&a_dev,size*size*sizeof(DTYPE));
   hipMalloc((void**)&x_dev,size*sizeof(DTYPE));
   hipMalloc((void**)&buff_dev,xblocks*size*sizeof(DTYPE));

	hipMemcpy(a_dev,a_host,size*size*sizeof(DTYPE),hipMemcpyHostToDevice);
	hipMemcpy(x_dev,x_host,size*sizeof(DTYPE),hipMemcpyHostToDevice);
   
   dim3 block(sx,sy);
   dim3 grid(size/block.x,size/block.y);
   
   int toreduce=size;//Anzahl der noch zu reduzierenden Wert, im ersten Kernel=size, da noch nichts reduziert wurde und man somit im ersten Berechnungsschritt size viel Skalare erhällt
   bool doComputation=true;

	//cache Konfiguration
	if(argc>4){
		if(atoi(argv[4])==1){//L1 Prefered
			std::cout<<"16 kB shared, 48kB L1"<<std::endl;
			hipFuncSetCacheConfig(reinterpret_cast<const void*>(kernel), hipFuncCachePreferL1);
		} else if(atoi(argv[4])==2){
			std::cout<<"48kB shared, 16kb L1"<<std::endl;
			hipFuncSetCacheConfig(reinterpret_cast<const void*>(kernel), hipFuncCachePreferShared);
		}else{
			std::cout<<"32kB shared, 32kB L1"<<std::endl;
			hipFuncSetCacheConfig(reinterpret_cast<const void*>(kernel), hipFuncCachePreferNone);
		}
	}
   //Start Zeitmessung, da dies nur der erste Aufruf ist und danach der Kernel nochmal aufgerufen wird, mit einem anderen Argument als du Computation, endet die Zeitmessung erst nach der while Schleife, inder der Kernel noch einmal aufgerufen wird.
   hipEventCreate(&start);
   hipEventCreate(&end);
   hipEventRecord(start,0);
   kernel<<<grid,block>>>(a_dev,x_dev,buff_dev,xblocks,size,doComputation,toreduce);//,y_dev,size);

	//Überprüfung des Zwischenergebnisses
	/*cudaMemcpy(buff_host,buff_dev,xblocks*size*sizeof(DTYPE),cudaMemcpyDeviceToHost);
   for(int lj=0;lj<10;lj++){
   	for(int li=0;li<10;li++)
   		std::cout<<buff_host[li+xblocks*lj]<<" ; ";
   	std::cout<<std::endl;
   }*/
   //Vorbereitung für weitere Durchführung
   doComputation=false; //Matrix Vektor Multiplikation musst nur einmal durchgeführt werden
   if(toreduce%sx==0)
   	toreduce/=sx;
   else
   	toreduce=toreduce/sx+1;
   while(toreduce>1){//toreduce=1 bedeut, dass auf einen Wert reduziert wurde und somit keine weitere Reduzierung notwendig ist
   	int xgrid;
   	int ygrid=size/sy;
   	if(toreduce%sx==0)
   		xgrid=toreduce/sx;
   	else
	   	xgrid=toreduce/sx+1;
   	dim3 itgrid(xgrid,ygrid);
   	kernel<<<itgrid,block>>>(a_dev,x_dev,buff_dev,xblocks,size,doComputation,toreduce);//,y_dev,size);
   	if(toreduce%sx==0)
   	toreduce/=sx;
	else
   	toreduce=toreduce/sx+1;	
   }
   hipEventRecord(end,0);
   hipEventSynchronize(end);
   hipEventElapsedTime(&kernelA_time,start,end);
   if(argc>4){
   	hipMemcpy(buff_host,buff_dev,xblocks*size*sizeof(DTYPE),hipMemcpyDeviceToHost);
   	std::cout<<"Das Endergebnis ist"<<std::endl;
   	for(int lj=0;lj<10;lj++){
   		//for(int li=0;li<10;li++)
   		//	std::cout<<buff_host[li+xblocks*lj]<<" ; ";
   //   		std::cout<<std::endl;
   		std::cout<<buff_host[xblocks*lj]<<std::endl;
   	}
   }
   std::cout<<"Computation time: "<<kernelA_time<<std::endl;
   //Bestimmung der Performance
   //float time_s=kernelA_time/1000;
   //float gflops=pow(10,-9)*size*size*2/time_s;
   float gflops=pow(10,-6)*size*size*2/kernelA_time;
   std::cout<<"Computation Performance in GFLOPs: "<<gflops<<std::endl;
   return 0;
}


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

                                        # End of file scope inline assembly
	.globl	_Z21__device_stub__kernelPfS_S_iibi # -- Begin function _Z21__device_stub__kernelPfS_S_iibi
	.type	_Z21__device_stub__kernelPfS_S_iibi,@function
_Z21__device_stub__kernelPfS_S_iibi:    # @_Z21__device_stub__kernelPfS_S_iibi
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$160, %rsp
	.cfi_def_cfa_offset 208
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	56(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	48(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	20(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	16(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	15(%rsp), %r8
	movb	%r9b, (%r8)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	208(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	32(%rsp), %r12
	leaq	24(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z6kernelPfS_S_iibi, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$176, %rsp
	.cfi_adjust_cfa_offset -176
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z21__device_stub__kernelPfS_S_iibi, .Lfunc_end0-_Z21__device_stub__kernelPfS_S_iibi
	.cfi_endproc
                                        # -- End function
	.globl	_Z5fillAPfi                     # -- Begin function _Z5fillAPfi
	.type	_Z5fillAPfi,@function
_Z5fillAPfi:                            # @_Z5fillAPfi
	.cfi_startproc
# %bb.0:
                                        # kill: def $esi killed $esi def $rsi
	testl	%esi, %esi
	je	.LBB1_3
# %bb.1:                                # %.lr.ph.preheader
	imull	%esi, %esi
	cmpl	$1, %esi
	adcl	$0, %esi
	xorl	%eax, %eax
.LBB1_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	$1065353216, (%rdi,%rax,4)      # imm = 0x3F800000
	incq	%rax
	cmpq	%rax, %rsi
	jne	.LBB1_2
.LBB1_3:                                # %._crit_edge
	retq
.Lfunc_end1:
	.size	_Z5fillAPfi, .Lfunc_end1-_Z5fillAPfi
	.cfi_endproc
                                        # -- End function
	.globl	_Z5fillXPfi                     # -- Begin function _Z5fillXPfi
	.type	_Z5fillXPfi,@function
_Z5fillXPfi:                            # @_Z5fillXPfi
	.cfi_startproc
# %bb.0:
	testl	%esi, %esi
	jle	.LBB2_3
# %bb.1:                                # %.lr.ph.preheader
	movl	%esi, %eax
	xorl	%ecx, %ecx
.LBB2_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	$1065353216, (%rdi,%rcx,4)      # imm = 0x3F800000
	incq	%rcx
	cmpq	%rcx, %rax
	jne	.LBB2_2
.LBB2_3:                                # %._crit_edge
	retq
.Lfunc_end2:
	.size	_Z5fillXPfi, .Lfunc_end2-_Z5fillXPfi
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI3_0:
	.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	$120, %rsp
	.cfi_def_cfa_offset 176
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	cmpl	$2, %edi
	movl	%edi, 28(%rsp)                  # 4-byte Spill
	movq	%rsi, 88(%rsp)                  # 8-byte Spill
	jge	.LBB3_1
# %bb.29:
	movl	$_ZSt4cout, %ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str, %esi
	movl	$36, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.1, %esi
	movl	$3, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$32, %ebp
	movl	$_ZSt4cout, %edi
	movl	$32, %esi
	callq	_ZNSolsEi
	movq	%rax, %rbx
	movl	$.L.str.2, %esi
	movl	$4, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%rbx, %rdi
	movl	$32, %esi
	callq	_ZNSolsEi
	movq	%rax, %rbx
	movl	$.L.str.3, %esi
	movl	$11, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%rbx, %rdi
	movl	$1, %esi
	callq	_ZNSolsEi
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$1024, %r15d                    # imm = 0x400
	movl	$32, %eax
	movq	%rax, 72(%rsp)                  # 8-byte Spill
	jmp	.LBB3_5
.LBB3_1:
	movq	%rsi, %r12
	movl	%edi, %ebx
	movq	8(%rsi), %rdi
	callq	atoi
	movl	%eax, %ebp
	movl	$32, %r14d
	movl	$1, %r15d
	cmpl	$2, %ebx
	je	.LBB3_4
# %bb.2:
	movq	16(%r12), %rdi
	callq	atoi
	movl	%eax, %r14d
	cmpl	$4, %ebx
	jb	.LBB3_4
# %bb.3:
	movq	24(%r12), %rdi
	callq	atoi
	movl	%eax, %r15d
.LBB3_4:
	movl	$_ZSt4cout, %ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.4, %esi
	movl	$38, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.1, %esi
	movl	$3, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movl	%ebp, %esi
	callq	_ZNSolsEi
	movq	%rax, %rbx
	movl	$.L.str.2, %esi
	movl	$4, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%rbx, %rdi
	movq	%r14, 72(%rsp)                  # 8-byte Spill
	movl	%r14d, %esi
	callq	_ZNSolsEi
	movq	%rax, %rbx
	movl	$.L.str.3, %esi
	movl	$11, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%rbx, %rdi
	movl	%r15d, %esi
	callq	_ZNSolsEi
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	shll	$10, %r15d
.LBB3_5:
	movl	%r15d, %eax
	cltd
	movl	%ebp, 20(%rsp)                  # 4-byte Spill
	idivl	%ebp
	movl	%eax, %ebp
	movl	%edx, 68(%rsp)                  # 4-byte Spill
	movl	$0, 24(%rsp)
	movl	%r15d, %r14d
	imull	%r14d, %r14d
	leaq	(,%r14,4), %rdi
	movq	%rdi, 104(%rsp)                 # 8-byte Spill
	callq	malloc
	movq	%rax, %r13
	movslq	%r15d, %r12
	leaq	(,%r12,4), %rdi
	movq	%rdi, 96(%rsp)                  # 8-byte Spill
	callq	malloc
	movq	%rax, %rbx
	movl	%ebp, 16(%rsp)                  # 4-byte Spill
	movl	%r15d, 12(%rsp)                 # 4-byte Spill
	imull	%r15d, %ebp
	movslq	%ebp, %rbp
	shlq	$2, %rbp
	movq	%rbp, %rdi
	callq	malloc
	movq	%rax, %r15
	testl	%r12d, %r12d
	je	.LBB3_11
# %bb.6:                                # %.lr.ph.preheader.i
	cmpl	$1, %r14d
	adcl	$0, %r14d
	xorl	%eax, %eax
.LBB3_7:                                # %.lr.ph.i
                                        # =>This Inner Loop Header: Depth=1
	movl	$1065353216, (%r13,%rax,4)      # imm = 0x3F800000
	incq	%rax
	cmpq	%rax, %r14
	jne	.LBB3_7
# %bb.8:                                # %_Z5fillAPfi.exit
	cmpl	$0, 12(%rsp)                    # 4-byte Folded Reload
	jle	.LBB3_11
# %bb.9:                                # %.lr.ph.preheader.i109
	movl	12(%rsp), %eax                  # 4-byte Reload
	xorl	%ecx, %ecx
.LBB3_10:                               # %.lr.ph.i111
                                        # =>This Inner Loop Header: Depth=1
	movl	$1065353216, (%rbx,%rcx,4)      # imm = 0x3F800000
	incq	%rcx
	cmpq	%rcx, %rax
	jne	.LBB3_10
.LBB3_11:                               # %_Z5fillXPfi.exit
	leaq	56(%rsp), %rdi
	movq	104(%rsp), %r12                 # 8-byte Reload
	movq	%r12, %rsi
	callq	hipMalloc
	leaq	40(%rsp), %rdi
	movq	96(%rsp), %r14                  # 8-byte Reload
	movq	%r14, %rsi
	callq	hipMalloc
	leaq	48(%rsp), %rdi
	movq	%rbp, 80(%rsp)                  # 8-byte Spill
	movq	%rbp, %rsi
	callq	hipMalloc
	leaq	56(%rsp), %rax
	movq	(%rax), %rdi
	movq	%r13, %rsi
	movq	%r12, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	40(%rsp), %rax
	movq	(%rax), %rdi
	movq	%rbx, %rsi
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movl	20(%rsp), %r14d                 # 4-byte Reload
	movl	%r14d, %eax
	movq	72(%rsp), %r13                  # 8-byte Reload
	movq	%r13, %r12
	shlq	$32, %r12
	orq	%rax, %r12
	movl	12(%rsp), %esi                  # 4-byte Reload
	movl	%esi, %eax
	xorl	%edx, %edx
	divl	%r14d
	movl	%eax, %ecx
	movl	%esi, %eax
	xorl	%edx, %edx
	divl	%r13d
	movl	%eax, %ebp
	shlq	$32, %rbp
	orq	%rcx, %rbp
	cmpl	$5, 28(%rsp)                    # 4-byte Folded Reload
	jl	.LBB3_18
# %bb.12:
	movq	88(%rsp), %rax                  # 8-byte Reload
	movq	32(%rax), %rdi
	callq	atoi
	movl	$_ZSt4cout, %edi
	cmpl	$2, %eax
	je	.LBB3_15
# %bb.13:
	cmpl	$1, %eax
	jne	.LBB3_16
# %bb.14:
	movl	$.L.str.5, %esi
	movl	$21, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$2, %ebx
	jmp	.LBB3_17
.LBB3_15:
	movl	$.L.str.6, %esi
	movl	$20, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$1, %ebx
	jmp	.LBB3_17
.LBB3_16:
	movl	$.L.str.7, %esi
	movl	$20, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	xorl	%ebx, %ebx
.LBB3_17:                               # %.sink.split
	movq	_ZSt4cout(%rip), %rax
	movl	$_ZSt4cout, %ecx
	movq	-24(%rax), %rdi
	addq	%rcx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_Z6kernelPfS_S_iibi, %edi
	movl	%ebx, %esi
	callq	hipFuncSetCacheConfig
.LBB3_18:
	leaq	112(%rsp), %rbx
	movq	%rbx, %rdi
	callq	hipEventCreate
	leaq	32(%rsp), %rdi
	callq	hipEventCreate
	movq	(%rbx), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	%rbp, %rdi
	movl	$1, %esi
	movq	%r12, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_20
# %bb.19:
	movq	56(%rsp), %rdi
	movq	40(%rsp), %rsi
	movq	48(%rsp), %rdx
	movl	12(%rsp), %r8d                  # 4-byte Reload
	movl	%r8d, (%rsp)
	movl	16(%rsp), %ecx                  # 4-byte Reload
	movl	$1, %r9d
	callq	_Z21__device_stub__kernelPfS_S_iibi
.LBB3_20:
	cmpl	$1, 68(%rsp)                    # 4-byte Folded Reload
	movl	16(%rsp), %ebp                  # 4-byte Reload
	sbbl	$-1, %ebp
	cmpl	$2, %ebp
	jl	.LBB3_25
# %bb.21:                               # %.lr.ph
	movl	12(%rsp), %eax                  # 4-byte Reload
	cltd
	idivl	%r13d
	movl	%eax, %ebx
	shlq	$32, %rbx
.LBB3_22:                               # =>This Inner Loop Header: Depth=1
	movl	%ebp, %eax
	cltd
	idivl	%r14d
	movl	%eax, %r13d
	movl	%edx, %r14d
	cmpl	$1, %edx
	movl	%eax, %edi
	sbbl	$-1, %edi
	orq	%rbx, %rdi
	movl	$1, %esi
	movq	%r12, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_24
# %bb.23:                               #   in Loop: Header=BB3_22 Depth=1
	movq	56(%rsp), %rdi
	movq	40(%rsp), %rsi
	movq	48(%rsp), %rdx
	movl	%ebp, (%rsp)
	movl	16(%rsp), %ecx                  # 4-byte Reload
	movl	12(%rsp), %r8d                  # 4-byte Reload
	xorl	%r9d, %r9d
	callq	_Z21__device_stub__kernelPfS_S_iibi
.LBB3_24:                               #   in Loop: Header=BB3_22 Depth=1
	cmpl	$1, %r14d
	sbbl	$-1, %r13d
	movl	%r13d, %ebp
	cmpl	$1, %r13d
	movl	20(%rsp), %r14d                 # 4-byte Reload
	jg	.LBB3_22
.LBB3_25:                               # %._crit_edge
	movq	32(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	32(%rsp), %rdi
	callq	hipEventSynchronize
	movq	112(%rsp), %rsi
	movq	32(%rsp), %rdx
	leaq	24(%rsp), %rdi
	callq	hipEventElapsedTime
	cmpl	$5, 28(%rsp)                    # 4-byte Folded Reload
	jl	.LBB3_28
# %bb.26:
	movq	48(%rsp), %rsi
	movq	%r15, %rdi
	movq	80(%rsp), %rdx                  # 8-byte Reload
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$_ZSt4cout, %ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.8, %esi
	movl	$19, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movslq	16(%rsp), %r14                  # 4-byte Folded Reload
	shlq	$2, %r14
	movl	$10, %r12d
.LBB3_27:                               # =>This Inner Loop Header: Depth=1
	xorps	%xmm0, %xmm0
	cvtss2sd	(%r15), %xmm0
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	addq	%r14, %r15
	decq	%r12
	jne	.LBB3_27
.LBB3_28:                               # %.loopexit
	movl	$_ZSt4cout, %edi
	movl	$.L.str.9, %esi
	movl	$18, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	xorps	%xmm0, %xmm0
	cvtss2sd	24(%rsp), %xmm0
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	xorps	%xmm0, %xmm0
	cvtsi2sdl	12(%rsp), %xmm0         # 4-byte Folded Reload
	movsd	.LCPI3_0(%rip), %xmm1           # xmm1 = mem[0],zero
	mulsd	%xmm0, %xmm1
	mulsd	%xmm0, %xmm1
	xorps	%xmm0, %xmm0
	cvtss2sd	24(%rsp), %xmm0
	addsd	%xmm1, %xmm1
	divsd	%xmm0, %xmm1
	xorps	%xmm0, %xmm0
	cvtsd2ss	%xmm1, %xmm0
	movss	%xmm0, 20(%rsp)                 # 4-byte Spill
	movl	$_ZSt4cout, %edi
	movl	$.L.str.10, %esi
	movl	$35, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	xorps	%xmm0, %xmm0
	cvtss2sd	20(%rsp), %xmm0                 # 4-byte Folded Reload
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	xorl	%eax, %eax
	addq	$120, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end3:
	.size	main, .Lfunc_end3-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB4_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB4_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z6kernelPfS_S_iibi, %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_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	_Z6kernelPfS_S_iibi,@object     # @_Z6kernelPfS_S_iibi
	.section	.rodata,"a",@progbits
	.globl	_Z6kernelPfS_S_iibi
	.p2align	3, 0x0
_Z6kernelPfS_S_iibi:
	.quad	_Z21__device_stub__kernelPfS_S_iibi
	.size	_Z6kernelPfS_S_iibi, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Do experiment with standard settings"
	.size	.L.str, 37

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"Sx="
	.size	.L.str.1, 4

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"\nSy="
	.size	.L.str.2, 5

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"\nSize=1024*"
	.size	.L.str.3, 12

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"Do experiment with individual settings"
	.size	.L.str.4, 39

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"16 kB shared, 48kB L1"
	.size	.L.str.5, 22

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"48kB shared, 16kb L1"
	.size	.L.str.6, 21

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"32kB shared, 32kB L1"
	.size	.L.str.7, 21

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"Das Endergebnis ist"
	.size	.L.str.8, 20

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"Computation time: "
	.size	.L.str.9, 19

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"Computation Performance in GFLOPs: "
	.size	.L.str.10, 36

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z6kernelPfS_S_iibi"
	.size	.L__unnamed_1, 20

	.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 _Z21__device_stub__kernelPfS_S_iibi
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z6kernelPfS_S_iibi
	.addrsig_sym _ZSt4cout
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z6kernelPfS_S_iibi     ; -- Begin function _Z6kernelPfS_S_iibi
	.globl	_Z6kernelPfS_S_iibi
	.p2align	8
	.type	_Z6kernelPfS_S_iibi,@function
_Z6kernelPfS_S_iibi:                    ; @_Z6kernelPfS_S_iibi
; %bb.0:
	s_clause 0x2
	s_load_b32 s8, s[0:1], 0x34
	s_load_b128 s[4:7], s[0:1], 0x18
	s_load_b64 s[2:3], s[0:1], 0x10
	v_bfe_u32 v6, v0, 10, 10
	v_and_b32_e32 v1, 0x3ff, v0
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s9, s8, 16
	s_and_b32 s8, s8, 0xffff
	s_and_b32 s6, s6, 1
	v_mul_u32_u24_e32 v0, s8, v6
	v_mad_u64_u32 v[4:5], null, s14, s8, v[1:2]
	v_mad_u64_u32 v[2:3], null, s15, s9, v[6:7]
	s_delay_alu instid0(VALU_DEP_3)
	v_add_lshl_u32 v3, v0, v1, 2
	s_cmp_eq_u32 s6, 0
	s_mov_b32 s6, -1
	s_cbranch_scc0 .LBB0_2
; %bb.1:
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[5:6], null, v2, s4, v[4:5]
	s_mov_b32 s6, 0
	v_ashrrev_i32_e32 v6, 31, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[5:6], 2, v[5:6]
	v_add_co_u32 v5, vcc_lo, s2, v5
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v6, vcc_lo, s3, v6, vcc_lo
	global_load_b32 v5, v[5:6], off
	s_waitcnt vmcnt(0)
	ds_store_b32 v3, v5
.LBB0_2:                                ; %Flow59
	s_and_not1_b32 vcc_lo, exec_lo, s6
	s_cbranch_vccnz .LBB0_4
; %bb.3:
	s_load_b128 s[16:19], s[0:1], 0x0
	v_mad_u64_u32 v[6:7], null, v2, s5, v[4:5]
	v_ashrrev_i32_e32 v5, 31, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[8:9], 2, v[4:5]
	v_ashrrev_i32_e32 v7, 31, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[6:7], 2, v[6:7]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v5, vcc_lo, s16, v6
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v6, vcc_lo, s17, v7, vcc_lo
	v_add_co_u32 v7, vcc_lo, s18, v8
	v_add_co_ci_u32_e32 v8, vcc_lo, s19, v9, vcc_lo
	global_load_b32 v5, v[5:6], off
	global_load_b32 v6, v[7:8], off
	s_waitcnt vmcnt(0)
	v_mul_f32_e32 v5, v5, v6
	ds_store_b32 v3, v5
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
.LBB0_4:
	v_cmp_gt_i32_e32 vcc_lo, s7, v4
	s_cmp_gt_u32 s8, 1
	s_cselect_b32 s0, -1, 0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s0, vcc_lo, s0
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_8
.LBB0_5:                                ; =>This Inner Loop Header: Depth=1
	s_lshr_b32 s5, s8, 1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v5, s5, v4
	v_cmp_gt_u32_e32 vcc_lo, s5, v1
	v_cmp_gt_i32_e64 s0, s7, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s6, vcc_lo, s0
	s_and_saveexec_b32 s0, s6
	s_cbranch_execz .LBB0_7
; %bb.6:                                ;   in Loop: Header=BB0_5 Depth=1
	v_lshl_add_u32 v5, s5, 2, v3
	ds_load_b32 v5, v5
	ds_load_b32 v6, v3
	s_waitcnt lgkmcnt(0)
	v_add_f32_e32 v5, v5, v6
	ds_store_b32 v3, v5
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
.LBB0_7:                                ;   in Loop: Header=BB0_5 Depth=1
	s_or_b32 exec_lo, exec_lo, s0
	s_cmp_gt_u32 s8, 3
	s_mov_b32 s8, s5
	s_cbranch_scc1 .LBB0_5
.LBB0_8:                                ; %Flow58
	s_or_b32 exec_lo, exec_lo, s1
	v_lshlrev_b32_e32 v3, 2, v0
	v_mad_u64_u32 v[0:1], null, v2, s4, s[14:15]
	v_mov_b32_e32 v1, 0
	ds_load_b32 v2, v3
	v_lshlrev_b64 v[0:1], 2, v[0:1]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s2, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s3, v1, vcc_lo
	s_waitcnt lgkmcnt(0)
	global_store_b32 v[0:1], v2, off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z6kernelPfS_S_iibi
		.amdhsa_group_segment_fixed_size 4096
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 296
		.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 10
		.amdhsa_next_free_sgpr 20
		.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	_Z6kernelPfS_S_iibi, .Lfunc_end0-_Z6kernelPfS_S_iibi
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 528
; NumSgprs: 22
; NumVgprs: 10
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 4096 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 22
; NumVGPRsForWavesPerEU: 10
; 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:           1
        .value_kind:     by_value
      - .offset:         36
        .size:           4
        .value_kind:     by_value
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         44
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         52
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         54
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         56
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         58
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         60
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         62
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         104
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 4096
    .kernarg_segment_align: 8
    .kernarg_segment_size: 296
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z6kernelPfS_S_iibi
    .private_segment_fixed_size: 0
    .sgpr_count:     22
    .sgpr_spill_count: 0
    .symbol:         _Z6kernelPfS_S_iibi.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     10
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
