/* 
 * Example of using reducing (tree) type algorithms to parallelize finding the sum of
 * a set of numbers. On a GF 8600 GT the two parallel algorithms (sumControl = 0 or 1) 
 * are about 35 times faster than the serial algorithm also running on the GPU but using 
 * global memory (sumControl=2), for an array of 512 floats.  This is both because
 * the parallel algorithms scale as ln N while the serial algorithm scales as N, and because the
 * parallel algorithms use the shared memory while the serial one uses the (generally slower)
 * global memory in these tests. If we do the serial algorithm on the same computer but use
 * shared memory (set sumControl=3), the parallel algorithms are only 3.5 times faster.  Thus a
 * factor of 10 in the speed up is because of using the shared memory in this example. Note, 
 * however, that shared memory can only be shared among the threads within a single block.
 * 
 */

#include <stdio.h>

#define BLOCKSIZE 512 

// Define some GPU timing utilities. These are invoked from the host program. Usage:
//     START_GPU;
//         kernelFunction <<< numBlocks, threadsPerBlock >>> (args)
//     STOP_GPU;
//     PRINT_GPU
// in the host code. This estimates the time for the kernel kernelFunction to run on the GPU.
// For a more extensive discusion, see Section 5.1.2 of the CUDA Best Practices Guide at
// http://developer.download.nvidia.com/compute/DevZone/docs/html/C/doc/CUDA_C_Best_Practices_Guide.pdf

float timeGPU;
cudaEvent_t start, stop;
#define START_GPU cudaEventCreate(&start); cudaEventCreate(&stop); cudaEventRecord(start, 0);
#define STOP_GPU cudaEventRecord(stop, 0); cudaEventSynchronize(stop);\
   cudaEventElapsedTime(&timeGPU, start, stop);\
   cudaEventDestroy(start);cudaEventDestroy(stop);
#define PRINT_GPU printf("\n\nTime to compute on GPU: %f ms \n", timeGPU);

// Define a utility to check for CUDA errors.  Place it immediately after a CUDA kernel
// call in the host code. The initial cudaDeviceSynchronize() command ensures that the device
// has completed all preceding requested tasks.

#define CUDA_ERROR_CHECK cudaDeviceSynchronize(); cudaError_t error = cudaGetLastError();\
   if(error != cudaSuccess){printf("***CUDA error: %s\n", cudaGetErrorString(error)); exit(-1);}\
   else{printf("\nNo CUDA errors detected\n" );}
   

// Device code.  Sums the elements of the array Array and puts the result in Sum

__global__ void SumKernel(float* Array, float* Sum, int arraySize)
{
	__device__ float reductionSum(int, float *);

	*Sum = reductionSum(arraySize, Array);
}

/*
  Function to do parallel reduction sum. This should scale as ln N.  The parallel butterfly
  algorithm taken from the literature works generally.  My homegrown parallel version
  works as written for even number of entries in the array, so this algorithm can be used 
  for an odd number by padding the array with an added zero entry. Note that this version 
  assumes that all summations are within one block, so a max of 512 threads on 1.1 devices
  (presently blocksize is set to 256).  One option for larger sums is to break the array up 
  onto multiple blocks, use this algorithm on each block to get a block sum, and then sum
  the block sums.
*/

__device__ float reductionSum(int length, float *array)
{
	float sum = 0.0f;
	
	// = 0 or 1 for parallel with shared memory, 2 for serial with global, 3 for serial with shared
	int sumControl = 0;
	
	// Copy the array to be summed into shared memory and initialize
	__shared__ float sarray[BLOCKSIZE];
	int i = threadIdx.x;
	sarray[i] = 0.0f;
	if(i<length) sarray[i] = array[i];
	__syncthreads();
	
	
	if(sumControl == 0)
	{
		// Parallel butterfly sum
		// see http://cs.anu.edu.au/files/systems/GPUWksp/PDFs/02_CUDAParallelProgrammingModel.pdf
		
		for(int bit=BLOCKSIZE/2; bit>0; bit /= 2)
		{
			if(i<length)
			{
				float t=sarray[i] + sarray[i^bit];
				__syncthreads();
				sarray[i] = t;
				__syncthreads();
			}
		}
		
		// The array entries sarray[i] with i<length/2 now all contain the sum
		
		sum = sarray[0];
	
	}
	else if(sumControl == 1)
	{
		// Another home-made parallel version of a reduction sum. As written, this requires an even
		// number of entries in the array to be summed, so pad with a zero to handle odd number
		// (or rewrite to deal with odd number).
		
		int L=length;
		int steps = 0;
		int inc = 1;
		float t = 0;
		
		while(L > 2 )
		{
			steps ++;
			if(i < (length-inc))
				t = sarray[i] + sarray[i+inc];
				__syncthreads();
				sarray[i] = t;
				__syncthreads();
			inc *= 2;
			L /= 2;
		}
		
		sum = sarray[0] + sarray[inc];  // This contains the sum
	}
	else if(sumControl == 2)
	{	
		// Serial version of sum accessing global (not shared) memory
	
		sum = 0.0f;
		for(int i=0; i<length; i++)
		{
			sum += array[i];
		}
	}
	else
	{
		// Serial version of sum accessing shared memory
	
		sum = 0.0f;
		for(int i=0; i<length; i++)
		{
			sum += sarray[i];
		}
	}
	
	return sum;
}



// Host code
int main(void)
{
	
	int arraySize = 512;
	float* Array;
	
	Array = (float*) malloc(sizeof(float) * arraySize);
	
	// Fill array with some numbers
	
	for(int i=0; i<arraySize; i++)
	{
		Array[i] = 2.0f * (float) i;
	}
	
	// Set up device pointers
	float *devPtrArray;
	
	float Sum;
	float* devPtrSum;                            
	cudaMalloc((void**)&devPtrSum, sizeof(float));
	
	// Allocate device memory
	cudaMalloc((void**)&devPtrArray, arraySize * sizeof(float));
	
	// Copy array to device
	cudaMemcpy(devPtrArray, Array, arraySize * sizeof(float), cudaMemcpyHostToDevice);
	
	// Launch the kernel.
	
	START_GPU;     // Start timer for device code
	
	SumKernel<<<1, 512>>>(devPtrArray, devPtrSum, arraySize);
	
	STOP_GPU;           // Stop timer for device code
	PRINT_GPU;          // Print timing for device code
	CUDA_ERROR_CHECK
	
	// Copy the sum back from the GPU to the host
	
	cudaMemcpy(&Sum, devPtrSum, sizeof(float), cudaMemcpyDeviceToHost);
	
	printf("\nSum array elements multiplied by 2 (GPU) = %6.2f\n\n", Sum);
	 
    // Free the memory allocated on the device   
	cudaFree(devPtrSum);
	cudaFree(devPtrArray);
	
	// Free the memory allocated on the CPU
	free(Array);
	
    return 0;
	
}


	.file	"tmpxft_0025657c_00000000-6_reducedSum.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
	.globl	_Z12reductionSumiPf
	.type	_Z12reductionSumiPf, @function
_Z12reductionSumiPf:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2027:
	.size	_Z12reductionSumiPf, .-_Z12reductionSumiPf
	.globl	_Z31__device_stub__Z9SumKernelPfS_iPfS_i
	.type	_Z31__device_stub__Z9SumKernelPfS_iPfS_i, @function
_Z31__device_stub__Z9SumKernelPfS_iPfS_i:
.LFB2053:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L4
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z9SumKernelPfS_i(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L4:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L6
	call	__stack_chk_fail@PLT
.L6:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2053:
	.size	_Z31__device_stub__Z9SumKernelPfS_iPfS_i, .-_Z31__device_stub__Z9SumKernelPfS_iPfS_i
	.globl	_Z9SumKernelPfS_i
	.type	_Z9SumKernelPfS_i, @function
_Z9SumKernelPfS_i:
.LFB2054:
	.cfi_startproc
	endbr64
	jmp	_Z31__device_stub__Z9SumKernelPfS_iPfS_i
	.cfi_endproc
.LFE2054:
	.size	_Z9SumKernelPfS_i, .-_Z9SumKernelPfS_i
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"\n\nTime to compute on GPU: %f ms \n"
.LC1:
	.string	"***CUDA error: %s\n"
.LC2:
	.string	"\nNo CUDA errors detected\n"
.LC3:
	.string	"\nSum array elements multiplied by 2 (GPU) = %6.2f\n\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movl	$2048, %edi
	subq	$48, %rsp
	.cfi_def_cfa_offset 64
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	call	malloc@PLT
	movq	%rax, %rbx
	xorl	%eax, %eax
.L11:
	cvtsi2ssl	%eax, %xmm0
	addss	%xmm0, %xmm0
	movss	%xmm0, (%rbx,%rax,4)
	incq	%rax
	cmpq	$512, %rax
	jne	.L11
	leaq	8(%rsp), %rdi
	movl	$4, %esi
	call	cudaMalloc@PLT
	movq	%rsp, %rdi
	movl	$2048, %esi
	call	cudaMalloc@PLT
	movq	(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rsi
	movl	$2048, %edx
	call	cudaMemcpy@PLT
	leaq	start(%rip), %rdi
	call	cudaEventCreate@PLT
	leaq	stop(%rip), %rdi
	call	cudaEventCreate@PLT
	movq	start(%rip), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	$8388609, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	salq	$9, %rdx
	movl	$1, %ecx
	movabsq	$4294967297, %rdi
	movl	$1, %esi
	movq	%rdx, 28(%rsp)
	movl	$1, 36(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L12
	movq	8(%rsp), %rsi
	movq	(%rsp), %rdi
	movl	$512, %edx
	call	_Z31__device_stub__Z9SumKernelPfS_iPfS_i
.L12:
	movq	stop(%rip), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	stop(%rip), %rdi
	call	cudaEventSynchronize@PLT
	movq	stop(%rip), %rdx
	movq	start(%rip), %rsi
	leaq	timeGPU(%rip), %rdi
	call	cudaEventElapsedTime@PLT
	movq	start(%rip), %rdi
	call	cudaEventDestroy@PLT
	movq	stop(%rip), %rdi
	call	cudaEventDestroy@PLT
	movl	$2, %edi
	movb	$1, %al
	leaq	.LC0(%rip), %rsi
	cvtss2sd	timeGPU(%rip), %xmm0
	call	__printf_chk@PLT
	call	cudaDeviceSynchronize@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L13
	call	cudaGetErrorString@PLT
	movl	$2, %edi
	leaq	.LC1(%rip), %rsi
	movq	%rax, %rdx
	xorl	%eax, %eax
	call	__printf_chk@PLT
	orl	$-1, %edi
	call	exit@PLT
.L13:
	leaq	.LC2(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	8(%rsp), %rsi
	movl	$2, %ecx
	movl	$4, %edx
	leaq	28(%rsp), %rdi
	call	cudaMemcpy@PLT
	leaq	.LC3(%rip), %rsi
	movl	$2, %edi
	movb	$1, %al
	cvtss2sd	28(%rsp), %xmm0
	call	__printf_chk@PLT
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	(%rsp), %rdi
	call	cudaFree@PLT
	movq	%rbx, %rdi
	call	free@PLT
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L14
	call	__stack_chk_fail@PLT
.L14:
	addq	$48, %rsp
	.cfi_def_cfa_offset 16
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2028:
	.size	main, .-main
	.section	.rodata.str1.1
.LC4:
	.string	"_Z9SumKernelPfS_i"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2056:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC4(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z9SumKernelPfS_i(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2056:
	.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
	.globl	stop
	.bss
	.align 8
	.type	stop, @object
	.size	stop, 8
stop:
	.zero	8
	.globl	start
	.align 8
	.type	start, @object
	.size	start, 8
start:
	.zero	8
	.globl	timeGPU
	.align 4
	.type	timeGPU, @object
	.size	timeGPU, 4
timeGPU:
	.zero	4
	.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:

/* 
 * Example of using reducing (tree) type algorithms to parallelize finding the sum of
 * a set of numbers. On a GF 8600 GT the two parallel algorithms (sumControl = 0 or 1) 
 * are about 35 times faster than the serial algorithm also running on the GPU but using 
 * global memory (sumControl=2), for an array of 512 floats.  This is both because
 * the parallel algorithms scale as ln N while the serial algorithm scales as N, and because the
 * parallel algorithms use the shared memory while the serial one uses the (generally slower)
 * global memory in these tests. If we do the serial algorithm on the same computer but use
 * shared memory (set sumControl=3), the parallel algorithms are only 3.5 times faster.  Thus a
 * factor of 10 in the speed up is because of using the shared memory in this example. Note, 
 * however, that shared memory can only be shared among the threads within a single block.
 * 
 */


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

#define BLOCKSIZE 512 

// Define some GPU timing utilities. These are invoked from the host program. Usage:
//     START_GPU;
//         kernelFunction <<< numBlocks, threadsPerBlock >>> (args)
//     STOP_GPU;
//     PRINT_GPU
// in the host code. This estimates the time for the kernel kernelFunction to run on the GPU.
// For a more extensive discusion, see Section 5.1.2 of the CUDA Best Practices Guide at
// http://developer.download.nvidia.com/compute/DevZone/docs/html/C/doc/CUDA_C_Best_Practices_Guide.pdf

float timeGPU;
hipEvent_t start, stop;
#define START_GPU hipEventCreate(&start); hipEventCreate(&stop); hipEventRecord(start, 0);
#define STOP_GPU hipEventRecord(stop, 0); hipEventSynchronize(stop);\
   hipEventElapsedTime(&timeGPU, start, stop);\
   hipEventDestroy(start);hipEventDestroy(stop);
#define PRINT_GPU printf("\n\nTime to compute on GPU: %f ms \n", timeGPU);

// Define a utility to check for CUDA errors.  Place it immediately after a CUDA kernel
// call in the host code. The initial cudaDeviceSynchronize() command ensures that the device
// has completed all preceding requested tasks.

#define CUDA_ERROR_CHECK hipDeviceSynchronize(); hipError_t error = hipGetLastError();\
   if(error != hipSuccess){printf("***CUDA error: %s\n", hipGetErrorString(error)); exit(-1);}\
   else{printf("\nNo CUDA errors detected\n" );}
   

// Device code.  Sums the elements of the array Array and puts the result in Sum

__global__ void SumKernel(float* Array, float* Sum, int arraySize)
{
	__device__ float reductionSum(int, float *);

	*Sum = reductionSum(arraySize, Array);
}

/*
  Function to do parallel reduction sum. This should scale as ln N.  The parallel butterfly
  algorithm taken from the literature works generally.  My homegrown parallel version
  works as written for even number of entries in the array, so this algorithm can be used 
  for an odd number by padding the array with an added zero entry. Note that this version 
  assumes that all summations are within one block, so a max of 512 threads on 1.1 devices
  (presently blocksize is set to 256).  One option for larger sums is to break the array up 
  onto multiple blocks, use this algorithm on each block to get a block sum, and then sum
  the block sums.
*/

__device__ float reductionSum(int length, float *array)
{
	float sum = 0.0f;
	
	// = 0 or 1 for parallel with shared memory, 2 for serial with global, 3 for serial with shared
	int sumControl = 0;
	
	// Copy the array to be summed into shared memory and initialize
	__shared__ float sarray[BLOCKSIZE];
	int i = threadIdx.x;
	sarray[i] = 0.0f;
	if(i<length) sarray[i] = array[i];
	__syncthreads();
	
	
	if(sumControl == 0)
	{
		// Parallel butterfly sum
		// see http://cs.anu.edu.au/files/systems/GPUWksp/PDFs/02_CUDAParallelProgrammingModel.pdf
		
		for(int bit=BLOCKSIZE/2; bit>0; bit /= 2)
		{
			if(i<length)
			{
				float t=sarray[i] + sarray[i^bit];
				__syncthreads();
				sarray[i] = t;
				__syncthreads();
			}
		}
		
		// The array entries sarray[i] with i<length/2 now all contain the sum
		
		sum = sarray[0];
	
	}
	else if(sumControl == 1)
	{
		// Another home-made parallel version of a reduction sum. As written, this requires an even
		// number of entries in the array to be summed, so pad with a zero to handle odd number
		// (or rewrite to deal with odd number).
		
		int L=length;
		int steps = 0;
		int inc = 1;
		float t = 0;
		
		while(L > 2 )
		{
			steps ++;
			if(i < (length-inc))
				t = sarray[i] + sarray[i+inc];
				__syncthreads();
				sarray[i] = t;
				__syncthreads();
			inc *= 2;
			L /= 2;
		}
		
		sum = sarray[0] + sarray[inc];  // This contains the sum
	}
	else if(sumControl == 2)
	{	
		// Serial version of sum accessing global (not shared) memory
	
		sum = 0.0f;
		for(int i=0; i<length; i++)
		{
			sum += array[i];
		}
	}
	else
	{
		// Serial version of sum accessing shared memory
	
		sum = 0.0f;
		for(int i=0; i<length; i++)
		{
			sum += sarray[i];
		}
	}
	
	return sum;
}



// Host code
int main(void)
{
	
	int arraySize = 512;
	float* Array;
	
	Array = (float*) malloc(sizeof(float) * arraySize);
	
	// Fill array with some numbers
	
	for(int i=0; i<arraySize; i++)
	{
		Array[i] = 2.0f * (float) i;
	}
	
	// Set up device pointers
	float *devPtrArray;
	
	float Sum;
	float* devPtrSum;                            
	hipMalloc((void**)&devPtrSum, sizeof(float));
	
	// Allocate device memory
	hipMalloc((void**)&devPtrArray, arraySize * sizeof(float));
	
	// Copy array to device
	hipMemcpy(devPtrArray, Array, arraySize * sizeof(float), hipMemcpyHostToDevice);
	
	// Launch the kernel.
	
	START_GPU;     // Start timer for device code
	
	SumKernel<<<1, 512>>>(devPtrArray, devPtrSum, arraySize);
	
	STOP_GPU;           // Stop timer for device code
	PRINT_GPU;          // Print timing for device code
	CUDA_ERROR_CHECK
	
	// Copy the sum back from the GPU to the host
	
	hipMemcpy(&Sum, devPtrSum, sizeof(float), hipMemcpyDeviceToHost);
	
	printf("\nSum array elements multiplied by 2 (GPU) = %6.2f\n\n", Sum);
	 
    // Free the memory allocated on the device   
	hipFree(devPtrSum);
	hipFree(devPtrArray);
	
	// Free the memory allocated on the CPU
	free(Array);
	
    return 0;
	
}


	.text
	.file	"reducedSum.hip"
	.globl	_Z24__device_stub__SumKernelPfS_i # -- Begin function _Z24__device_stub__SumKernelPfS_i
	.type	_Z24__device_stub__SumKernelPfS_i,@function
_Z24__device_stub__SumKernelPfS_i:      # @_Z24__device_stub__SumKernelPfS_i
	.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	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%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	$_Z9SumKernelPfS_i, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	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	_Z24__device_stub__SumKernelPfS_i, .Lfunc_end0-_Z24__device_stub__SumKernelPfS_i
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	subq	$24, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	movl	$2048, %edi                     # imm = 0x800
	callq	malloc
	movq	%rax, %rbx
	xorl	%eax, %eax
.LBB1_1:                                # =>This Inner Loop Header: Depth=1
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	addss	%xmm0, %xmm0
	movss	%xmm0, (%rbx,%rax,4)
	incq	%rax
	cmpq	$512, %rax                      # imm = 0x200
	jne	.LBB1_1
# %bb.2:
	movq	%rsp, %rdi
	movl	$4, %esi
	callq	hipMalloc
	leaq	8(%rsp), %r14
	movl	$2048, %esi                     # imm = 0x800
	movq	%r14, %rdi
	callq	hipMalloc
	movq	(%r14), %rdi
	movl	$2048, %edx                     # imm = 0x800
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$start, %edi
	callq	hipEventCreate
	movl	$stop, %edi
	callq	hipEventCreate
	movq	start(%rip), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	leaq	511(%rdi), %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_4
# %bb.3:
	movq	8(%rsp), %rdi
	movq	(%rsp), %rsi
	movl	$512, %edx                      # imm = 0x200
	callq	_Z24__device_stub__SumKernelPfS_i
.LBB1_4:
	movq	stop(%rip), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	stop(%rip), %rdi
	callq	hipEventSynchronize
	movq	start(%rip), %rsi
	movq	stop(%rip), %rdx
	movl	$timeGPU, %edi
	callq	hipEventElapsedTime
	movq	start(%rip), %rdi
	callq	hipEventDestroy
	movq	stop(%rip), %rdi
	callq	hipEventDestroy
	xorps	%xmm0, %xmm0
	cvtss2sd	timeGPU(%rip), %xmm0
	movl	$.L.str, %edi
	movb	$1, %al
	callq	printf
	callq	hipDeviceSynchronize
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB1_6
# %bb.5:
	movl	$.Lstr, %edi
	callq	puts@PLT
	movq	(%rsp), %rsi
	leaq	20(%rsp), %r14
	movl	$4, %edx
	movq	%r14, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	xorps	%xmm0, %xmm0
	cvtss2sd	(%r14), %xmm0
	movl	$.L.str.3, %edi
	movb	$1, %al
	callq	printf
	movq	(%rsp), %rdi
	callq	hipFree
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	%rbx, %rdi
	callq	free
	xorl	%eax, %eax
	addq	$24, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.LBB1_6:
	.cfi_def_cfa_offset 48
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.1, %edi
	movq	%rax, %rsi
	xorl	%eax, %eax
	callq	printf
	movl	$-1, %edi
	callq	exit
.Lfunc_end1:
	.size	main, .Lfunc_end1-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB2_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB2_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z9SumKernelPfS_i, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end2:
	.size	__hip_module_ctor, .Lfunc_end2-__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	.LBB3_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
.LBB3_2:
	retq
.Lfunc_end3:
	.size	__hip_module_dtor, .Lfunc_end3-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	timeGPU,@object                 # @timeGPU
	.bss
	.globl	timeGPU
	.p2align	2, 0x0
timeGPU:
	.long	0x00000000                      # float 0
	.size	timeGPU, 4

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

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

	.type	_Z9SumKernelPfS_i,@object       # @_Z9SumKernelPfS_i
	.section	.rodata,"a",@progbits
	.globl	_Z9SumKernelPfS_i
	.p2align	3, 0x0
_Z9SumKernelPfS_i:
	.quad	_Z24__device_stub__SumKernelPfS_i
	.size	_Z9SumKernelPfS_i, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"\n\nTime to compute on GPU: %f ms \n"
	.size	.L.str, 34

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"***CUDA error: %s\n"
	.size	.L.str.1, 19

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"\nSum array elements multiplied by 2 (GPU) = %6.2f\n\n"
	.size	.L.str.3, 52

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z9SumKernelPfS_i"
	.size	.L__unnamed_1, 18

	.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	"\nNo CUDA errors detected"
	.size	.Lstr, 25

	.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 _Z24__device_stub__SumKernelPfS_i
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym timeGPU
	.addrsig_sym start
	.addrsig_sym stop
	.addrsig_sym _Z9SumKernelPfS_i
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z9SumKernelPfS_i       ; -- Begin function _Z9SumKernelPfS_i
	.globl	_Z9SumKernelPfS_i
	.p2align	8
	.type	_Z9SumKernelPfS_i,@function
_Z9SumKernelPfS_i:                      ; @_Z9SumKernelPfS_i
; %bb.0:
	s_clause 0x1
	s_load_b32 s4, s[0:1], 0x10
	s_load_b128 s[0:3], s[0:1], 0x0
	v_dual_mov_b32 v2, 0 :: v_dual_lshlrev_b32 v1, 2, v0
	ds_store_b32 v1, v2
	s_waitcnt lgkmcnt(0)
	v_cmp_gt_i32_e32 vcc_lo, s4, v0
	s_and_saveexec_b32 s4, vcc_lo
	s_cbranch_execz .LBB0_2
; %bb.1:
	global_load_b32 v2, v1, s[0:1]
	s_waitcnt vmcnt(0)
	ds_store_b32 v1, v2
.LBB0_2:
	s_or_b32 exec_lo, exec_lo, s4
	s_movk_i32 s0, 0x100
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
.LBB0_3:                                ; =>This Inner Loop Header: Depth=1
	s_and_saveexec_b32 s1, vcc_lo
	s_cbranch_execz .LBB0_5
; %bb.4:                                ;   in Loop: Header=BB0_3 Depth=1
	v_xor_b32_e32 v2, s0, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_lshlrev_b32_e32 v2, 2, v2
	ds_load_b32 v3, v1
	ds_load_b32 v2, v2
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	v_add_f32_e32 v2, v3, v2
	ds_store_b32 v1, v2
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
.LBB0_5:                                ;   in Loop: Header=BB0_3 Depth=1
	s_or_b32 exec_lo, exec_lo, s1
	s_lshr_b32 s1, s0, 1
	s_cmp_gt_u32 s0, 1
	s_mov_b32 s0, s1
	s_cbranch_scc1 .LBB0_3
; %bb.6:                                ; %_Z12reductionSumiPf.exit
	v_mov_b32_e32 v0, 0
	ds_load_b32 v1, v0
	s_waitcnt lgkmcnt(0)
	global_store_b32 v0, v1, s[2:3]
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z9SumKernelPfS_i
		.amdhsa_group_segment_fixed_size 2048
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 20
		.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 4
		.amdhsa_next_free_sgpr 5
		.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	_Z9SumKernelPfS_i, .Lfunc_end0-_Z9SumKernelPfS_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 232
; NumSgprs: 7
; NumVgprs: 4
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 2048 bytes/workgroup (compile time only)
; SGPRBlocks: 0
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 7
; NumVGPRsForWavesPerEU: 4
; 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
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
    .group_segment_fixed_size: 2048
    .kernarg_segment_align: 8
    .kernarg_segment_size: 20
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z9SumKernelPfS_i
    .private_segment_fixed_size: 0
    .sgpr_count:     7
    .sgpr_spill_count: 0
    .symbol:         _Z9SumKernelPfS_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     4
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
