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

	
  Source Code :  sharedMemoryReadingSameWord.cu

  Objective   :  Example code to demonstrate the different access patterns of float3 array 
                 in the global memory the corresponding advantages in terms of the bandwidth 
                 that is achievable

  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

  output:	 The different bandwidths of the shared memory that are achieved while 
                 accessing the same word and the access pattern without any bank conflicts

 Modified     :  Aug 2011
                 

 Author	      :  RarchK		

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

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

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

#define NO_OF_PATTERNS 2 
#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 reading the same word by all the threads  
// the access pattern is conflict-free 
// since the word gets broadcasted to all the threads
//
/////////////////////////////////////////////////////////////////////////////////////////////////

__global__ void sharedMemReadSameWord(void)
{
	__shared__  float array[BLOCK_SIZE];
	int idx = 3;
	float __attribute__((unused)) scalar ;
	scalar = array[idx];
}

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






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 two kernals 
// finding the corresponding 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);
			sharedMemReadSameWord <<< gridSize,BLOCK_SIZE >>> ();
			cudaEventRecord(stop,0);
			cudaEventSynchronize(stop);
			cudaEventElapsedTime(&elapsedTimes[1][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 taken by the two kernels
// the associated maximum bandwidth achieved by the different kernels
//
///////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////

void printResults()
{
	printf("Demonstrating that reading same 32-bit word by all threads will not cause bank conflicts\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("Reading the same 32-bit word");
					  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_0033d7be_00000000-6_sharedMemoryReadingSameWord.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 that reading same 32-bit word by all threads will not cause bank conflicts\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   "
.LC7:
	.string	"\t %.6f \t\t %f \t\t  %f \t\t  %f\n"
.LC8:
	.string	"Reading the same 32-bit word"
.LC9:
	.string	"\n ------------------------------------------------------------------------------------------------------------------------------\n"
	.text
	.globl	_Z12printResultsv
	.type	_Z12printResultsv, @function
_Z12printResultsv:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	.LC0(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC7(%rip), %rbx
	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
	leaq	.LC5(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	%rbx, %rsi
	movl	$2, %edi
	movb	$4, %al
	movsd	_ZL7maxtime(%rip), %xmm3
	movsd	_ZL7mintime(%rip), %xmm2
	movsd	_ZL7avgtime(%rip), %xmm1
	movss	_ZL10bandWidths(%rip), %xmm0
	divss	.LC6(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
	leaq	.LC8(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	%rbx, %rsi
	movl	$2, %edi
	movb	$4, %al
	movsd	8+_ZL7maxtime(%rip), %xmm3
	movsd	8+_ZL7mintime(%rip), %xmm2
	movsd	8+_ZL7avgtime(%rip), %xmm1
	movss	4+_ZL10bandWidths(%rip), %xmm0
	divss	.LC6(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
	leaq	.LC9(%rip), %rsi
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 8
	movl	$2, %edi
	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	.L4
	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
.L4:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L6
	call	__stack_chk_fail@PLT
.L6:
	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	_Z40__device_stub__Z21sharedMemReadSameWordvv
	.type	_Z40__device_stub__Z21sharedMemReadSameWordvv, @function
_Z40__device_stub__Z21sharedMemReadSameWordvv:
.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	.L10
	pushq	8(%rsp)
	.cfi_def_cfa_offset 104
	leaq	_Z21sharedMemReadSameWordv(%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
.L10:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L12
	call	__stack_chk_fail@PLT
.L12:
	addq	$88, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2055:
	.size	_Z40__device_stub__Z21sharedMemReadSameWordvv, .-_Z40__device_stub__Z21sharedMemReadSameWordvv
	.globl	_Z21sharedMemReadSameWordv
	.type	_Z21sharedMemReadSameWordv, @function
_Z21sharedMemReadSameWordv:
.LFB2056:
	.cfi_startproc
	endbr64
	jmp	_Z40__device_stub__Z21sharedMemReadSameWordvv
	.cfi_endproc
.LFE2056:
	.size	_Z21sharedMemReadSameWordv, .-_Z21sharedMemReadSameWordv
	.section	.rodata.str1.1
.LC13:
	.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	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	movabsq	$4294982921, %r13
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	movabsq	$4294967424, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	addq	$-128, %rsp
	.cfi_def_cfa_offset 176
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	movq	%rsp, %rdi
	leaq	40(%rsp), %rbx
	leaq	80(%rsp), %r14
	call	cudaEventCreate@PLT
	leaq	8(%rsp), %rdi
	movq	%rbx, %rbp
	call	cudaEventCreate@PLT
.L18:
	movq	(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%r12, %rdx
	movl	$1, %ecx
	movq	%r13, %rdi
	movl	$1, %esi
	movq	%r12, 28(%rsp)
	movl	$1, 36(%rsp)
	movq	%r13, 16(%rsp)
	movl	$1, 24(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L16
	call	_Z45__device_stub__Z26sharedMemAccessWithStride1vv
.L16:
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	8(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	8(%rsp), %rdx
	movq	(%rsp), %rsi
	movq	%rbp, %rdi
	call	cudaEventElapsedTime@PLT
	movq	(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%r12, %rdx
	movl	$1, %ecx
	movq	%r13, %rdi
	movl	$1, %esi
	movq	%r12, 28(%rsp)
	movl	$1, 36(%rsp)
	movq	%r13, 16(%rsp)
	movl	$1, 24(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L17
	call	_Z40__device_stub__Z21sharedMemReadSameWordvv
.L17:
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	8(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	8(%rsp), %rdx
	movq	(%rsp), %rsi
	leaq	40(%rbp), %rdi
	addq	$4, %rbp
	call	cudaEventElapsedTime@PLT
	cmpq	%r14, %rbp
	jne	.L18
	movsd	.LC10(%rip), %xmm2
	movsd	_ZL7avgtime(%rip), %xmm0
	leaq	36(%rbx), %rax
	movsd	_ZL7maxtime(%rip), %xmm5
	movsd	8+_ZL7avgtime(%rip), %xmm1
	movsd	8+_ZL7maxtime(%rip), %xmm4
	movaps	%xmm2, %xmm3
.L19:
	cvtss2sd	4(%rbx), %xmm6
	addq	$4, %rbx
	addsd	%xmm6, %xmm0
	minsd	%xmm6, %xmm2
	maxsd	%xmm6, %xmm5
	cvtss2sd	40(%rbx), %xmm6
	addsd	%xmm6, %xmm1
	minsd	%xmm6, %xmm3
	maxsd	%xmm6, %xmm4
	cmpq	%rax, %rbx
	jne	.L19
	movsd	%xmm5, _ZL7maxtime(%rip)
	movsd	.LC11(%rip), %xmm5
	movsd	%xmm4, 8+_ZL7maxtime(%rip)
	divsd	%xmm5, %xmm0
	movsd	%xmm2, _ZL7mintime(%rip)
	movsd	%xmm3, 8+_ZL7mintime(%rip)
	divsd	%xmm5, %xmm1
	movsd	%xmm0, _ZL7avgtime(%rip)
	movsd	.LC12(%rip), %xmm0
	movaps	%xmm0, %xmm4
	divsd	%xmm2, %xmm4
	movsd	%xmm1, 8+_ZL7avgtime(%rip)
	divsd	%xmm3, %xmm0
	cvtsd2ss	%xmm4, %xmm4
	movss	%xmm4, _ZL10bandWidths(%rip)
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, 4+_ZL10bandWidths(%rip)
	call	_Z12printResultsv
	leaq	.LC13(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movq	8(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L27
	call	__stack_chk_fail@PLT
.L27:
	subq	$-128, %rsp
	.cfi_def_cfa_offset 48
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	main, .-main
	.section	.rodata.str1.1
.LC14:
	.string	"_Z21sharedMemReadSameWordv"
.LC15:
	.string	"_Z26sharedMemAccessWithStride1v"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2058:
	.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	.LC14(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z21sharedMemReadSameWordv(%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	.LC15(%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
.LFE2058:
	.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,8,8
	.local	_ZL7mintime
	.comm	_ZL7mintime,16,16
	.local	_ZL7maxtime
	.comm	_ZL7maxtime,16,16
	.local	_ZL7avgtime
	.comm	_ZL7avgtime,16,16
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC6:
	.long	1232348160
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC10:
	.long	-536870912
	.long	1206910975
	.align 8
.LC11:
	.long	0
	.long	1075970048
	.align 8
.LC12:
	.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 :  sharedMemoryReadingSameWord.cu

  Objective   :  Example code to demonstrate the different access patterns of float3 array 
                 in the global memory the corresponding advantages in terms of the bandwidth 
                 that is achievable

  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

  output:	 The different bandwidths of the shared memory that are achieved while 
                 accessing the same word and the access pattern without any bank conflicts

 Modified     :  Aug 2011
                 

 Author	      :  RarchK		

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

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

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

#define NO_OF_PATTERNS 2 
#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 reading the same word by all the threads  
// the access pattern is conflict-free 
// since the word gets broadcasted to all the threads
//
/////////////////////////////////////////////////////////////////////////////////////////////////

__global__ void sharedMemReadSameWord(void)
{
	__shared__  float array[BLOCK_SIZE];
	int idx = 3;
	float __attribute__((unused)) scalar ;
	scalar = array[idx];
}

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






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 two kernals 
// finding the corresponding 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);
			sharedMemReadSameWord <<< gridSize,BLOCK_SIZE >>> ();
			hipEventRecord(stop,0);
			hipEventSynchronize(stop);
			hipEventElapsedTime(&elapsedTimes[1][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 taken by the two kernels
// the associated maximum bandwidth achieved by the different kernels
//
///////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////

void printResults()
{
	printf("Demonstrating that reading same 32-bit word by all threads will not cause bank conflicts\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("Reading the same 32-bit word");
					  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	"sharedMemoryReadingSameWord.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	_Z36__device_stub__sharedMemReadSameWordv # -- Begin function _Z36__device_stub__sharedMemReadSameWordv
	.type	_Z36__device_stub__sharedMemReadSameWordv,@function
_Z36__device_stub__sharedMemReadSameWordv: # @_Z36__device_stub__sharedMemReadSameWordv
	.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	$_Z21sharedMemReadSameWordv, %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	_Z36__device_stub__sharedMemReadSameWordv, .Lfunc_end1-_Z36__device_stub__sharedMemReadSameWordv
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI2_0:
	.quad	0x4022000000000000              # double 9
.LCPI2_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	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	subq	$104, %rsp
	.cfi_def_cfa_offset 144
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -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
	leaq	16(%rsp), %r14
	movl	$10, %r12d
	leaq	15497(%rbx), %r15
.LBB2_1:                                # =>This Inner Loop Header: Depth=1
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	%r15, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_3
# %bb.2:                                #   in Loop: Header=BB2_1 Depth=1
	callq	_Z41__device_stub__sharedMemAccessWithStride1v
.LBB2_3:                                #   in Loop: Header=BB2_1 Depth=1
	movq	(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	(%rsp), %rdi
	callq	hipEventSynchronize
	movq	8(%rsp), %rsi
	movq	(%rsp), %rdx
	movq	%r14, %rdi
	callq	hipEventElapsedTime
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	%r15, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_5
# %bb.4:                                #   in Loop: Header=BB2_1 Depth=1
	callq	_Z36__device_stub__sharedMemReadSameWordv
.LBB2_5:                                #   in Loop: Header=BB2_1 Depth=1
	movq	(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	(%rsp), %rdi
	callq	hipEventSynchronize
	leaq	40(%r14), %rdi
	movq	8(%rsp), %rsi
	movq	(%rsp), %rdx
	callq	hipEventElapsedTime
	addq	$4, %r14
	decq	%r12
	jne	.LBB2_1
# %bb.6:                                # %.preheader64.preheader
	movabsq	$5183643170566569984, %rax      # imm = 0x47EFFFFFE0000000
	movq	%rax, _ZL7mintime(%rip)
	movq	%rax, _ZL7mintime+8(%rip)
	movl	$1, %eax
.LBB2_7:                                # %.preheader62
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_8 Depth 2
	movb	$1, %cl
	xorl	%edx, %edx
.LBB2_8:                                #   Parent Loop BB2_7 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	leaq	(%rdx,%rdx,4), %rsi
	leaq	(%rsp,%rsi,8), %rsi
	addq	$16, %rsi
	xorps	%xmm0, %xmm0
	cvtss2sd	(%rsi,%rax,4), %xmm0
	movsd	_ZL7avgtime(,%rdx,8), %xmm1     # xmm1 = mem[0],zero
	addsd	%xmm0, %xmm1
	movsd	%xmm1, _ZL7avgtime(,%rdx,8)
	movsd	_ZL7mintime(,%rdx,8), %xmm1     # xmm1 = mem[0],zero
	minsd	%xmm0, %xmm1
	movsd	%xmm1, _ZL7mintime(,%rdx,8)
	movsd	_ZL7maxtime(,%rdx,8), %xmm1     # xmm1 = mem[0],zero
	maxsd	%xmm0, %xmm1
	movsd	%xmm1, _ZL7maxtime(,%rdx,8)
	movl	$1, %edx
	testb	$1, %cl
	movl	$0, %ecx
	jne	.LBB2_8
# %bb.9:                                #   in Loop: Header=BB2_7 Depth=1
	incq	%rax
	cmpq	$10, %rax
	jne	.LBB2_7
# %bb.10:                               # %.preheader
	movsd	_ZL7avgtime(%rip), %xmm0        # xmm0 = mem[0],zero
	movsd	.LCPI2_0(%rip), %xmm1           # xmm1 = mem[0],zero
	divsd	%xmm1, %xmm0
	movsd	.LCPI2_1(%rip), %xmm2           # xmm2 = mem[0],zero
	movapd	%xmm2, %xmm3
	divsd	_ZL7mintime(%rip), %xmm3
	movsd	%xmm0, _ZL7avgtime(%rip)
	xorps	%xmm0, %xmm0
	cvtsd2ss	%xmm3, %xmm0
	movss	%xmm0, _ZL10bandWidths(%rip)
	movsd	_ZL7avgtime+8(%rip), %xmm0      # xmm0 = mem[0],zero
	divsd	%xmm1, %xmm0
	divsd	_ZL7mintime+8(%rip), %xmm2
	movsd	%xmm0, _ZL7avgtime+8(%rip)
	xorps	%xmm0, %xmm0
	cvtsd2ss	%xmm2, %xmm0
	movss	%xmm0, _ZL10bandWidths+4(%rip)
	callq	_Z12printResultsv
	movl	$.Lstr, %edi
	callq	puts@PLT
	movq	8(%rsp), %rdi
	callq	hipEventDestroy
	movq	(%rsp), %rdi
	callq	hipEventDestroy
	xorl	%eax, %eax
	addq	$104, %rsp
	.cfi_def_cfa_offset 40
	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	main, .Lfunc_end2-main
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _Z12printResultsv
.LCPI3_0:
	.long	0x49742400                      # float 1.0E+6
	.text
	.globl	_Z12printResultsv
	.type	_Z12printResultsv,@function
_Z12printResultsv:                      # @_Z12printResultsv
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %rbp, -16
	movl	$.Lstr.1, %edi
	callq	puts@PLT
	xorl	%ebx, %ebx
	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
	movb	$1, %bpl
	movl	$.L.str.6, %r14d
.LBB3_1:                                # =>This Inner Loop Header: Depth=1
	testb	$1, %bpl
	movl	$.L.str.7, %edi
	cmovneq	%r14, %rdi
	xorl	%eax, %eax
	callq	printf
	movss	_ZL10bandWidths(,%rbx,4), %xmm0 # xmm0 = mem[0],zero,zero,zero
	divss	.LCPI3_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.8, %edi
	movb	$4, %al
	callq	printf
	movl	$1, %ebx
	testb	$1, %bpl
	movl	$0, %ebp
	jne	.LBB3_1
# %bb.2:
	movl	$.Lstr.5, %edi
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	puts@PLT                        # TAILCALL
.Lfunc_end3:
	.size	_Z12printResultsv, .Lfunc_end3-_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	.LBB4_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB4_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_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	$_Z21sharedMemReadSameWordv, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end4:
	.size	__hip_module_ctor, .Lfunc_end4-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB5_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB5_2:
	retq
.Lfunc_end5:
	.size	__hip_module_dtor, .Lfunc_end5-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z26sharedMemAccessWithStride1v,@object # @_Z26sharedMemAccessWithStride1v
	.section	.rodata,"a",@progbits
	.globl	_Z26sharedMemAccessWithStride1v
	.p2align	3, 0x0
_Z26sharedMemAccessWithStride1v:
	.quad	_Z41__device_stub__sharedMemAccessWithStride1v
	.size	_Z26sharedMemAccessWithStride1v, 8

	.type	_Z21sharedMemReadSameWordv,@object # @_Z21sharedMemReadSameWordv
	.globl	_Z21sharedMemReadSameWordv
	.p2align	3, 0x0
_Z21sharedMemReadSameWordv:
	.quad	_Z36__device_stub__sharedMemReadSameWordv
	.size	_Z21sharedMemReadSameWordv, 8

	.type	_ZL7mintime,@object             # @_ZL7mintime
	.local	_ZL7mintime
	.comm	_ZL7mintime,16,16
	.type	_ZL7avgtime,@object             # @_ZL7avgtime
	.local	_ZL7avgtime
	.comm	_ZL7avgtime,16,16
	.type	_ZL7maxtime,@object             # @_ZL7maxtime
	.local	_ZL7maxtime
	.comm	_ZL7maxtime,16,16
	.type	_ZL10bandWidths,@object         # @_ZL10bandWidths
	.local	_ZL10bandWidths
	.comm	_ZL10bandWidths,8,4
	.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, 29

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"Reading the same 32-bit word"
	.size	.L.str.7, 29

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"\t %.6f \t\t %f \t\t  %f \t\t  %f\n"
	.size	.L.str.8, 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	"_Z21sharedMemReadSameWordv"
	.size	.L__unnamed_2, 27

	.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 that reading same 32-bit word by all threads will not cause bank conflicts"
	.size	.Lstr.1, 89

	.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

	.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 _Z36__device_stub__sharedMemReadSameWordv
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z26sharedMemAccessWithStride1v
	.addrsig_sym _Z21sharedMemReadSameWordv
	.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	_Z21sharedMemReadSameWordv ; -- Begin function _Z21sharedMemReadSameWordv
	.globl	_Z21sharedMemReadSameWordv
	.p2align	8
	.type	_Z21sharedMemReadSameWordv,@function
_Z21sharedMemReadSameWordv:             ; @_Z21sharedMemReadSameWordv
; %bb.0:
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z21sharedMemReadSameWordv
		.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	_Z21sharedMemReadSameWordv, .Lfunc_end1-_Z21sharedMemReadSameWordv
                                        ; -- 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:           _Z21sharedMemReadSameWordv
    .private_segment_fixed_size: 0
    .sgpr_count:     0
    .sgpr_spill_count: 0
    .symbol:         _Z21sharedMemReadSameWordv.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
