/*written by Cheng Chen
parallel computing second project part 1 12/4/2016*/
#include <iostream>
#include <stdlib.h>

#include <algorithm>
#include <ctime>
#include "kernels.cuh"
using namespace std;


int randNum() {
	double ran = (double)rand();
	return ran;
}

// double timeTransfer(struct timeval start, struct timeval end){
//   double startInuSec = start.tv_sec * 1000 * 1000 + start.tv_usec;
//   double endInuSec = end.tv_sec * 1000 * 1000 + end.tv_usec;
//   return (endInuSec - startInuSec)/1000000.0;
// }

void concurrent_all(double* array, int n){
	double *h_max;
	double *h_std;
	double *h_mean;
	double *h_min;

	h_max = (double*)malloc(sizeof(double));
	h_std = (double*)malloc(sizeof(double));
	h_mean = (double*)malloc(sizeof(double));
	h_min = (double*)malloc(sizeof(double));
	dim3 gridSize = 256;
	dim3 blockSize = 256;


	double *d_max;
	double *d_min;
	double *d_mean;
	double *d_std;
	double *d_array;
	int *d_mutex,*d_mutex1, *d_mutex2;
	//cudaMalloc((void**)&d_array, n*sizeof(double));
	cudaMalloc((void**)&d_max, sizeof(double));
	cudaMalloc((void**)&d_min, sizeof(double));
	cudaMalloc((void**)&d_mean, sizeof(double));
	cudaMalloc((void**)&d_std, sizeof(double));
	cudaMalloc((void**)&d_array, n*sizeof(double));



	cudaMemset(d_max, 0, sizeof(double));
	cudaMemset(d_min, 0, sizeof(double));
	cudaMemset(d_std, 0, sizeof(double));
	cudaMemset(d_mean, 0, sizeof(double));
	cudaMalloc((void**)&d_mutex, sizeof(int));
	cudaMemset(d_mutex, 0, sizeof(int));
	cudaMalloc((void**)&d_mutex1, sizeof(int));
	cudaMemset(d_mutex1, 0, sizeof(int));
	cudaMalloc((void**)&d_mutex2, sizeof(int));
	cudaMemset(d_mutex2, 0, sizeof(int));

	cudaMemcpy(d_array, array, n*sizeof(double), cudaMemcpyHostToDevice);

	float gpu_elapsed_time;
	cudaEvent_t gpu_start, gpu_stop;
	cudaEventCreate(&gpu_start);
	cudaEventCreate(&gpu_stop);


	cudaEventRecord(gpu_start, 0);


	cout<<"--------GPU version concurrent------------"<<endl;
	find_maximum_kernel<<<gridSize, blockSize>>>(d_array, d_max, d_mutex, n);

	find_minimum_kernel<<<gridSize, blockSize>>>(d_array, d_min, d_mutex1, n);

	mean_kernel<<<gridSize, blockSize>>>(d_array, d_mean, d_mutex2, n);

	cudaMemcpy(h_mean, d_mean, sizeof(double), cudaMemcpyDeviceToHost);

	*h_mean = *h_mean / n;

	std_kernel<<<gridSize, blockSize>>>(d_array, d_std, d_mutex2, n, *h_mean);

	cudaMemcpy(h_std, d_std, sizeof(double), cudaMemcpyDeviceToHost);

	*h_std = sqrt(*h_std/n);
	//concurrent_kernel<<<gridSize, blockSize>>>(d_array, d_max, d_min, d_mean, d_mutex, n);
	//cudaMemcpy(h_mean, d_mean, sizeof(double), cudaMemcpyDeviceToHost);
	cudaEventRecord(gpu_stop, 0);
	cudaEventSynchronize(gpu_stop);
	cudaEventElapsedTime(&gpu_elapsed_time, gpu_start, gpu_stop);
	cudaEventDestroy(gpu_start);
	cudaEventDestroy(gpu_stop);


  // cudaStreamDestroy(stream1);
	// cudaStreamDestroy(stream2);
	// cudaStreamDestroy(stream3);

	cudaMemcpy(h_max, d_max, sizeof(double), cudaMemcpyDeviceToHost);
	cudaMemcpy(h_min, d_min, sizeof(double), cudaMemcpyDeviceToHost);
	//cudaMemcpy(h_std, d_std, sizeof(double), cudaMemcpyDeviceToHost);

	std::cout<<"max "<<h_max<<" min "<<h_min<<" mean "<<h_mean << " Stand_d: "<< h_std<<endl;
	std::cout<<"array size "<< n <<" time "<<gpu_elapsed_time<<endl;

	//cpu version
	cout<<"-----------cpu Version concurrent---------------------"<<endl;
	double max=array[0], min=array[0], mean = 0, standDeviation = 0;
	for(int i = 0; i < n; ++i){
		if(max < array[i]){
			max = array[i];
		}
		if(min < array[i]){
			min = array[i];
		}
		mean += array[i];
	}
	mean = mean / n;
	clock_t cpu_start = clock();
	for(int i = 0; i < n; i++){
		standDeviation+=(array[i] - mean)*(array[i] - mean);
	}
	standDeviation =sqrt( standDeviation/n);
	clock_t cpu_stop = clock();
	double cpu_elapsed_time = 1000*(cpu_stop - cpu_start)/CLOCKS_PER_SEC;

	std::cout<<"max "<<max<<" min "<<min<<" mean "<<mean << " Stand d "<<standDeviation<<endl;
	std::cout<<"array size "<< n<<" time "<<cpu_elapsed_time<<endl;
	cudaFree(d_array);
	cudaFree(d_max);
	cudaFree(d_mutex);
	free(array);

}

int main(int argc, char *argv[]) {
	//set the total number
	int totalNum = 0;
	if(argc < 2){
		totalNum = 50 * 1000 * 1000;
	}else {
		totalNum =atoi(argv[1]);
	}
	double *d_array;
	double *d_max;
	int *d_mutex;
	double *h_max;


	//allocate memory space for the random numbers
	double* nums = (double*)malloc(totalNum * sizeof(double));
	h_max = (double*)malloc(sizeof(double));
	cudaMalloc((void**)&d_array, totalNum * sizeof(double));
	cudaMalloc((void**)&d_max, sizeof(double));
	cudaMalloc((void**)&d_mutex, sizeof(int));
	cudaMemset(d_max, 0, sizeof(float));
	cudaMemset(d_mutex, 0, sizeof(int));

	//fill array with data
	srand(time(NULL));
	for (int i = 0; i < totalNum; ++i) {
		nums[i] = randNum();
	}

	//set up timing variables
	float gpu_elapased_time;
	cudaEvent_t gpu_start, gpu_stop;
	cudaEventCreate(&gpu_start);
	cudaEventCreate(&gpu_stop);

	//copy from host to device
	cudaMemcpy(d_array, nums, totalNum * sizeof(double), cudaMemcpyHostToDevice);
	dim3 gridSize = 256;
	dim3 blockSize = 256;


	//call kernel
	printf("---------find max-------------------\n");

	cudaEventRecord(gpu_start, 0);
	find_maximum_kernel <<<gridSize, blockSize>>>(d_array, d_max, d_mutex, totalNum);


	//copy from GPU to host
	cudaMemcpy(h_max, d_max, sizeof(double), cudaMemcpyDeviceToHost);
	cudaEventRecord(gpu_stop, 0);
	cudaEventSynchronize(gpu_stop);
	cudaEventElapsedTime(&gpu_elapased_time, gpu_start, gpu_stop);
	cudaEventDestroy(gpu_start);
	cudaEventDestroy(gpu_stop);
	std::cout << "size of arrays " << totalNum <<endl;
	std::cout << "Maximum number on gpu was: " << *h_max << std::endl;
	cout<<"GPU time " << gpu_elapased_time << endl;




	//cpu version
	//printf("------------CPU version find the maximum---------------\n");
	clock_t cpu_start = clock();
	double maxNum = nums[0];
	for (int i = 0; i < totalNum; i++) {
		if (maxNum < nums[i]) {
			maxNum = nums[i];
		}
	}
	clock_t cpu_stop = clock();
	clock_t cpu_elapsed_time = 1000*(cpu_stop - cpu_start)/CLOCKS_PER_SEC;
	cout << "CPU time " << cpu_elapsed_time << endl;
	cout << "the cpu max number is " << maxNum << endl;



	printf("---------find min-------------------\n");
	//allocate memory space for the random numbers
	double* h_min;
	h_min = (double*)malloc(sizeof(double));
	double* d_min;
	cudaMalloc((void**)&d_min, sizeof(double));
	cudaMemset(d_min, 0, sizeof(float));
	cudaMemset(d_mutex, 0, sizeof(int));



	cudaEventCreate(&gpu_start);
	cudaEventCreate(&gpu_stop);



	cudaEventRecord(gpu_start, 0);
	find_minimum_kernel <<<gridSize, blockSize>>>(d_array, d_min, d_mutex, totalNum);


	//copy from GPU to host
	cudaMemcpy(h_min, d_min, sizeof(double), cudaMemcpyDeviceToHost);
	cudaEventRecord(gpu_stop, 0);
	cudaEventSynchronize(gpu_stop);
	cudaEventElapsedTime(&gpu_elapased_time, gpu_start, gpu_stop);

	std::cout << "size of arrays " << totalNum << endl;
	std::cout << "the gpu min number was: " << *h_min << std::endl;

	cout<<" gpu time " << gpu_elapased_time << endl;




	//cpu version
	//printf("------------CPU version find the minimun---------------\n");
	//gettimeofday(&start, NULL);
	double minimum = nums[0];
	cpu_start = clock();
	for (int i = 0; i < totalNum; i++) {
		if (minimum < nums[i]) {
			minimum = nums[i];
		}
	}
	cpu_stop = clock();
	cpu_elapsed_time = 1000*(cpu_stop - cpu_start)/CLOCKS_PER_SEC;
	cout << "cpu time " << cpu_elapsed_time << endl;
	cout << "the cpu min number is " << minimum << endl;


		printf("--------- arithmetic mean-------------------\n");
		//allocate memory space for the random numbers
		double* h_mean;
		h_mean = (double*)malloc(sizeof(double));
		double* d_mean;
		cudaMalloc((void**)&d_mean, sizeof(double));
		cudaMemset(d_mean, 0, sizeof(float));
		cudaMemset(d_mutex, 0, sizeof(int));


		//set up timing variables
		cudaEventCreate(&gpu_start);
		cudaEventCreate(&gpu_stop);



		cudaEventRecord(gpu_start, 0);
		mean_kernel <<<gridSize, blockSize>>> (d_array, d_mean, d_mutex, totalNum);


		//copy from GPU to host
		cudaMemcpy(h_mean, d_mean, sizeof(double), cudaMemcpyDeviceToHost);
		cudaEventRecord(gpu_stop, 0);
		cudaEventSynchronize(gpu_stop);
		cudaEventElapsedTime(&gpu_elapased_time, gpu_start, gpu_stop);
		cudaEventDestroy(gpu_start);
		cudaEventDestroy(gpu_stop);
		*h_mean = *h_mean/totalNum;
		std::cout << "mean of the array calculate by gpu was: " << *h_mean<< std::endl;
		std::cout << "size of arrays " << totalNum << endl;
		cout<<"Gpu time " << gpu_elapased_time << endl;




		//cpu version
		//printf("------------CPU version find the arithmetic mean---------------\n");
		//gettimeofday(&start, NULL);
		double mean = nums[0];
		cpu_start = clock();
		for (int i = 1; i < totalNum; i++) {
				mean += nums[i];
		}
		mean = mean/totalNum;
		cpu_stop = clock();
		cpu_elapsed_time = 1000*(cpu_stop - cpu_start)/CLOCKS_PER_SEC;
		cout << "cpu time " << cpu_elapsed_time << endl;
		cout << "the  mean is " << mean << endl;

		//GPU version std
		printf("-----------------GPU version STD--------------------\n");
		double *h_std = (double*)malloc(sizeof(double));
		double *d_std;
		cudaMalloc((void**)&d_std, sizeof(double));
		cudaMemset(d_std, 0, sizeof(double));
		cudaEventCreate(&gpu_start);
		cudaEventCreate(&gpu_stop);

		cudaEventRecord(gpu_start, 0);
		std_kernel<<<gridSize, blockSize>>>(d_array, d_std, d_mutex, totalNum, *h_mean);
		cudaEventRecord(gpu_stop,0);
		cudaEventSynchronize(gpu_stop);
		cudaEventElapsedTime(&gpu_elapased_time, gpu_start, gpu_stop);
		cudaMemcpy(h_std, d_std, sizeof(double), cudaMemcpyDeviceToHost);
		cudaEventDestroy(gpu_start);
		cudaEventDestroy(gpu_stop);
		*h_std = sqrt(*h_std/totalNum);
		std::cout<<"GPU std is " << *h_std<<endl;
		std::cout<<"GPU time "<<gpu_elapased_time<<endl;

		//run the cpu version std
		//printf("------------CPU version find the STD---------------\n");
		//gettimeofday(&start, NULL);
		double stand_d= 0;
		cpu_start = clock();
		for (int i = 0; i < totalNum; i++) {
				stand_d += (nums[i]-mean)*(nums[i]-mean);
		}
		stand_d = stand_d/totalNum;
		stand_d = sqrt(stand_d);
		cpu_stop = clock();
		cpu_elapsed_time = 1000*(cpu_stop - cpu_start)/CLOCKS_PER_SEC;
		cout << "CPU time " << cpu_elapsed_time << endl;
		cout << "the  stand deviation is " << stand_d << endl;

		//free the numbers
		free(h_max);
		free(h_min);
		free(h_std);
		free(h_mean);
		//free gpu


	cudaFree(d_max);
	cudaFree(d_min);
	cudaFree(d_std);
	cudaFree(d_mean);
	cudaFree(d_array);



	concurrent_all(nums, totalNum);

}

	.file	"tmpxft_00272caf_00000000-6_main.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3894:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3894:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.align 2
	.type	_ZNSolsEPFRSoS_E.isra.0, @function
_ZNSolsEPFRSoS_E.isra.0:
.LFB4549:
	.cfi_startproc
	jmp	*%rsi
	.cfi_endproc
.LFE4549:
	.size	_ZNSolsEPFRSoS_E.isra.0, .-_ZNSolsEPFRSoS_E.isra.0
	.globl	_Z7randNumv
	.type	_Z7randNumv, @function
_Z7randNumv:
.LFB3889:
	.cfi_startproc
	endbr64
	jmp	rand@PLT
	.cfi_endproc
.LFE3889:
	.size	_Z7randNumv, .-_Z7randNumv
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"--------GPU version concurrent------------"
.LC2:
	.string	"max "
.LC3:
	.string	" min "
.LC4:
	.string	" mean "
.LC5:
	.string	" Stand_d: "
.LC6:
	.string	"array size "
.LC7:
	.string	" time "
.LC8:
	.string	"-----------cpu Version concurrent---------------------"
.LC9:
	.string	" Stand d "
	.text
	.globl	_Z14concurrent_allPdi
	.type	_Z14concurrent_allPdi, @function
_Z14concurrent_allPdi:
.LFB3890:
	.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
	movq	%rdi, %rbp
	movl	$8, %edi
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%esi, %ebx
	subq	$168, %rsp
	.cfi_def_cfa_offset 224
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	call	malloc@PLT
	movl	$8, %edi
	movq	%rax, %r15
	call	malloc@PLT
	movl	$8, %edi
	movq	%rax, %r13
	call	malloc@PLT
	movl	$8, %edi
	movq	%rax, %r12
	call	malloc@PLT
	leaq	48(%rsp), %rdi
	movl	$8, %esi
	movl	$1, 136(%rsp)
	movl	$1, 148(%rsp)
	movq	%rax, %r14
	movl	$16777217, %eax
	salq	$8, %rax
	movq	%rax, 128(%rsp)
	movq	%rax, 140(%rsp)
	call	cudaMalloc@PLT
	leaq	56(%rsp), %rdi
	movl	$8, %esi
	call	cudaMalloc@PLT
	leaq	64(%rsp), %rdi
	movl	$8, %esi
	call	cudaMalloc@PLT
	leaq	72(%rsp), %rdi
	movl	$8, %esi
	call	cudaMalloc@PLT
	movslq	%ebx, %rax
	leaq	80(%rsp), %rdi
	salq	$3, %rax
	movq	%rax, %rsi
	movq	%rax, 8(%rsp)
	call	cudaMalloc@PLT
	movq	48(%rsp), %rdi
	xorl	%esi, %esi
	movl	$8, %edx
	call	cudaMemset@PLT
	movq	56(%rsp), %rdi
	xorl	%esi, %esi
	movl	$8, %edx
	call	cudaMemset@PLT
	movq	72(%rsp), %rdi
	xorl	%esi, %esi
	movl	$8, %edx
	call	cudaMemset@PLT
	movq	64(%rsp), %rdi
	movl	$8, %edx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	leaq	88(%rsp), %rdi
	movl	$4, %esi
	call	cudaMalloc@PLT
	movq	88(%rsp), %rdi
	movl	$4, %edx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	leaq	96(%rsp), %rdi
	movl	$4, %esi
	call	cudaMalloc@PLT
	movq	96(%rsp), %rdi
	movl	$4, %edx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	leaq	104(%rsp), %rdi
	movl	$4, %esi
	call	cudaMalloc@PLT
	movq	104(%rsp), %rdi
	xorl	%esi, %esi
	movl	$4, %edx
	call	cudaMemset@PLT
	movq	8(%rsp), %rdx
	movq	80(%rsp), %rdi
	movq	%rbp, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	leaq	112(%rsp), %rdi
	call	cudaEventCreate@PLT
	leaq	120(%rsp), %rdi
	call	cudaEventCreate@PLT
	movq	112(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	leaq	.LC1(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movl	148(%rsp), %ecx
	xorl	%r9d, %r9d
	movq	140(%rsp), %rdx
	movq	128(%rsp), %rdi
	movl	136(%rsp), %esi
	xorl	%r8d, %r8d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L5
	movq	88(%rsp), %rdx
	movq	48(%rsp), %rsi
	movl	%ebx, %ecx
	movq	80(%rsp), %rdi
	call	_Z19find_maximum_kernelPdS_Pii@PLT
.L5:
	movl	148(%rsp), %ecx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	140(%rsp), %rdx
	movq	128(%rsp), %rdi
	movl	136(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L6
	movq	96(%rsp), %rdx
	movq	56(%rsp), %rsi
	movl	%ebx, %ecx
	movq	80(%rsp), %rdi
	call	_Z19find_minimum_kernelPdS_Pii@PLT
.L6:
	movl	148(%rsp), %ecx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	140(%rsp), %rdx
	movq	128(%rsp), %rdi
	movl	136(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L7
	movq	104(%rsp), %rdx
	movq	64(%rsp), %rsi
	movl	%ebx, %ecx
	movq	80(%rsp), %rdi
	call	_Z11mean_kernelPdS_Pii@PLT
.L7:
	movq	64(%rsp), %rsi
	movl	$2, %ecx
	movl	$8, %edx
	movq	%r12, %rdi
	call	cudaMemcpy@PLT
	cvtsi2sdl	%ebx, %xmm3
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movsd	(%r12), %xmm0
	movl	148(%rsp), %ecx
	movq	140(%rsp), %rdx
	movq	128(%rsp), %rdi
	movl	136(%rsp), %esi
	divsd	%xmm3, %xmm0
	movsd	%xmm3, 8(%rsp)
	movsd	%xmm0, (%r12)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L8
	movsd	(%r12), %xmm0
	movq	104(%rsp), %rdx
	movl	%ebx, %ecx
	movq	72(%rsp), %rsi
	movq	80(%rsp), %rdi
	call	_Z10std_kernelPdS_Piid@PLT
.L8:
	movq	72(%rsp), %rsi
	movl	$2, %ecx
	movl	$8, %edx
	movq	%r13, %rdi
	call	cudaMemcpy@PLT
	movsd	0(%r13), %xmm0
	divsd	8(%rsp), %xmm0
	call	sqrt@PLT
	movq	120(%rsp), %rdi
	xorl	%esi, %esi
	movsd	%xmm0, 0(%r13)
	call	cudaEventRecord@PLT
	movq	120(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	120(%rsp), %rdx
	movq	112(%rsp), %rsi
	leaq	44(%rsp), %rdi
	call	cudaEventElapsedTime@PLT
	movq	112(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movq	120(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movq	48(%rsp), %rsi
	movq	%r15, %rdi
	movl	$2, %ecx
	movl	$8, %edx
	call	cudaMemcpy@PLT
	movq	56(%rsp), %rsi
	movl	$2, %ecx
	movq	%r14, %rdi
	movl	$8, %edx
	call	cudaMemcpy@PLT
	leaq	.LC2(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%r15, %rsi
	leaq	.LC3(%rip), %r15
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIPKvEERSoT_@PLT
	movq	%r15, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%r14, %rsi
	leaq	.LC4(%rip), %r14
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIPKvEERSoT_@PLT
	movq	%r14, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%r12, %rsi
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIPKvEERSoT_@PLT
	leaq	.LC5(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%r13, %rsi
	leaq	.LC6(%rip), %r13
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIPKvEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	%r13, %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC7(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cvtss2sd	44(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC8(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movsd	0(%rbp), %xmm6
	xorl	%eax, %eax
	xorps	%xmm0, %xmm0
	movsd	%xmm6, 16(%rsp)
.L9:
	cmpl	%eax, %ebx
	jle	.L19
	movsd	0(%rbp,%rax,8), %xmm1
	incq	%rax
	movaps	%xmm1, %xmm2
	maxsd	16(%rsp), %xmm2
	addsd	%xmm1, %xmm0
	movsd	%xmm2, 16(%rsp)
	jmp	.L9
.L19:
	divsd	8(%rsp), %xmm0
	movsd	%xmm0, 24(%rsp)
	call	clock@PLT
	xorl	%edx, %edx
	xorps	%xmm0, %xmm0
	movq	%rax, %r12
.L12:
	cmpl	%edx, %ebx
	jle	.L20
	movsd	0(%rbp,%rdx,8), %xmm1
	subsd	24(%rsp), %xmm1
	incq	%rdx
	mulsd	%xmm1, %xmm1
	addsd	%xmm1, %xmm0
	jmp	.L12
.L20:
	divsd	8(%rsp), %xmm0
	call	sqrt@PLT
	movsd	%xmm0, 8(%rsp)
	call	clock@PLT
	movl	$1000, %ecx
	leaq	.LC2(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	subq	%r12, %rax
	cqto
	idivq	%rcx
	cvtsi2sdq	%rax, %xmm6
	movq	%xmm6, %r12
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	16(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	%r15, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	16(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	%r14, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	24(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	leaq	.LC9(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	8(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	%r13, %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC7(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%r12, %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	80(%rsp), %rdi
	call	cudaFree@PLT
	movq	48(%rsp), %rdi
	call	cudaFree@PLT
	movq	88(%rsp), %rdi
	call	cudaFree@PLT
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L14
	call	__stack_chk_fail@PLT
.L14:
	addq	$168, %rsp
	.cfi_def_cfa_offset 56
	movq	%rbp, %rdi
	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
	jmp	free@PLT
	.cfi_endproc
.LFE3890:
	.size	_Z14concurrent_allPdi, .-_Z14concurrent_allPdi
	.section	.rodata.str1.1
.LC10:
	.string	"---------find max-------------------\n"
.LC11:
	.string	"size of arrays "
.LC12:
	.string	"Maximum number on gpu was: "
.LC13:
	.string	"GPU time "
.LC14:
	.string	"CPU time "
.LC15:
	.string	"the cpu max number is "
.LC16:
	.string	"---------find min-------------------\n"
.LC17:
	.string	"the gpu min number was: "
.LC18:
	.string	" gpu time "
.LC19:
	.string	"cpu time "
.LC20:
	.string	"the cpu min number is "
.LC21:
	.string	"--------- arithmetic mean-------------------\n"
.LC22:
	.string	"mean of the array calculate by gpu was: "
.LC23:
	.string	"Gpu time "
.LC24:
	.string	"the  mean is "
.LC25:
	.string	"-----------------GPU version STD--------------------\n"
.LC26:
	.string	"GPU std is "
.LC27:
	.string	"the  stand deviation is "
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3891:
	.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
	movl	$50000000, %ebx
	subq	$184, %rsp
	.cfi_def_cfa_offset 240
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	decl	%edi
	jle	.L22
	movq	8(%rsi), %rdi
	call	atoi@PLT
	movl	%eax, %ebx
.L22:
	movslq	%ebx, %r12
	xorl	%r13d, %r13d
	salq	$3, %r12
	movq	%r12, %rdi
	call	malloc@PLT
	movl	$8, %edi
	movq	%rax, %rbp
	call	malloc@PLT
	leaq	80(%rsp), %rdi
	movq	%r12, %rsi
	movq	%rax, 40(%rsp)
	call	cudaMalloc@PLT
	leaq	88(%rsp), %rdi
	movl	$8, %esi
	call	cudaMalloc@PLT
	leaq	96(%rsp), %rdi
	movl	$4, %esi
	call	cudaMalloc@PLT
	movq	88(%rsp), %rdi
	movl	$4, %edx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	movq	96(%rsp), %rdi
	movl	$4, %edx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	xorl	%edi, %edi
	call	time@PLT
	movl	%eax, %edi
	call	srand@PLT
.L23:
	cmpl	%r13d, %ebx
	jle	.L43
	call	rand@PLT
	cvtsi2sdl	%eax, %xmm0
	movsd	%xmm0, 0(%rbp,%r13,8)
	incq	%r13
	jmp	.L23
.L43:
	leaq	104(%rsp), %rax
	movq	%rax, %rdi
	movq	%rax, 16(%rsp)
	call	cudaEventCreate@PLT
	leaq	112(%rsp), %rax
	movq	%rax, %rdi
	movq	%rax, 24(%rsp)
	call	cudaEventCreate@PLT
	movq	80(%rsp), %rdi
	movq	%r12, %rdx
	movq	%rbp, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movl	$16777217, %eax
	movl	$2, %edi
	leaq	.LC10(%rip), %rsi
	movl	$1, 152(%rsp)
	salq	$8, %rax
	movq	%rax, 144(%rsp)
	movq	%rax, 156(%rsp)
	xorl	%eax, %eax
	movl	$1, 164(%rsp)
	call	__printf_chk@PLT
	movq	104(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	164(%rsp), %ecx
	xorl	%r9d, %r9d
	movq	156(%rsp), %rdx
	movq	144(%rsp), %rdi
	movl	152(%rsp), %esi
	xorl	%r8d, %r8d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L25
	movq	96(%rsp), %rdx
	movq	88(%rsp), %rsi
	movl	%ebx, %ecx
	movq	80(%rsp), %rdi
	call	_Z19find_maximum_kernelPdS_Pii@PLT
.L25:
	movq	88(%rsp), %rsi
	movq	40(%rsp), %rdi
	movl	$2, %ecx
	movl	$8, %edx
	leaq	_ZSt4cout(%rip), %r12
	leaq	.LC11(%rip), %r14
	call	cudaMemcpy@PLT
	movq	112(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	112(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	112(%rsp), %rdx
	movq	104(%rsp), %rsi
	leaq	76(%rsp), %rax
	movq	%rax, %rdi
	movq	%rax, 32(%rsp)
	call	cudaEventElapsedTime@PLT
	movq	104(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movq	112(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movq	%r14, %rsi
	movq	%r12, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC12(%rip), %rsi
	movq	%r12, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	movq	40(%rsp), %rax
	movsd	(%rax), %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC13(%rip), %rsi
	movq	%r12, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cvtss2sd	76(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	call	clock@PLT
	movsd	0(%rbp), %xmm0
	movq	%rax, %r15
	xorl	%eax, %eax
.L26:
	cmpl	%eax, %ebx
	jle	.L44
	movsd	0(%rbp,%rax,8), %xmm1
	incq	%rax
	maxsd	%xmm0, %xmm1
	movaps	%xmm1, %xmm0
	jmp	.L26
.L44:
	movsd	%xmm0, 8(%rsp)
	call	clock@PLT
	leaq	.LC14(%rip), %rsi
	movq	%r12, %rdi
	movq	%rax, %r13
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	$1000, %ecx
	movq	%rax, %rdi
	movq	%r13, %rax
	subq	%r15, %rax
	cqto
	idivq	%rcx
	movq	%rax, %rsi
	call	_ZNSo9_M_insertIlEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC15(%rip), %rsi
	movq	%r12, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	8(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC16(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$8, %edi
	call	malloc@PLT
	leaq	120(%rsp), %rdi
	movl	$8, %esi
	movq	%rax, 48(%rsp)
	call	cudaMalloc@PLT
	movq	120(%rsp), %rdi
	xorl	%esi, %esi
	movl	$4, %edx
	call	cudaMemset@PLT
	movq	96(%rsp), %rdi
	movl	$4, %edx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	movq	16(%rsp), %rdi
	call	cudaEventCreate@PLT
	movq	24(%rsp), %rdi
	call	cudaEventCreate@PLT
	movq	104(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	164(%rsp), %ecx
	xorl	%r9d, %r9d
	movq	156(%rsp), %rdx
	movq	144(%rsp), %rdi
	movl	152(%rsp), %esi
	xorl	%r8d, %r8d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L29
	movq	96(%rsp), %rdx
	movq	120(%rsp), %rsi
	movl	%ebx, %ecx
	movq	80(%rsp), %rdi
	call	_Z19find_minimum_kernelPdS_Pii@PLT
.L29:
	movq	120(%rsp), %rsi
	movq	48(%rsp), %rdi
	movl	$2, %ecx
	movl	$8, %edx
	call	cudaMemcpy@PLT
	movq	112(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	112(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	112(%rsp), %rdx
	movq	104(%rsp), %rsi
	movq	32(%rsp), %rdi
	call	cudaEventElapsedTime@PLT
	movq	%r14, %rsi
	movq	%r12, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC17(%rip), %rsi
	movq	%r12, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	movq	48(%rsp), %rax
	movsd	(%rax), %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC18(%rip), %rsi
	movq	%r12, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cvtss2sd	76(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movsd	0(%rbp), %xmm0
	movsd	%xmm0, 8(%rsp)
	call	clock@PLT
	movsd	8(%rsp), %xmm0
	movq	%rax, %r15
	xorl	%eax, %eax
.L30:
	cmpl	%eax, %ebx
	jle	.L45
	movsd	0(%rbp,%rax,8), %xmm1
	incq	%rax
	maxsd	%xmm0, %xmm1
	movaps	%xmm1, %xmm0
	jmp	.L30
.L45:
	leaq	_ZSt4cout(%rip), %r13
	movsd	%xmm0, 8(%rsp)
	call	clock@PLT
	leaq	.LC19(%rip), %rsi
	movq	%r13, %rdi
	movq	%rax, %r12
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	$1000, %ecx
	movq	%rax, %rdi
	movq	%r12, %rax
	subq	%r15, %rax
	cqto
	idivq	%rcx
	movq	%rax, %rsi
	call	_ZNSo9_M_insertIlEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC20(%rip), %rsi
	movq	%r13, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	8(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC21(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$8, %edi
	call	malloc@PLT
	leaq	128(%rsp), %rdi
	movl	$8, %esi
	movq	%rax, %r12
	call	cudaMalloc@PLT
	movq	128(%rsp), %rdi
	xorl	%esi, %esi
	movl	$4, %edx
	call	cudaMemset@PLT
	movq	96(%rsp), %rdi
	movl	$4, %edx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	movq	16(%rsp), %rdi
	call	cudaEventCreate@PLT
	movq	24(%rsp), %rdi
	call	cudaEventCreate@PLT
	movq	104(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	164(%rsp), %ecx
	xorl	%r9d, %r9d
	movq	156(%rsp), %rdx
	movq	144(%rsp), %rdi
	movl	152(%rsp), %esi
	xorl	%r8d, %r8d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L33
	movq	96(%rsp), %rdx
	movq	128(%rsp), %rsi
	movl	%ebx, %ecx
	movq	80(%rsp), %rdi
	call	_Z11mean_kernelPdS_Pii@PLT
.L33:
	movq	128(%rsp), %rsi
	movl	$2, %ecx
	movl	$8, %edx
	movq	%r12, %rdi
	call	cudaMemcpy@PLT
	movq	112(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	112(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	112(%rsp), %rdx
	movq	104(%rsp), %rsi
	movq	32(%rsp), %rdi
	call	cudaEventElapsedTime@PLT
	movq	104(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movq	112(%rsp), %rdi
	call	cudaEventDestroy@PLT
	cvtsi2sdl	%ebx, %xmm2
	movsd	(%r12), %xmm0
	movq	%r13, %rdi
	leaq	.LC22(%rip), %rsi
	divsd	%xmm2, %xmm0
	movsd	%xmm2, 8(%rsp)
	movsd	%xmm0, (%r12)
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	(%r12), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	%r14, %rsi
	movq	%r13, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC23(%rip), %rsi
	movq	%r13, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cvtss2sd	76(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movsd	0(%rbp), %xmm0
	movsd	%xmm0, 56(%rsp)
	call	clock@PLT
	movsd	56(%rsp), %xmm0
	movq	%rax, %r14
	xorl	%eax, %eax
.L34:
	incq	%rax
	cmpl	%eax, %ebx
	jle	.L46
	addsd	0(%rbp,%rax,8), %xmm0
	jmp	.L34
.L46:
	divsd	8(%rsp), %xmm0
	movsd	%xmm0, 56(%rsp)
	call	clock@PLT
	leaq	.LC19(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	movq	%rax, %r13
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	$1000, %ecx
	movq	%rax, %rdi
	movq	%r13, %rax
	subq	%r14, %rax
	cqto
	idivq	%rcx
	movq	%rax, %rsi
	call	_ZNSo9_M_insertIlEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC24(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	56(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC25(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$8, %edi
	call	malloc@PLT
	leaq	136(%rsp), %rdi
	movl	$8, %esi
	movq	%rax, %r14
	call	cudaMalloc@PLT
	movq	136(%rsp), %rdi
	movl	$8, %edx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	movq	16(%rsp), %rdi
	call	cudaEventCreate@PLT
	movq	24(%rsp), %rdi
	call	cudaEventCreate@PLT
	movq	104(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	164(%rsp), %ecx
	xorl	%r9d, %r9d
	movq	156(%rsp), %rdx
	movq	144(%rsp), %rdi
	movl	152(%rsp), %esi
	xorl	%r8d, %r8d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L36
	movsd	(%r12), %xmm0
	movq	96(%rsp), %rdx
	movl	%ebx, %ecx
	movq	136(%rsp), %rsi
	movq	80(%rsp), %rdi
	call	_Z10std_kernelPdS_Piid@PLT
.L36:
	movq	112(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	112(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	112(%rsp), %rdx
	movq	104(%rsp), %rsi
	movq	32(%rsp), %rdi
	call	cudaEventElapsedTime@PLT
	movq	136(%rsp), %rsi
	movl	$2, %ecx
	movq	%r14, %rdi
	movl	$8, %edx
	call	cudaMemcpy@PLT
	movq	104(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movq	112(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movsd	(%r14), %xmm0
	divsd	8(%rsp), %xmm0
	call	sqrt@PLT
	leaq	.LC26(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	movsd	%xmm0, (%r14)
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	(%r14), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC13(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cvtss2sd	76(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	call	clock@PLT
	xorps	%xmm0, %xmm0
	movq	%rax, %r15
	xorl	%eax, %eax
.L37:
	cmpl	%eax, %ebx
	jle	.L47
	movsd	0(%rbp,%rax,8), %xmm1
	subsd	56(%rsp), %xmm1
	incq	%rax
	mulsd	%xmm1, %xmm1
	addsd	%xmm1, %xmm0
	jmp	.L37
.L47:
	divsd	8(%rsp), %xmm0
	call	sqrt@PLT
	movsd	%xmm0, 8(%rsp)
	call	clock@PLT
	leaq	.LC14(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	movq	%rax, %r13
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	$1000, %ecx
	movq	%rax, %rdi
	movq	%r13, %rax
	subq	%r15, %rax
	cqto
	idivq	%rcx
	movq	%rax, %rsi
	call	_ZNSo9_M_insertIlEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC27(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	8(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	40(%rsp), %rdi
	call	free@PLT
	movq	48(%rsp), %rdi
	call	free@PLT
	movq	%r14, %rdi
	call	free@PLT
	movq	%r12, %rdi
	call	free@PLT
	movq	88(%rsp), %rdi
	call	cudaFree@PLT
	movq	120(%rsp), %rdi
	call	cudaFree@PLT
	movq	136(%rsp), %rdi
	call	cudaFree@PLT
	movq	128(%rsp), %rdi
	call	cudaFree@PLT
	movq	80(%rsp), %rdi
	call	cudaFree@PLT
	movl	%ebx, %esi
	movq	%rbp, %rdi
	call	_Z14concurrent_allPdi
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L39
	call	__stack_chk_fail@PLT
.L39:
	addq	$184, %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
.LFE3891:
	.size	main, .-main
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3917:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	movq	%rax, %rdi
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE3917:
	.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
	.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:

/*written by Cheng Chen
parallel computing second project part 1 12/4/2016*/

#include <hip/hip_runtime.h>
#include <iostream>
#include <stdlib.h>

#include <algorithm>
#include <ctime>
#include "kernels.cuh"
using namespace std;


int randNum() {
	double ran = (double)rand();
	return ran;
}

// double timeTransfer(struct timeval start, struct timeval end){
//   double startInuSec = start.tv_sec * 1000 * 1000 + start.tv_usec;
//   double endInuSec = end.tv_sec * 1000 * 1000 + end.tv_usec;
//   return (endInuSec - startInuSec)/1000000.0;
// }

void concurrent_all(double* array, int n){
	double *h_max;
	double *h_std;
	double *h_mean;
	double *h_min;

	h_max = (double*)malloc(sizeof(double));
	h_std = (double*)malloc(sizeof(double));
	h_mean = (double*)malloc(sizeof(double));
	h_min = (double*)malloc(sizeof(double));
	dim3 gridSize = 256;
	dim3 blockSize = 256;


	double *d_max;
	double *d_min;
	double *d_mean;
	double *d_std;
	double *d_array;
	int *d_mutex,*d_mutex1, *d_mutex2;
	//cudaMalloc((void**)&d_array, n*sizeof(double));
	hipMalloc((void**)&d_max, sizeof(double));
	hipMalloc((void**)&d_min, sizeof(double));
	hipMalloc((void**)&d_mean, sizeof(double));
	hipMalloc((void**)&d_std, sizeof(double));
	hipMalloc((void**)&d_array, n*sizeof(double));



	hipMemset(d_max, 0, sizeof(double));
	hipMemset(d_min, 0, sizeof(double));
	hipMemset(d_std, 0, sizeof(double));
	hipMemset(d_mean, 0, sizeof(double));
	hipMalloc((void**)&d_mutex, sizeof(int));
	hipMemset(d_mutex, 0, sizeof(int));
	hipMalloc((void**)&d_mutex1, sizeof(int));
	hipMemset(d_mutex1, 0, sizeof(int));
	hipMalloc((void**)&d_mutex2, sizeof(int));
	hipMemset(d_mutex2, 0, sizeof(int));

	hipMemcpy(d_array, array, n*sizeof(double), hipMemcpyHostToDevice);

	float gpu_elapsed_time;
	hipEvent_t gpu_start, gpu_stop;
	hipEventCreate(&gpu_start);
	hipEventCreate(&gpu_stop);


	hipEventRecord(gpu_start, 0);


	cout<<"--------GPU version concurrent------------"<<endl;
	find_maximum_kernel<<<gridSize, blockSize>>>(d_array, d_max, d_mutex, n);

	find_minimum_kernel<<<gridSize, blockSize>>>(d_array, d_min, d_mutex1, n);

	mean_kernel<<<gridSize, blockSize>>>(d_array, d_mean, d_mutex2, n);

	hipMemcpy(h_mean, d_mean, sizeof(double), hipMemcpyDeviceToHost);

	*h_mean = *h_mean / n;

	std_kernel<<<gridSize, blockSize>>>(d_array, d_std, d_mutex2, n, *h_mean);

	hipMemcpy(h_std, d_std, sizeof(double), hipMemcpyDeviceToHost);

	*h_std = sqrt(*h_std/n);
	//concurrent_kernel<<<gridSize, blockSize>>>(d_array, d_max, d_min, d_mean, d_mutex, n);
	//cudaMemcpy(h_mean, d_mean, sizeof(double), cudaMemcpyDeviceToHost);
	hipEventRecord(gpu_stop, 0);
	hipEventSynchronize(gpu_stop);
	hipEventElapsedTime(&gpu_elapsed_time, gpu_start, gpu_stop);
	hipEventDestroy(gpu_start);
	hipEventDestroy(gpu_stop);


  // cudaStreamDestroy(stream1);
	// cudaStreamDestroy(stream2);
	// cudaStreamDestroy(stream3);

	hipMemcpy(h_max, d_max, sizeof(double), hipMemcpyDeviceToHost);
	hipMemcpy(h_min, d_min, sizeof(double), hipMemcpyDeviceToHost);
	//cudaMemcpy(h_std, d_std, sizeof(double), cudaMemcpyDeviceToHost);

	std::cout<<"max "<<h_max<<" min "<<h_min<<" mean "<<h_mean << " Stand_d: "<< h_std<<endl;
	std::cout<<"array size "<< n <<" time "<<gpu_elapsed_time<<endl;

	//cpu version
	cout<<"-----------cpu Version concurrent---------------------"<<endl;
	double max=array[0], min=array[0], mean = 0, standDeviation = 0;
	for(int i = 0; i < n; ++i){
		if(max < array[i]){
			max = array[i];
		}
		if(min < array[i]){
			min = array[i];
		}
		mean += array[i];
	}
	mean = mean / n;
	clock_t cpu_start = clock();
	for(int i = 0; i < n; i++){
		standDeviation+=(array[i] - mean)*(array[i] - mean);
	}
	standDeviation =sqrt( standDeviation/n);
	clock_t cpu_stop = clock();
	double cpu_elapsed_time = 1000*(cpu_stop - cpu_start)/CLOCKS_PER_SEC;

	std::cout<<"max "<<max<<" min "<<min<<" mean "<<mean << " Stand d "<<standDeviation<<endl;
	std::cout<<"array size "<< n<<" time "<<cpu_elapsed_time<<endl;
	hipFree(d_array);
	hipFree(d_max);
	hipFree(d_mutex);
	free(array);

}

int main(int argc, char *argv[]) {
	//set the total number
	int totalNum = 0;
	if(argc < 2){
		totalNum = 50 * 1000 * 1000;
	}else {
		totalNum =atoi(argv[1]);
	}
	double *d_array;
	double *d_max;
	int *d_mutex;
	double *h_max;


	//allocate memory space for the random numbers
	double* nums = (double*)malloc(totalNum * sizeof(double));
	h_max = (double*)malloc(sizeof(double));
	hipMalloc((void**)&d_array, totalNum * sizeof(double));
	hipMalloc((void**)&d_max, sizeof(double));
	hipMalloc((void**)&d_mutex, sizeof(int));
	hipMemset(d_max, 0, sizeof(float));
	hipMemset(d_mutex, 0, sizeof(int));

	//fill array with data
	srand(time(NULL));
	for (int i = 0; i < totalNum; ++i) {
		nums[i] = randNum();
	}

	//set up timing variables
	float gpu_elapased_time;
	hipEvent_t gpu_start, gpu_stop;
	hipEventCreate(&gpu_start);
	hipEventCreate(&gpu_stop);

	//copy from host to device
	hipMemcpy(d_array, nums, totalNum * sizeof(double), hipMemcpyHostToDevice);
	dim3 gridSize = 256;
	dim3 blockSize = 256;


	//call kernel
	printf("---------find max-------------------\n");

	hipEventRecord(gpu_start, 0);
	find_maximum_kernel <<<gridSize, blockSize>>>(d_array, d_max, d_mutex, totalNum);


	//copy from GPU to host
	hipMemcpy(h_max, d_max, sizeof(double), hipMemcpyDeviceToHost);
	hipEventRecord(gpu_stop, 0);
	hipEventSynchronize(gpu_stop);
	hipEventElapsedTime(&gpu_elapased_time, gpu_start, gpu_stop);
	hipEventDestroy(gpu_start);
	hipEventDestroy(gpu_stop);
	std::cout << "size of arrays " << totalNum <<endl;
	std::cout << "Maximum number on gpu was: " << *h_max << std::endl;
	cout<<"GPU time " << gpu_elapased_time << endl;




	//cpu version
	//printf("------------CPU version find the maximum---------------\n");
	clock_t cpu_start = clock();
	double maxNum = nums[0];
	for (int i = 0; i < totalNum; i++) {
		if (maxNum < nums[i]) {
			maxNum = nums[i];
		}
	}
	clock_t cpu_stop = clock();
	clock_t cpu_elapsed_time = 1000*(cpu_stop - cpu_start)/CLOCKS_PER_SEC;
	cout << "CPU time " << cpu_elapsed_time << endl;
	cout << "the cpu max number is " << maxNum << endl;



	printf("---------find min-------------------\n");
	//allocate memory space for the random numbers
	double* h_min;
	h_min = (double*)malloc(sizeof(double));
	double* d_min;
	hipMalloc((void**)&d_min, sizeof(double));
	hipMemset(d_min, 0, sizeof(float));
	hipMemset(d_mutex, 0, sizeof(int));



	hipEventCreate(&gpu_start);
	hipEventCreate(&gpu_stop);



	hipEventRecord(gpu_start, 0);
	find_minimum_kernel <<<gridSize, blockSize>>>(d_array, d_min, d_mutex, totalNum);


	//copy from GPU to host
	hipMemcpy(h_min, d_min, sizeof(double), hipMemcpyDeviceToHost);
	hipEventRecord(gpu_stop, 0);
	hipEventSynchronize(gpu_stop);
	hipEventElapsedTime(&gpu_elapased_time, gpu_start, gpu_stop);

	std::cout << "size of arrays " << totalNum << endl;
	std::cout << "the gpu min number was: " << *h_min << std::endl;

	cout<<" gpu time " << gpu_elapased_time << endl;




	//cpu version
	//printf("------------CPU version find the minimun---------------\n");
	//gettimeofday(&start, NULL);
	double minimum = nums[0];
	cpu_start = clock();
	for (int i = 0; i < totalNum; i++) {
		if (minimum < nums[i]) {
			minimum = nums[i];
		}
	}
	cpu_stop = clock();
	cpu_elapsed_time = 1000*(cpu_stop - cpu_start)/CLOCKS_PER_SEC;
	cout << "cpu time " << cpu_elapsed_time << endl;
	cout << "the cpu min number is " << minimum << endl;


		printf("--------- arithmetic mean-------------------\n");
		//allocate memory space for the random numbers
		double* h_mean;
		h_mean = (double*)malloc(sizeof(double));
		double* d_mean;
		hipMalloc((void**)&d_mean, sizeof(double));
		hipMemset(d_mean, 0, sizeof(float));
		hipMemset(d_mutex, 0, sizeof(int));


		//set up timing variables
		hipEventCreate(&gpu_start);
		hipEventCreate(&gpu_stop);



		hipEventRecord(gpu_start, 0);
		mean_kernel <<<gridSize, blockSize>>> (d_array, d_mean, d_mutex, totalNum);


		//copy from GPU to host
		hipMemcpy(h_mean, d_mean, sizeof(double), hipMemcpyDeviceToHost);
		hipEventRecord(gpu_stop, 0);
		hipEventSynchronize(gpu_stop);
		hipEventElapsedTime(&gpu_elapased_time, gpu_start, gpu_stop);
		hipEventDestroy(gpu_start);
		hipEventDestroy(gpu_stop);
		*h_mean = *h_mean/totalNum;
		std::cout << "mean of the array calculate by gpu was: " << *h_mean<< std::endl;
		std::cout << "size of arrays " << totalNum << endl;
		cout<<"Gpu time " << gpu_elapased_time << endl;




		//cpu version
		//printf("------------CPU version find the arithmetic mean---------------\n");
		//gettimeofday(&start, NULL);
		double mean = nums[0];
		cpu_start = clock();
		for (int i = 1; i < totalNum; i++) {
				mean += nums[i];
		}
		mean = mean/totalNum;
		cpu_stop = clock();
		cpu_elapsed_time = 1000*(cpu_stop - cpu_start)/CLOCKS_PER_SEC;
		cout << "cpu time " << cpu_elapsed_time << endl;
		cout << "the  mean is " << mean << endl;

		//GPU version std
		printf("-----------------GPU version STD--------------------\n");
		double *h_std = (double*)malloc(sizeof(double));
		double *d_std;
		hipMalloc((void**)&d_std, sizeof(double));
		hipMemset(d_std, 0, sizeof(double));
		hipEventCreate(&gpu_start);
		hipEventCreate(&gpu_stop);

		hipEventRecord(gpu_start, 0);
		std_kernel<<<gridSize, blockSize>>>(d_array, d_std, d_mutex, totalNum, *h_mean);
		hipEventRecord(gpu_stop,0);
		hipEventSynchronize(gpu_stop);
		hipEventElapsedTime(&gpu_elapased_time, gpu_start, gpu_stop);
		hipMemcpy(h_std, d_std, sizeof(double), hipMemcpyDeviceToHost);
		hipEventDestroy(gpu_start);
		hipEventDestroy(gpu_stop);
		*h_std = sqrt(*h_std/totalNum);
		std::cout<<"GPU std is " << *h_std<<endl;
		std::cout<<"GPU time "<<gpu_elapased_time<<endl;

		//run the cpu version std
		//printf("------------CPU version find the STD---------------\n");
		//gettimeofday(&start, NULL);
		double stand_d= 0;
		cpu_start = clock();
		for (int i = 0; i < totalNum; i++) {
				stand_d += (nums[i]-mean)*(nums[i]-mean);
		}
		stand_d = stand_d/totalNum;
		stand_d = sqrt(stand_d);
		cpu_stop = clock();
		cpu_elapsed_time = 1000*(cpu_stop - cpu_start)/CLOCKS_PER_SEC;
		cout << "CPU time " << cpu_elapsed_time << endl;
		cout << "the  stand deviation is " << stand_d << endl;

		//free the numbers
		free(h_max);
		free(h_min);
		free(h_std);
		free(h_mean);
		//free gpu


	hipFree(d_max);
	hipFree(d_min);
	hipFree(d_std);
	hipFree(d_mean);
	hipFree(d_array);



	concurrent_all(nums, totalNum);

}

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

                                        # End of file scope inline assembly
	.globl	_Z7randNumv                     # -- Begin function _Z7randNumv
	.type	_Z7randNumv,@function
_Z7randNumv:                            # @_Z7randNumv
	.cfi_startproc
# %bb.0:
	jmp	rand                            # TAILCALL
.Lfunc_end0:
	.size	_Z7randNumv, .Lfunc_end0-_Z7randNumv
	.cfi_endproc
                                        # -- End function
	.globl	_Z14concurrent_allPdi           # -- Begin function _Z14concurrent_allPdi
	.type	_Z14concurrent_allPdi,@function
_Z14concurrent_allPdi:                  # @_Z14concurrent_allPdi
	.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	$136, %rsp
	.cfi_def_cfa_offset 192
	.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	%esi, %ebp
	movq	%rdi, %rbx
	movl	$8, %edi
	callq	malloc
	movq	%rax, 24(%rsp)                  # 8-byte Spill
	movl	$8, %edi
	callq	malloc
	movq	%rax, 16(%rsp)                  # 8-byte Spill
	movl	$8, %edi
	callq	malloc
	movq	%rax, 8(%rsp)                   # 8-byte Spill
	movl	$8, %edi
	callq	malloc
	movq	%rax, 48(%rsp)                  # 8-byte Spill
	leaq	72(%rsp), %r14
	movl	$8, %esi
	movq	%r14, %rdi
	callq	hipMalloc
	leaq	112(%rsp), %r15
	movl	$8, %esi
	movq	%r15, %rdi
	callq	hipMalloc
	leaq	64(%rsp), %rdi
	movl	$8, %esi
	callq	hipMalloc
	leaq	104(%rsp), %r12
	movl	$8, %esi
	movq	%r12, %rdi
	callq	hipMalloc
	movslq	%ebp, %r13
	shlq	$3, %r13
	leaq	40(%rsp), %rdi
	movq	%r13, %rsi
	callq	hipMalloc
	movq	(%r14), %rdi
	movl	$8, %edx
	xorl	%esi, %esi
	callq	hipMemset
	movq	(%r15), %rdi
	movabsq	$4294967552, %r15               # imm = 0x100000100
	movl	$8, %edx
	xorl	%esi, %esi
	callq	hipMemset
	movq	(%r12), %rdi
	movl	$8, %edx
	xorl	%esi, %esi
	callq	hipMemset
	leaq	64(%rsp), %rax
	movq	(%rax), %rdi
	movl	$8, %edx
	xorl	%esi, %esi
	callq	hipMemset
	leaq	96(%rsp), %r14
	movl	$4, %esi
	movq	%r14, %rdi
	callq	hipMalloc
	movq	(%r14), %rdi
	movl	$4, %edx
	xorl	%esi, %esi
	callq	hipMemset
	leaq	120(%rsp), %r14
	movl	$4, %esi
	movq	%r14, %rdi
	callq	hipMalloc
	movq	(%r14), %rdi
	movl	$4, %edx
	xorl	%esi, %esi
	callq	hipMemset
	leaq	88(%rsp), %r14
	movl	$4, %esi
	movq	%r14, %rdi
	callq	hipMalloc
	movq	(%r14), %rdi
	movl	$4, %edx
	xorl	%esi, %esi
	callq	hipMemset
	leaq	40(%rsp), %rax
	movq	(%rax), %rdi
	movq	%rbx, %rsi
	movq	%r13, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	80(%rsp), %r14
	movq	%r14, %rdi
	callq	hipEventCreate
	leaq	56(%rsp), %rdi
	callq	hipEventCreate
	movq	(%r14), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movl	$_ZSt4cout, %r14d
	movl	$_ZSt4cout, %edi
	movl	$.L.str, %esi
	movl	$42, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r14, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	%r15, %rdi
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_2
# %bb.1:
	movq	40(%rsp), %rdi
	movq	72(%rsp), %rsi
	movq	96(%rsp), %rdx
	movl	%ebp, %ecx
	callq	_Z34__device_stub__find_maximum_kernelPdS_Pii
.LBB1_2:
	movq	%r15, %rdi
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_4
# %bb.3:
	movq	40(%rsp), %rdi
	movq	112(%rsp), %rsi
	movq	120(%rsp), %rdx
	movl	%ebp, %ecx
	callq	_Z34__device_stub__find_minimum_kernelPdS_Pii
.LBB1_4:
	movq	%r15, %rdi
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	movq	16(%rsp), %r13                  # 8-byte Reload
	jne	.LBB1_6
# %bb.5:
	movq	40(%rsp), %rdi
	movq	64(%rsp), %rsi
	movq	88(%rsp), %rdx
	movl	%ebp, %ecx
	callq	_Z26__device_stub__mean_kernelPdS_Pii
.LBB1_6:
	movq	64(%rsp), %rsi
	movl	$8, %edx
	movq	8(%rsp), %r14                   # 8-byte Reload
	movq	%r14, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movsd	(%r14), %xmm0                   # xmm0 = mem[0],zero
	cvtsi2sd	%ebp, %xmm1
	movsd	%xmm1, 32(%rsp)                 # 8-byte Spill
	divsd	%xmm1, %xmm0
	movsd	%xmm0, (%r14)
	movq	%r15, %rdi
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_8
# %bb.7:
	movq	40(%rsp), %rdi
	movq	104(%rsp), %rsi
	movq	88(%rsp), %rdx
	movq	8(%rsp), %rax                   # 8-byte Reload
	movsd	(%rax), %xmm0                   # xmm0 = mem[0],zero
	movl	%ebp, %ecx
	callq	_Z25__device_stub__std_kernelPdS_Piid
.LBB1_8:
	movq	104(%rsp), %rsi
	movl	$8, %edx
	movq	%r13, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movsd	(%r13), %xmm0                   # xmm0 = mem[0],zero
	divsd	32(%rsp), %xmm0                 # 8-byte Folded Reload
	xorpd	%xmm1, %xmm1
	ucomisd	%xmm1, %xmm0
	jb	.LBB1_10
# %bb.9:
	sqrtsd	%xmm0, %xmm0
	jmp	.LBB1_11
.LBB1_10:                               # %call.sqrt
	callq	sqrt
.LBB1_11:                               # %.split
	movsd	%xmm0, (%r13)
	movq	56(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	56(%rsp), %rdi
	callq	hipEventSynchronize
	movq	80(%rsp), %rsi
	movq	56(%rsp), %rdx
	leaq	132(%rsp), %r12
	movq	%r12, %rdi
	callq	hipEventElapsedTime
	movq	80(%rsp), %rdi
	callq	hipEventDestroy
	movq	56(%rsp), %rdi
	callq	hipEventDestroy
	movq	72(%rsp), %rsi
	movl	$8, %edx
	movq	24(%rsp), %r14                  # 8-byte Reload
	movq	%r14, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	112(%rsp), %rsi
	movl	$8, %edx
	movq	48(%rsp), %r15                  # 8-byte Reload
	movq	%r15, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$_ZSt4cout, %edi
	movl	$.L.str.1, %esi
	movl	$4, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movq	%r14, %rsi
	callq	_ZNSo9_M_insertIPKvEERSoT_
	movq	%rax, %r14
	movl	$.L.str.2, %esi
	movl	$5, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%r14, %rdi
	movq	%r15, %rsi
	callq	_ZNSo9_M_insertIPKvEERSoT_
	movq	%rax, %r14
	movl	$.L.str.3, %esi
	movl	$6, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%r14, %rdi
	movq	8(%rsp), %rsi                   # 8-byte Reload
	callq	_ZNSo9_M_insertIPKvEERSoT_
	movq	%rax, %r14
	movl	$.L.str.4, %esi
	movl	$10, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%r14, %rdi
	movq	%r13, %rsi
	callq	_ZNSo9_M_insertIPKvEERSoT_
	movq	%rax, %r14
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r14, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r14, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.5, %esi
	movl	$11, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movl	%ebp, %esi
	callq	_ZNSolsEi
	movq	%rax, %r14
	movl	$.L.str.6, %esi
	movl	$6, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	cvtss2sd	(%r12), %xmm0
	movq	%r14, %rdi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r14
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r14, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r14, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.7, %esi
	movl	$54, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	movl	$_ZSt4cout, %eax
	addq	%rax, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movsd	(%rbx), %xmm3                   # xmm3 = mem[0],zero
	testl	%ebp, %ebp
	jle	.LBB1_12
# %bb.19:                               # %.lr.ph.preheader
	movl	%ebp, %r15d
	xorpd	%xmm2, %xmm2
	xorl	%eax, %eax
	movapd	%xmm3, %xmm4
.LBB1_20:                               # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movapd	%xmm4, %xmm0
	movapd	%xmm3, %xmm1
	movsd	(%rbx,%rax,8), %xmm4            # xmm4 = mem[0],zero
	movapd	%xmm4, %xmm3
	maxsd	%xmm1, %xmm3
	addsd	%xmm4, %xmm2
	maxsd	%xmm0, %xmm4
	incq	%rax
	cmpq	%rax, %r15
	jne	.LBB1_20
# %bb.15:                               # %._crit_edge
	movsd	%xmm4, 24(%rsp)                 # 8-byte Spill
	movsd	%xmm3, 16(%rsp)                 # 8-byte Spill
	movsd	%xmm2, 8(%rsp)                  # 8-byte Spill
	callq	clock
	movsd	8(%rsp), %xmm2                  # 8-byte Reload
                                        # xmm2 = mem[0],zero
	movq	%rax, %r14
	movsd	32(%rsp), %xmm1                 # 8-byte Reload
                                        # xmm1 = mem[0],zero
	divsd	%xmm1, %xmm2
	xorps	%xmm0, %xmm0
	testl	%ebp, %ebp
	jle	.LBB1_18
# %bb.16:                               # %.lr.ph128.preheader
	xorl	%eax, %eax
.LBB1_17:                               # %.lr.ph128
                                        # =>This Inner Loop Header: Depth=1
	movsd	(%rbx,%rax,8), %xmm3            # xmm3 = mem[0],zero
	subsd	%xmm2, %xmm3
	mulsd	%xmm3, %xmm3
	addsd	%xmm3, %xmm0
	incq	%rax
	cmpq	%rax, %r15
	jne	.LBB1_17
.LBB1_18:
	movsd	16(%rsp), %xmm4                 # 8-byte Reload
                                        # xmm4 = mem[0],zero
	movsd	24(%rsp), %xmm3                 # 8-byte Reload
                                        # xmm3 = mem[0],zero
	jmp	.LBB1_13
.LBB1_12:                               # %._crit_edge.thread
	xorps	%xmm0, %xmm0
	divsd	32(%rsp), %xmm0                 # 8-byte Folded Reload
	movsd	%xmm0, 8(%rsp)                  # 8-byte Spill
	movsd	%xmm3, 16(%rsp)                 # 8-byte Spill
	callq	clock
	xorpd	%xmm0, %xmm0
	movsd	16(%rsp), %xmm4                 # 8-byte Reload
                                        # xmm4 = mem[0],zero
	movsd	8(%rsp), %xmm2                  # 8-byte Reload
                                        # xmm2 = mem[0],zero
	movsd	32(%rsp), %xmm1                 # 8-byte Reload
                                        # xmm1 = mem[0],zero
	movq	%rax, %r14
	movaps	%xmm4, %xmm3
.LBB1_13:                               # %._crit_edge129
	movsd	%xmm3, 24(%rsp)                 # 8-byte Spill
	movsd	%xmm4, 16(%rsp)                 # 8-byte Spill
	movsd	%xmm2, 8(%rsp)                  # 8-byte Spill
	divsd	%xmm1, %xmm0
	xorpd	%xmm1, %xmm1
	ucomisd	%xmm1, %xmm0
	jb	.LBB1_21
# %bb.14:
	sqrtsd	%xmm0, %xmm0
	jmp	.LBB1_22
.LBB1_21:                               # %call.sqrt148
	callq	sqrt
.LBB1_22:                               # %._crit_edge129.split
	movsd	%xmm0, 48(%rsp)                 # 8-byte Spill
	callq	clock
	subq	%r14, %rax
	movabsq	$2361183241434822607, %rcx      # imm = 0x20C49BA5E353F7CF
	imulq	%rcx
	movq	%rdx, %rax
	shrq	$63, %rax
	sarq	$7, %rdx
	addq	%rax, %rdx
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rdx, %xmm0
	movsd	%xmm0, 32(%rsp)                 # 8-byte Spill
	movl	$_ZSt4cout, %edi
	movl	$.L.str.1, %esi
	movl	$4, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movsd	16(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r14
	movl	$.L.str.2, %esi
	movl	$5, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%r14, %rdi
	movsd	24(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r14
	movl	$.L.str.3, %esi
	movl	$6, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%r14, %rdi
	movsd	8(%rsp), %xmm0                  # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r14
	movl	$.L.str.8, %esi
	movl	$9, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%r14, %rdi
	movsd	48(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r14
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r14, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r14, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.5, %esi
	movl	$11, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movl	%ebp, %esi
	callq	_ZNSolsEi
	movq	%rax, %r14
	movl	$.L.str.6, %esi
	movl	$6, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%r14, %rdi
	movsd	32(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r14
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r14, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r14, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	40(%rsp), %rdi
	callq	hipFree
	movq	72(%rsp), %rdi
	callq	hipFree
	movq	96(%rsp), %rdi
	callq	hipFree
	movq	%rbx, %rdi
	callq	free
	addq	$136, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z14concurrent_allPdi, .Lfunc_end1-_Z14concurrent_allPdi
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$120, %rsp
	.cfi_def_cfa_offset 176
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	$50000000, %ebx                 # imm = 0x2FAF080
	cmpl	$2, %edi
	jl	.LBB2_2
# %bb.1:
	movq	8(%rsi), %rdi
	callq	atoi
	movl	%eax, %ebx
.LBB2_2:
	movslq	%ebx, %rbp
	leaq	(,%rbp,8), %r15
	movq	%r15, %rdi
	callq	malloc
	movq	%rax, %r14
	movl	$8, %edi
	callq	malloc
	movq	%rax, 104(%rsp)                 # 8-byte Spill
	leaq	56(%rsp), %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	leaq	96(%rsp), %r12
	movl	$8, %esi
	movq	%r12, %rdi
	callq	hipMalloc
	leaq	48(%rsp), %r13
	movl	$4, %esi
	movq	%r13, %rdi
	callq	hipMalloc
	movq	(%r12), %rdi
	movl	$4, %edx
	xorl	%esi, %esi
	callq	hipMemset
	movq	(%r13), %rdi
	movl	$4, %edx
	xorl	%esi, %esi
	callq	hipMemset
	xorl	%edi, %edi
	callq	time
	movl	%eax, %edi
	callq	srand
	movl	%ebx, 12(%rsp)                  # 4-byte Spill
	movl	%ebx, %ebx
	testl	%ebp, %ebp
	jle	.LBB2_5
# %bb.3:                                # %.lr.ph.preheader
	xorl	%r12d, %r12d
.LBB2_4:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	callq	rand
	xorps	%xmm0, %xmm0
	cvtsi2sd	%eax, %xmm0
	movsd	%xmm0, (%r14,%r12,8)
	incq	%r12
	cmpq	%r12, %rbx
	jne	.LBB2_4
.LBB2_5:                                # %._crit_edge
	movabsq	$4294967552, %r12               # imm = 0x100000100
	leaq	32(%rsp), %r13
	movq	%r13, %rdi
	callq	hipEventCreate
	leaq	16(%rsp), %rdi
	callq	hipEventCreate
	movq	56(%rsp), %rdi
	movq	%r14, %rsi
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$.Lstr, %edi
	callq	puts@PLT
	movq	(%r13), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	%r12, %rdi
	movl	$1, %esi
	movq	%r12, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	movl	12(%rsp), %ebp                  # 4-byte Reload
	jne	.LBB2_7
# %bb.6:
	movq	56(%rsp), %rdi
	movq	96(%rsp), %rsi
	movq	48(%rsp), %rdx
	movl	%ebp, %ecx
	callq	_Z34__device_stub__find_maximum_kernelPdS_Pii
.LBB2_7:
	movq	96(%rsp), %rsi
	movl	$8, %edx
	movq	104(%rsp), %r12                 # 8-byte Reload
	movq	%r12, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	16(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	16(%rsp), %rdi
	callq	hipEventSynchronize
	movq	32(%rsp), %rsi
	movq	16(%rsp), %rdx
	leaq	44(%rsp), %r15
	movq	%r15, %rdi
	callq	hipEventElapsedTime
	movq	32(%rsp), %rdi
	callq	hipEventDestroy
	movq	16(%rsp), %rdi
	callq	hipEventDestroy
	movl	$_ZSt4cout, %edi
	movl	$.L.str.10, %esi
	movl	$15, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movl	%ebp, %esi
	callq	_ZNSolsEi
	movq	%rax, %r13
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r13, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r13, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.11, %esi
	movl	$27, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movsd	(%r12), %xmm0                   # xmm0 = mem[0],zero
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r13
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r13, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r13, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.12, %esi
	movl	$9, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	xorps	%xmm0, %xmm0
	cvtss2sd	(%r15), %xmm0
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r15
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	callq	clock
	movq	%rax, %r15
	movsd	(%r14), %xmm0                   # xmm0 = mem[0],zero
	cmpl	$0, 12(%rsp)                    # 4-byte Folded Reload
	jle	.LBB2_8
# %bb.13:                               # %.lr.ph161.preheader
	xorl	%eax, %eax
	movabsq	$4294967552, %r12               # imm = 0x100000100
.LBB2_14:                               # %.lr.ph161
                                        # =>This Inner Loop Header: Depth=1
	movsd	(%r14,%rax,8), %xmm1            # xmm1 = mem[0],zero
	maxsd	%xmm0, %xmm1
	incq	%rax
	movapd	%xmm1, %xmm0
	cmpq	%rax, %rbx
	jne	.LBB2_14
	jmp	.LBB2_9
.LBB2_8:
	movapd	%xmm0, %xmm1
	movabsq	$4294967552, %r12               # imm = 0x100000100
.LBB2_9:                                # %._crit_edge162
	movsd	%xmm1, 24(%rsp)                 # 8-byte Spill
	callq	clock
	subq	%r15, %rax
	movabsq	$2361183241434822607, %rcx      # imm = 0x20C49BA5E353F7CF
	imulq	%rcx
	movq	%rdx, %r15
	shrq	$63, %r15
	sarq	$7, %rdx
	addq	%rdx, %r15
	movl	$_ZSt4cout, %edi
	movl	$.L.str.13, %esi
	movl	$9, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movq	%r15, %rsi
	callq	_ZNSo9_M_insertIlEERSoT_
	movq	%rax, %r15
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.14, %esi
	movl	$22, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movsd	24(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r15
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$.Lstr.1, %edi
	callq	puts@PLT
	movl	$8, %edi
	callq	malloc
	movq	%rax, %rbp
	leaq	88(%rsp), %r15
	movl	$8, %esi
	movq	%r15, %rdi
	callq	hipMalloc
	movq	(%r15), %rdi
	movl	$4, %edx
	xorl	%esi, %esi
	callq	hipMemset
	movq	48(%rsp), %rdi
	movl	$4, %edx
	xorl	%esi, %esi
	callq	hipMemset
	leaq	32(%rsp), %r15
	movq	%r15, %rdi
	callq	hipEventCreate
	leaq	16(%rsp), %rdi
	callq	hipEventCreate
	movq	(%r15), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	%r12, %rdi
	movl	$1, %esi
	movq	%r12, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_11
# %bb.10:
	movq	56(%rsp), %rdi
	movq	88(%rsp), %rsi
	movq	48(%rsp), %rdx
	movl	12(%rsp), %ecx                  # 4-byte Reload
	callq	_Z34__device_stub__find_minimum_kernelPdS_Pii
.LBB2_11:
	movq	88(%rsp), %rsi
	movl	$8, %edx
	movq	%rbp, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	16(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	16(%rsp), %rdi
	callq	hipEventSynchronize
	movq	32(%rsp), %rsi
	movq	16(%rsp), %rdx
	leaq	44(%rsp), %r15
	movq	%r15, %rdi
	callq	hipEventElapsedTime
	movl	$_ZSt4cout, %edi
	movl	$.L.str.10, %esi
	movl	$15, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movl	12(%rsp), %r13d                 # 4-byte Reload
	movl	%r13d, %esi
	callq	_ZNSolsEi
	movq	%rbp, %r12
	movq	%rax, %rbp
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbp, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbp, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.16, %esi
	movl	$24, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%r12, 112(%rsp)                 # 8-byte Spill
	movsd	(%r12), %xmm0                   # xmm0 = mem[0],zero
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %rbp
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbp, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbp, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.17, %esi
	movl	$10, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	xorps	%xmm0, %xmm0
	cvtss2sd	(%r15), %xmm0
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r15
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movsd	(%r14), %xmm0                   # xmm0 = mem[0],zero
	movsd	%xmm0, 24(%rsp)                 # 8-byte Spill
	callq	clock
	movq	%rax, %r15
	testl	%r13d, %r13d
	jle	.LBB2_12
# %bb.15:                               # %.lr.ph166.preheader
	xorl	%eax, %eax
	movabsq	$4294967552, %r12               # imm = 0x100000100
	movsd	24(%rsp), %xmm1                 # 8-byte Reload
                                        # xmm1 = mem[0],zero
.LBB2_16:                               # %.lr.ph166
                                        # =>This Inner Loop Header: Depth=1
	movsd	(%r14,%rax,8), %xmm0            # xmm0 = mem[0],zero
	maxsd	%xmm1, %xmm0
	incq	%rax
	movapd	%xmm0, %xmm1
	cmpq	%rax, %rbx
	jne	.LBB2_16
	jmp	.LBB2_17
.LBB2_12:
	movsd	24(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	movabsq	$4294967552, %r12               # imm = 0x100000100
.LBB2_17:                               # %._crit_edge167
	movsd	%xmm0, 24(%rsp)                 # 8-byte Spill
	callq	clock
	subq	%r15, %rax
	movabsq	$2361183241434822607, %rcx      # imm = 0x20C49BA5E353F7CF
	imulq	%rcx
	movq	%rdx, %r15
	shrq	$63, %r15
	sarq	$7, %rdx
	addq	%rdx, %r15
	movl	$_ZSt4cout, %edi
	movl	$.L.str.18, %esi
	movl	$9, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movq	%r15, %rsi
	callq	_ZNSo9_M_insertIlEERSoT_
	movq	%rax, %r15
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.19, %esi
	movl	$22, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movsd	24(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r15
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$.Lstr.2, %edi
	callq	puts@PLT
	movl	$8, %edi
	callq	malloc
	movq	%rax, %rbp
	leaq	80(%rsp), %r15
	movl	$8, %esi
	movq	%r15, %rdi
	callq	hipMalloc
	movq	(%r15), %rdi
	movl	$4, %edx
	xorl	%esi, %esi
	callq	hipMemset
	movq	48(%rsp), %rdi
	movl	$4, %edx
	xorl	%esi, %esi
	callq	hipMemset
	leaq	32(%rsp), %r15
	movq	%r15, %rdi
	callq	hipEventCreate
	leaq	16(%rsp), %rdi
	callq	hipEventCreate
	movq	(%r15), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	%r12, %rdi
	movl	$1, %esi
	movq	%r12, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_19
# %bb.18:
	movq	56(%rsp), %rdi
	movq	80(%rsp), %rsi
	movq	48(%rsp), %rdx
	movl	12(%rsp), %ecx                  # 4-byte Reload
	callq	_Z26__device_stub__mean_kernelPdS_Pii
.LBB2_19:
	movq	80(%rsp), %rsi
	movl	$8, %edx
	movq	%rbp, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	16(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	16(%rsp), %rdi
	callq	hipEventSynchronize
	movq	32(%rsp), %rsi
	movq	16(%rsp), %rdx
	leaq	44(%rsp), %r15
	movq	%r15, %rdi
	callq	hipEventElapsedTime
	movq	32(%rsp), %rdi
	callq	hipEventDestroy
	movq	16(%rsp), %rdi
	callq	hipEventDestroy
	movsd	(%rbp), %xmm0                   # xmm0 = mem[0],zero
	cvtsi2sdl	12(%rsp), %xmm1         # 4-byte Folded Reload
	movsd	%xmm1, 64(%rsp)                 # 8-byte Spill
	divsd	%xmm1, %xmm0
	movsd	%xmm0, (%rbp)
	movl	$_ZSt4cout, %edi
	movl	$.L.str.21, %esi
	movl	$40, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movsd	(%rbp), %xmm0                   # xmm0 = mem[0],zero
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r13
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r13, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r13, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.10, %esi
	movl	$15, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movl	12(%rsp), %esi                  # 4-byte Reload
	callq	_ZNSolsEi
	movq	%rax, %r13
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r13, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r13, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.22, %esi
	movl	$9, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	xorps	%xmm0, %xmm0
	cvtss2sd	(%r15), %xmm0
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r15
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movsd	(%r14), %xmm0                   # xmm0 = mem[0],zero
	movsd	%xmm0, 24(%rsp)                 # 8-byte Spill
	callq	clock
	movq	%rax, %r15
	cmpl	$2, 12(%rsp)                    # 4-byte Folded Reload
	jl	.LBB2_20
# %bb.25:                               # %.lr.ph172.preheader
	movl	$1, %eax
	movsd	24(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
.LBB2_26:                               # %.lr.ph172
                                        # =>This Inner Loop Header: Depth=1
	addsd	(%r14,%rax,8), %xmm0
	incq	%rax
	cmpq	%rax, %rbx
	jne	.LBB2_26
	jmp	.LBB2_21
.LBB2_20:
	movsd	24(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
.LBB2_21:                               # %._crit_edge173
	divsd	64(%rsp), %xmm0                 # 8-byte Folded Reload
	movsd	%xmm0, 24(%rsp)                 # 8-byte Spill
	callq	clock
	subq	%r15, %rax
	movabsq	$2361183241434822607, %rcx      # imm = 0x20C49BA5E353F7CF
	imulq	%rcx
	movq	%rdx, %r15
	shrq	$63, %r15
	sarq	$7, %rdx
	addq	%rdx, %r15
	movl	$_ZSt4cout, %edi
	movl	$.L.str.18, %esi
	movl	$9, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movq	%r15, %rsi
	callq	_ZNSo9_M_insertIlEERSoT_
	movq	%rax, %r15
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.23, %esi
	movl	$13, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movsd	24(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r15
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$.Lstr.3, %edi
	callq	puts@PLT
	movl	$8, %edi
	callq	malloc
	movq	%rax, %r15
	leaq	72(%rsp), %r13
	movl	$8, %esi
	movq	%r13, %rdi
	callq	hipMalloc
	movq	(%r13), %rdi
	movl	$8, %edx
	xorl	%esi, %esi
	callq	hipMemset
	leaq	32(%rsp), %r13
	movq	%r13, %rdi
	callq	hipEventCreate
	leaq	16(%rsp), %rdi
	callq	hipEventCreate
	movq	(%r13), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	%r12, %rdi
	movl	$1, %esi
	movq	%r12, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	movl	12(%rsp), %r13d                 # 4-byte Reload
	jne	.LBB2_23
# %bb.22:
	movq	56(%rsp), %rdi
	movq	72(%rsp), %rsi
	movq	48(%rsp), %rdx
	movsd	(%rbp), %xmm0                   # xmm0 = mem[0],zero
	movl	%r13d, %ecx
	callq	_Z25__device_stub__std_kernelPdS_Piid
.LBB2_23:
	movq	16(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	16(%rsp), %rdi
	callq	hipEventSynchronize
	movq	32(%rsp), %rsi
	movq	16(%rsp), %rdx
	leaq	44(%rsp), %rdi
	callq	hipEventElapsedTime
	movq	72(%rsp), %rsi
	movl	$8, %edx
	movq	%r15, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	32(%rsp), %rdi
	callq	hipEventDestroy
	movq	16(%rsp), %rdi
	callq	hipEventDestroy
	movsd	(%r15), %xmm0                   # xmm0 = mem[0],zero
	divsd	64(%rsp), %xmm0                 # 8-byte Folded Reload
	xorpd	%xmm1, %xmm1
	ucomisd	%xmm1, %xmm0
	jb	.LBB2_27
# %bb.24:
	sqrtsd	%xmm0, %xmm0
	jmp	.LBB2_28
.LBB2_27:                               # %call.sqrt
	callq	sqrt
.LBB2_28:                               # %.split
	movsd	%xmm0, (%r15)
	movl	$_ZSt4cout, %edi
	movl	$.L.str.25, %esi
	movl	$11, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movsd	(%r15), %xmm0                   # xmm0 = mem[0],zero
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r12
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r12, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r12, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.12, %esi
	movl	$9, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	xorps	%xmm0, %xmm0
	cvtss2sd	44(%rsp), %xmm0
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r12
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r12, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r12, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	callq	clock
	movq	%rax, %r12
	testl	%r13d, %r13d
	movsd	24(%rsp), %xmm1                 # 8-byte Reload
                                        # xmm1 = mem[0],zero
	xorps	%xmm0, %xmm0
	jle	.LBB2_31
# %bb.29:                               # %.lr.ph178.preheader
	xorl	%eax, %eax
.LBB2_30:                               # %.lr.ph178
                                        # =>This Inner Loop Header: Depth=1
	movsd	(%r14,%rax,8), %xmm2            # xmm2 = mem[0],zero
	subsd	%xmm1, %xmm2
	mulsd	%xmm2, %xmm2
	addsd	%xmm2, %xmm0
	incq	%rax
	cmpq	%rax, %rbx
	jne	.LBB2_30
.LBB2_31:                               # %._crit_edge179
	divsd	64(%rsp), %xmm0                 # 8-byte Folded Reload
	xorpd	%xmm1, %xmm1
	ucomisd	%xmm1, %xmm0
	jb	.LBB2_33
# %bb.32:
	sqrtsd	%xmm0, %xmm0
	jmp	.LBB2_34
.LBB2_33:                               # %call.sqrt211
	callq	sqrt
.LBB2_34:                               # %._crit_edge179.split
	movsd	%xmm0, 24(%rsp)                 # 8-byte Spill
	callq	clock
	subq	%r12, %rax
	movabsq	$2361183241434822607, %rcx      # imm = 0x20C49BA5E353F7CF
	imulq	%rcx
	movq	%rdx, %r12
	shrq	$63, %r12
	sarq	$7, %rdx
	addq	%rdx, %r12
	movl	$_ZSt4cout, %edi
	movl	$.L.str.13, %esi
	movl	$9, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movq	%r12, %rsi
	callq	_ZNSo9_M_insertIlEERSoT_
	movq	%rax, %r12
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r12, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r12, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.26, %esi
	movl	$24, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movsd	24(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r12
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r12, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r12, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	104(%rsp), %rdi                 # 8-byte Reload
	callq	free
	movq	112(%rsp), %rdi                 # 8-byte Reload
	callq	free
	movq	%r15, %rdi
	callq	free
	movq	%rbp, %rdi
	callq	free
	movq	96(%rsp), %rdi
	callq	hipFree
	movq	88(%rsp), %rdi
	callq	hipFree
	movq	72(%rsp), %rdi
	callq	hipFree
	movq	80(%rsp), %rdi
	callq	hipFree
	movq	56(%rsp), %rdi
	callq	hipFree
	movq	%r14, %rdi
	movl	%r13d, %esi
	callq	_Z14concurrent_allPdi
	xorl	%eax, %eax
	addq	$120, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	main, .Lfunc_end2-main
	.cfi_endproc
                                        # -- End function
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"--------GPU version concurrent------------"
	.size	.L.str, 43

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"max "
	.size	.L.str.1, 5

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	" min "
	.size	.L.str.2, 6

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	" mean "
	.size	.L.str.3, 7

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	" Stand_d: "
	.size	.L.str.4, 11

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"array size "
	.size	.L.str.5, 12

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	" time "
	.size	.L.str.6, 7

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"-----------cpu Version concurrent---------------------"
	.size	.L.str.7, 55

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	" Stand d "
	.size	.L.str.8, 10

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"size of arrays "
	.size	.L.str.10, 16

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"Maximum number on gpu was: "
	.size	.L.str.11, 28

	.type	.L.str.12,@object               # @.str.12
.L.str.12:
	.asciz	"GPU time "
	.size	.L.str.12, 10

	.type	.L.str.13,@object               # @.str.13
.L.str.13:
	.asciz	"CPU time "
	.size	.L.str.13, 10

	.type	.L.str.14,@object               # @.str.14
.L.str.14:
	.asciz	"the cpu max number is "
	.size	.L.str.14, 23

	.type	.L.str.16,@object               # @.str.16
.L.str.16:
	.asciz	"the gpu min number was: "
	.size	.L.str.16, 25

	.type	.L.str.17,@object               # @.str.17
.L.str.17:
	.asciz	" gpu time "
	.size	.L.str.17, 11

	.type	.L.str.18,@object               # @.str.18
.L.str.18:
	.asciz	"cpu time "
	.size	.L.str.18, 10

	.type	.L.str.19,@object               # @.str.19
.L.str.19:
	.asciz	"the cpu min number is "
	.size	.L.str.19, 23

	.type	.L.str.21,@object               # @.str.21
.L.str.21:
	.asciz	"mean of the array calculate by gpu was: "
	.size	.L.str.21, 41

	.type	.L.str.22,@object               # @.str.22
.L.str.22:
	.asciz	"Gpu time "
	.size	.L.str.22, 10

	.type	.L.str.23,@object               # @.str.23
.L.str.23:
	.asciz	"the  mean is "
	.size	.L.str.23, 14

	.type	.L.str.25,@object               # @.str.25
.L.str.25:
	.asciz	"GPU std is "
	.size	.L.str.25, 12

	.type	.L.str.26,@object               # @.str.26
.L.str.26:
	.asciz	"the  stand deviation is "
	.size	.L.str.26, 25

	.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	"---------find max-------------------"
	.size	.Lstr, 37

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"---------find min-------------------"
	.size	.Lstr.1, 37

	.type	.Lstr.2,@object                 # @str.2
.Lstr.2:
	.asciz	"--------- arithmetic mean-------------------"
	.size	.Lstr.2, 45

	.type	.Lstr.3,@object                 # @str.3
.Lstr.3:
	.asciz	"-----------------GPU version STD--------------------"
	.size	.Lstr.3, 53

	.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 _ZSt4cout
	.addrsig_sym __hip_cuid_

	.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_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
