

/********************************************************************************************

	
  Source Code :  sharedMemoryStridedAccessPatterns.cu

  Objective   :	 This code demonstrates bank conflicts that can occur while accessing 
                 the shared memory

  Descritpion :  Example code to demonstrate that while reading the same word by all the 
                 threads, there will not be any serialization since all threads are accessing 
                 from the same bank the 32-bit word gets broadcasted to all the threads -- 
		 hence bandwidth can be comparable to the value got when there were no 
                 bank conflicts

  Description :  Example code to demonstrate the different strided access patterns in the 
                 shared memory the corresponding advantages in terms of the bandwidth of the 
                 shared memory that is achievable. 

                 The following strided accesses are demonstrated: 

		    stride of one 32-bit word      : resulting in no bank conflicts
	            stride of two 32-bit words     : causes 2-way bank conflicts 
	            stride of three 32-bit words   : causes no bank conflicts
		    stride of eight 32-bit words   : causes 8-way bank conflicts
	     	    stride of sixteen 32-bit words : causes 16-way bank conflicts

  Input       :  none

  output      :	 The different bandwidths that are achieved for the different strided accesses 
                  of the shared memory

 Modified     :  Aug 2011
                 

 Author	     :   RarchK		

****************************************************************************************/

#include <cuda.h>
#include <stdio.h>
#include <float.h>

#define BLOCK_SIZE 128   // 128 threads per block
#define TRANSFERED_DATA_SIZE 2000000   //2 MB

#define NO_OF_PATTERNS 5 
#define NTIMES 10

#define MIN(x,y) ((x)<(y)?(x):(y))
#define MAX(x,y) ((x)>(y)?(x):(y))

void printResults(void);


////////////////////////////////////////////////////////////////////////////////////////////////
//
// access pattern involving accessing the elements with a stride of one 32-bit word
// the access pattern causes no bank conflicts
//
/////////////////////////////////////////////////////////////////////////////////////////////////

__global__ void sharedMemAccessWithStride1(void)
{
	__shared__ __attribute__((unused)) float array[BLOCK_SIZE];
	int idx = threadIdx.x;
	float scalar = 5.0f;
	array[idx] = scalar;
}

////////////////////////////////////////////////////////////////////////////////////////////////


////////////////////////////////////////////////////////////////////////////////////////////////
//
// access pattern involving accessing the elements with a stride of three 32-bit words 
// the access pattern causes no bank conflicts
//
/////////////////////////////////////////////////////////////////////////////////////////////////

__global__ void sharedMemAccessWithStride3(void)
{
	__shared__ __attribute__((unused)) float array[BLOCK_SIZE];
	int idx = (threadIdx.x * 3) % BLOCK_SIZE;
	float scalar = 5.0f;
	array[idx] = scalar;
}

////////////////////////////////////////////////////////////////////////////////////////////////


////////////////////////////////////////////////////////////////////////////////////////////////
//
// access pattern involving accessing the elements with a stride of two 32-bit words 
// the access pattern causes 2-way bank conflicts. 
//
/////////////////////////////////////////////////////////////////////////////////////////////////

__global__ void sharedMemAccessWithStride2(void)
{
	__shared__ __attribute__((unused)) float array[BLOCK_SIZE];
	int idx = (threadIdx.x * 2) % BLOCK_SIZE;
	float scalar = 5.0f;
	array[idx] = scalar;
}

////////////////////////////////////////////////////////////////////////////////////////////////

////////////////////////////////////////////////////////////////////////////////////////////////
//
// access pattern involving accessing the elements with a stride of eight 32-bit words 
// the access pattern causes 8-way bank conflicts. 
//
/////////////////////////////////////////////////////////////////////////////////////////////////

__global__ void sharedMemAccessWithStride8(void)
{
	__shared__ __attribute__((unused)) float array[BLOCK_SIZE];
	int idx = (threadIdx.x * 8) % BLOCK_SIZE;
	float scalar = 5.0f;
	array[idx] = scalar;
}

////////////////////////////////////////////////////////////////////////////////////////////////



////////////////////////////////////////////////////////////////////////////////////////////////
//
// access pattern involving accessing the elements with a stride of sixteen 32-bit words 
// the access pattern causes 16-way bank conflicts. 
//
/////////////////////////////////////////////////////////////////////////////////////////////////

__global__ void sharedMemAccessWithStride16(void)
{
	__shared__ __attribute__((unused)) float array[BLOCK_SIZE];
	int idx = (threadIdx.x * 16) % BLOCK_SIZE;
	float scalar = 5.0f;
	array[idx] = scalar;
}

///////////////////////////////////////////////////////////////////////////////////////////////





static double  avgtime[NO_OF_PATTERNS] = {0}, maxtime[NO_OF_PATTERNS] = {0}, mintime[NO_OF_PATTERNS];
static float bandWidths[NO_OF_PATTERNS] = {0}; 

////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
//
// the main routene 
// for timing the different strided access patterns 
// finding the band widths
//
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////

int main(int argc, char* argv[])
{
	float elapsedTimes[NO_OF_PATTERNS][NTIMES];
	cudaEvent_t start,stop;
	
	double bytes = sizeof(float) * TRANSFERED_DATA_SIZE;
	
	// event creation, which will be used for timing the code
	cudaEventCreate(&start);
	cudaEventCreate(&stop);
	
	//finding the 1D grid size 
	int gridSize = TRANSFERED_DATA_SIZE/BLOCK_SIZE;
	
	if( TRANSFERED_DATA_SIZE % BLOCK_SIZE != 0 ) 
			gridSize += 1;
	
	// running each pattern NTIMES
	for(int k=0; k < NTIMES; k++)
	{
		// timing the kernels corresponding to different access patterns
			
			// PATTERN 1
			cudaEventRecord(start,0);
			sharedMemAccessWithStride1 <<< gridSize,BLOCK_SIZE >>> ();
			cudaEventRecord(stop,0);
			cudaEventSynchronize(stop);
			cudaEventElapsedTime(&elapsedTimes[0][k],start,stop);

			// PATTERN 2
			cudaEventRecord(start,0);
			sharedMemAccessWithStride3 <<< gridSize,BLOCK_SIZE >>> ();
			cudaEventRecord(stop,0);
			cudaEventSynchronize(stop);
			cudaEventElapsedTime(&elapsedTimes[1][k],start,stop);
		
			// PATTERN 3
			cudaEventRecord(start,0);
			sharedMemAccessWithStride2 <<< gridSize,BLOCK_SIZE >>> ();
			cudaEventRecord(stop,0);
			cudaEventSynchronize(stop);
			cudaEventElapsedTime(&elapsedTimes[2][k],start,stop);
			
			// PATTERN 4
			cudaEventRecord(start,0);
			sharedMemAccessWithStride8 <<< gridSize,BLOCK_SIZE >>> ();
			cudaEventRecord(stop,0);
			cudaEventSynchronize(stop);
			cudaEventElapsedTime(&elapsedTimes[3][k],start,stop);

			//PATTERN 5
			cudaEventRecord(start,0);
			sharedMemAccessWithStride16 <<< gridSize,BLOCK_SIZE >>> ();
			cudaEventRecord(stop,0);
			cudaEventSynchronize(stop);
			cudaEventElapsedTime(&elapsedTimes[4][k],start,stop);


	} // end of the for loop involving NTIMES


	// intializing the mintime array
	for(int i=0; i < NO_OF_PATTERNS;i++)
		mintime[i] = FLT_MAX;


   for (int k=1; k < NTIMES; k++) // skiping the first iteration
	{
	   for (int i=0; i < NO_OF_PATTERNS; i++)
	   {
	       avgtime[i] = avgtime[i] + elapsedTimes[i][k];
	       mintime[i] = MIN(mintime[i],elapsedTimes[i][k]);
	       maxtime[i] = MAX(maxtime[i], elapsedTimes[i][k]);
	   }
	}


	// calculation of the different band widths that are achieved by different access patterns  
	for(int i=0; i < NO_OF_PATTERNS; i++)
	{
		avgtime[i] = avgtime[i]/(double)(NTIMES-1); // finding the average time
		bandWidths[i] = bytes/mintime[i];
	}

   printResults();

	printf("\n\n**** successful termination of the program  ****\n\n");

	cudaEventDestroy(start);
	cudaEventDestroy(stop);

	return 0;
}

////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////



////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
//
// prints the results containig the minimum, maximum, average times that the different strided access pattern have taken
// the associated maximum bandwidth of the different patterns
//
///////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////

void printResults()
{

	printf("Demonstrating the different strided accesses of the shared memory\n");
		
   printf("The transfered data size (Bytes): %ld\n",TRANSFERED_DATA_SIZE*sizeof(float));

   printf("\n-------------------------------------------------------------------------------------------------------------------------------\n");
   printf("Pattern \t\t\t Bandwidth (GB/sec) \t Avg time (ms)  \t Min time (ms)  \t Max time (ms)\n");
   printf("-------------------------------------------------------------------------------------------------------------------------------\n");
	// printing the results for different access patterns
	for(int i=0; i < NO_OF_PATTERNS; i++)
	{
		
		switch(i)
		{
			case 0: printf("Stride of one 32-bit word     ");
					  break;

			case 1: printf("Stride of three 32-bit words  ");
					  break;
			
			case 2: printf("Stride of two 32-bit words    ");
					  break;

			case 3: printf("Stride of eight 32-bit words  ");
					  break;

			case 4: printf("Stride of sixteen 32-bit words");
					  break;

		}
		
		printf("\t %.6f \t\t %f \t\t  %f \t\t  %f\n",bandWidths[i]/1000000,avgtime[i],mintime[i],maxtime[i]);
  }
	
   printf("\n ------------------------------------------------------------------------------------------------------------------------------\n");

}

///////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////



	.file	"tmpxft_00204a53_00000000-6_sharedMemoryStridedAccessPatterns.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2031:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2031:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"Demonstrating the different strided accesses of the shared memory\n"
.LC1:
	.string	"The transfered data size (Bytes): %ld\n"
.LC2:
	.string	"\n-------------------------------------------------------------------------------------------------------------------------------\n"
.LC3:
	.string	"Pattern \t\t\t Bandwidth (GB/sec) \t Avg time (ms)  \t Min time (ms)  \t Max time (ms)\n"
.LC4:
	.string	"-------------------------------------------------------------------------------------------------------------------------------\n"
.LC5:
	.string	"Stride of one 32-bit word     "
.LC6:
	.string	"Stride of three 32-bit words  "
.LC7:
	.string	"Stride of two 32-bit words    "
.LC8:
	.string	"Stride of eight 32-bit words  "
.LC9:
	.string	"Stride of sixteen 32-bit words"
.LC11:
	.string	"\t %.6f \t\t %f \t\t  %f \t\t  %f\n"
.LC12:
	.string	"\n ------------------------------------------------------------------------------------------------------------------------------\n"
	.text
	.globl	_Z12printResultsv
	.type	_Z12printResultsv, @function
_Z12printResultsv:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	leaq	.LC0(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	leaq	_ZL7maxtime(%rip), %rbp
	xorl	%ebx, %ebx
	pushq	%rcx
	.cfi_def_cfa_offset 32
	call	__printf_chk@PLT
	movl	$8000000, %edx
	leaq	.LC1(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	leaq	.LC2(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC3(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L8:
	leaq	.LC8(%rip), %rsi
	cmpq	$3, %rbx
	je	.L11
	leaq	.LC9(%rip), %rsi
	cmpl	$4, %ebx
	je	.L11
	leaq	.LC6(%rip), %rsi
	cmpq	$1, %rbx
	je	.L11
	leaq	.LC5(%rip), %rsi
	cmpq	$2, %rbx
	jne	.L11
	leaq	.LC7(%rip), %rsi
.L11:
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movsd	0(%rbp,%rbx,8), %xmm3
	movl	$2, %edi
	leaq	_ZL10bandWidths(%rip), %rax
	movss	(%rax,%rbx,4), %xmm0
	leaq	_ZL7mintime(%rip), %rax
	leaq	.LC11(%rip), %rsi
	divss	.LC10(%rip), %xmm0
	movsd	(%rax,%rbx,8), %xmm2
	leaq	_ZL7avgtime(%rip), %rax
	cvtss2sd	%xmm0, %xmm0
	movsd	(%rax,%rbx,8), %xmm1
	movb	$4, %al
	incq	%rbx
	call	__printf_chk@PLT
	cmpq	$5, %rbx
	jne	.L8
	popq	%rdx
	.cfi_def_cfa_offset 24
	leaq	.LC12(%rip), %rsi
	popq	%rbx
	.cfi_def_cfa_offset 16
	movl	$2, %edi
	xorl	%eax, %eax
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	__printf_chk@PLT
	.cfi_endproc
.LFE2028:
	.size	_Z12printResultsv, .-_Z12printResultsv
	.globl	_Z45__device_stub__Z26sharedMemAccessWithStride1vv
	.type	_Z45__device_stub__Z26sharedMemAccessWithStride1vv, @function
_Z45__device_stub__Z26sharedMemAccessWithStride1vv:
.LFB2053:
	.cfi_startproc
	endbr64
	subq	$88, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	movabsq	$4294967297, %rax
	leaq	8(%rsp), %rcx
	movq	%rsp, %rdx
	leaq	28(%rsp), %rsi
	leaq	16(%rsp), %rdi
	movl	$1, 24(%rsp)
	movl	$1, 36(%rsp)
	movq	%rax, 16(%rsp)
	movq	%rax, 28(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L13
	pushq	8(%rsp)
	.cfi_def_cfa_offset 104
	leaq	_Z26sharedMemAccessWithStride1v(%rip), %rdi
	pushq	8(%rsp)
	.cfi_def_cfa_offset 112
	movq	44(%rsp), %rcx
	movl	52(%rsp), %r8d
	movq	32(%rsp), %rsi
	movl	40(%rsp), %edx
	leaq	80(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 104
	popq	%rdx
	.cfi_def_cfa_offset 96
.L13:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L15
	call	__stack_chk_fail@PLT
.L15:
	addq	$88, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2053:
	.size	_Z45__device_stub__Z26sharedMemAccessWithStride1vv, .-_Z45__device_stub__Z26sharedMemAccessWithStride1vv
	.globl	_Z26sharedMemAccessWithStride1v
	.type	_Z26sharedMemAccessWithStride1v, @function
_Z26sharedMemAccessWithStride1v:
.LFB2054:
	.cfi_startproc
	endbr64
	jmp	_Z45__device_stub__Z26sharedMemAccessWithStride1vv
	.cfi_endproc
.LFE2054:
	.size	_Z26sharedMemAccessWithStride1v, .-_Z26sharedMemAccessWithStride1v
	.globl	_Z45__device_stub__Z26sharedMemAccessWithStride3vv
	.type	_Z45__device_stub__Z26sharedMemAccessWithStride3vv, @function
_Z45__device_stub__Z26sharedMemAccessWithStride3vv:
.LFB2055:
	.cfi_startproc
	endbr64
	subq	$88, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	movabsq	$4294967297, %rax
	leaq	8(%rsp), %rcx
	movq	%rsp, %rdx
	leaq	28(%rsp), %rsi
	leaq	16(%rsp), %rdi
	movl	$1, 24(%rsp)
	movl	$1, 36(%rsp)
	movq	%rax, 16(%rsp)
	movq	%rax, 28(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L18
	pushq	8(%rsp)
	.cfi_def_cfa_offset 104
	leaq	_Z26sharedMemAccessWithStride3v(%rip), %rdi
	pushq	8(%rsp)
	.cfi_def_cfa_offset 112
	movq	44(%rsp), %rcx
	movl	52(%rsp), %r8d
	movq	32(%rsp), %rsi
	movl	40(%rsp), %edx
	leaq	80(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 104
	popq	%rdx
	.cfi_def_cfa_offset 96
.L18:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L20
	call	__stack_chk_fail@PLT
.L20:
	addq	$88, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2055:
	.size	_Z45__device_stub__Z26sharedMemAccessWithStride3vv, .-_Z45__device_stub__Z26sharedMemAccessWithStride3vv
	.globl	_Z26sharedMemAccessWithStride3v
	.type	_Z26sharedMemAccessWithStride3v, @function
_Z26sharedMemAccessWithStride3v:
.LFB2056:
	.cfi_startproc
	endbr64
	jmp	_Z45__device_stub__Z26sharedMemAccessWithStride3vv
	.cfi_endproc
.LFE2056:
	.size	_Z26sharedMemAccessWithStride3v, .-_Z26sharedMemAccessWithStride3v
	.globl	_Z45__device_stub__Z26sharedMemAccessWithStride2vv
	.type	_Z45__device_stub__Z26sharedMemAccessWithStride2vv, @function
_Z45__device_stub__Z26sharedMemAccessWithStride2vv:
.LFB2057:
	.cfi_startproc
	endbr64
	subq	$88, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	movabsq	$4294967297, %rax
	leaq	8(%rsp), %rcx
	movq	%rsp, %rdx
	leaq	28(%rsp), %rsi
	leaq	16(%rsp), %rdi
	movl	$1, 24(%rsp)
	movl	$1, 36(%rsp)
	movq	%rax, 16(%rsp)
	movq	%rax, 28(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L23
	pushq	8(%rsp)
	.cfi_def_cfa_offset 104
	leaq	_Z26sharedMemAccessWithStride2v(%rip), %rdi
	pushq	8(%rsp)
	.cfi_def_cfa_offset 112
	movq	44(%rsp), %rcx
	movl	52(%rsp), %r8d
	movq	32(%rsp), %rsi
	movl	40(%rsp), %edx
	leaq	80(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 104
	popq	%rdx
	.cfi_def_cfa_offset 96
.L23:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L25
	call	__stack_chk_fail@PLT
.L25:
	addq	$88, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2057:
	.size	_Z45__device_stub__Z26sharedMemAccessWithStride2vv, .-_Z45__device_stub__Z26sharedMemAccessWithStride2vv
	.globl	_Z26sharedMemAccessWithStride2v
	.type	_Z26sharedMemAccessWithStride2v, @function
_Z26sharedMemAccessWithStride2v:
.LFB2058:
	.cfi_startproc
	endbr64
	jmp	_Z45__device_stub__Z26sharedMemAccessWithStride2vv
	.cfi_endproc
.LFE2058:
	.size	_Z26sharedMemAccessWithStride2v, .-_Z26sharedMemAccessWithStride2v
	.globl	_Z45__device_stub__Z26sharedMemAccessWithStride8vv
	.type	_Z45__device_stub__Z26sharedMemAccessWithStride8vv, @function
_Z45__device_stub__Z26sharedMemAccessWithStride8vv:
.LFB2059:
	.cfi_startproc
	endbr64
	subq	$88, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	movabsq	$4294967297, %rax
	leaq	8(%rsp), %rcx
	movq	%rsp, %rdx
	leaq	28(%rsp), %rsi
	leaq	16(%rsp), %rdi
	movl	$1, 24(%rsp)
	movl	$1, 36(%rsp)
	movq	%rax, 16(%rsp)
	movq	%rax, 28(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L28
	pushq	8(%rsp)
	.cfi_def_cfa_offset 104
	leaq	_Z26sharedMemAccessWithStride8v(%rip), %rdi
	pushq	8(%rsp)
	.cfi_def_cfa_offset 112
	movq	44(%rsp), %rcx
	movl	52(%rsp), %r8d
	movq	32(%rsp), %rsi
	movl	40(%rsp), %edx
	leaq	80(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 104
	popq	%rdx
	.cfi_def_cfa_offset 96
.L28:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L30
	call	__stack_chk_fail@PLT
.L30:
	addq	$88, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2059:
	.size	_Z45__device_stub__Z26sharedMemAccessWithStride8vv, .-_Z45__device_stub__Z26sharedMemAccessWithStride8vv
	.globl	_Z26sharedMemAccessWithStride8v
	.type	_Z26sharedMemAccessWithStride8v, @function
_Z26sharedMemAccessWithStride8v:
.LFB2060:
	.cfi_startproc
	endbr64
	jmp	_Z45__device_stub__Z26sharedMemAccessWithStride8vv
	.cfi_endproc
.LFE2060:
	.size	_Z26sharedMemAccessWithStride8v, .-_Z26sharedMemAccessWithStride8v
	.globl	_Z46__device_stub__Z27sharedMemAccessWithStride16vv
	.type	_Z46__device_stub__Z27sharedMemAccessWithStride16vv, @function
_Z46__device_stub__Z27sharedMemAccessWithStride16vv:
.LFB2061:
	.cfi_startproc
	endbr64
	subq	$88, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	movabsq	$4294967297, %rax
	leaq	8(%rsp), %rcx
	movq	%rsp, %rdx
	leaq	28(%rsp), %rsi
	leaq	16(%rsp), %rdi
	movl	$1, 24(%rsp)
	movl	$1, 36(%rsp)
	movq	%rax, 16(%rsp)
	movq	%rax, 28(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L33
	pushq	8(%rsp)
	.cfi_def_cfa_offset 104
	leaq	_Z27sharedMemAccessWithStride16v(%rip), %rdi
	pushq	8(%rsp)
	.cfi_def_cfa_offset 112
	movq	44(%rsp), %rcx
	movl	52(%rsp), %r8d
	movq	32(%rsp), %rsi
	movl	40(%rsp), %edx
	leaq	80(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 104
	popq	%rdx
	.cfi_def_cfa_offset 96
.L33:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L35
	call	__stack_chk_fail@PLT
.L35:
	addq	$88, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2061:
	.size	_Z46__device_stub__Z27sharedMemAccessWithStride16vv, .-_Z46__device_stub__Z27sharedMemAccessWithStride16vv
	.globl	_Z27sharedMemAccessWithStride16v
	.type	_Z27sharedMemAccessWithStride16v, @function
_Z27sharedMemAccessWithStride16v:
.LFB2062:
	.cfi_startproc
	endbr64
	jmp	_Z46__device_stub__Z27sharedMemAccessWithStride16vv
	.cfi_endproc
.LFE2062:
	.size	_Z27sharedMemAccessWithStride16v, .-_Z27sharedMemAccessWithStride16v
	.section	.rodata.str1.1
.LC16:
	.string	"\n\n**** successful termination of the program  ****\n\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movabsq	$4294982921, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movl	$33554433, %ebx
	salq	$7, %rbx
	subq	$264, %rsp
	.cfi_def_cfa_offset 304
	movq	%fs:40, %rax
	movq	%rax, 248(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rdi
	leaq	48(%rsp), %r12
	call	cudaEventCreate@PLT
	leaq	16(%rsp), %rdi
	leaq	88(%rsp), %r13
	call	cudaEventCreate@PLT
.L44:
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%rbx, %rdx
	movl	$1, %ecx
	movq	%rbp, %rdi
	movl	$1, %esi
	movq	%rbx, 36(%rsp)
	movl	$1, 44(%rsp)
	movq	%rbp, 24(%rsp)
	movl	$1, 32(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L39
	call	_Z45__device_stub__Z26sharedMemAccessWithStride1vv
.L39:
	movq	16(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	16(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	16(%rsp), %rdx
	movq	8(%rsp), %rsi
	movq	%r12, %rdi
	call	cudaEventElapsedTime@PLT
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%rbx, %rdx
	movl	$1, %ecx
	movq	%rbp, %rdi
	movl	$1, %esi
	movq	%rbx, 36(%rsp)
	movl	$1, 44(%rsp)
	movq	%rbp, 24(%rsp)
	movl	$1, 32(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L40
	call	_Z45__device_stub__Z26sharedMemAccessWithStride3vv
.L40:
	movq	16(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	16(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	16(%rsp), %rdx
	movq	8(%rsp), %rsi
	leaq	40(%r12), %rdi
	call	cudaEventElapsedTime@PLT
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%rbx, %rdx
	movl	$1, %ecx
	movq	%rbp, %rdi
	movl	$1, %esi
	movq	%rbx, 36(%rsp)
	movl	$1, 44(%rsp)
	movq	%rbp, 24(%rsp)
	movl	$1, 32(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	call	_Z45__device_stub__Z26sharedMemAccessWithStride2vv
.L41:
	movq	16(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	16(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	16(%rsp), %rdx
	movq	8(%rsp), %rsi
	leaq	80(%r12), %rdi
	call	cudaEventElapsedTime@PLT
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%rbx, %rdx
	movl	$1, %ecx
	movq	%rbp, %rdi
	movl	$1, %esi
	movq	%rbx, 36(%rsp)
	movl	$1, 44(%rsp)
	movq	%rbp, 24(%rsp)
	movl	$1, 32(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L42
	call	_Z45__device_stub__Z26sharedMemAccessWithStride8vv
.L42:
	movq	16(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	16(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	16(%rsp), %rdx
	movq	8(%rsp), %rsi
	leaq	120(%r12), %rdi
	call	cudaEventElapsedTime@PLT
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%rbx, %rdx
	movl	$1, %ecx
	movq	%rbp, %rdi
	movl	$1, %esi
	movq	%rbx, 36(%rsp)
	movl	$1, 44(%rsp)
	movq	%rbp, 24(%rsp)
	movl	$1, 32(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L43
	call	_Z46__device_stub__Z27sharedMemAccessWithStride16vv
.L43:
	movq	16(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	16(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	16(%rsp), %rdx
	movq	8(%rsp), %rsi
	leaq	160(%r12), %rdi
	addq	$4, %r12
	call	cudaEventElapsedTime@PLT
	cmpq	%r13, %r12
	jne	.L44
	movsd	.LC13(%rip), %xmm0
	leaq	52(%rsp), %rcx
	leaq	_ZL7avgtime(%rip), %rdx
	leaq	_ZL7mintime(%rip), %rsi
	leaq	_ZL7maxtime(%rip), %r8
	movsd	%xmm0, _ZL7mintime(%rip)
	movsd	%xmm0, 8+_ZL7mintime(%rip)
	movsd	%xmm0, 16+_ZL7mintime(%rip)
	movsd	%xmm0, 24+_ZL7mintime(%rip)
	movsd	%xmm0, 32+_ZL7mintime(%rip)
.L45:
	movq	%rcx, %rdi
	xorl	%eax, %eax
.L50:
	movsd	(%rdx,%rax), %xmm1
	cvtss2sd	(%rdi), %xmm0
	addq	$40, %rdi
	addsd	%xmm0, %xmm1
	movsd	%xmm1, (%rdx,%rax)
	movsd	(%rsi,%rax), %xmm1
	minsd	%xmm0, %xmm1
	movsd	%xmm1, (%rsi,%rax)
	movsd	(%r8,%rax), %xmm1
	maxsd	%xmm0, %xmm1
	movsd	%xmm1, (%r8,%rax)
	addq	$8, %rax
	cmpq	$40, %rax
	jne	.L50
	addq	$4, %rcx
	cmpq	%r13, %rcx
	jne	.L45
	movsd	.LC14(%rip), %xmm2
	xorl	%eax, %eax
	movsd	.LC15(%rip), %xmm1
	leaq	_ZL10bandWidths(%rip), %rcx
.L51:
	movsd	(%rdx,%rax,8), %xmm0
	divsd	%xmm2, %xmm0
	movsd	%xmm0, (%rdx,%rax,8)
	movaps	%xmm1, %xmm0
	divsd	(%rsi,%rax,8), %xmm0
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, (%rcx,%rax,4)
	incq	%rax
	cmpq	$5, %rax
	jne	.L51
	call	_Z12printResultsv
	leaq	.LC16(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	8(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movq	16(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movq	248(%rsp), %rax
	subq	%fs:40, %rax
	je	.L52
	call	__stack_chk_fail@PLT
.L52:
	addq	$264, %rsp
	.cfi_def_cfa_offset 40
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	main, .-main
	.section	.rodata.str1.1
.LC17:
	.string	"_Z27sharedMemAccessWithStride16v"
.LC18:
	.string	"_Z26sharedMemAccessWithStride8v"
.LC19:
	.string	"_Z26sharedMemAccessWithStride2v"
.LC20:
	.string	"_Z26sharedMemAccessWithStride3v"
.LC21:
	.string	"_Z26sharedMemAccessWithStride1v"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2064:
	.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	_Z27sharedMemAccessWithStride16v(%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	_Z26sharedMemAccessWithStride8v(%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	_Z26sharedMemAccessWithStride2v(%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	.LC20(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z26sharedMemAccessWithStride3v(%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	.LC21(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z26sharedMemAccessWithStride1v(%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
.LFE2064:
	.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
	.local	_ZL10bandWidths
	.comm	_ZL10bandWidths,20,16
	.local	_ZL7mintime
	.comm	_ZL7mintime,40,32
	.local	_ZL7maxtime
	.comm	_ZL7maxtime,40,32
	.local	_ZL7avgtime
	.comm	_ZL7avgtime,40,32
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC10:
	.long	1232348160
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC13:
	.long	-536870912
	.long	1206910975
	.align 8
.LC14:
	.long	0
	.long	1075970048
	.align 8
.LC15:
	.long	0
	.long	1096713344
	.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:



/********************************************************************************************

	
  Source Code :  sharedMemoryStridedAccessPatterns.cu

  Objective   :	 This code demonstrates bank conflicts that can occur while accessing 
                 the shared memory

  Descritpion :  Example code to demonstrate that while reading the same word by all the 
                 threads, there will not be any serialization since all threads are accessing 
                 from the same bank the 32-bit word gets broadcasted to all the threads -- 
		 hence bandwidth can be comparable to the value got when there were no 
                 bank conflicts

  Description :  Example code to demonstrate the different strided access patterns in the 
                 shared memory the corresponding advantages in terms of the bandwidth of the 
                 shared memory that is achievable. 

                 The following strided accesses are demonstrated: 

		    stride of one 32-bit word      : resulting in no bank conflicts
	            stride of two 32-bit words     : causes 2-way bank conflicts 
	            stride of three 32-bit words   : causes no bank conflicts
		    stride of eight 32-bit words   : causes 8-way bank conflicts
	     	    stride of sixteen 32-bit words : causes 16-way bank conflicts

  Input       :  none

  output      :	 The different bandwidths that are achieved for the different strided accesses 
                  of the shared memory

 Modified     :  Aug 2011
                 

 Author	     :   RarchK		

****************************************************************************************/

#include <hip/hip_runtime.h>
#include <stdio.h>
#include <float.h>

#define BLOCK_SIZE 128   // 128 threads per block
#define TRANSFERED_DATA_SIZE 2000000   //2 MB

#define NO_OF_PATTERNS 5 
#define NTIMES 10

#define MIN(x,y) ((x)<(y)?(x):(y))
#define MAX(x,y) ((x)>(y)?(x):(y))

void printResults(void);


////////////////////////////////////////////////////////////////////////////////////////////////
//
// access pattern involving accessing the elements with a stride of one 32-bit word
// the access pattern causes no bank conflicts
//
/////////////////////////////////////////////////////////////////////////////////////////////////

__global__ void sharedMemAccessWithStride1(void)
{
	__shared__ __attribute__((unused)) float array[BLOCK_SIZE];
	int idx = threadIdx.x;
	float scalar = 5.0f;
	array[idx] = scalar;
}

////////////////////////////////////////////////////////////////////////////////////////////////


////////////////////////////////////////////////////////////////////////////////////////////////
//
// access pattern involving accessing the elements with a stride of three 32-bit words 
// the access pattern causes no bank conflicts
//
/////////////////////////////////////////////////////////////////////////////////////////////////

__global__ void sharedMemAccessWithStride3(void)
{
	__shared__ __attribute__((unused)) float array[BLOCK_SIZE];
	int idx = (threadIdx.x * 3) % BLOCK_SIZE;
	float scalar = 5.0f;
	array[idx] = scalar;
}

////////////////////////////////////////////////////////////////////////////////////////////////


////////////////////////////////////////////////////////////////////////////////////////////////
//
// access pattern involving accessing the elements with a stride of two 32-bit words 
// the access pattern causes 2-way bank conflicts. 
//
/////////////////////////////////////////////////////////////////////////////////////////////////

__global__ void sharedMemAccessWithStride2(void)
{
	__shared__ __attribute__((unused)) float array[BLOCK_SIZE];
	int idx = (threadIdx.x * 2) % BLOCK_SIZE;
	float scalar = 5.0f;
	array[idx] = scalar;
}

////////////////////////////////////////////////////////////////////////////////////////////////

////////////////////////////////////////////////////////////////////////////////////////////////
//
// access pattern involving accessing the elements with a stride of eight 32-bit words 
// the access pattern causes 8-way bank conflicts. 
//
/////////////////////////////////////////////////////////////////////////////////////////////////

__global__ void sharedMemAccessWithStride8(void)
{
	__shared__ __attribute__((unused)) float array[BLOCK_SIZE];
	int idx = (threadIdx.x * 8) % BLOCK_SIZE;
	float scalar = 5.0f;
	array[idx] = scalar;
}

////////////////////////////////////////////////////////////////////////////////////////////////



////////////////////////////////////////////////////////////////////////////////////////////////
//
// access pattern involving accessing the elements with a stride of sixteen 32-bit words 
// the access pattern causes 16-way bank conflicts. 
//
/////////////////////////////////////////////////////////////////////////////////////////////////

__global__ void sharedMemAccessWithStride16(void)
{
	__shared__ __attribute__((unused)) float array[BLOCK_SIZE];
	int idx = (threadIdx.x * 16) % BLOCK_SIZE;
	float scalar = 5.0f;
	array[idx] = scalar;
}

///////////////////////////////////////////////////////////////////////////////////////////////





static double  avgtime[NO_OF_PATTERNS] = {0}, maxtime[NO_OF_PATTERNS] = {0}, mintime[NO_OF_PATTERNS];
static float bandWidths[NO_OF_PATTERNS] = {0}; 

////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
//
// the main routene 
// for timing the different strided access patterns 
// finding the band widths
//
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////

int main(int argc, char* argv[])
{
	float elapsedTimes[NO_OF_PATTERNS][NTIMES];
	hipEvent_t start,stop;
	
	double bytes = sizeof(float) * TRANSFERED_DATA_SIZE;
	
	// event creation, which will be used for timing the code
	hipEventCreate(&start);
	hipEventCreate(&stop);
	
	//finding the 1D grid size 
	int gridSize = TRANSFERED_DATA_SIZE/BLOCK_SIZE;
	
	if( TRANSFERED_DATA_SIZE % BLOCK_SIZE != 0 ) 
			gridSize += 1;
	
	// running each pattern NTIMES
	for(int k=0; k < NTIMES; k++)
	{
		// timing the kernels corresponding to different access patterns
			
			// PATTERN 1
			hipEventRecord(start,0);
			sharedMemAccessWithStride1 <<< gridSize,BLOCK_SIZE >>> ();
			hipEventRecord(stop,0);
			hipEventSynchronize(stop);
			hipEventElapsedTime(&elapsedTimes[0][k],start,stop);

			// PATTERN 2
			hipEventRecord(start,0);
			sharedMemAccessWithStride3 <<< gridSize,BLOCK_SIZE >>> ();
			hipEventRecord(stop,0);
			hipEventSynchronize(stop);
			hipEventElapsedTime(&elapsedTimes[1][k],start,stop);
		
			// PATTERN 3
			hipEventRecord(start,0);
			sharedMemAccessWithStride2 <<< gridSize,BLOCK_SIZE >>> ();
			hipEventRecord(stop,0);
			hipEventSynchronize(stop);
			hipEventElapsedTime(&elapsedTimes[2][k],start,stop);
			
			// PATTERN 4
			hipEventRecord(start,0);
			sharedMemAccessWithStride8 <<< gridSize,BLOCK_SIZE >>> ();
			hipEventRecord(stop,0);
			hipEventSynchronize(stop);
			hipEventElapsedTime(&elapsedTimes[3][k],start,stop);

			//PATTERN 5
			hipEventRecord(start,0);
			sharedMemAccessWithStride16 <<< gridSize,BLOCK_SIZE >>> ();
			hipEventRecord(stop,0);
			hipEventSynchronize(stop);
			hipEventElapsedTime(&elapsedTimes[4][k],start,stop);


	} // end of the for loop involving NTIMES


	// intializing the mintime array
	for(int i=0; i < NO_OF_PATTERNS;i++)
		mintime[i] = FLT_MAX;


   for (int k=1; k < NTIMES; k++) // skiping the first iteration
	{
	   for (int i=0; i < NO_OF_PATTERNS; i++)
	   {
	       avgtime[i] = avgtime[i] + elapsedTimes[i][k];
	       mintime[i] = MIN(mintime[i],elapsedTimes[i][k]);
	       maxtime[i] = MAX(maxtime[i], elapsedTimes[i][k]);
	   }
	}


	// calculation of the different band widths that are achieved by different access patterns  
	for(int i=0; i < NO_OF_PATTERNS; i++)
	{
		avgtime[i] = avgtime[i]/(double)(NTIMES-1); // finding the average time
		bandWidths[i] = bytes/mintime[i];
	}

   printResults();

	printf("\n\n**** successful termination of the program  ****\n\n");

	hipEventDestroy(start);
	hipEventDestroy(stop);

	return 0;
}

////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////



////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
//
// prints the results containig the minimum, maximum, average times that the different strided access pattern have taken
// the associated maximum bandwidth of the different patterns
//
///////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////

void printResults()
{

	printf("Demonstrating the different strided accesses of the shared memory\n");
		
   printf("The transfered data size (Bytes): %ld\n",TRANSFERED_DATA_SIZE*sizeof(float));

   printf("\n-------------------------------------------------------------------------------------------------------------------------------\n");
   printf("Pattern \t\t\t Bandwidth (GB/sec) \t Avg time (ms)  \t Min time (ms)  \t Max time (ms)\n");
   printf("-------------------------------------------------------------------------------------------------------------------------------\n");
	// printing the results for different access patterns
	for(int i=0; i < NO_OF_PATTERNS; i++)
	{
		
		switch(i)
		{
			case 0: printf("Stride of one 32-bit word     ");
					  break;

			case 1: printf("Stride of three 32-bit words  ");
					  break;
			
			case 2: printf("Stride of two 32-bit words    ");
					  break;

			case 3: printf("Stride of eight 32-bit words  ");
					  break;

			case 4: printf("Stride of sixteen 32-bit words");
					  break;

		}
		
		printf("\t %.6f \t\t %f \t\t  %f \t\t  %f\n",bandWidths[i]/1000000,avgtime[i],mintime[i],maxtime[i]);
  }
	
   printf("\n ------------------------------------------------------------------------------------------------------------------------------\n");

}

///////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////



	.text
	.file	"sharedMemoryStridedAccessPatterns.hip"
	.globl	_Z41__device_stub__sharedMemAccessWithStride1v # -- Begin function _Z41__device_stub__sharedMemAccessWithStride1v
	.type	_Z41__device_stub__sharedMemAccessWithStride1v,@function
_Z41__device_stub__sharedMemAccessWithStride1v: # @_Z41__device_stub__sharedMemAccessWithStride1v
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	subq	$56, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rbx
	leaq	24(%rsp), %r14
	leaq	16(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%rbx, %rdi
	movq	%r14, %rsi
	movq	%r15, %rdx
	movq	%r12, %rcx
	callq	__hipPopCallConfiguration
	movq	(%rbx), %rsi
	movl	8(%rbx), %edx
	movq	(%r14), %rcx
	movl	8(%r14), %r8d
	movq	%rsp, %r9
	movl	$_Z26sharedMemAccessWithStride1v, %edi
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	pushq	(%r15)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$72, %rsp
	.cfi_adjust_cfa_offset -72
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z41__device_stub__sharedMemAccessWithStride1v, .Lfunc_end0-_Z41__device_stub__sharedMemAccessWithStride1v
	.cfi_endproc
                                        # -- End function
	.globl	_Z41__device_stub__sharedMemAccessWithStride3v # -- Begin function _Z41__device_stub__sharedMemAccessWithStride3v
	.type	_Z41__device_stub__sharedMemAccessWithStride3v,@function
_Z41__device_stub__sharedMemAccessWithStride3v: # @_Z41__device_stub__sharedMemAccessWithStride3v
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	subq	$56, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rbx
	leaq	24(%rsp), %r14
	leaq	16(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%rbx, %rdi
	movq	%r14, %rsi
	movq	%r15, %rdx
	movq	%r12, %rcx
	callq	__hipPopCallConfiguration
	movq	(%rbx), %rsi
	movl	8(%rbx), %edx
	movq	(%r14), %rcx
	movl	8(%r14), %r8d
	movq	%rsp, %r9
	movl	$_Z26sharedMemAccessWithStride3v, %edi
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	pushq	(%r15)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$72, %rsp
	.cfi_adjust_cfa_offset -72
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z41__device_stub__sharedMemAccessWithStride3v, .Lfunc_end1-_Z41__device_stub__sharedMemAccessWithStride3v
	.cfi_endproc
                                        # -- End function
	.globl	_Z41__device_stub__sharedMemAccessWithStride2v # -- Begin function _Z41__device_stub__sharedMemAccessWithStride2v
	.type	_Z41__device_stub__sharedMemAccessWithStride2v,@function
_Z41__device_stub__sharedMemAccessWithStride2v: # @_Z41__device_stub__sharedMemAccessWithStride2v
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	subq	$56, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rbx
	leaq	24(%rsp), %r14
	leaq	16(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%rbx, %rdi
	movq	%r14, %rsi
	movq	%r15, %rdx
	movq	%r12, %rcx
	callq	__hipPopCallConfiguration
	movq	(%rbx), %rsi
	movl	8(%rbx), %edx
	movq	(%r14), %rcx
	movl	8(%r14), %r8d
	movq	%rsp, %r9
	movl	$_Z26sharedMemAccessWithStride2v, %edi
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	pushq	(%r15)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$72, %rsp
	.cfi_adjust_cfa_offset -72
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	_Z41__device_stub__sharedMemAccessWithStride2v, .Lfunc_end2-_Z41__device_stub__sharedMemAccessWithStride2v
	.cfi_endproc
                                        # -- End function
	.globl	_Z41__device_stub__sharedMemAccessWithStride8v # -- Begin function _Z41__device_stub__sharedMemAccessWithStride8v
	.type	_Z41__device_stub__sharedMemAccessWithStride8v,@function
_Z41__device_stub__sharedMemAccessWithStride8v: # @_Z41__device_stub__sharedMemAccessWithStride8v
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	subq	$56, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rbx
	leaq	24(%rsp), %r14
	leaq	16(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%rbx, %rdi
	movq	%r14, %rsi
	movq	%r15, %rdx
	movq	%r12, %rcx
	callq	__hipPopCallConfiguration
	movq	(%rbx), %rsi
	movl	8(%rbx), %edx
	movq	(%r14), %rcx
	movl	8(%r14), %r8d
	movq	%rsp, %r9
	movl	$_Z26sharedMemAccessWithStride8v, %edi
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	pushq	(%r15)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$72, %rsp
	.cfi_adjust_cfa_offset -72
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end3:
	.size	_Z41__device_stub__sharedMemAccessWithStride8v, .Lfunc_end3-_Z41__device_stub__sharedMemAccessWithStride8v
	.cfi_endproc
                                        # -- End function
	.globl	_Z42__device_stub__sharedMemAccessWithStride16v # -- Begin function _Z42__device_stub__sharedMemAccessWithStride16v
	.type	_Z42__device_stub__sharedMemAccessWithStride16v,@function
_Z42__device_stub__sharedMemAccessWithStride16v: # @_Z42__device_stub__sharedMemAccessWithStride16v
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	subq	$56, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rbx
	leaq	24(%rsp), %r14
	leaq	16(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%rbx, %rdi
	movq	%r14, %rsi
	movq	%r15, %rdx
	movq	%r12, %rcx
	callq	__hipPopCallConfiguration
	movq	(%rbx), %rsi
	movl	8(%rbx), %edx
	movq	(%r14), %rcx
	movl	8(%r14), %r8d
	movq	%rsp, %r9
	movl	$_Z27sharedMemAccessWithStride16v, %edi
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	pushq	(%r15)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$72, %rsp
	.cfi_adjust_cfa_offset -72
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end4:
	.size	_Z42__device_stub__sharedMemAccessWithStride16v, .Lfunc_end4-_Z42__device_stub__sharedMemAccessWithStride16v
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI5_0:
	.quad	0x4022000000000000              # double 9
.LCPI5_1:
	.quad	0x415e848000000000              # double 8.0E+6
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	subq	$224, %rsp
	.cfi_def_cfa_offset 256
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movabsq	$4294967424, %rbx               # imm = 0x100000080
	leaq	8(%rsp), %rdi
	callq	hipEventCreate
	movq	%rsp, %rdi
	callq	hipEventCreate
	xorl	%r15d, %r15d
	leaq	15497(%rbx), %r14
.LBB5_1:                                # =>This Inner Loop Header: Depth=1
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_3
# %bb.2:                                #   in Loop: Header=BB5_1 Depth=1
	callq	_Z41__device_stub__sharedMemAccessWithStride1v
.LBB5_3:                                #   in Loop: Header=BB5_1 Depth=1
	movq	(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	(%rsp), %rdi
	callq	hipEventSynchronize
	leaq	(%rsp,%r15), %rdi
	addq	$16, %rdi
	movq	8(%rsp), %rsi
	movq	(%rsp), %rdx
	callq	hipEventElapsedTime
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_5
# %bb.4:                                #   in Loop: Header=BB5_1 Depth=1
	callq	_Z41__device_stub__sharedMemAccessWithStride3v
.LBB5_5:                                #   in Loop: Header=BB5_1 Depth=1
	movq	(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	(%rsp), %rdi
	callq	hipEventSynchronize
	leaq	(%rsp,%r15), %rdi
	addq	$56, %rdi
	movq	8(%rsp), %rsi
	movq	(%rsp), %rdx
	callq	hipEventElapsedTime
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_7
# %bb.6:                                #   in Loop: Header=BB5_1 Depth=1
	callq	_Z41__device_stub__sharedMemAccessWithStride2v
.LBB5_7:                                #   in Loop: Header=BB5_1 Depth=1
	movq	(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	(%rsp), %rdi
	callq	hipEventSynchronize
	leaq	(%rsp,%r15), %rdi
	addq	$96, %rdi
	movq	8(%rsp), %rsi
	movq	(%rsp), %rdx
	callq	hipEventElapsedTime
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_9
# %bb.8:                                #   in Loop: Header=BB5_1 Depth=1
	callq	_Z41__device_stub__sharedMemAccessWithStride8v
.LBB5_9:                                #   in Loop: Header=BB5_1 Depth=1
	movq	(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	(%rsp), %rdi
	callq	hipEventSynchronize
	leaq	(%rsp,%r15), %rdi
	addq	$136, %rdi
	movq	8(%rsp), %rsi
	movq	(%rsp), %rdx
	callq	hipEventElapsedTime
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_11
# %bb.10:                               #   in Loop: Header=BB5_1 Depth=1
	callq	_Z42__device_stub__sharedMemAccessWithStride16v
.LBB5_11:                               #   in Loop: Header=BB5_1 Depth=1
	movq	(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	(%rsp), %rdi
	callq	hipEventSynchronize
	leaq	(%rsp,%r15), %rdi
	addq	$176, %rdi
	movq	8(%rsp), %rsi
	movq	(%rsp), %rdx
	callq	hipEventElapsedTime
	addq	$4, %r15
	cmpq	$40, %r15
	jne	.LBB5_1
# %bb.12:                               # %.preheader103.preheader
	movq	$-40, %rax
	movabsq	$5183643170566569984, %rcx      # imm = 0x47EFFFFFE0000000
.LBB5_13:                               # %.preheader103
                                        # =>This Inner Loop Header: Depth=1
	movq	%rcx, _ZL7mintime+40(%rax)
	addq	$8, %rax
	jne	.LBB5_13
# %bb.14:                               # %.preheader101.preheader
	movl	$1, %eax
	leaq	20(%rsp), %rcx
.LBB5_15:                               # %.preheader101
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB5_16 Depth 2
	movq	$-40, %rdx
	movq	%rcx, %rsi
.LBB5_16:                               #   Parent Loop BB5_15 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	xorps	%xmm0, %xmm0
	cvtss2sd	(%rsi), %xmm0
	movsd	_ZL7avgtime+40(%rdx), %xmm1     # xmm1 = mem[0],zero
	addsd	%xmm0, %xmm1
	movsd	%xmm1, _ZL7avgtime+40(%rdx)
	movsd	_ZL7mintime+40(%rdx), %xmm1     # xmm1 = mem[0],zero
	minsd	%xmm0, %xmm1
	movsd	%xmm1, _ZL7mintime+40(%rdx)
	movsd	_ZL7maxtime+40(%rdx), %xmm1     # xmm1 = mem[0],zero
	maxsd	%xmm0, %xmm1
	movsd	%xmm1, _ZL7maxtime+40(%rdx)
	addq	$40, %rsi
	addq	$8, %rdx
	jne	.LBB5_16
# %bb.17:                               #   in Loop: Header=BB5_15 Depth=1
	incq	%rax
	addq	$4, %rcx
	cmpq	$10, %rax
	jne	.LBB5_15
# %bb.18:                               # %.preheader.preheader
	movq	$-20, %rax
	movsd	.LCPI5_0(%rip), %xmm0           # xmm0 = mem[0],zero
	movsd	.LCPI5_1(%rip), %xmm1           # xmm1 = mem[0],zero
.LBB5_19:                               # %.preheader
                                        # =>This Inner Loop Header: Depth=1
	movsd	_ZL7avgtime+40(%rax,%rax), %xmm2 # xmm2 = mem[0],zero
	divsd	%xmm0, %xmm2
	movsd	%xmm2, _ZL7avgtime+40(%rax,%rax)
	movapd	%xmm1, %xmm2
	divsd	_ZL7mintime+40(%rax,%rax), %xmm2
	cvtsd2ss	%xmm2, %xmm2
	movss	%xmm2, _ZL10bandWidths+20(%rax)
	addq	$4, %rax
	jne	.LBB5_19
# %bb.20:
	callq	_Z12printResultsv
	movl	$.Lstr, %edi
	callq	puts@PLT
	movq	8(%rsp), %rdi
	callq	hipEventDestroy
	movq	(%rsp), %rdi
	callq	hipEventDestroy
	xorl	%eax, %eax
	addq	$224, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end5:
	.size	main, .Lfunc_end5-main
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _Z12printResultsv
.LCPI6_0:
	.long	0x49742400                      # float 1.0E+6
	.text
	.globl	_Z12printResultsv
	.type	_Z12printResultsv,@function
_Z12printResultsv:                      # @_Z12printResultsv
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movl	$.Lstr.1, %edi
	callq	puts@PLT
	movl	$.L.str.2, %edi
	movl	$8000000, %esi                  # imm = 0x7A1200
	xorl	%eax, %eax
	callq	printf
	movl	$.Lstr.2, %edi
	callq	puts@PLT
	movl	$.Lstr.3, %edi
	callq	puts@PLT
	movl	$.Lstr.4, %edi
	callq	puts@PLT
	xorl	%ebx, %ebx
.LBB6_1:                                # =>This Inner Loop Header: Depth=1
	cmpl	$4, %ebx
	ja	.LBB6_3
# %bb.2:                                # %switch.lookup
                                        #   in Loop: Header=BB6_1 Depth=1
	movq	.Lswitch.table._Z12printResultsv(,%rbx,8), %rdi
	xorl	%eax, %eax
	callq	printf
.LBB6_3:                                #   in Loop: Header=BB6_1 Depth=1
	movss	_ZL10bandWidths(,%rbx,4), %xmm0 # xmm0 = mem[0],zero,zero,zero
	divss	.LCPI6_0(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	movsd	_ZL7avgtime(,%rbx,8), %xmm1     # xmm1 = mem[0],zero
	movsd	_ZL7mintime(,%rbx,8), %xmm2     # xmm2 = mem[0],zero
	movsd	_ZL7maxtime(,%rbx,8), %xmm3     # xmm3 = mem[0],zero
	movl	$.L.str.11, %edi
	movb	$4, %al
	callq	printf
	incq	%rbx
	cmpq	$5, %rbx
	jne	.LBB6_1
# %bb.4:
	movl	$.Lstr.5, %edi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	puts@PLT                        # TAILCALL
.Lfunc_end6:
	.size	_Z12printResultsv, .Lfunc_end6-_Z12printResultsv
	.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	.LBB7_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB7_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z26sharedMemAccessWithStride1v, %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	$_Z26sharedMemAccessWithStride3v, %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	$_Z26sharedMemAccessWithStride2v, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z26sharedMemAccessWithStride8v, %esi
	movl	$.L__unnamed_4, %edx
	movl	$.L__unnamed_4, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z27sharedMemAccessWithStride16v, %esi
	movl	$.L__unnamed_5, %edx
	movl	$.L__unnamed_5, %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_end7:
	.size	__hip_module_ctor, .Lfunc_end7-__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	.LBB8_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
.LBB8_2:
	retq
.Lfunc_end8:
	.size	__hip_module_dtor, .Lfunc_end8-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z26sharedMemAccessWithStride1v,@object # @_Z26sharedMemAccessWithStride1v
	.section	.rodata,"a",@progbits
	.globl	_Z26sharedMemAccessWithStride1v
	.p2align	3, 0x0
_Z26sharedMemAccessWithStride1v:
	.quad	_Z41__device_stub__sharedMemAccessWithStride1v
	.size	_Z26sharedMemAccessWithStride1v, 8

	.type	_Z26sharedMemAccessWithStride3v,@object # @_Z26sharedMemAccessWithStride3v
	.globl	_Z26sharedMemAccessWithStride3v
	.p2align	3, 0x0
_Z26sharedMemAccessWithStride3v:
	.quad	_Z41__device_stub__sharedMemAccessWithStride3v
	.size	_Z26sharedMemAccessWithStride3v, 8

	.type	_Z26sharedMemAccessWithStride2v,@object # @_Z26sharedMemAccessWithStride2v
	.globl	_Z26sharedMemAccessWithStride2v
	.p2align	3, 0x0
_Z26sharedMemAccessWithStride2v:
	.quad	_Z41__device_stub__sharedMemAccessWithStride2v
	.size	_Z26sharedMemAccessWithStride2v, 8

	.type	_Z26sharedMemAccessWithStride8v,@object # @_Z26sharedMemAccessWithStride8v
	.globl	_Z26sharedMemAccessWithStride8v
	.p2align	3, 0x0
_Z26sharedMemAccessWithStride8v:
	.quad	_Z41__device_stub__sharedMemAccessWithStride8v
	.size	_Z26sharedMemAccessWithStride8v, 8

	.type	_Z27sharedMemAccessWithStride16v,@object # @_Z27sharedMemAccessWithStride16v
	.globl	_Z27sharedMemAccessWithStride16v
	.p2align	3, 0x0
_Z27sharedMemAccessWithStride16v:
	.quad	_Z42__device_stub__sharedMemAccessWithStride16v
	.size	_Z27sharedMemAccessWithStride16v, 8

	.type	_ZL7mintime,@object             # @_ZL7mintime
	.local	_ZL7mintime
	.comm	_ZL7mintime,40,16
	.type	_ZL7avgtime,@object             # @_ZL7avgtime
	.local	_ZL7avgtime
	.comm	_ZL7avgtime,40,16
	.type	_ZL7maxtime,@object             # @_ZL7maxtime
	.local	_ZL7maxtime
	.comm	_ZL7maxtime,40,16
	.type	_ZL10bandWidths,@object         # @_ZL10bandWidths
	.local	_ZL10bandWidths
	.comm	_ZL10bandWidths,20,16
	.type	.L.str.2,@object                # @.str.2
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.2:
	.asciz	"The transfered data size (Bytes): %ld\n"
	.size	.L.str.2, 39

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"Stride of one 32-bit word     "
	.size	.L.str.6, 31

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"Stride of three 32-bit words  "
	.size	.L.str.7, 31

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"Stride of two 32-bit words    "
	.size	.L.str.8, 31

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"Stride of eight 32-bit words  "
	.size	.L.str.9, 31

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"Stride of sixteen 32-bit words"
	.size	.L.str.10, 31

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"\t %.6f \t\t %f \t\t  %f \t\t  %f\n"
	.size	.L.str.11, 28

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

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z26sharedMemAccessWithStride3v"
	.size	.L__unnamed_2, 32

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"_Z26sharedMemAccessWithStride2v"
	.size	.L__unnamed_3, 32

	.type	.L__unnamed_4,@object           # @3
.L__unnamed_4:
	.asciz	"_Z26sharedMemAccessWithStride8v"
	.size	.L__unnamed_4, 32

	.type	.L__unnamed_5,@object           # @4
.L__unnamed_5:
	.asciz	"_Z27sharedMemAccessWithStride16v"
	.size	.L__unnamed_5, 33

	.type	__hip_fatbin_wrapper,@object    # @__hip_fatbin_wrapper
	.section	.hipFatBinSegment,"a",@progbits
	.p2align	3, 0x0
__hip_fatbin_wrapper:
	.long	1212764230                      # 0x48495046
	.long	1                               # 0x1
	.quad	__hip_fatbin
	.quad	0
	.size	__hip_fatbin_wrapper, 24

	.type	__hip_gpubin_handle,@object     # @__hip_gpubin_handle
	.local	__hip_gpubin_handle
	.comm	__hip_gpubin_handle,8,8
	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	__hip_module_ctor
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.type	.Lstr,@object                   # @str
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr:
	.asciz	"\n\n**** successful termination of the program  ****\n"
	.size	.Lstr, 52

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"Demonstrating the different strided accesses of the shared memory"
	.size	.Lstr.1, 66

	.type	.Lstr.2,@object                 # @str.2
.Lstr.2:
	.asciz	"\n-------------------------------------------------------------------------------------------------------------------------------"
	.size	.Lstr.2, 129

	.type	.Lstr.3,@object                 # @str.3
.Lstr.3:
	.asciz	"Pattern \t\t\t Bandwidth (GB/sec) \t Avg time (ms)  \t Min time (ms)  \t Max time (ms)"
	.size	.Lstr.3, 81

	.type	.Lstr.4,@object                 # @str.4
.Lstr.4:
	.asciz	"-------------------------------------------------------------------------------------------------------------------------------"
	.size	.Lstr.4, 128

	.type	.Lstr.5,@object                 # @str.5
.Lstr.5:
	.asciz	"\n ------------------------------------------------------------------------------------------------------------------------------"
	.size	.Lstr.5, 129

	.type	.Lswitch.table._Z12printResultsv,@object # @switch.table._Z12printResultsv
	.section	.rodata,"a",@progbits
	.p2align	3, 0x0
.Lswitch.table._Z12printResultsv:
	.quad	.L.str.6
	.quad	.L.str.7
	.quad	.L.str.8
	.quad	.L.str.9
	.quad	.L.str.10
	.size	.Lswitch.table._Z12printResultsv, 40

	.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 _Z41__device_stub__sharedMemAccessWithStride1v
	.addrsig_sym _Z41__device_stub__sharedMemAccessWithStride3v
	.addrsig_sym _Z41__device_stub__sharedMemAccessWithStride2v
	.addrsig_sym _Z41__device_stub__sharedMemAccessWithStride8v
	.addrsig_sym _Z42__device_stub__sharedMemAccessWithStride16v
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z26sharedMemAccessWithStride1v
	.addrsig_sym _Z26sharedMemAccessWithStride3v
	.addrsig_sym _Z26sharedMemAccessWithStride2v
	.addrsig_sym _Z26sharedMemAccessWithStride8v
	.addrsig_sym _Z27sharedMemAccessWithStride16v
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z26sharedMemAccessWithStride1v ; -- Begin function _Z26sharedMemAccessWithStride1v
	.globl	_Z26sharedMemAccessWithStride1v
	.p2align	8
	.type	_Z26sharedMemAccessWithStride1v,@function
_Z26sharedMemAccessWithStride1v:        ; @_Z26sharedMemAccessWithStride1v
; %bb.0:
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z26sharedMemAccessWithStride1v
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 0
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 0
		.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 1
		.amdhsa_next_free_sgpr 1
		.amdhsa_reserve_vcc 0
		.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	_Z26sharedMemAccessWithStride1v, .Lfunc_end0-_Z26sharedMemAccessWithStride1v
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 4
; NumSgprs: 0
; NumVgprs: 0
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 0
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 1
; NumVGPRsForWavesPerEU: 1
; 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	_Z26sharedMemAccessWithStride3v ; -- Begin function _Z26sharedMemAccessWithStride3v
	.globl	_Z26sharedMemAccessWithStride3v
	.p2align	8
	.type	_Z26sharedMemAccessWithStride3v,@function
_Z26sharedMemAccessWithStride3v:        ; @_Z26sharedMemAccessWithStride3v
; %bb.0:
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z26sharedMemAccessWithStride3v
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 0
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 0
		.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 1
		.amdhsa_next_free_sgpr 1
		.amdhsa_reserve_vcc 0
		.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	_Z26sharedMemAccessWithStride3v, .Lfunc_end1-_Z26sharedMemAccessWithStride3v
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 4
; NumSgprs: 0
; NumVgprs: 0
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 0
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 1
; NumVGPRsForWavesPerEU: 1
; 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	_Z26sharedMemAccessWithStride2v ; -- Begin function _Z26sharedMemAccessWithStride2v
	.globl	_Z26sharedMemAccessWithStride2v
	.p2align	8
	.type	_Z26sharedMemAccessWithStride2v,@function
_Z26sharedMemAccessWithStride2v:        ; @_Z26sharedMemAccessWithStride2v
; %bb.0:
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z26sharedMemAccessWithStride2v
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 0
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 0
		.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 1
		.amdhsa_next_free_sgpr 1
		.amdhsa_reserve_vcc 0
		.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	_Z26sharedMemAccessWithStride2v, .Lfunc_end2-_Z26sharedMemAccessWithStride2v
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 4
; NumSgprs: 0
; NumVgprs: 0
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 0
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 1
; NumVGPRsForWavesPerEU: 1
; 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	_Z26sharedMemAccessWithStride8v ; -- Begin function _Z26sharedMemAccessWithStride8v
	.globl	_Z26sharedMemAccessWithStride8v
	.p2align	8
	.type	_Z26sharedMemAccessWithStride8v,@function
_Z26sharedMemAccessWithStride8v:        ; @_Z26sharedMemAccessWithStride8v
; %bb.0:
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z26sharedMemAccessWithStride8v
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 0
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 0
		.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 1
		.amdhsa_next_free_sgpr 1
		.amdhsa_reserve_vcc 0
		.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_end3:
	.size	_Z26sharedMemAccessWithStride8v, .Lfunc_end3-_Z26sharedMemAccessWithStride8v
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 4
; NumSgprs: 0
; NumVgprs: 0
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 0
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 1
; NumVGPRsForWavesPerEU: 1
; 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	_Z27sharedMemAccessWithStride16v ; -- Begin function _Z27sharedMemAccessWithStride16v
	.globl	_Z27sharedMemAccessWithStride16v
	.p2align	8
	.type	_Z27sharedMemAccessWithStride16v,@function
_Z27sharedMemAccessWithStride16v:       ; @_Z27sharedMemAccessWithStride16v
; %bb.0:
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z27sharedMemAccessWithStride16v
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 0
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 0
		.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 1
		.amdhsa_next_free_sgpr 1
		.amdhsa_reserve_vcc 0
		.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_end4:
	.size	_Z27sharedMemAccessWithStride16v, .Lfunc_end4-_Z27sharedMemAccessWithStride16v
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 4
; NumSgprs: 0
; NumVgprs: 0
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 0
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 1
; NumVGPRsForWavesPerEU: 1
; 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
	.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:           []
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 4
    .kernarg_segment_size: 0
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z26sharedMemAccessWithStride1v
    .private_segment_fixed_size: 0
    .sgpr_count:     0
    .sgpr_spill_count: 0
    .symbol:         _Z26sharedMemAccessWithStride1v.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     0
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:           []
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 4
    .kernarg_segment_size: 0
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z26sharedMemAccessWithStride3v
    .private_segment_fixed_size: 0
    .sgpr_count:     0
    .sgpr_spill_count: 0
    .symbol:         _Z26sharedMemAccessWithStride3v.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     0
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:           []
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 4
    .kernarg_segment_size: 0
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z26sharedMemAccessWithStride2v
    .private_segment_fixed_size: 0
    .sgpr_count:     0
    .sgpr_spill_count: 0
    .symbol:         _Z26sharedMemAccessWithStride2v.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     0
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:           []
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 4
    .kernarg_segment_size: 0
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z26sharedMemAccessWithStride8v
    .private_segment_fixed_size: 0
    .sgpr_count:     0
    .sgpr_spill_count: 0
    .symbol:         _Z26sharedMemAccessWithStride8v.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     0
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:           []
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 4
    .kernarg_segment_size: 0
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z27sharedMemAccessWithStride16v
    .private_segment_fixed_size: 0
    .sgpr_count:     0
    .sgpr_spill_count: 0
    .symbol:         _Z27sharedMemAccessWithStride16v.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     0
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
