

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

	
  Source Code :  SOAvsAOS.cu

  Objective   : Example code to demonstrate the advantage of having Stucture of arrays rather 
                than array of structures in the application while representing data the 
                corresponding advantages in terms of the bandwidth of the global memory that 
                is achievable

  Description : This example takes "Triangle" as structure with three arrays of three floating 
                points each representing the three vertices of a triangle the same information 
                is also is stored using a structure "Triangles" which has arrays for each field 
                of each vertex  both the representations are intialised generating typical 
                access patterns that will be present while accessing those structures

  Input       :	none

  Output      :	The different bandwidths that are achieved while accessing data from different 
                data representations

 Modified     : Aug 2011
                

 Author	      : RarchK		

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

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

#define NO_OF_TRIANGLES 1000000 
#define BLOCK_SIZE 128


#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);

// Triangle structure -- which is the structure in "Array of Structures"
struct Triangle {
						float A[3];
						float B[3];
						float C[3];
					 };

// Triangles structure -- which is the structure in "Structure of Arrays"
struct Triangles {
						float *Ax, *Ay, *Az;
						float *Bx, *By, *Bz;
						float *Cx, *Cy, *Cz;
					  };


/////////////////////////////////////////////////////////////////////////////////////
//
// initializing the array of Triangle types.
// each element is intialized to -- (0,0,0),(1,1,1),(2,2,2)
//
////////////////////////////////////////////////////////////////////////////////////

__global__ void setTriangles(Triangle *myArrayOfTriangles)
{
	int idx = threadIdx.x + blockIdx.x*blockDim.x;

	if(idx < NO_OF_TRIANGLES)
	{
		myArrayOfTriangles[idx].A[0] = 0;
		myArrayOfTriangles[idx].A[1] = 0;
		myArrayOfTriangles[idx].A[2] = 0;
		myArrayOfTriangles[idx].B[0] = 1;
		myArrayOfTriangles[idx].B[1] = 1;
		myArrayOfTriangles[idx].B[2] = 1;
		myArrayOfTriangles[idx].C[0] = 2;
		myArrayOfTriangles[idx].C[1] = 2;
		myArrayOfTriangles[idx].C[2] = 2;
	}
}

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




/////////////////////////////////////////////////////////////////////////////////////
//
// initializing the Triangles structure
// each triangle's vertices are initialized to -- (0,0,0),(1,1,1),(2,2,2)
//
////////////////////////////////////////////////////////////////////////////////////

__global__ void setTriangles2(Triangles myTrianglesStructure)
{
	int idx = threadIdx.x + blockIdx.x*blockDim.x;

	if(idx < NO_OF_TRIANGLES)
	{
		myTrianglesStructure.Ax[idx] = 0;
		myTrianglesStructure.Ay[idx] = 0;
		myTrianglesStructure.Az[idx] = 0;
		myTrianglesStructure.Bx[idx] = 1;
		myTrianglesStructure.By[idx] = 1;
		myTrianglesStructure.Bz[idx] = 1;
		myTrianglesStructure.Cx[idx] = 2;
		myTrianglesStructure.Cy[idx] = 2;
		myTrianglesStructure.Cz[idx] = 2;
	}
}

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



///////////////////////////////////////////////////////////////////////////////////////////////////////
//
// checking for the equivalence of both the initialising kernals
// 
//////////////////////////////////////////////////////////////////////////////////////////////////////

__global__ void Test(Triangle *myArrayOfTriangles, Triangles myTrianglesStructure, int* dCorrectBool)
{

	int idx = threadIdx.x + blockIdx.x*blockDim.x;
	if(idx < NO_OF_TRIANGLES)
	if(myTrianglesStructure.Ax[idx] != myArrayOfTriangles[idx].A[0]
		||	myTrianglesStructure.Ay[idx] !=  myArrayOfTriangles[idx].A[1]
		|| myTrianglesStructure.Az[idx] !=  myArrayOfTriangles[idx].A[2]
		|| myTrianglesStructure.Bx[idx] !=  myArrayOfTriangles[idx].B[0]
		|| myTrianglesStructure.By[idx] !=  myArrayOfTriangles[idx].B[1]
		|| myTrianglesStructure.Bz[idx] !=  myArrayOfTriangles[idx].B[2]
		|| myTrianglesStructure.Cx[idx] !=  myArrayOfTriangles[idx].C[0]
		|| myTrianglesStructure.Cy[idx] !=  myArrayOfTriangles[idx].C[1]
		|| myTrianglesStructure.Cz[idx] !=  myArrayOfTriangles[idx].C[2] )
	
		*dCorrectBool = 0;
}

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




static double  avgtime[NO_OF_PATTERNS] = {0}, maxtime[NO_OF_PATTERNS] = {0}, mintime[NO_OF_PATTERNS];
static float bandWidths[NO_OF_PATTERNS] = {0}; 
static double bytes = sizeof(float) * NO_OF_TRIANGLES * 9;

////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
//
// the main routene 
// for timing the different kernals which are accessing array of structures and structure of arrays
// finding the bandwidths
//
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////

int main(int argc, char* argv[])
{
	Triangle *myArrayOfTriangles;          // Array of Structures -- array of struct Triangle type
	Triangles myTrianglesStructure;        // Structure of Arrays
	int* dCorrectBool;                     // on the device variable
	int hCorrectBool = 1;

	float elapsedTimes[NO_OF_PATTERNS][NTIMES];
	cudaEvent_t start,stop;
	cudaError_t err = cudaSuccess;
	
	
	// event creation, which will be used for timing the code
	cudaEventCreate(&start);
	cudaEventCreate(&stop);


	// memory allocation on the device

	err = cudaMalloc((void **)&myArrayOfTriangles,sizeof(Triangle)*NO_OF_TRIANGLES);
	if(err == cudaErrorMemoryAllocation)
	{
		printf("error in device memory allocation for - myArrayOfTriangles\n exiting out of the program.....\n");
		exit(-1);
	}
	
	err = cudaMalloc((void **)&myTrianglesStructure.Ax,sizeof(float)*NO_OF_TRIANGLES);
	if(err == cudaErrorMemoryAllocation)
	{
		printf("error in device memory allocation for - myTrianglesStructure.Ax\n exiting out of the program.....\n");
		exit(-1);
	}

	err = cudaMalloc((void **)&myTrianglesStructure.Ay,sizeof(float)*NO_OF_TRIANGLES);
	if(err == cudaErrorMemoryAllocation)
	{
		printf("error in device memory allocation for - myTrianglesStructure.Ay\n exiting out of the program.....\n");
		exit(-1);
	}
	
	err = cudaMalloc((void **)&myTrianglesStructure.Az,sizeof(float)*NO_OF_TRIANGLES);
	if(err == cudaErrorMemoryAllocation)
	{
		printf("error in device memory allocation for - myTrianglesStructure.Az\n exiting out of the program.....\n");
		exit(-1);
	}
	
	err = cudaMalloc((void **)&myTrianglesStructure.Bx,sizeof(float)*NO_OF_TRIANGLES);
	if(err == cudaErrorMemoryAllocation)
	{
		printf("error in device memory allocation for - myTrianglesStructure.Bx\n exiting out of the program.....\n");
		exit(-1);
	}
	
	err = cudaMalloc((void **)&myTrianglesStructure.By,sizeof(float)*NO_OF_TRIANGLES);
	if(err == cudaErrorMemoryAllocation)
	{
		printf("error in device memory allocation for - myTrianglesStructure.By\n exiting out of the program.....\n");
		exit(-1);
	}
	
	err = cudaMalloc((void **)&myTrianglesStructure.Bz,sizeof(float)*NO_OF_TRIANGLES);
	if(err == cudaErrorMemoryAllocation)
	{
		printf("error in device memory allocation for - myTrianglesStructure.Bz\n exiting out of the program.....\n");
		exit(-1);
	}
	
	err = cudaMalloc((void **)&myTrianglesStructure.Cx,sizeof(float)*NO_OF_TRIANGLES);
	if(err == cudaErrorMemoryAllocation)
	{
		printf("error in device memory allocation for - myTrianglesStructure.Cx\n exiting out of the program.....\n");
		exit(-1);
	}
	
	err = cudaMalloc((void **)&myTrianglesStructure.Cy,sizeof(float)*NO_OF_TRIANGLES);
	if(err == cudaErrorMemoryAllocation)
	{
		printf("error in device memory allocation for - myTrianglesStructure.Cy\n exiting out of the program.....\n");
		exit(-1);
	}
	
	err = cudaMalloc((void **)&myTrianglesStructure.Cz,sizeof(float)*NO_OF_TRIANGLES);
	if(err == cudaErrorMemoryAllocation)
	{
		printf("error in device memory allocation for - myTrianglesStructure.Cz\n exiting out of the program.....\n");
		exit(-1);
	}
	
	err = cudaMalloc((void**)&dCorrectBool, sizeof(int));
	if(err == cudaErrorMemoryAllocation)
	{
		printf("error in device memory allocation for dCorrectBool\n exiting out of the program.....\n");
		exit(-1);
	}

	// copying hCorrectBool into dCorrectBool
	cudaMemcpy(dCorrectBool,&hCorrectBool,sizeof(int),cudaMemcpyHostToDevice);


	//finding the 1D grid size 
	int gridSize = NO_OF_TRIANGLES/BLOCK_SIZE;
	
	if( NO_OF_TRIANGLES % 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);
			setTriangles <<< gridSize , BLOCK_SIZE >>> (myArrayOfTriangles);
			cudaEventRecord(stop,0);
			cudaEventSynchronize(stop);
			cudaEventElapsedTime(&elapsedTimes[0][k],start,stop);

			// PATTERN 2
			cudaEventRecord(start,0);
			setTriangles2 <<< gridSize , BLOCK_SIZE >>> (myTrianglesStructure);
			cudaEventRecord(stop,0);
			cudaEventSynchronize(stop);
			cudaEventElapsedTime(&elapsedTimes[1][k],start,stop);


	} // end of the for loop involving NTIMES

	Test  <<< gridSize, BLOCK_SIZE >>>  (myArrayOfTriangles,myTrianglesStructure,dCorrectBool); // testing the equivalence of both the intializing kernals
	
	// copying back the value of dCorrectBool into hCorrectBool
	cudaMemcpy(&hCorrectBool,dCorrectBool,sizeof(int),cudaMemcpyDeviceToHost);

	if(hCorrectBool != 1) // if the kernels are not equivalent
	{
		printf("the kernel executions are not equivalent\n");
		printf("exiting out of the program\n");
		
	}
 
   else // if kernals are equivalent 
	{
		// 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");
   }

	//destroying the events
	cudaEventDestroy(start);
	cudaEventDestroy(stop);
	
	//freeing the allocated memory on the device 
	cudaFree(myArrayOfTriangles);
	cudaFree(myTrianglesStructure.Ax);
	cudaFree(myTrianglesStructure.Ay);
	cudaFree(myTrianglesStructure.Az);
	cudaFree(myTrianglesStructure.Bx);
	cudaFree(myTrianglesStructure.By);
	cudaFree(myTrianglesStructure.Bz);
	cudaFree(myTrianglesStructure.Cx);
	cudaFree(myTrianglesStructure.Cy);
	cudaFree(myTrianglesStructure.Cz);

	return 0;

}

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



////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
//
// prints the results containig the minimum, maximum, average times taken by the two initializing kernels
// the associated maximum bandwidth of global memory achieved by the kernels
//
///////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////

void printResults()
{

	printf("Demonstrating the advantages of Structure of arrays over Array of structures in data representation\n");
		
   printf("The transfered data size (Bytes): %ld\n",NO_OF_TRIANGLES*sizeof(float)*9);

   printf("\n-------------------------------------------------------------------------------------------------------------------------------\n");
   printf("Pattern \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("Array of Structures");
					  break;

			case 1: printf("Structure of Arrays");
					  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_0029a5ad_00000000-6_SOAvsAOS.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 advantages of Structure of arrays over Array of structures in data representation\n"
.LC1:
	.string	"The transfered data size (Bytes): %ld\n"
.LC2:
	.string	"\n-------------------------------------------------------------------------------------------------------------------------------\n"
.LC3:
	.string	"Pattern \t\t Bandwidth (GB/sec) \t Avg time (ms)  \t Min time (ms)  \t Max time (ms)\n"
.LC4:
	.string	"-------------------------------------------------------------------------------------------------------------------------------\n"
.LC5:
	.string	"Array of Structures"
.LC7:
	.string	"\t %.6f \t\t %f \t\t  %f \t\t  %f\n"
.LC8:
	.string	"Structure of Arrays"
.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	$36000000, %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	_Z40__device_stub__Z12setTrianglesP8TriangleP8Triangle
	.type	_Z40__device_stub__Z12setTrianglesP8TriangleP8Triangle, @function
_Z40__device_stub__Z12setTrianglesP8TriangleP8Triangle:
.LFB2053:
	.cfi_startproc
	endbr64
	subq	$104, %rsp
	.cfi_def_cfa_offset 112
	movq	%rdi, 8(%rsp)
	leaq	24(%rsp), %rcx
	leaq	16(%rsp), %rdx
	leaq	44(%rsp), %rsi
	leaq	32(%rsp), %rdi
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 40(%rsp)
	movq	%rax, 80(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L4
	pushq	24(%rsp)
	.cfi_def_cfa_offset 120
	leaq	_Z12setTrianglesP8Triangle(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 128
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	96(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 120
	popq	%rdx
	.cfi_def_cfa_offset 112
.L4:
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L6
	call	__stack_chk_fail@PLT
.L6:
	addq	$104, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2053:
	.size	_Z40__device_stub__Z12setTrianglesP8TriangleP8Triangle, .-_Z40__device_stub__Z12setTrianglesP8TriangleP8Triangle
	.globl	_Z12setTrianglesP8Triangle
	.type	_Z12setTrianglesP8Triangle, @function
_Z12setTrianglesP8Triangle:
.LFB2054:
	.cfi_startproc
	endbr64
	jmp	_Z40__device_stub__Z12setTrianglesP8TriangleP8Triangle
	.cfi_endproc
.LFE2054:
	.size	_Z12setTrianglesP8Triangle, .-_Z12setTrianglesP8Triangle
	.globl	_Z41__device_stub__Z13setTriangles29TrianglesR9Triangles
	.type	_Z41__device_stub__Z13setTriangles29TrianglesR9Triangles, @function
_Z41__device_stub__Z13setTriangles29TrianglesR9Triangles:
.LFB2055:
	.cfi_startproc
	endbr64
	subq	$88, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rcx
	movq	%rdi, 64(%rsp)
	movq	%rsp, %rdx
	movabsq	$4294967297, %rax
	leaq	28(%rsp), %rsi
	leaq	16(%rsp), %rdi
	movl	$1, 24(%rsp)
	movq	%rax, 16(%rsp)
	movq	%rax, 28(%rsp)
	movl	$1, 36(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L10
	pushq	8(%rsp)
	.cfi_def_cfa_offset 104
	leaq	_Z13setTriangles29Triangles(%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	_Z41__device_stub__Z13setTriangles29TrianglesR9Triangles, .-_Z41__device_stub__Z13setTriangles29TrianglesR9Triangles
	.globl	_Z13setTriangles29Triangles
	.type	_Z13setTriangles29Triangles, @function
_Z13setTriangles29Triangles:
.LFB2056:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	16(%rsp), %rdi
	call	_Z41__device_stub__Z13setTriangles29TrianglesR9Triangles
	popq	%rdx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2056:
	.size	_Z13setTriangles29Triangles, .-_Z13setTriangles29Triangles
	.globl	_Z43__device_stub__Z4TestP8Triangle9TrianglesPiP8TriangleR9TrianglesPi
	.type	_Z43__device_stub__Z4TestP8Triangle9TrianglesPiP8TriangleR9TrianglesPi, @function
_Z43__device_stub__Z4TestP8Triangle9TrianglesPiP8TriangleR9TrianglesPi:
.LFB2057:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	24(%rsp), %rcx
	leaq	32(%rsp), %rdi
	movq	%rdx, (%rsp)
	leaq	16(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movq	%rsi, 88(%rsp)
	leaq	44(%rsp), %rsi
	movq	%rax, 80(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movl	$1, 40(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L16
	pushq	24(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z4TestP8Triangle9TrianglesPi(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 144
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	96(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L16:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L18
	call	__stack_chk_fail@PLT
.L18:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2057:
	.size	_Z43__device_stub__Z4TestP8Triangle9TrianglesPiP8TriangleR9TrianglesPi, .-_Z43__device_stub__Z4TestP8Triangle9TrianglesPiP8TriangleR9TrianglesPi
	.globl	_Z4TestP8Triangle9TrianglesPi
	.type	_Z4TestP8Triangle9TrianglesPi, @function
_Z4TestP8Triangle9TrianglesPi:
.LFB2058:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	movq	%rsi, %rdx
	leaq	16(%rsp), %rsi
	call	_Z43__device_stub__Z4TestP8Triangle9TrianglesPiP8TriangleR9TrianglesPi
	popq	%rdx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2058:
	.size	_Z4TestP8Triangle9TrianglesPi, .-_Z4TestP8Triangle9TrianglesPi
	.section	.rodata.str1.1
.LC11:
	.string	"error in device memory allocation for - myArrayOfTriangles\n exiting out of the program.....\n"
.LC12:
	.string	"error in device memory allocation for - myTrianglesStructure.Ax\n exiting out of the program.....\n"
.LC13:
	.string	"error in device memory allocation for - myTrianglesStructure.Ay\n exiting out of the program.....\n"
.LC14:
	.string	"error in device memory allocation for - myTrianglesStructure.Az\n exiting out of the program.....\n"
.LC15:
	.string	"error in device memory allocation for - myTrianglesStructure.Bx\n exiting out of the program.....\n"
.LC16:
	.string	"error in device memory allocation for - myTrianglesStructure.By\n exiting out of the program.....\n"
.LC17:
	.string	"error in device memory allocation for - myTrianglesStructure.Bz\n exiting out of the program.....\n"
.LC18:
	.string	"error in device memory allocation for - myTrianglesStructure.Cx\n exiting out of the program.....\n"
.LC19:
	.string	"error in device memory allocation for - myTrianglesStructure.Cy\n exiting out of the program.....\n"
.LC20:
	.string	"error in device memory allocation for - myTrianglesStructure.Cz\n exiting out of the program.....\n"
.LC21:
	.string	"error in device memory allocation for dCorrectBool\n exiting out of the program.....\n"
.LC22:
	.string	"the kernel executions are not equivalent\n"
.LC23:
	.string	"exiting out of the program\n"
.LC26:
	.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	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$328, %rsp
	.cfi_def_cfa_offset 384
	movq	%fs:40, %rax
	movq	%rax, 312(%rsp)
	xorl	%eax, %eax
	leaq	48(%rsp), %rdi
	movl	$1, 28(%rsp)
	call	cudaEventCreate@PLT
	leaq	56(%rsp), %rdi
	call	cudaEventCreate@PLT
	movl	$36000000, %esi
	leaq	32(%rsp), %rdi
	call	cudaMalloc@PLT
	leaq	.LC11(%rip), %rsi
	cmpl	$2, %eax
	je	.L53
	leaq	88(%rsp), %rbp
	movl	$4000000, %esi
	movq	%rbp, %rdi
	call	cudaMalloc@PLT
	cmpl	$2, %eax
	jne	.L24
	leaq	.LC12(%rip), %rsi
.L53:
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	orl	$-1, %edi
	call	exit@PLT
.L24:
	movl	$4000000, %esi
	leaq	96(%rsp), %rdi
	call	cudaMalloc@PLT
	leaq	.LC13(%rip), %rsi
	cmpl	$2, %eax
	je	.L53
	movl	$4000000, %esi
	leaq	104(%rsp), %rdi
	call	cudaMalloc@PLT
	leaq	.LC14(%rip), %rsi
	cmpl	$2, %eax
	je	.L53
	movl	$4000000, %esi
	leaq	112(%rsp), %rdi
	call	cudaMalloc@PLT
	leaq	.LC15(%rip), %rsi
	cmpl	$2, %eax
	je	.L53
	movl	$4000000, %esi
	leaq	120(%rsp), %rdi
	call	cudaMalloc@PLT
	leaq	.LC16(%rip), %rsi
	cmpl	$2, %eax
	je	.L53
	movl	$4000000, %esi
	leaq	128(%rsp), %rdi
	call	cudaMalloc@PLT
	leaq	.LC17(%rip), %rsi
	cmpl	$2, %eax
	je	.L53
	movl	$4000000, %esi
	leaq	136(%rsp), %rdi
	call	cudaMalloc@PLT
	leaq	.LC18(%rip), %rsi
	cmpl	$2, %eax
	je	.L53
	movl	$4000000, %esi
	leaq	144(%rsp), %rdi
	call	cudaMalloc@PLT
	leaq	.LC19(%rip), %rsi
	cmpl	$2, %eax
	je	.L53
	movl	$4000000, %esi
	leaq	152(%rsp), %rdi
	call	cudaMalloc@PLT
	leaq	.LC20(%rip), %rsi
	cmpl	$2, %eax
	je	.L53
	movl	$4, %esi
	leaq	40(%rsp), %rdi
	call	cudaMalloc@PLT
	leaq	.LC21(%rip), %rsi
	cmpl	$2, %eax
	je	.L53
	leaq	28(%rsp), %rax
	movq	40(%rsp), %rdi
	movl	$1, %ecx
	movl	$4, %edx
	movq	%rax, %rsi
	leaq	232(%rsp), %rbx
	movq	%rax, 8(%rsp)
	leaq	272(%rsp), %r15
	call	cudaMemcpy@PLT
	movq	%rbx, %r12
	movabsq	$4294967424, %r13
	movabsq	$4294975109, %r14
.L36:
	movq	48(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%r13, %rdx
	movl	$1, %ecx
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%r13, 160(%rsp)
	movl	$1, 168(%rsp)
	movq	%r14, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L34
	movq	32(%rsp), %rdi
	call	_Z40__device_stub__Z12setTrianglesP8TriangleP8Triangle
.L34:
	movq	56(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	56(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	56(%rsp), %rdx
	movq	48(%rsp), %rsi
	movq	%r12, %rdi
	call	cudaEventElapsedTime@PLT
	movq	48(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%r13, %rdx
	movl	$1, %ecx
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%r13, 76(%rsp)
	movl	$1, 84(%rsp)
	movq	%r14, 64(%rsp)
	movl	$1, 72(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L35
	leaq	160(%rsp), %rdi
	movl	$18, %ecx
	movq	%rbp, %rsi
	rep movsl
	leaq	160(%rsp), %rdi
	call	_Z41__device_stub__Z13setTriangles29TrianglesR9Triangles
.L35:
	movq	56(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	56(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	56(%rsp), %rdx
	movq	48(%rsp), %rsi
	leaq	40(%r12), %rdi
	addq	$4, %r12
	call	cudaEventElapsedTime@PLT
	cmpq	%r15, %r12
	jne	.L36
	movl	$33554433, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	salq	$7, %rdx
	movl	$1, %esi
	movabsq	$4294975109, %rdi
	movl	$1, 84(%rsp)
	movq	%rdx, 76(%rsp)
	movq	%rdi, 64(%rsp)
	movl	$1, 72(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L37
	leaq	160(%rsp), %rdi
	movl	$18, %ecx
	movq	%rbp, %rsi
	movq	40(%rsp), %rdx
	rep movsl
	movq	32(%rsp), %rdi
	leaq	160(%rsp), %rsi
	call	_Z43__device_stub__Z4TestP8Triangle9TrianglesPiP8TriangleR9TrianglesPi
.L37:
	movq	40(%rsp), %rsi
	movq	8(%rsp), %rdi
	movl	$2, %ecx
	movl	$4, %edx
	call	cudaMemcpy@PLT
	cmpl	$1, 28(%rsp)
	jne	.L38
	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
	jmp	.L39
.L38:
	leaq	.LC22(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC23(%rip), %rsi
	jmp	.L52
.L39:
	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	.L39
	movsd	%xmm5, _ZL7maxtime(%rip)
	movsd	.LC24(%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	.LC25(%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	.LC26(%rip), %rsi
.L52:
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	48(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movq	56(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movq	32(%rsp), %rdi
	call	cudaFree@PLT
	movq	88(%rsp), %rdi
	call	cudaFree@PLT
	movq	96(%rsp), %rdi
	call	cudaFree@PLT
	movq	104(%rsp), %rdi
	call	cudaFree@PLT
	movq	112(%rsp), %rdi
	call	cudaFree@PLT
	movq	120(%rsp), %rdi
	call	cudaFree@PLT
	movq	128(%rsp), %rdi
	call	cudaFree@PLT
	movq	136(%rsp), %rdi
	call	cudaFree@PLT
	movq	144(%rsp), %rdi
	call	cudaFree@PLT
	movq	152(%rsp), %rdi
	call	cudaFree@PLT
	movq	312(%rsp), %rax
	subq	%fs:40, %rax
	je	.L48
	call	__stack_chk_fail@PLT
.L48:
	addq	$328, %rsp
	.cfi_def_cfa_offset 56
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	main, .-main
	.section	.rodata.str1.1
.LC27:
	.string	"_Z4TestP8Triangle9TrianglesPi"
.LC28:
	.string	"_Z13setTriangles29Triangles"
.LC29:
	.string	"_Z12setTrianglesP8Triangle"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2060:
	.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	.LC27(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z4TestP8Triangle9TrianglesPi(%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	.LC28(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z13setTriangles29Triangles(%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	.LC29(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z12setTrianglesP8Triangle(%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
.LFE2060:
	.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
.LC24:
	.long	0
	.long	1075970048
	.align 8
.LC25:
	.long	0
	.long	1098984072
	.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 :  SOAvsAOS.cu

  Objective   : Example code to demonstrate the advantage of having Stucture of arrays rather 
                than array of structures in the application while representing data the 
                corresponding advantages in terms of the bandwidth of the global memory that 
                is achievable

  Description : This example takes "Triangle" as structure with three arrays of three floating 
                points each representing the three vertices of a triangle the same information 
                is also is stored using a structure "Triangles" which has arrays for each field 
                of each vertex  both the representations are intialised generating typical 
                access patterns that will be present while accessing those structures

  Input       :	none

  Output      :	The different bandwidths that are achieved while accessing data from different 
                data representations

 Modified     : Aug 2011
                

 Author	      : RarchK		

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

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

#define NO_OF_TRIANGLES 1000000 
#define BLOCK_SIZE 128


#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);

// Triangle structure -- which is the structure in "Array of Structures"
struct Triangle {
						float A[3];
						float B[3];
						float C[3];
					 };

// Triangles structure -- which is the structure in "Structure of Arrays"
struct Triangles {
						float *Ax, *Ay, *Az;
						float *Bx, *By, *Bz;
						float *Cx, *Cy, *Cz;
					  };


/////////////////////////////////////////////////////////////////////////////////////
//
// initializing the array of Triangle types.
// each element is intialized to -- (0,0,0),(1,1,1),(2,2,2)
//
////////////////////////////////////////////////////////////////////////////////////

__global__ void setTriangles(Triangle *myArrayOfTriangles)
{
	int idx = threadIdx.x + blockIdx.x*blockDim.x;

	if(idx < NO_OF_TRIANGLES)
	{
		myArrayOfTriangles[idx].A[0] = 0;
		myArrayOfTriangles[idx].A[1] = 0;
		myArrayOfTriangles[idx].A[2] = 0;
		myArrayOfTriangles[idx].B[0] = 1;
		myArrayOfTriangles[idx].B[1] = 1;
		myArrayOfTriangles[idx].B[2] = 1;
		myArrayOfTriangles[idx].C[0] = 2;
		myArrayOfTriangles[idx].C[1] = 2;
		myArrayOfTriangles[idx].C[2] = 2;
	}
}

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




/////////////////////////////////////////////////////////////////////////////////////
//
// initializing the Triangles structure
// each triangle's vertices are initialized to -- (0,0,0),(1,1,1),(2,2,2)
//
////////////////////////////////////////////////////////////////////////////////////

__global__ void setTriangles2(Triangles myTrianglesStructure)
{
	int idx = threadIdx.x + blockIdx.x*blockDim.x;

	if(idx < NO_OF_TRIANGLES)
	{
		myTrianglesStructure.Ax[idx] = 0;
		myTrianglesStructure.Ay[idx] = 0;
		myTrianglesStructure.Az[idx] = 0;
		myTrianglesStructure.Bx[idx] = 1;
		myTrianglesStructure.By[idx] = 1;
		myTrianglesStructure.Bz[idx] = 1;
		myTrianglesStructure.Cx[idx] = 2;
		myTrianglesStructure.Cy[idx] = 2;
		myTrianglesStructure.Cz[idx] = 2;
	}
}

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



///////////////////////////////////////////////////////////////////////////////////////////////////////
//
// checking for the equivalence of both the initialising kernals
// 
//////////////////////////////////////////////////////////////////////////////////////////////////////

__global__ void Test(Triangle *myArrayOfTriangles, Triangles myTrianglesStructure, int* dCorrectBool)
{

	int idx = threadIdx.x + blockIdx.x*blockDim.x;
	if(idx < NO_OF_TRIANGLES)
	if(myTrianglesStructure.Ax[idx] != myArrayOfTriangles[idx].A[0]
		||	myTrianglesStructure.Ay[idx] !=  myArrayOfTriangles[idx].A[1]
		|| myTrianglesStructure.Az[idx] !=  myArrayOfTriangles[idx].A[2]
		|| myTrianglesStructure.Bx[idx] !=  myArrayOfTriangles[idx].B[0]
		|| myTrianglesStructure.By[idx] !=  myArrayOfTriangles[idx].B[1]
		|| myTrianglesStructure.Bz[idx] !=  myArrayOfTriangles[idx].B[2]
		|| myTrianglesStructure.Cx[idx] !=  myArrayOfTriangles[idx].C[0]
		|| myTrianglesStructure.Cy[idx] !=  myArrayOfTriangles[idx].C[1]
		|| myTrianglesStructure.Cz[idx] !=  myArrayOfTriangles[idx].C[2] )
	
		*dCorrectBool = 0;
}

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




static double  avgtime[NO_OF_PATTERNS] = {0}, maxtime[NO_OF_PATTERNS] = {0}, mintime[NO_OF_PATTERNS];
static float bandWidths[NO_OF_PATTERNS] = {0}; 
static double bytes = sizeof(float) * NO_OF_TRIANGLES * 9;

////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
//
// the main routene 
// for timing the different kernals which are accessing array of structures and structure of arrays
// finding the bandwidths
//
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////

int main(int argc, char* argv[])
{
	Triangle *myArrayOfTriangles;          // Array of Structures -- array of struct Triangle type
	Triangles myTrianglesStructure;        // Structure of Arrays
	int* dCorrectBool;                     // on the device variable
	int hCorrectBool = 1;

	float elapsedTimes[NO_OF_PATTERNS][NTIMES];
	hipEvent_t start,stop;
	hipError_t err = hipSuccess;
	
	
	// event creation, which will be used for timing the code
	hipEventCreate(&start);
	hipEventCreate(&stop);


	// memory allocation on the device

	err = hipMalloc((void **)&myArrayOfTriangles,sizeof(Triangle)*NO_OF_TRIANGLES);
	if(err == hipErrorOutOfMemory)
	{
		printf("error in device memory allocation for - myArrayOfTriangles\n exiting out of the program.....\n");
		exit(-1);
	}
	
	err = hipMalloc((void **)&myTrianglesStructure.Ax,sizeof(float)*NO_OF_TRIANGLES);
	if(err == hipErrorOutOfMemory)
	{
		printf("error in device memory allocation for - myTrianglesStructure.Ax\n exiting out of the program.....\n");
		exit(-1);
	}

	err = hipMalloc((void **)&myTrianglesStructure.Ay,sizeof(float)*NO_OF_TRIANGLES);
	if(err == hipErrorOutOfMemory)
	{
		printf("error in device memory allocation for - myTrianglesStructure.Ay\n exiting out of the program.....\n");
		exit(-1);
	}
	
	err = hipMalloc((void **)&myTrianglesStructure.Az,sizeof(float)*NO_OF_TRIANGLES);
	if(err == hipErrorOutOfMemory)
	{
		printf("error in device memory allocation for - myTrianglesStructure.Az\n exiting out of the program.....\n");
		exit(-1);
	}
	
	err = hipMalloc((void **)&myTrianglesStructure.Bx,sizeof(float)*NO_OF_TRIANGLES);
	if(err == hipErrorOutOfMemory)
	{
		printf("error in device memory allocation for - myTrianglesStructure.Bx\n exiting out of the program.....\n");
		exit(-1);
	}
	
	err = hipMalloc((void **)&myTrianglesStructure.By,sizeof(float)*NO_OF_TRIANGLES);
	if(err == hipErrorOutOfMemory)
	{
		printf("error in device memory allocation for - myTrianglesStructure.By\n exiting out of the program.....\n");
		exit(-1);
	}
	
	err = hipMalloc((void **)&myTrianglesStructure.Bz,sizeof(float)*NO_OF_TRIANGLES);
	if(err == hipErrorOutOfMemory)
	{
		printf("error in device memory allocation for - myTrianglesStructure.Bz\n exiting out of the program.....\n");
		exit(-1);
	}
	
	err = hipMalloc((void **)&myTrianglesStructure.Cx,sizeof(float)*NO_OF_TRIANGLES);
	if(err == hipErrorOutOfMemory)
	{
		printf("error in device memory allocation for - myTrianglesStructure.Cx\n exiting out of the program.....\n");
		exit(-1);
	}
	
	err = hipMalloc((void **)&myTrianglesStructure.Cy,sizeof(float)*NO_OF_TRIANGLES);
	if(err == hipErrorOutOfMemory)
	{
		printf("error in device memory allocation for - myTrianglesStructure.Cy\n exiting out of the program.....\n");
		exit(-1);
	}
	
	err = hipMalloc((void **)&myTrianglesStructure.Cz,sizeof(float)*NO_OF_TRIANGLES);
	if(err == hipErrorOutOfMemory)
	{
		printf("error in device memory allocation for - myTrianglesStructure.Cz\n exiting out of the program.....\n");
		exit(-1);
	}
	
	err = hipMalloc((void**)&dCorrectBool, sizeof(int));
	if(err == hipErrorOutOfMemory)
	{
		printf("error in device memory allocation for dCorrectBool\n exiting out of the program.....\n");
		exit(-1);
	}

	// copying hCorrectBool into dCorrectBool
	hipMemcpy(dCorrectBool,&hCorrectBool,sizeof(int),hipMemcpyHostToDevice);


	//finding the 1D grid size 
	int gridSize = NO_OF_TRIANGLES/BLOCK_SIZE;
	
	if( NO_OF_TRIANGLES % 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);
			setTriangles <<< gridSize , BLOCK_SIZE >>> (myArrayOfTriangles);
			hipEventRecord(stop,0);
			hipEventSynchronize(stop);
			hipEventElapsedTime(&elapsedTimes[0][k],start,stop);

			// PATTERN 2
			hipEventRecord(start,0);
			setTriangles2 <<< gridSize , BLOCK_SIZE >>> (myTrianglesStructure);
			hipEventRecord(stop,0);
			hipEventSynchronize(stop);
			hipEventElapsedTime(&elapsedTimes[1][k],start,stop);


	} // end of the for loop involving NTIMES

	Test  <<< gridSize, BLOCK_SIZE >>>  (myArrayOfTriangles,myTrianglesStructure,dCorrectBool); // testing the equivalence of both the intializing kernals
	
	// copying back the value of dCorrectBool into hCorrectBool
	hipMemcpy(&hCorrectBool,dCorrectBool,sizeof(int),hipMemcpyDeviceToHost);

	if(hCorrectBool != 1) // if the kernels are not equivalent
	{
		printf("the kernel executions are not equivalent\n");
		printf("exiting out of the program\n");
		
	}
 
   else // if kernals are equivalent 
	{
		// 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");
   }

	//destroying the events
	hipEventDestroy(start);
	hipEventDestroy(stop);
	
	//freeing the allocated memory on the device 
	hipFree(myArrayOfTriangles);
	hipFree(myTrianglesStructure.Ax);
	hipFree(myTrianglesStructure.Ay);
	hipFree(myTrianglesStructure.Az);
	hipFree(myTrianglesStructure.Bx);
	hipFree(myTrianglesStructure.By);
	hipFree(myTrianglesStructure.Bz);
	hipFree(myTrianglesStructure.Cx);
	hipFree(myTrianglesStructure.Cy);
	hipFree(myTrianglesStructure.Cz);

	return 0;

}

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



////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
//
// prints the results containig the minimum, maximum, average times taken by the two initializing kernels
// the associated maximum bandwidth of global memory achieved by the kernels
//
///////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////

void printResults()
{

	printf("Demonstrating the advantages of Structure of arrays over Array of structures in data representation\n");
		
   printf("The transfered data size (Bytes): %ld\n",NO_OF_TRIANGLES*sizeof(float)*9);

   printf("\n-------------------------------------------------------------------------------------------------------------------------------\n");
   printf("Pattern \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("Array of Structures");
					  break;

			case 1: printf("Structure of Arrays");
					  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	"SOAvsAOS.hip"
	.globl	_Z27__device_stub__setTrianglesP8Triangle # -- Begin function _Z27__device_stub__setTrianglesP8Triangle
	.type	_Z27__device_stub__setTrianglesP8Triangle,@function
_Z27__device_stub__setTrianglesP8Triangle: # @_Z27__device_stub__setTrianglesP8Triangle
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$64, %rsp
	.cfi_def_cfa_offset 112
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	movq	%rsp, %rbx
	movq	%rax, (%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z12setTrianglesP8Triangle, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$80, %rsp
	.cfi_adjust_cfa_offset -80
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z27__device_stub__setTrianglesP8Triangle, .Lfunc_end0-_Z27__device_stub__setTrianglesP8Triangle
	.cfi_endproc
                                        # -- End function
	.globl	_Z28__device_stub__setTriangles29Triangles # -- Begin function _Z28__device_stub__setTriangles29Triangles
	.type	_Z28__device_stub__setTriangles29Triangles,@function
_Z28__device_stub__setTriangles29Triangles: # @_Z28__device_stub__setTriangles29Triangles
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$64, %rsp
	.cfi_def_cfa_offset 112
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	112(%rsp), %rax
	movq	%rsp, %rbx
	movq	%rax, (%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z13setTriangles29Triangles, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$80, %rsp
	.cfi_adjust_cfa_offset -80
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z28__device_stub__setTriangles29Triangles, .Lfunc_end1-_Z28__device_stub__setTriangles29Triangles
	.cfi_endproc
                                        # -- End function
	.globl	_Z19__device_stub__TestP8Triangle9TrianglesPi # -- Begin function _Z19__device_stub__TestP8Triangle9TrianglesPi
	.type	_Z19__device_stub__TestP8Triangle9TrianglesPi,@function
_Z19__device_stub__TestP8Triangle9TrianglesPi: # @_Z19__device_stub__TestP8Triangle9TrianglesPi
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$96, %rsp
	.cfi_def_cfa_offset 144
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	16(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	leaq	144(%rsp), %rax
	movq	%rax, 8(%rbx)
	movq	%rcx, 16(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%rsp, %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z4TestP8Triangle9TrianglesPi, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$112, %rsp
	.cfi_adjust_cfa_offset -112
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	_Z19__device_stub__TestP8Triangle9TrianglesPi, .Lfunc_end2-_Z19__device_stub__TestP8Triangle9TrianglesPi
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI3_0:
	.quad	0x4022000000000000              # double 9
.LCPI3_1:
	.quad	0x41812a8800000000              # double 3.6E+7
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$408, %rsp                      # imm = 0x198
	.cfi_def_cfa_offset 464
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	$1, 100(%rsp)
	leaq	88(%rsp), %rdi
	callq	hipEventCreate
	leaq	80(%rsp), %rdi
	callq	hipEventCreate
	leaq	112(%rsp), %rdi
	movl	$36000000, %esi                 # imm = 0x2255100
	callq	hipMalloc
	cmpl	$2, %eax
	je	.LBB3_1
# %bb.3:
	leaq	120(%rsp), %rdi
	movl	$4000000, %esi                  # imm = 0x3D0900
	callq	hipMalloc
	cmpl	$2, %eax
	je	.LBB3_4
# %bb.5:
	leaq	128(%rsp), %rdi
	movl	$4000000, %esi                  # imm = 0x3D0900
	callq	hipMalloc
	cmpl	$2, %eax
	je	.LBB3_6
# %bb.7:
	leaq	136(%rsp), %rdi
	movl	$4000000, %esi                  # imm = 0x3D0900
	callq	hipMalloc
	cmpl	$2, %eax
	je	.LBB3_8
# %bb.9:
	leaq	144(%rsp), %rdi
	movl	$4000000, %esi                  # imm = 0x3D0900
	callq	hipMalloc
	cmpl	$2, %eax
	je	.LBB3_10
# %bb.11:
	leaq	152(%rsp), %rdi
	movl	$4000000, %esi                  # imm = 0x3D0900
	callq	hipMalloc
	cmpl	$2, %eax
	je	.LBB3_12
# %bb.13:
	leaq	160(%rsp), %rdi
	movl	$4000000, %esi                  # imm = 0x3D0900
	callq	hipMalloc
	cmpl	$2, %eax
	je	.LBB3_14
# %bb.15:
	leaq	168(%rsp), %rdi
	movl	$4000000, %esi                  # imm = 0x3D0900
	callq	hipMalloc
	cmpl	$2, %eax
	je	.LBB3_16
# %bb.17:
	leaq	176(%rsp), %rdi
	movl	$4000000, %esi                  # imm = 0x3D0900
	callq	hipMalloc
	cmpl	$2, %eax
	je	.LBB3_18
# %bb.19:
	leaq	184(%rsp), %rdi
	movl	$4000000, %esi                  # imm = 0x3D0900
	callq	hipMalloc
	cmpl	$2, %eax
	je	.LBB3_20
# %bb.21:
	leaq	104(%rsp), %rdi
	movl	$4, %esi
	callq	hipMalloc
	cmpl	$2, %eax
	je	.LBB3_22
# %bb.23:
	movabsq	$4294967424, %rbx               # imm = 0x100000080
	movq	104(%rsp), %rdi
	leaq	100(%rsp), %rsi
	movl	$4, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	256(%rsp), %r15
	movl	$10, %r13d
	leaq	7685(%rbx), %r14
	leaq	336(%rsp), %r12
	leaq	192(%rsp), %rbp
.LBB3_24:                               # =>This Inner Loop Header: Depth=1
	movq	88(%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	.LBB3_26
# %bb.25:                               #   in Loop: Header=BB3_24 Depth=1
	movq	112(%rsp), %rdi
	callq	_Z27__device_stub__setTrianglesP8Triangle
.LBB3_26:                               #   in Loop: Header=BB3_24 Depth=1
	movq	80(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	80(%rsp), %rdi
	callq	hipEventSynchronize
	movq	88(%rsp), %rsi
	movq	80(%rsp), %rdx
	movq	%r15, %rdi
	callq	hipEventElapsedTime
	movq	88(%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	.LBB3_28
# %bb.27:                               #   in Loop: Header=BB3_24 Depth=1
	movl	$9, %ecx
	movq	%r12, %rdi
	leaq	120(%rsp), %rsi
	rep;movsq (%rsi), %es:(%rdi)
	movq	%r12, 192(%rsp)
	leaq	240(%rsp), %rdi
	leaq	224(%rsp), %rsi
	leaq	216(%rsp), %rdx
	leaq	208(%rsp), %rcx
	callq	__hipPopCallConfiguration
	movq	216(%rsp), %rax
	movq	208(%rsp), %rdi
	movq	240(%rsp), %rsi
	movl	248(%rsp), %edx
	movq	224(%rsp), %rcx
	movl	232(%rsp), %r8d
	movq	%rdi, 8(%rsp)
	movq	%rax, (%rsp)
	movl	$_Z13setTriangles29Triangles, %edi
	movq	%rbp, %r9
	callq	hipLaunchKernel
.LBB3_28:                               #   in Loop: Header=BB3_24 Depth=1
	movq	80(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	80(%rsp), %rdi
	callq	hipEventSynchronize
	leaq	40(%r15), %rdi
	movq	88(%rsp), %rsi
	movq	80(%rsp), %rdx
	callq	hipEventElapsedTime
	addq	$4, %r15
	decq	%r13
	jne	.LBB3_24
# %bb.29:
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_31
# %bb.30:
	movq	112(%rsp), %rax
	movq	104(%rsp), %rdx
	leaq	120(%rsp), %rsi
	movl	$9, %ecx
	movq	%rsp, %rdi
	rep;movsq (%rsi), %es:(%rdi)
	movq	%rax, %rdi
	movq	%rdx, %rsi
	callq	_Z19__device_stub__TestP8Triangle9TrianglesPi
.LBB3_31:
	movq	104(%rsp), %rsi
	leaq	100(%rsp), %rbx
	movl	$4, %edx
	movq	%rbx, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	cmpl	$1, (%rbx)
	jne	.LBB3_38
# %bb.32:                               # %.preheader102.preheader
	movabsq	$5183643170566569984, %rax      # imm = 0x47EFFFFFE0000000
	movq	%rax, _ZL7mintime(%rip)
	movq	%rax, _ZL7mintime+8(%rip)
	movl	$1, %eax
.LBB3_33:                               # %.preheader100
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB3_34 Depth 2
	movb	$1, %cl
	xorl	%edx, %edx
.LBB3_34:                               #   Parent Loop BB3_33 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	leaq	(%rdx,%rdx,4), %rsi
	leaq	(%rsp,%rsi,8), %rsi
	addq	$256, %rsi                      # imm = 0x100
	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	.LBB3_34
# %bb.35:                               #   in Loop: Header=BB3_33 Depth=1
	incq	%rax
	cmpq	$10, %rax
	jne	.LBB3_33
# %bb.36:                               # %.preheader
	movsd	_ZL7avgtime(%rip), %xmm0        # xmm0 = mem[0],zero
	movsd	.LCPI3_0(%rip), %xmm1           # xmm1 = mem[0],zero
	divsd	%xmm1, %xmm0
	movsd	.LCPI3_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
	jmp	.LBB3_37
.LBB3_38:
	movl	$.Lstr.1, %edi
	callq	puts@PLT
	movl	$.Lstr.2, %edi
.LBB3_37:
	callq	puts@PLT
	movq	88(%rsp), %rdi
	callq	hipEventDestroy
	movq	80(%rsp), %rdi
	callq	hipEventDestroy
	movq	112(%rsp), %rdi
	callq	hipFree
	movq	120(%rsp), %rdi
	callq	hipFree
	movq	128(%rsp), %rdi
	callq	hipFree
	movq	136(%rsp), %rdi
	callq	hipFree
	movq	144(%rsp), %rdi
	callq	hipFree
	movq	152(%rsp), %rdi
	callq	hipFree
	movq	160(%rsp), %rdi
	callq	hipFree
	movq	168(%rsp), %rdi
	callq	hipFree
	movq	176(%rsp), %rdi
	callq	hipFree
	movq	184(%rsp), %rdi
	callq	hipFree
	xorl	%eax, %eax
	addq	$408, %rsp                      # imm = 0x198
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.LBB3_1:
	.cfi_def_cfa_offset 464
	movl	$.Lstr.13, %edi
	jmp	.LBB3_2
.LBB3_4:
	movl	$.Lstr.12, %edi
	jmp	.LBB3_2
.LBB3_6:
	movl	$.Lstr.11, %edi
	jmp	.LBB3_2
.LBB3_8:
	movl	$.Lstr.10, %edi
	jmp	.LBB3_2
.LBB3_10:
	movl	$.Lstr.9, %edi
	jmp	.LBB3_2
.LBB3_12:
	movl	$.Lstr.8, %edi
	jmp	.LBB3_2
.LBB3_14:
	movl	$.Lstr.7, %edi
	jmp	.LBB3_2
.LBB3_16:
	movl	$.Lstr.6, %edi
	jmp	.LBB3_2
.LBB3_18:
	movl	$.Lstr.5, %edi
	jmp	.LBB3_2
.LBB3_20:
	movl	$.Lstr.4, %edi
	jmp	.LBB3_2
.LBB3_22:
	movl	$.Lstr.3, %edi
.LBB3_2:
	callq	puts@PLT
	movl	$-1, %edi
	callq	exit
.Lfunc_end3:
	.size	main, .Lfunc_end3-main
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _Z12printResultsv
.LCPI4_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.14, %edi
	callq	puts@PLT
	xorl	%ebx, %ebx
	movl	$.L.str.15, %edi
	movl	$36000000, %esi                 # imm = 0x2255100
	xorl	%eax, %eax
	callq	printf
	movl	$.Lstr.15, %edi
	callq	puts@PLT
	movl	$.Lstr.16, %edi
	callq	puts@PLT
	movl	$.Lstr.17, %edi
	callq	puts@PLT
	movb	$1, %bpl
	movl	$.L.str.19, %r14d
.LBB4_1:                                # =>This Inner Loop Header: Depth=1
	testb	$1, %bpl
	movl	$.L.str.20, %edi
	cmovneq	%r14, %rdi
	xorl	%eax, %eax
	callq	printf
	movss	_ZL10bandWidths(,%rbx,4), %xmm0 # xmm0 = mem[0],zero,zero,zero
	divss	.LCPI4_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.21, %edi
	movb	$4, %al
	callq	printf
	movl	$1, %ebx
	testb	$1, %bpl
	movl	$0, %ebp
	jne	.LBB4_1
# %bb.2:
	movl	$.Lstr.18, %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_end4:
	.size	_Z12printResultsv, .Lfunc_end4-_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	.LBB5_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB5_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z12setTrianglesP8Triangle, %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	$_Z13setTriangles29Triangles, %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	$_Z4TestP8Triangle9TrianglesPi, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end5:
	.size	__hip_module_ctor, .Lfunc_end5-__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	.LBB6_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
.LBB6_2:
	retq
.Lfunc_end6:
	.size	__hip_module_dtor, .Lfunc_end6-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z12setTrianglesP8Triangle,@object # @_Z12setTrianglesP8Triangle
	.section	.rodata,"a",@progbits
	.globl	_Z12setTrianglesP8Triangle
	.p2align	3, 0x0
_Z12setTrianglesP8Triangle:
	.quad	_Z27__device_stub__setTrianglesP8Triangle
	.size	_Z12setTrianglesP8Triangle, 8

	.type	_Z13setTriangles29Triangles,@object # @_Z13setTriangles29Triangles
	.globl	_Z13setTriangles29Triangles
	.p2align	3, 0x0
_Z13setTriangles29Triangles:
	.quad	_Z28__device_stub__setTriangles29Triangles
	.size	_Z13setTriangles29Triangles, 8

	.type	_Z4TestP8Triangle9TrianglesPi,@object # @_Z4TestP8Triangle9TrianglesPi
	.globl	_Z4TestP8Triangle9TrianglesPi
	.p2align	3, 0x0
_Z4TestP8Triangle9TrianglesPi:
	.quad	_Z19__device_stub__TestP8Triangle9TrianglesPi
	.size	_Z4TestP8Triangle9TrianglesPi, 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.15,@object               # @.str.15
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.15:
	.asciz	"The transfered data size (Bytes): %ld\n"
	.size	.L.str.15, 39

	.type	.L.str.19,@object               # @.str.19
.L.str.19:
	.asciz	"Array of Structures"
	.size	.L.str.19, 20

	.type	.L.str.20,@object               # @.str.20
.L.str.20:
	.asciz	"Structure of Arrays"
	.size	.L.str.20, 20

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

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z12setTrianglesP8Triangle"
	.size	.L__unnamed_1, 27

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z13setTriangles29Triangles"
	.size	.L__unnamed_2, 28

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

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

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

	.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	"the kernel executions are not equivalent"
	.size	.Lstr.1, 41

	.type	.Lstr.2,@object                 # @str.2
.Lstr.2:
	.asciz	"exiting out of the program"
	.size	.Lstr.2, 27

	.type	.Lstr.3,@object                 # @str.3
.Lstr.3:
	.asciz	"error in device memory allocation for dCorrectBool\n exiting out of the program....."
	.size	.Lstr.3, 84

	.type	.Lstr.4,@object                 # @str.4
.Lstr.4:
	.asciz	"error in device memory allocation for - myTrianglesStructure.Cz\n exiting out of the program....."
	.size	.Lstr.4, 97

	.type	.Lstr.5,@object                 # @str.5
.Lstr.5:
	.asciz	"error in device memory allocation for - myTrianglesStructure.Cy\n exiting out of the program....."
	.size	.Lstr.5, 97

	.type	.Lstr.6,@object                 # @str.6
.Lstr.6:
	.asciz	"error in device memory allocation for - myTrianglesStructure.Cx\n exiting out of the program....."
	.size	.Lstr.6, 97

	.type	.Lstr.7,@object                 # @str.7
.Lstr.7:
	.asciz	"error in device memory allocation for - myTrianglesStructure.Bz\n exiting out of the program....."
	.size	.Lstr.7, 97

	.type	.Lstr.8,@object                 # @str.8
.Lstr.8:
	.asciz	"error in device memory allocation for - myTrianglesStructure.By\n exiting out of the program....."
	.size	.Lstr.8, 97

	.type	.Lstr.9,@object                 # @str.9
.Lstr.9:
	.asciz	"error in device memory allocation for - myTrianglesStructure.Bx\n exiting out of the program....."
	.size	.Lstr.9, 97

	.type	.Lstr.10,@object                # @str.10
.Lstr.10:
	.asciz	"error in device memory allocation for - myTrianglesStructure.Az\n exiting out of the program....."
	.size	.Lstr.10, 97

	.type	.Lstr.11,@object                # @str.11
.Lstr.11:
	.asciz	"error in device memory allocation for - myTrianglesStructure.Ay\n exiting out of the program....."
	.size	.Lstr.11, 97

	.type	.Lstr.12,@object                # @str.12
.Lstr.12:
	.asciz	"error in device memory allocation for - myTrianglesStructure.Ax\n exiting out of the program....."
	.size	.Lstr.12, 97

	.type	.Lstr.13,@object                # @str.13
.Lstr.13:
	.asciz	"error in device memory allocation for - myArrayOfTriangles\n exiting out of the program....."
	.size	.Lstr.13, 92

	.type	.Lstr.14,@object                # @str.14
.Lstr.14:
	.asciz	"Demonstrating the advantages of Structure of arrays over Array of structures in data representation"
	.size	.Lstr.14, 100

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

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

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

	.type	.Lstr.18,@object                # @str.18
.Lstr.18:
	.asciz	"\n ------------------------------------------------------------------------------------------------------------------------------"
	.size	.Lstr.18, 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 _Z27__device_stub__setTrianglesP8Triangle
	.addrsig_sym _Z28__device_stub__setTriangles29Triangles
	.addrsig_sym _Z19__device_stub__TestP8Triangle9TrianglesPi
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z12setTrianglesP8Triangle
	.addrsig_sym _Z13setTriangles29Triangles
	.addrsig_sym _Z4TestP8Triangle9TrianglesPi
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z12setTrianglesP8Triangle ; -- Begin function _Z12setTrianglesP8Triangle
	.globl	_Z12setTrianglesP8Triangle
	.p2align	8
	.type	_Z12setTrianglesP8Triangle,@function
_Z12setTrianglesP8Triangle:             ; @_Z12setTrianglesP8Triangle
; %bb.0:
	s_load_b32 s2, s[0:1], 0x14
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_i32_e32 0xf4240, v1
	s_cbranch_execz .LBB0_2
; %bb.1:
	s_load_b64 s[0:1], s[0:1], 0x0
	v_dual_mov_b32 v2, 0 :: v_dual_mov_b32 v5, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_dual_mov_b32 v8, 2.0 :: v_dual_mov_b32 v3, v2
	v_dual_mov_b32 v4, v2 :: v_dual_mov_b32 v7, v5
	s_delay_alu instid0(VALU_DEP_2)
	v_dual_mov_b32 v6, v5 :: v_dual_mov_b32 v9, v8
	s_waitcnt lgkmcnt(0)
	v_mad_i64_i32 v[10:11], null, v1, 36, s[0:1]
	s_clause 0x2
	global_store_b128 v[10:11], v[2:5], off
	global_store_b128 v[10:11], v[6:9], off offset:16
	global_store_b32 v[10:11], v8, off offset:32
.LBB0_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z12setTrianglesP8Triangle
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 264
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 12
		.amdhsa_next_free_sgpr 16
		.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	_Z12setTrianglesP8Triangle, .Lfunc_end0-_Z12setTrianglesP8Triangle
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 148
; NumSgprs: 16
; NumVgprs: 12
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 1
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 16
; NumVGPRsForWavesPerEU: 12
; 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	_Z13setTriangles29Triangles ; -- Begin function _Z13setTriangles29Triangles
	.globl	_Z13setTriangles29Triangles
	.p2align	8
	.type	_Z13setTriangles29Triangles,@function
_Z13setTriangles29Triangles:            ; @_Z13setTriangles29Triangles
; %bb.0:
	s_load_b32 s2, s[0:1], 0x54
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_i32_e32 0xf4240, v1
	s_cbranch_execz .LBB1_2
; %bb.1:
	s_load_b512 s[4:19], s[0:1], 0x0
	v_ashrrev_i32_e32 v2, 31, v1
	s_load_b64 s[0:1], s[0:1], 0x40
	v_dual_mov_b32 v10, 0 :: v_dual_mov_b32 v11, 1.0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s4, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v3, vcc_lo, s5, v1, vcc_lo
	v_add_co_u32 v4, vcc_lo, s6, v0
	v_add_co_ci_u32_e32 v5, vcc_lo, s7, v1, vcc_lo
	v_add_co_u32 v6, vcc_lo, s8, v0
	v_add_co_ci_u32_e32 v7, vcc_lo, s9, v1, vcc_lo
	v_add_co_u32 v8, vcc_lo, s10, v0
	v_add_co_ci_u32_e32 v9, vcc_lo, s11, v1, vcc_lo
	global_store_b32 v[2:3], v10, off
	global_store_b32 v[4:5], v10, off
	global_store_b32 v[6:7], v10, off
	global_store_b32 v[8:9], v11, off
	v_add_co_u32 v2, vcc_lo, s12, v0
	v_add_co_ci_u32_e32 v3, vcc_lo, s13, v1, vcc_lo
	v_add_co_u32 v4, vcc_lo, s14, v0
	v_add_co_ci_u32_e32 v5, vcc_lo, s15, v1, vcc_lo
	v_add_co_u32 v6, vcc_lo, s16, v0
	v_add_co_ci_u32_e32 v7, vcc_lo, s17, v1, vcc_lo
	v_add_co_u32 v8, vcc_lo, s18, v0
	v_add_co_ci_u32_e32 v9, vcc_lo, s19, v1, vcc_lo
	v_mov_b32_e32 v10, 2.0
	v_add_co_u32 v0, vcc_lo, s0, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_store_b32 v[2:3], v11, off
	global_store_b32 v[4:5], v11, off
	global_store_b32 v[6:7], v10, off
	global_store_b32 v[8:9], v10, off
	global_store_b32 v[0:1], v10, off
.LBB1_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z13setTriangles29Triangles
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 328
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 12
		.amdhsa_next_free_sgpr 20
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end1:
	.size	_Z13setTriangles29Triangles, .Lfunc_end1-_Z13setTriangles29Triangles
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 292
; NumSgprs: 22
; NumVgprs: 12
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 22
; NumVGPRsForWavesPerEU: 12
; 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	_Z4TestP8Triangle9TrianglesPi ; -- Begin function _Z4TestP8Triangle9TrianglesPi
	.globl	_Z4TestP8Triangle9TrianglesPi
	.p2align	8
	.type	_Z4TestP8Triangle9TrianglesPi,@function
_Z4TestP8Triangle9TrianglesPi:          ; @_Z4TestP8Triangle9TrianglesPi
; %bb.0:
	s_load_b32 s2, s[0:1], 0x64
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[4:5], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_i32_e32 0xf4240, v4
	s_cbranch_execz .LBB2_19
; %bb.1:
	s_load_b128 s[20:23], s[0:1], 0x0
	v_ashrrev_i32_e32 v5, 31, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[0:1], 2, v[4:5]
	s_waitcnt lgkmcnt(0)
	v_mad_i64_i32 v[2:3], null, v4, 36, s[20:21]
	v_add_co_u32 v5, vcc_lo, s22, v0
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v6, vcc_lo, s23, v1, vcc_lo
	s_mov_b32 s23, exec_lo
	global_load_b32 v5, v[5:6], off
	global_load_b32 v6, v[2:3], off
	s_waitcnt vmcnt(0)
	v_cmp_neq_f32_e64 s22, v5, v6
	v_cmpx_eq_f32_e32 v5, v6
	s_cbranch_execz .LBB2_17
; %bb.2:
	s_load_b512 s[4:19], s[0:1], 0x10
	s_mov_b32 s24, exec_lo
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v5, vcc_lo, s4, v0
	v_add_co_ci_u32_e32 v6, vcc_lo, s5, v1, vcc_lo
	global_load_b32 v7, v[2:3], off offset:4
	global_load_b32 v5, v[5:6], off
	s_waitcnt vmcnt(0)
	v_cmp_neq_f32_e32 vcc_lo, v5, v7
	v_cmpx_eq_f32_e32 v5, v7
	s_cbranch_execz .LBB2_16
; %bb.3:
	v_add_co_u32 v5, s2, s6, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v6, s2, s7, v1, s2
	s_mov_b32 s7, exec_lo
	global_load_b32 v2, v[2:3], off offset:8
	global_load_b32 v3, v[5:6], off
	s_waitcnt vmcnt(0)
	v_cmp_neq_f32_e64 s2, v3, v2
	v_cmpx_eq_f32_e32 v3, v2
	s_cbranch_execz .LBB2_15
; %bb.4:
	v_mad_i64_i32 v[2:3], null, v4, 36, s[20:21]
	v_add_co_u32 v4, s3, s8, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v5, s3, s9, v1, s3
	global_load_b32 v4, v[4:5], off
	global_load_b32 v5, v[2:3], off offset:12
	s_waitcnt vmcnt(0)
	v_cmp_neq_f32_e64 s3, v4, v5
	v_cmp_eq_f32_e64 s4, v4, v5
	s_delay_alu instid0(VALU_DEP_1)
	s_and_saveexec_b32 s8, s4
	s_cbranch_execz .LBB2_14
; %bb.5:
	v_add_co_u32 v6, s4, s10, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e64 v7, s4, s11, v1, s4
	v_add_co_u32 v4, s4, v2, 12
	v_add_co_ci_u32_e64 v5, s4, 0, v3, s4
	global_load_b32 v6, v[6:7], off
	global_load_b32 v7, v[4:5], off offset:4
	s_waitcnt vmcnt(0)
	v_cmp_neq_f32_e64 s4, v6, v7
	v_cmp_eq_f32_e64 s5, v6, v7
	s_delay_alu instid0(VALU_DEP_1)
	s_and_saveexec_b32 s9, s5
	s_cbranch_execz .LBB2_13
; %bb.6:
	v_add_co_u32 v6, s5, s12, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v7, s5, s13, v1, s5
	global_load_b32 v4, v[4:5], off offset:8
	global_load_b32 v5, v[6:7], off
	s_waitcnt vmcnt(0)
	v_cmp_neq_f32_e64 s5, v5, v4
	v_cmp_eq_f32_e64 s6, v5, v4
	s_delay_alu instid0(VALU_DEP_1)
	s_and_saveexec_b32 s10, s6
	s_cbranch_execz .LBB2_12
; %bb.7:
	v_add_co_u32 v4, s6, s14, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v5, s6, s15, v1, s6
	s_mov_b32 s12, exec_lo
	global_load_b32 v6, v[2:3], off offset:24
	global_load_b32 v4, v[4:5], off
	s_waitcnt vmcnt(0)
	v_cmp_neq_f32_e64 s11, v4, v6
	v_cmpx_eq_f32_e32 v4, v6
	s_cbranch_execz .LBB2_11
; %bb.8:
	v_add_co_u32 v4, s6, s16, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e64 v5, s6, s17, v1, s6
	v_add_co_u32 v2, s6, v2, 24
	v_add_co_ci_u32_e64 v3, s6, 0, v3, s6
	s_mov_b32 s14, exec_lo
	global_load_b32 v4, v[4:5], off
	global_load_b32 v5, v[2:3], off offset:4
	s_waitcnt vmcnt(0)
	v_cmp_neq_f32_e64 s13, v4, v5
	v_cmpx_eq_f32_e32 v4, v5
	s_cbranch_execz .LBB2_10
; %bb.9:
	v_add_co_u32 v0, s6, s18, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v1, s6, s19, v1, s6
	s_and_not1_b32 s13, s13, exec_lo
	global_load_b32 v0, v[0:1], off
	global_load_b32 v1, v[2:3], off offset:8
	s_waitcnt vmcnt(0)
	v_cmp_neq_f32_e64 s6, v0, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s6, s6, exec_lo
	s_or_b32 s13, s13, s6
.LBB2_10:                               ; %Flow54
	s_or_b32 exec_lo, exec_lo, s14
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	s_and_not1_b32 s6, s11, exec_lo
	s_and_b32 s11, s13, exec_lo
	s_or_b32 s11, s6, s11
.LBB2_11:                               ; %Flow53
	s_or_b32 exec_lo, exec_lo, s12
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	s_and_not1_b32 s5, s5, exec_lo
	s_and_b32 s6, s11, exec_lo
	s_or_b32 s5, s5, s6
.LBB2_12:                               ; %Flow52
	s_or_b32 exec_lo, exec_lo, s10
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	s_and_not1_b32 s4, s4, exec_lo
	s_and_b32 s5, s5, exec_lo
	s_or_b32 s4, s4, s5
.LBB2_13:                               ; %Flow51
	s_or_b32 exec_lo, exec_lo, s9
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	s_and_not1_b32 s3, s3, exec_lo
	s_and_b32 s4, s4, exec_lo
	s_or_b32 s3, s3, s4
.LBB2_14:                               ; %Flow50
	s_or_b32 exec_lo, exec_lo, s8
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	s_and_not1_b32 s2, s2, exec_lo
	s_and_b32 s3, s3, exec_lo
	s_or_b32 s2, s2, s3
.LBB2_15:                               ; %Flow49
	s_or_b32 exec_lo, exec_lo, s7
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	s_and_not1_b32 s3, vcc_lo, exec_lo
	s_and_b32 s2, s2, exec_lo
	s_or_b32 vcc_lo, s3, s2
.LBB2_16:                               ; %Flow48
	s_or_b32 exec_lo, exec_lo, s24
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	s_and_not1_b32 s2, s22, exec_lo
	s_and_b32 s3, vcc_lo, exec_lo
	s_or_b32 s22, s2, s3
.LBB2_17:                               ; %Flow
	s_or_b32 exec_lo, exec_lo, s23
	s_delay_alu instid0(VALU_DEP_2) | instid1(SALU_CYCLE_1)
	s_and_b32 exec_lo, exec_lo, s22
	s_cbranch_execz .LBB2_19
; %bb.18:
	s_load_b64 s[0:1], s[0:1], 0x50
	v_mov_b32_e32 v0, 0
	s_waitcnt lgkmcnt(0)
	global_store_b32 v0, v0, s[0:1]
.LBB2_19:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z4TestP8Triangle9TrianglesPi
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 344
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 8
		.amdhsa_next_free_sgpr 25
		.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	_Z4TestP8Triangle9TrianglesPi, .Lfunc_end2-_Z4TestP8Triangle9TrianglesPi
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 880
; NumSgprs: 27
; NumVgprs: 8
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 3
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 27
; NumVGPRsForWavesPerEU: 8
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .offset:         8
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         12
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         16
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         20
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         22
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         24
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         26
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         28
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         30
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         48
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         56
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         72
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 264
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z12setTrianglesP8Triangle
    .private_segment_fixed_size: 0
    .sgpr_count:     16
    .sgpr_spill_count: 0
    .symbol:         _Z12setTrianglesP8Triangle.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     12
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .offset:         0
        .size:           72
        .value_kind:     by_value
      - .offset:         72
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         76
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         80
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         84
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         86
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         88
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         90
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         92
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         94
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         112
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         120
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         128
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         136
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 328
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z13setTriangles29Triangles
    .private_segment_fixed_size: 0
    .sgpr_count:     22
    .sgpr_spill_count: 0
    .symbol:         _Z13setTriangles29Triangles.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     12
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .offset:         8
        .size:           72
        .value_kind:     by_value
      - .address_space:  global
        .offset:         80
        .size:           8
        .value_kind:     global_buffer
      - .offset:         88
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         92
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         96
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         100
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         102
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         104
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         106
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         108
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         110
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         128
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         136
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         144
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         152
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 344
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z4TestP8Triangle9TrianglesPi
    .private_segment_fixed_size: 0
    .sgpr_count:     27
    .sgpr_spill_count: 0
    .symbol:         _Z4TestP8Triangle9TrianglesPi.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     8
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
