/*
Copyright 2021 Fixstars Corporation
Licensed under the Apache License, Version 2.0 (the "License");
you may not use this file except in compliance with the License.
You may obtain a copy of the License at
http ://www.apache.org/licenses/LICENSE-2.0
Unless required by applicable law or agreed to in writing, software
distributed under the License is distributed on an "AS IS" BASIS,
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and
limitations under the License.
*/

#include <iostream>
#include <algorithm>
#include <random>
#include <chrono>
#include <cmath>


static constexpr int NUM_TRIALS = 11;


// CPU版ReLUカーネル
void relu_cpu(float *dst, const float *src, int n){
	for(int i = 0; i < n; ++i){
		dst[i] = std::max(0.0f, src[i]);
	}
}

// CPU版処理時間計測
// NUM_TRIALS 回計測して中央値を求める
double relu_cpu_benchmark(float *dst, const float *src, int n){
	std::vector<double> durations(NUM_TRIALS);
	for(int i = 0; i < NUM_TRIALS; ++i){
		const auto begin = std::chrono::steady_clock::now();
		relu_cpu(dst, src, n);
		const auto end = std::chrono::steady_clock::now();
		const auto duration = std::chrono::duration_cast<std::chrono::microseconds>(end - begin);
		durations[i] = duration.count() * 1e-3;
	}
	std::sort(durations.begin(), durations.end());
	return durations[NUM_TRIALS / 2];
}


// GPU版ReLUカーネル
__global__ void relu_gpu(float *dst, const float *src, int n){
	const int i = blockIdx.x * blockDim.x + threadIdx.x;
	if(i < n){ dst[i] = max(0.0f, src[i]); }
}

void call_relu_gpu(float *dst, const float *src, int n){
	const int bdim = 256, gdim = (n + bdim - 1) / bdim;
	relu_gpu<<<gdim, bdim>>>(dst, src, n);
}

// GPU版処理時間計測
// NUM_TRIALS 回計測して中央値を求める
double relu_gpu_benchmark(float *h_dst, const float *h_src, int n){
	// デバイスメモリの確保
	float *d_dst = nullptr, *d_src = nullptr;
	cudaMalloc(&d_src, sizeof(float) * n);
	cudaMalloc(&d_dst, sizeof(float) * n);
	// 入力データの転送
	cudaMemcpy(d_src, h_src, sizeof(float) * n, cudaMemcpyHostToDevice);

	std::vector<double> durations(NUM_TRIALS);
	for(int i = 0; i < NUM_TRIALS; ++i){
		const auto begin = std::chrono::steady_clock::now();
		call_relu_gpu(d_dst, d_src, n);
		cudaDeviceSynchronize();  // GPUカーネルの終了を待つ
		const auto end = std::chrono::steady_clock::now();
		const auto duration = std::chrono::duration_cast<std::chrono::microseconds>(end - begin);
		durations[i] = duration.count() * 1e-3;
	}

	// 出力データの転送
	cudaMemcpy(h_dst, d_dst, sizeof(float) * n, cudaMemcpyDeviceToHost);
	// デバイスメモリの開放
	cudaFree(d_src);
	cudaFree(d_dst);

	// 中央値を求める
	std::sort(durations.begin(), durations.end());
	return durations[NUM_TRIALS / 2];
}


// 検算
bool validate(const float *expect, const float *actual, int n){
	bool valid = true;
	for(int i = 0; i < n; ++i){
		if(std::fabs(expect[i] - actual[i]) > 1e-4){
			valid = false;
		}
	}
	return valid;
}


int main(int argc, char *argv[]){
	if(argc < 2){
		std::cerr << "Usage: " << argv[0] << " n" << std::endl;
		return 0;
	}

	const int n = atoi(argv[1]);
	std::cout << "n = " << n << std::endl;

	std::default_random_engine engine;
	std::uniform_real_distribution<float> dist(-1.0f, 1.0f);

	std::vector<float> src(n), cpu_dst(n), gpu_dst(n);
	for(int i = 0; i < n; ++i){
		src[i] = dist(engine);
	}

	const auto cpu_duration =
		relu_cpu_benchmark(cpu_dst.data(), src.data(), n);
	std::cout << "CPU: " << cpu_duration << " [ms]" << std::endl;

	const auto gpu_duration =
		relu_gpu_benchmark(gpu_dst.data(), src.data(), n);
	std::cout << "GPU: " << gpu_duration << " [ms]" << std::endl;

	const auto valid = validate(cpu_dst.data(), gpu_dst.data(), n);
	std::cout << "Validation: " << (valid ? "Success" : "Failed") << std::endl;

	return 0;
}

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

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB5046:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE5046:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_ZSt13__adjust_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEEldNS0_5__ops15_Iter_less_iterEEvT_T0_SA_T1_T2_.isra.0, @function
_ZSt13__adjust_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEEldNS0_5__ops15_Iter_less_iterEEvT_T0_SA_T1_T2_.isra.0:
.LFB6019:
	.cfi_startproc
	leaq	-1(%rdx), %rax
	movq	%rsi, %r9
	movl	$2, %ecx
	movq	%rdx, %rsi
	cqto
	idivq	%rcx
	movq	%r9, %rcx
.L3:
	cmpq	%rax, %rcx
	jge	.L13
	leaq	1(%rcx), %rdx
	addq	%rdx, %rdx
	leaq	-1(%rdx), %r8
	movsd	(%rdi,%r8,8), %xmm1
	comisd	(%rdi,%rdx,8), %xmm1
	cmovbe	%rdx, %r8
	movsd	(%rdi,%r8,8), %xmm1
	movsd	%xmm1, (%rdi,%rcx,8)
	movq	%r8, %rcx
	jmp	.L3
.L13:
	testl	$1, %esi
	jne	.L6
	leaq	-2(%rsi), %rax
	movl	$2, %esi
	cqto
	idivq	%rsi
	cmpq	%rax, %rcx
	jne	.L6
	leaq	1(%rcx,%rcx), %rax
	movsd	(%rdi,%rax,8), %xmm1
	movsd	%xmm1, (%rdi,%rcx,8)
	movq	%rax, %rcx
.L6:
	leaq	-1(%rcx), %rax
	movl	$2, %esi
.L11:
	cqto
	idivq	%rsi
	leaq	(%rdi,%rcx,8), %rdx
	cmpq	%rcx, %r9
	jge	.L7
	movsd	(%rdi,%rax,8), %xmm1
	comisd	%xmm1, %xmm0
	jbe	.L7
	movsd	%xmm1, (%rdx)
	leaq	-1(%rax), %rdx
	movq	%rax, %rcx
	movq	%rdx, %rax
	jmp	.L11
.L7:
	movsd	%xmm0, (%rdx)
	ret
	.cfi_endproc
.LFE6019:
	.size	_ZSt13__adjust_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEEldNS0_5__ops15_Iter_less_iterEEvT_T0_SA_T1_T2_.isra.0, .-_ZSt13__adjust_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEEldNS0_5__ops15_Iter_less_iterEEvT_T0_SA_T1_T2_.isra.0
	.align 2
	.type	_ZNSolsEPFRSoS_E.isra.0, @function
_ZNSolsEPFRSoS_E.isra.0:
.LFB6021:
	.cfi_startproc
	jmp	*%rsi
	.cfi_endproc
.LFE6021:
	.size	_ZNSolsEPFRSoS_E.isra.0, .-_ZNSolsEPFRSoS_E.isra.0
	.type	_ZSt16__insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_.isra.0, @function
_ZSt16__insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_.isra.0:
.LFB6022:
	.cfi_startproc
	cmpq	%rsi, %rdi
	jne	.L30
	ret
.L30:
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movl	$8, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movq	%rsi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	leaq	8(%rdi), %rbx
	subq	$24, %rsp
	.cfi_def_cfa_offset 64
.L17:
	cmpq	%rbx, %r12
	je	.L34
	movsd	(%rbx), %xmm0
	movsd	0(%rbp), %xmm1
	comisd	%xmm0, %xmm1
	jbe	.L29
	movq	%rbx, %rdx
	subq	%rbp, %rdx
	cmpq	$8, %rdx
	jle	.L20
	movq	%r13, %rdi
	movq	%rbp, %rsi
	movsd	%xmm0, 8(%rsp)
	subq	%rdx, %rdi
	addq	%rbx, %rdi
	call	memmove@PLT
	movsd	8(%rsp), %xmm0
	jmp	.L21
.L20:
	jne	.L21
	movsd	%xmm1, (%rbx)
.L21:
	movsd	%xmm0, 0(%rbp)
	jmp	.L22
.L29:
	movq	%rbx, %rax
.L18:
	movsd	-8(%rax), %xmm1
	comisd	%xmm0, %xmm1
	jbe	.L35
	movsd	%xmm1, (%rax)
	subq	$8, %rax
	jmp	.L18
.L35:
	movsd	%xmm0, (%rax)
.L22:
	addq	$8, %rbx
	jmp	.L17
.L34:
	addq	$24, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE6022:
	.size	_ZSt16__insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_.isra.0, .-_ZSt16__insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_.isra.0
	.type	_ZSt16__introsort_loopIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEElNS0_5__ops15_Iter_less_iterEEvT_S9_T0_T1_.isra.0, @function
_ZSt16__introsort_loopIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEElNS0_5__ops15_Iter_less_iterEEvT_S9_T0_T1_.isra.0:
.LFB6024:
	.cfi_startproc
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movq	%rsi, %r10
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	leaq	8(%rdi), %r12
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movq	%rdx, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movq	%rdi, %rbx
	pushq	%rcx
	.cfi_def_cfa_offset 48
.L37:
	movq	%r10, %rax
	subq	%rbx, %rax
	cmpq	$128, %rax
	jle	.L36
	movq	%rax, %r13
	sarq	$3, %r13
	testq	%rbp, %rbp
	jne	.L38
	leaq	-2(%r13), %r11
	sarq	%r11
.L40:
	movsd	(%rbx,%r11,8), %xmm0
	movq	%r13, %rdx
	movq	%r11, %rsi
	movq	%rbx, %rdi
	call	_ZSt13__adjust_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEEldNS0_5__ops15_Iter_less_iterEEvT_T0_SA_T1_T2_.isra.0
	testq	%r11, %r11
	je	.L41
	decq	%r11
	jmp	.L40
.L41:
	movq	%r10, %rax
	subq	$8, %r10
	subq	%rbx, %rax
	cmpq	$8, %rax
	jle	.L36
	movsd	(%rbx), %xmm1
	movq	%r10, %rdx
	movsd	(%r10), %xmm0
	xorl	%esi, %esi
	subq	%rbx, %rdx
	movq	%rbx, %rdi
	movsd	%xmm1, (%r10)
	sarq	$3, %rdx
	call	_ZSt13__adjust_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEEldNS0_5__ops15_Iter_less_iterEEvT_T0_SA_T1_T2_.isra.0
	jmp	.L41
.L38:
	sarq	$4, %rax
	movsd	8(%rbx), %xmm3
	movsd	(%rbx), %xmm0
	movq	%r12, %rdx
	leaq	(%rbx,%rax,8), %rax
	movsd	-8(%r10), %xmm1
	movsd	(%rax), %xmm2
	comisd	%xmm3, %xmm2
	jbe	.L67
	comisd	%xmm2, %xmm1
	ja	.L71
	comisd	%xmm3, %xmm1
	jbe	.L69
	jmp	.L73
.L67:
	comisd	%xmm3, %xmm1
	jbe	.L70
.L69:
	movsd	%xmm3, (%rbx)
	movsd	%xmm0, 8(%rbx)
.L48:
	movq	%r10, %rax
	jmp	.L51
.L70:
	comisd	%xmm2, %xmm1
	jbe	.L71
.L73:
	movsd	%xmm1, (%rbx)
	movsd	%xmm0, -8(%r10)
	jmp	.L48
.L71:
	movsd	%xmm2, (%rbx)
	movsd	%xmm0, (%rax)
	jmp	.L48
.L58:
	movsd	%xmm1, 0(%r13)
	movsd	%xmm0, (%rdx)
	leaq	8(%r13), %rdx
.L51:
	movsd	(%rbx), %xmm2
.L56:
	movq	%rdx, %r13
	movsd	(%rdx), %xmm0
	leaq	8(%rdx), %rdx
	comisd	%xmm0, %xmm2
	ja	.L56
	subq	$8, %rax
.L57:
	movq	%rax, %rdx
	movsd	(%rax), %xmm1
	subq	$8, %rax
	comisd	%xmm2, %xmm1
	ja	.L57
	movq	%rdx, %rax
	cmpq	%rdx, %r13
	jb	.L58
	decq	%rbp
	movq	%r10, %rsi
	movq	%r13, %rdi
	movq	%rbp, %rdx
	call	_ZSt16__introsort_loopIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEElNS0_5__ops15_Iter_less_iterEEvT_S9_T0_T1_.isra.0
	movq	%r13, %r10
	jmp	.L37
.L36:
	popq	%rax
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE6024:
	.size	_ZSt16__introsort_loopIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEElNS0_5__ops15_Iter_less_iterEEvT_S9_T0_T1_.isra.0, .-_ZSt16__introsort_loopIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEElNS0_5__ops15_Iter_less_iterEEvT_S9_T0_T1_.isra.0
	.type	_ZSt6__sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_.isra.0, @function
_ZSt6__sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_.isra.0:
.LFB6031:
	.cfi_startproc
	cmpq	%rsi, %rdi
	je	.L85
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rsi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rsi, %rbx
	subq	%rdi, %rbx
	movq	%rbx, %rax
	sarq	$3, %rax
	subq	$24, %rsp
	.cfi_def_cfa_offset 48
	bsrq	%rax, %rdx
	movq	%rdi, 8(%rsp)
	xorq	$63, %rdx
	testq	%rax, %rax
	movl	$64, %eax
	cmovne	%edx, %eax
	movl	$63, %edx
	subl	%eax, %edx
	movslq	%edx, %rdx
	addq	%rdx, %rdx
	call	_ZSt16__introsort_loopIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEElNS0_5__ops15_Iter_less_iterEEvT_S9_T0_T1_.isra.0
	cmpq	$128, %rbx
	movq	8(%rsp), %rdi
	jle	.L77
	leaq	128(%rdi), %rbx
	movq	%rbx, %rsi
	call	_ZSt16__insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_.isra.0
	movq	%rbx, %rax
.L78:
	cmpq	%rax, %rbp
	je	.L88
	movsd	(%rax), %xmm0
	movq	%rax, %rdx
.L79:
	movsd	-8(%rdx), %xmm1
	comisd	%xmm0, %xmm1
	jbe	.L89
	movsd	%xmm1, (%rdx)
	subq	$8, %rdx
	jmp	.L79
.L89:
	movsd	%xmm0, (%rdx)
	addq	$8, %rax
	jmp	.L78
.L77:
	addq	$24, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 24
	movq	%rbp, %rsi
	popq	%rbx
	.cfi_restore 3
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_restore 6
	.cfi_def_cfa_offset 8
	jmp	_ZSt16__insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_.isra.0
.L88:
	.cfi_restore_state
	addq	$24, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
.L85:
	.cfi_restore 3
	.cfi_restore 6
	ret
	.cfi_endproc
.LFE6031:
	.size	_ZSt6__sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_.isra.0, .-_ZSt6__sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_.isra.0
	.globl	_Z8relu_cpuPfPKfi
	.type	_Z8relu_cpuPfPKfi, @function
_Z8relu_cpuPfPKfi:
.LFB5035:
	.cfi_startproc
	endbr64
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	xorps	%xmm0, %xmm0
	movq	%fs:40, %rax
	movq	%rax, 8(%rsp)
	xorl	%eax, %eax
	leaq	4(%rsp), %r8
.L91:
	cmpl	%eax, %edx
	jle	.L96
	movss	(%rsi), %xmm1
	movq	%r8, %rcx
	movss	%xmm0, 4(%rsp)
	comiss	%xmm0, %xmm1
	cmova	%rsi, %rcx
	addq	$4, %rsi
	movss	(%rcx), %xmm1
	movss	%xmm1, (%rdi,%rax,4)
	incq	%rax
	jmp	.L91
.L96:
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	je	.L94
	call	__stack_chk_fail@PLT
.L94:
	addq	$24, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE5035:
	.size	_Z8relu_cpuPfPKfi, .-_Z8relu_cpuPfPKfi
	.globl	_Z8validatePKfS0_i
	.type	_Z8validatePKfS0_i, @function
_Z8validatePKfS0_i:
.LFB5042:
	.cfi_startproc
	endbr64
	movsd	.LC2(%rip), %xmm2
	xorl	%eax, %eax
	movb	$1, %cl
	xorl	%r8d, %r8d
	movss	.LC1(%rip), %xmm1
.L98:
	cmpl	%eax, %edx
	jle	.L102
	movss	(%rdi,%rax,4), %xmm0
	subss	(%rsi,%rax,4), %xmm0
	andps	%xmm1, %xmm0
	cvtss2sd	%xmm0, %xmm0
	ucomisd	%xmm2, %xmm0
	cmova	%r8d, %ecx
	incq	%rax
	jmp	.L98
.L102:
	movl	%ecx, %eax
	ret
	.cfi_endproc
.LFE5042:
	.size	_Z8validatePKfS0_i, .-_Z8validatePKfS0_i
	.globl	_Z31__device_stub__Z8relu_gpuPfPKfiPfPKfi
	.type	_Z31__device_stub__Z8relu_gpuPfPKfiPfPKfi, @function
_Z31__device_stub__Z8relu_gpuPfPKfiPfPKfi:
.LFB5068:
	.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	.L103
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z8relu_gpuPfPKfi(%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
.L103:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L105
	call	__stack_chk_fail@PLT
.L105:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE5068:
	.size	_Z31__device_stub__Z8relu_gpuPfPKfiPfPKfi, .-_Z31__device_stub__Z8relu_gpuPfPKfiPfPKfi
	.globl	_Z8relu_gpuPfPKfi
	.type	_Z8relu_gpuPfPKfi, @function
_Z8relu_gpuPfPKfi:
.LFB5069:
	.cfi_startproc
	endbr64
	jmp	_Z31__device_stub__Z8relu_gpuPfPKfiPfPKfi
	.cfi_endproc
.LFE5069:
	.size	_Z8relu_gpuPfPKfi, .-_Z8relu_gpuPfPKfi
	.globl	_Z13call_relu_gpuPfPKfi
	.type	_Z13call_relu_gpuPfPKfi, @function
_Z13call_relu_gpuPfPKfi:
.LFB5040:
	.cfi_startproc
	endbr64
	leal	255(%rdx), %eax
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movl	$256, %ecx
	movq	%rsi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movq	%rdi, %rbp
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movl	%edx, %ebx
	cltd
	idivl	%ecx
	subq	$32, %rsp
	.cfi_def_cfa_offset 64
	movl	$16777217, %edx
	movl	$1, %ecx
	salq	$8, %rdx
	movl	%eax, 8(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 12(%rsp)
	movq	8(%rsp), %rdi
	movl	16(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L108
	addq	$32, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 32
	movl	%ebx, %edx
	movq	%r12, %rsi
	movq	%rbp, %rdi
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	jmp	_Z31__device_stub__Z8relu_gpuPfPKfiPfPKfi
.L108:
	.cfi_restore_state
	addq	$32, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE5040:
	.size	_Z13call_relu_gpuPfPKfi, .-_Z13call_relu_gpuPfPKfi
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC3:
	.string	"_Z8relu_gpuPfPKfi"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB5071:
	.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	.LC3(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z8relu_gpuPfPKfi(%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
.LFE5071:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.text._ZNSt6vectorIdSaIdEED2Ev,"axG",@progbits,_ZNSt6vectorIdSaIdEED5Ev,comdat
	.align 2
	.weak	_ZNSt6vectorIdSaIdEED2Ev
	.type	_ZNSt6vectorIdSaIdEED2Ev, @function
_ZNSt6vectorIdSaIdEED2Ev:
.LFB5426:
	.cfi_startproc
	endbr64
	movq	(%rdi), %rax
	testq	%rax, %rax
	je	.L113
	movq	16(%rdi), %rsi
	movq	%rax, %rdi
	subq	%rax, %rsi
	jmp	_ZdlPvm@PLT
.L113:
	ret
	.cfi_endproc
.LFE5426:
	.size	_ZNSt6vectorIdSaIdEED2Ev, .-_ZNSt6vectorIdSaIdEED2Ev
	.weak	_ZNSt6vectorIdSaIdEED1Ev
	.set	_ZNSt6vectorIdSaIdEED1Ev,_ZNSt6vectorIdSaIdEED2Ev
	.text
	.globl	_Z18relu_cpu_benchmarkPfPKfi
	.type	_Z18relu_cpu_benchmarkPfPKfi, @function
_Z18relu_cpu_benchmarkPfPKfi:
.LFB5036:
	.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
	movl	%edx, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rsi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$56, %rsp
	.cfi_def_cfa_offset 112
	movq	%rdi, 8(%rsp)
	movl	$88, %edi
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	call	_Znwm@PLT
	leaq	88(%rax), %r15
	movq	%rax, 16(%rsp)
	movq	%rax, %rbx
	movq	%r15, 32(%rsp)
	movq	$0x000000000, (%rax)
	xorl	%eax, %eax
.L116:
	addq	$8, %rax
	cmpq	$88, %rax
	je	.L122
	movq	$0x000000000, (%rbx,%rax)
	jmp	.L116
.L122:
	movq	%r15, 24(%rsp)
	xorl	%r14d, %r14d
.L118:
	call	_ZNSt6chrono3_V212steady_clock3nowEv@PLT
	movq	8(%rsp), %rdi
	movl	%r13d, %edx
	movq	%r12, %rsi
	movq	%rax, %rbp
	call	_Z8relu_cpuPfPKfi
	call	_ZNSt6chrono3_V212steady_clock3nowEv@PLT
	movl	$1000, %ecx
	subq	%rbp, %rax
	cqto
	idivq	%rcx
	cvtsi2sdq	%rax, %xmm0
	mulsd	.LC5(%rip), %xmm0
	movsd	%xmm0, (%rbx,%r14,8)
	incq	%r14
	cmpq	$11, %r14
	jne	.L118
	movq	%r15, %rsi
	movq	%rbx, %rdi
	call	_ZSt6__sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_.isra.0
	movsd	40(%rbx), %xmm0
	leaq	16(%rsp), %rdi
	movsd	%xmm0, 8(%rsp)
	call	_ZNSt6vectorIdSaIdEED1Ev
	movsd	8(%rsp), %xmm0
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L119
	call	__stack_chk_fail@PLT
.L119:
	addq	$56, %rsp
	.cfi_def_cfa_offset 56
	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
.LFE5036:
	.size	_Z18relu_cpu_benchmarkPfPKfi, .-_Z18relu_cpu_benchmarkPfPKfi
	.globl	_Z18relu_gpu_benchmarkPfPKfi
	.type	_Z18relu_gpu_benchmarkPfPKfi, @function
_Z18relu_gpu_benchmarkPfPKfi:
.LFB5041:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA5041
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%rdi, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movl	%edx, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rsi, %rbx
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	movq	%fs:40, %rax
	movq	%rax, 56(%rsp)
	xorl	%eax, %eax
	xorl	%eax, %eax
	leaq	24(%rsp), %rdi
	movq	%rax, 16(%rsp)
	movq	%rax, 24(%rsp)
	movslq	%edx, %rax
	salq	$2, %rax
	movq	%rax, %rsi
	movq	%rax, 8(%rsp)
.LEHB0:
	call	cudaMalloc@PLT
	movq	8(%rsp), %rsi
	leaq	16(%rsp), %rdi
	call	cudaMalloc@PLT
	movq	8(%rsp), %rdx
	movq	24(%rsp), %rdi
	movq	%rbx, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movl	$88, %edi
	call	_Znwm@PLT
.LEHE0:
	leaq	88(%rax), %r13
	movq	%rax, 32(%rsp)
	movq	%rax, %rbx
	movq	%r13, 48(%rsp)
	movq	$0x000000000, (%rax)
	xorl	%eax, %eax
.L124:
	addq	$8, %rax
	cmpq	$88, %rax
	je	.L134
	movq	$0x000000000, (%rbx,%rax)
	jmp	.L124
.L134:
	movq	%r13, 40(%rsp)
	xorl	%ebp, %ebp
.L126:
	call	_ZNSt6chrono3_V212steady_clock3nowEv@PLT
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	%r12d, %edx
	movq	%rax, %r15
.LEHB1:
	call	_Z13call_relu_gpuPfPKfi
	call	cudaDeviceSynchronize@PLT
	call	_ZNSt6chrono3_V212steady_clock3nowEv@PLT
	movl	$1000, %ecx
	subq	%r15, %rax
	cqto
	idivq	%rcx
	cvtsi2sdq	%rax, %xmm0
	mulsd	.LC5(%rip), %xmm0
	movsd	%xmm0, (%rbx,%rbp,8)
	incq	%rbp
	cmpq	$11, %rbp
	jne	.L126
	movq	8(%rsp), %rdx
	movq	16(%rsp), %rsi
	movl	$2, %ecx
	movq	%r14, %rdi
	call	cudaMemcpy@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
.LEHE1:
	movq	%r13, %rsi
	movq	%rbx, %rdi
	call	_ZSt6__sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_.isra.0
	movsd	40(%rbx), %xmm0
	leaq	32(%rsp), %rdi
	movsd	%xmm0, 8(%rsp)
	call	_ZNSt6vectorIdSaIdEED1Ev
	movsd	8(%rsp), %xmm0
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L129
	jmp	.L132
.L130:
	endbr64
	movq	%rax, %rbx
.L127:
	leaq	32(%rsp), %rdi
	call	_ZNSt6vectorIdSaIdEED1Ev
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L132
	movq	%rbx, %rdi
.LEHB2:
	call	_Unwind_Resume@PLT
.LEHE2:
.L132:
	call	__stack_chk_fail@PLT
.L129:
	addq	$72, %rsp
	.cfi_def_cfa_offset 56
	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
.LFE5041:
	.globl	__gxx_personality_v0
	.section	.gcc_except_table,"a",@progbits
.LLSDA5041:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE5041-.LLSDACSB5041
.LLSDACSB5041:
	.uleb128 .LEHB0-.LFB5041
	.uleb128 .LEHE0-.LEHB0
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB1-.LFB5041
	.uleb128 .LEHE1-.LEHB1
	.uleb128 .L130-.LFB5041
	.uleb128 0
	.uleb128 .LEHB2-.LFB5041
	.uleb128 .LEHE2-.LEHB2
	.uleb128 0
	.uleb128 0
.LLSDACSE5041:
	.text
	.size	_Z18relu_gpu_benchmarkPfPKfi, .-_Z18relu_gpu_benchmarkPfPKfi
	.section	.rodata._ZNSt6vectorIfSaIfEEC2EmRKS0_.str1.1,"aMS",@progbits,1
.LC6:
	.string	"cannot create std::vector larger than max_size()"
	.section	.text._ZNSt6vectorIfSaIfEEC2EmRKS0_,"axG",@progbits,_ZNSt6vectorIfSaIfEEC5EmRKS0_,comdat
	.align 2
	.weak	_ZNSt6vectorIfSaIfEEC2EmRKS0_
	.type	_ZNSt6vectorIfSaIfEEC2EmRKS0_, @function
_ZNSt6vectorIfSaIfEEC2EmRKS0_:
.LFB5452:
	.cfi_startproc
	endbr64
	movabsq	$2305843009213693951, %rax
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	cmpq	%rsi, %rax
	jnb	.L136
	leaq	.LC6(%rip), %rdi
	call	_ZSt20__throw_length_errorPKc@PLT
.L136:
	xorl	%eax, %eax
	movq	%rdi, %rbx
	leaq	0(,%rsi,4), %r12
	movq	%rsi, %rbp
	movq	%rax, (%rdi)
	movq	%rax, 8(%rdi)
	movq	%rax, 16(%rdi)
	xorl	%eax, %eax
	testq	%rsi, %rsi
	je	.L137
	movq	%r12, %rdi
	call	_Znwm@PLT
.L137:
	addq	%rax, %r12
	movq	%rax, (%rbx)
	movq	%rax, 8(%rbx)
	movq	%r12, 16(%rbx)
	testq	%rbp, %rbp
	je	.L138
	movl	$0x00000000, (%rax)
	addq	$4, %rax
	decq	%rbp
	je	.L138
.L139:
	cmpq	%rax, %r12
	je	.L138
	movl	$0x00000000, (%rax)
	addq	$4, %rax
	jmp	.L139
.L138:
	movq	%rax, 8(%rbx)
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE5452:
	.size	_ZNSt6vectorIfSaIfEEC2EmRKS0_, .-_ZNSt6vectorIfSaIfEEC2EmRKS0_
	.weak	_ZNSt6vectorIfSaIfEEC1EmRKS0_
	.set	_ZNSt6vectorIfSaIfEEC1EmRKS0_,_ZNSt6vectorIfSaIfEEC2EmRKS0_
	.section	.text._ZNSt6vectorIfSaIfEED2Ev,"axG",@progbits,_ZNSt6vectorIfSaIfEED5Ev,comdat
	.align 2
	.weak	_ZNSt6vectorIfSaIfEED2Ev
	.type	_ZNSt6vectorIfSaIfEED2Ev, @function
_ZNSt6vectorIfSaIfEED2Ev:
.LFB5455:
	.cfi_startproc
	endbr64
	movq	(%rdi), %rax
	testq	%rax, %rax
	je	.L150
	movq	16(%rdi), %rsi
	movq	%rax, %rdi
	subq	%rax, %rsi
	jmp	_ZdlPvm@PLT
.L150:
	ret
	.cfi_endproc
.LFE5455:
	.size	_ZNSt6vectorIfSaIfEED2Ev, .-_ZNSt6vectorIfSaIfEED2Ev
	.weak	_ZNSt6vectorIfSaIfEED1Ev
	.set	_ZNSt6vectorIfSaIfEED1Ev,_ZNSt6vectorIfSaIfEED2Ev
	.section	.rodata.str1.1
.LC8:
	.string	"Success"
.LC9:
	.string	"Failed"
.LC10:
	.string	"Usage: "
.LC11:
	.string	" n"
.LC12:
	.string	"n = "
.LC15:
	.string	"CPU: "
.LC16:
	.string	" [ms]"
.LC17:
	.string	"GPU: "
.LC18:
	.string	"Validation: "
	.section	.text.startup
	.globl	main
	.type	main, @function
main:
.LFB5043:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA5043
	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
	movq	%rsi, %rbx
	subq	$136, %rsp
	.cfi_def_cfa_offset 192
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	decl	%edi
	jg	.L153
	leaq	.LC10(%rip), %rsi
	leaq	_ZSt4cerr(%rip), %rdi
.LEHB3:
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	(%rbx), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	leaq	.LC11(%rip), %rsi
	movq	%rax, %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
	jmp	.L154
.L153:
	movq	8(%rsi), %rdi
	leaq	_ZSt4cout(%rip), %rbp
	call	atoi@PLT
	leaq	.LC12(%rip), %rsi
	movq	%rbp, %rdi
	movslq	%eax, %r12
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%r12d, %esi
	movq	%r12, %rbx
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	96(%rsp), %rax
	movq	%r12, %rsi
	movq	%rax, 8(%rsp)
	leaq	48(%rsp), %rax
	movq	8(%rsp), %rdx
	movq	%rax, %rdi
	movq	%rax, 16(%rsp)
	call	_ZNSt6vectorIfSaIfEEC1EmRKS0_
.LEHE3:
	leaq	72(%rsp), %rax
	movq	8(%rsp), %rdx
	movq	%r12, %rsi
	movq	%rax, %rdi
	movq	%rax, 24(%rsp)
.LEHB4:
	call	_ZNSt6vectorIfSaIfEEC1EmRKS0_
.LEHE4:
	movq	8(%rsp), %rdi
	leaq	47(%rsp), %rdx
	movq	%r12, %rsi
.LEHB5:
	call	_ZNSt6vectorIfSaIfEEC1EmRKS0_
.LEHE5:
	movq	48(%rsp), %r12
	xorl	%esi, %esi
	movl	$1, %edx
	xorps	%xmm2, %xmm2
	movss	.LC13(%rip), %xmm3
	movss	.LC14(%rip), %xmm1
	movl	$2147483647, %edi
.L155:
	cmpl	%esi, %ebx
	jle	.L173
	imulq	$16807, %rdx, %rax
	xorl	%edx, %edx
	divq	%rdi
	movq	%rdx, %rcx
	decq	%rcx
	js	.L156
	cvtsi2ssq	%rcx, %xmm0
	jmp	.L157
.L156:
	movq	%rcx, %rax
	andl	$1, %ecx
	shrq	%rax
	orq	%rcx, %rax
	cvtsi2ssq	%rax, %xmm0
	addss	%xmm0, %xmm0
.L157:
	addss	%xmm2, %xmm0
	mulss	%xmm3, %xmm0
	comiss	%xmm1, %xmm0
	jb	.L158
	movss	.LC7(%rip), %xmm0
.L158:
	addss	%xmm0, %xmm0
	subss	%xmm1, %xmm0
	movss	%xmm0, (%r12,%rsi,4)
	incq	%rsi
	jmp	.L155
.L173:
	movq	72(%rsp), %r13
	movl	%ebx, %edx
	movq	%r12, %rsi
	movq	%r13, %rdi
.LEHB6:
	call	_Z18relu_cpu_benchmarkPfPKfi
	leaq	.LC15(%rip), %rsi
	movq	%rbp, %rdi
	movq	%xmm0, %r14
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	movq	%r14, %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	leaq	.LC16(%rip), %r15
	movq	%rax, %rdi
	movq	%r15, %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	96(%rsp), %r14
	movl	%ebx, %edx
	movq	%r12, %rsi
	movq	%r14, %rdi
	call	_Z18relu_gpu_benchmarkPfPKfi
	leaq	.LC17(%rip), %rsi
	movq	%rbp, %rdi
	movq	%xmm0, %r12
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	movq	%r12, %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	%rax, %rdi
	movq	%r15, %rsi
	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	%ebx, %edx
	movq	%r14, %rsi
	movq	%r13, %rdi
	call	_Z8validatePKfS0_i
	leaq	.LC18(%rip), %rsi
	movq	%rbp, %rdi
	movl	%eax, %ebx
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	testb	%bl, %bl
	leaq	.LC9(%rip), %rsi
	leaq	.LC8(%rip), %rax
	cmovne	%rax, %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
.LEHE6:
	movq	8(%rsp), %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	movq	24(%rsp), %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	movq	16(%rsp), %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
.L154:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L165
	jmp	.L171
.L170:
	endbr64
	movq	%rax, %rbx
.L161:
	movq	8(%rsp), %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	jmp	.L162
.L169:
	endbr64
	movq	%rax, %rbx
.L162:
	movq	24(%rsp), %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	jmp	.L163
.L168:
	endbr64
	movq	%rax, %rbx
.L163:
	movq	16(%rsp), %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L171
	movq	%rbx, %rdi
.LEHB7:
	call	_Unwind_Resume@PLT
.LEHE7:
.L171:
	call	__stack_chk_fail@PLT
.L165:
	addq	$136, %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
.LFE5043:
	.section	.gcc_except_table
.LLSDA5043:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE5043-.LLSDACSB5043
.LLSDACSB5043:
	.uleb128 .LEHB3-.LFB5043
	.uleb128 .LEHE3-.LEHB3
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB4-.LFB5043
	.uleb128 .LEHE4-.LEHB4
	.uleb128 .L168-.LFB5043
	.uleb128 0
	.uleb128 .LEHB5-.LFB5043
	.uleb128 .LEHE5-.LEHB5
	.uleb128 .L169-.LFB5043
	.uleb128 0
	.uleb128 .LEHB6-.LFB5043
	.uleb128 .LEHE6-.LEHB6
	.uleb128 .L170-.LFB5043
	.uleb128 0
	.uleb128 .LEHB7-.LFB5043
	.uleb128 .LEHE7-.LEHB7
	.uleb128 0
	.uleb128 0
.LLSDACSE5043:
	.section	.text.startup
	.size	main, .-main
	.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
	.section	.rodata.cst16,"aM",@progbits,16
	.align 16
.LC1:
	.long	2147483647
	.long	0
	.long	0
	.long	0
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC2:
	.long	-350469331
	.long	1058682594
	.align 8
.LC5:
	.long	-755914244
	.long	1062232653
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC7:
	.long	1065353215
	.align 4
.LC13:
	.long	805306368
	.align 4
.LC14:
	.long	1065353216
	.hidden	DW.ref.__gxx_personality_v0
	.weak	DW.ref.__gxx_personality_v0
	.section	.data.rel.local.DW.ref.__gxx_personality_v0,"awG",@progbits,DW.ref.__gxx_personality_v0,comdat
	.align 8
	.type	DW.ref.__gxx_personality_v0, @object
	.size	DW.ref.__gxx_personality_v0, 8
DW.ref.__gxx_personality_v0:
	.quad	__gxx_personality_v0
	.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:

/*
Copyright 2021 Fixstars Corporation
Licensed under the Apache License, Version 2.0 (the "License");
you may not use this file except in compliance with the License.
You may obtain a copy of the License at
http ://www.apache.org/licenses/LICENSE-2.0
Unless required by applicable law or agreed to in writing, software
distributed under the License is distributed on an "AS IS" BASIS,
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and
limitations under the License.
*/


#include <hip/hip_runtime.h>
#include <iostream>
#include <algorithm>
#include <random>
#include <chrono>
#include <cmath>


static constexpr int NUM_TRIALS = 11;


// CPU版ReLUカーネル
void relu_cpu(float *dst, const float *src, int n){
	for(int i = 0; i < n; ++i){
		dst[i] = std::max(0.0f, src[i]);
	}
}

// CPU版処理時間計測
// NUM_TRIALS 回計測して中央値を求める
double relu_cpu_benchmark(float *dst, const float *src, int n){
	std::vector<double> durations(NUM_TRIALS);
	for(int i = 0; i < NUM_TRIALS; ++i){
		const auto begin = std::chrono::steady_clock::now();
		relu_cpu(dst, src, n);
		const auto end = std::chrono::steady_clock::now();
		const auto duration = std::chrono::duration_cast<std::chrono::microseconds>(end - begin);
		durations[i] = duration.count() * 1e-3;
	}
	std::sort(durations.begin(), durations.end());
	return durations[NUM_TRIALS / 2];
}


// GPU版ReLUカーネル
__global__ void relu_gpu(float *dst, const float *src, int n){
	const int i = blockIdx.x * blockDim.x + threadIdx.x;
	if(i < n){ dst[i] = max(0.0f, src[i]); }
}

void call_relu_gpu(float *dst, const float *src, int n){
	const int bdim = 256, gdim = (n + bdim - 1) / bdim;
	relu_gpu<<<gdim, bdim>>>(dst, src, n);
}

// GPU版処理時間計測
// NUM_TRIALS 回計測して中央値を求める
double relu_gpu_benchmark(float *h_dst, const float *h_src, int n){
	// デバイスメモリの確保
	float *d_dst = nullptr, *d_src = nullptr;
	hipMalloc(&d_src, sizeof(float) * n);
	hipMalloc(&d_dst, sizeof(float) * n);
	// 入力データの転送
	hipMemcpy(d_src, h_src, sizeof(float) * n, hipMemcpyHostToDevice);

	std::vector<double> durations(NUM_TRIALS);
	for(int i = 0; i < NUM_TRIALS; ++i){
		const auto begin = std::chrono::steady_clock::now();
		call_relu_gpu(d_dst, d_src, n);
		hipDeviceSynchronize();  // GPUカーネルの終了を待つ
		const auto end = std::chrono::steady_clock::now();
		const auto duration = std::chrono::duration_cast<std::chrono::microseconds>(end - begin);
		durations[i] = duration.count() * 1e-3;
	}

	// 出力データの転送
	hipMemcpy(h_dst, d_dst, sizeof(float) * n, hipMemcpyDeviceToHost);
	// デバイスメモリの開放
	hipFree(d_src);
	hipFree(d_dst);

	// 中央値を求める
	std::sort(durations.begin(), durations.end());
	return durations[NUM_TRIALS / 2];
}


// 検算
bool validate(const float *expect, const float *actual, int n){
	bool valid = true;
	for(int i = 0; i < n; ++i){
		if(std::fabs(expect[i] - actual[i]) > 1e-4){
			valid = false;
		}
	}
	return valid;
}


int main(int argc, char *argv[]){
	if(argc < 2){
		std::cerr << "Usage: " << argv[0] << " n" << std::endl;
		return 0;
	}

	const int n = atoi(argv[1]);
	std::cout << "n = " << n << std::endl;

	std::default_random_engine engine;
	std::uniform_real_distribution<float> dist(-1.0f, 1.0f);

	std::vector<float> src(n), cpu_dst(n), gpu_dst(n);
	for(int i = 0; i < n; ++i){
		src[i] = dist(engine);
	}

	const auto cpu_duration =
		relu_cpu_benchmark(cpu_dst.data(), src.data(), n);
	std::cout << "CPU: " << cpu_duration << " [ms]" << std::endl;

	const auto gpu_duration =
		relu_gpu_benchmark(gpu_dst.data(), src.data(), n);
	std::cout << "GPU: " << gpu_duration << " [ms]" << std::endl;

	const auto valid = validate(cpu_dst.data(), gpu_dst.data(), n);
	std::cout << "Validation: " << (valid ? "Success" : "Failed") << std::endl;

	return 0;
}

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

                                        # End of file scope inline assembly
	.globl	_Z8relu_cpuPfPKfi               # -- Begin function _Z8relu_cpuPfPKfi
	.type	_Z8relu_cpuPfPKfi,@function
_Z8relu_cpuPfPKfi:                      # @_Z8relu_cpuPfPKfi
	.cfi_startproc
# %bb.0:
	testl	%edx, %edx
	jle	.LBB0_3
# %bb.1:                                # %.lr.ph.preheader
	movl	%edx, %eax
	xorl	%ecx, %ecx
	xorps	%xmm0, %xmm0
.LBB0_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movss	(%rsi,%rcx,4), %xmm1            # xmm1 = mem[0],zero,zero,zero
	maxss	%xmm0, %xmm1
	movss	%xmm1, (%rdi,%rcx,4)
	incq	%rcx
	cmpq	%rcx, %rax
	jne	.LBB0_2
.LBB0_3:                                # %._crit_edge
	retq
.Lfunc_end0:
	.size	_Z8relu_cpuPfPKfi, .Lfunc_end0-_Z8relu_cpuPfPKfi
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _Z18relu_cpu_benchmarkPfPKfi
.LCPI1_0:
	.quad	0x3f50624dd2f1a9fc              # double 0.001
	.text
	.globl	_Z18relu_cpu_benchmarkPfPKfi
	.type	_Z18relu_cpu_benchmarkPfPKfi,@function
_Z18relu_cpu_benchmarkPfPKfi:           # @_Z18relu_cpu_benchmarkPfPKfi
.Lfunc_begin0:
	.cfi_startproc
	.cfi_personality 3, __gxx_personality_v0
	.cfi_lsda 3, .Lexception0
# %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	$24, %rsp
	.cfi_def_cfa_offset 80
	.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	%edx, %ebx
	movq	%rsi, %r15
	movq	%rdi, %r12
	movl	$88, %edi
	callq	_Znwm
	movq	%rax, %rbp
	movl	%ebx, 8(%rsp)                   # 4-byte Spill
	movl	%ebx, %r14d
	xorps	%xmm0, %xmm0
	movups	%xmm0, (%rax)
	movups	%xmm0, 16(%rax)
	movups	%xmm0, 32(%rax)
	movups	%xmm0, 48(%rax)
	movups	%xmm0, 64(%rax)
	movq	%rax, 16(%rsp)                  # 8-byte Spill
	movq	$0, 80(%rax)
	xorl	%ebx, %ebx
.LBB1_1:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB1_3 Depth 2
	callq	_ZNSt6chrono3_V212steady_clock3nowEv
	movq	%rax, %r13
	cmpl	$0, 8(%rsp)                     # 4-byte Folded Reload
	xorps	%xmm1, %xmm1
	jle	.LBB1_4
# %bb.2:                                # %.lr.ph.i.preheader
                                        #   in Loop: Header=BB1_1 Depth=1
	xorl	%eax, %eax
.LBB1_3:                                # %.lr.ph.i
                                        #   Parent Loop BB1_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movss	(%r15,%rax,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
	maxss	%xmm1, %xmm0
	movss	%xmm0, (%r12,%rax,4)
	incq	%rax
	cmpq	%rax, %r14
	jne	.LBB1_3
.LBB1_4:                                # %.loopexit
                                        #   in Loop: Header=BB1_1 Depth=1
	callq	_ZNSt6chrono3_V212steady_clock3nowEv
	subq	%r13, %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
	mulsd	.LCPI1_0(%rip), %xmm0
	movsd	%xmm0, (%rbp,%rbx,8)
	incq	%rbx
	cmpq	$11, %rbx
	jne	.LBB1_1
# %bb.5:
	movq	16(%rsp), %rbx                  # 8-byte Reload
	addq	$88, %rbx
.Ltmp0:
	movl	$6, %edx
	movq	%rbp, %rdi
	movq	%rbx, %rsi
	callq	_ZSt16__introsort_loopIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEElNS0_5__ops15_Iter_less_iterEEvT_S9_T0_T1_
.Ltmp1:
# %bb.6:                                # %.noexc
.Ltmp2:
	movq	%rbp, %rdi
	movq	%rbx, %rsi
	callq	_ZSt22__final_insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_
.Ltmp3:
# %bb.7:                                # %_ZNSt6vectorIdSaIdEED2Ev.exit
	movsd	40(%rbp), %xmm0                 # xmm0 = mem[0],zero
	movsd	%xmm0, 8(%rsp)                  # 8-byte Spill
	movq	%rbp, %rdi
	callq	_ZdlPv
	movsd	8(%rsp), %xmm0                  # 8-byte Reload
                                        # xmm0 = mem[0],zero
	addq	$24, %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
.LBB1_8:                                # %_ZNSt6vectorIdSaIdEED2Ev.exit17
	.cfi_def_cfa_offset 80
.Ltmp4:
	movq	%rax, %r14
	movq	%rbp, %rdi
	callq	_ZdlPv
	movq	%r14, %rdi
	callq	_Unwind_Resume@PLT
.Lfunc_end1:
	.size	_Z18relu_cpu_benchmarkPfPKfi, .Lfunc_end1-_Z18relu_cpu_benchmarkPfPKfi
	.cfi_endproc
	.section	.gcc_except_table,"a",@progbits
	.p2align	2, 0x0
GCC_except_table1:
.Lexception0:
	.byte	255                             # @LPStart Encoding = omit
	.byte	255                             # @TType Encoding = omit
	.byte	1                               # Call site Encoding = uleb128
	.uleb128 .Lcst_end0-.Lcst_begin0
.Lcst_begin0:
	.uleb128 .Lfunc_begin0-.Lfunc_begin0    # >> Call Site 1 <<
	.uleb128 .Ltmp0-.Lfunc_begin0           #   Call between .Lfunc_begin0 and .Ltmp0
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp0-.Lfunc_begin0           # >> Call Site 2 <<
	.uleb128 .Ltmp3-.Ltmp0                  #   Call between .Ltmp0 and .Ltmp3
	.uleb128 .Ltmp4-.Lfunc_begin0           #     jumps to .Ltmp4
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp3-.Lfunc_begin0           # >> Call Site 3 <<
	.uleb128 .Lfunc_end1-.Ltmp3             #   Call between .Ltmp3 and .Lfunc_end1
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
.Lcst_end0:
	.p2align	2, 0x0
                                        # -- End function
	.text
	.globl	_Z23__device_stub__relu_gpuPfPKfi # -- Begin function _Z23__device_stub__relu_gpuPfPKfi
	.type	_Z23__device_stub__relu_gpuPfPKfi,@function
_Z23__device_stub__relu_gpuPfPKfi:      # @_Z23__device_stub__relu_gpuPfPKfi
	.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	$_Z8relu_gpuPfPKfi, %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_end2:
	.size	_Z23__device_stub__relu_gpuPfPKfi, .Lfunc_end2-_Z23__device_stub__relu_gpuPfPKfi
	.cfi_endproc
                                        # -- End function
	.globl	_Z13call_relu_gpuPfPKfi         # -- Begin function _Z13call_relu_gpuPfPKfi
	.type	_Z13call_relu_gpuPfPKfi,@function
_Z13call_relu_gpuPfPKfi:                # @_Z13call_relu_gpuPfPKfi
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movl	%edx, %ebx
	movq	%rsi, %r14
	movq	%rdi, %r15
	leal	255(%rbx), %eax
	leal	510(%rbx), %edi
	testl	%eax, %eax
	cmovnsl	%eax, %edi
	sarl	$8, %edi
	btsq	$32, %rdi
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$256, %rdx                      # imm = 0x100
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	je	.LBB3_2
# %bb.1:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.LBB3_2:
	.cfi_def_cfa_offset 32
	movq	%r15, %rdi
	movq	%r14, %rsi
	movl	%ebx, %edx
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	jmp	_Z23__device_stub__relu_gpuPfPKfi # TAILCALL
.Lfunc_end3:
	.size	_Z13call_relu_gpuPfPKfi, .Lfunc_end3-_Z13call_relu_gpuPfPKfi
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _Z18relu_gpu_benchmarkPfPKfi
.LCPI4_0:
	.quad	0x3f50624dd2f1a9fc              # double 0.001
	.text
	.globl	_Z18relu_gpu_benchmarkPfPKfi
	.type	_Z18relu_gpu_benchmarkPfPKfi,@function
_Z18relu_gpu_benchmarkPfPKfi:           # @_Z18relu_gpu_benchmarkPfPKfi
.Lfunc_begin1:
	.cfi_startproc
	.cfi_personality 3, __gxx_personality_v0
	.cfi_lsda 3, .Lexception1
# %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	$40, %rsp
	.cfi_def_cfa_offset 96
	.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	%edx, %ebp
	movq	%rsi, %rbx
	movq	%rdi, 32(%rsp)                  # 8-byte Spill
	xorl	%r15d, %r15d
	leaq	16(%rsp), %r14
	movq	%r15, (%r14)
	leaq	24(%rsp), %r13
	movq	%r15, (%r13)
	movslq	%edx, %r12
	shlq	$2, %r12
	movq	%r13, %rdi
	movq	%r12, %rsi
	callq	hipMalloc
	movq	%r14, %rdi
	movq	%r12, %rsi
	callq	hipMalloc
	movq	(%r13), %rdi
	movq	%rbx, %rsi
	movq	%r12, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$88, %edi
	callq	_Znwm
	movq	%rax, %rbx
	addq	$88, %rax
	movq	%rax, 8(%rsp)                   # 8-byte Spill
	xorpd	%xmm0, %xmm0
	movupd	%xmm0, (%rbx)
	movupd	%xmm0, 16(%rbx)
	movupd	%xmm0, 32(%rbx)
	movupd	%xmm0, 48(%rbx)
	movupd	%xmm0, 64(%rbx)
	movq	%r15, 80(%rbx)
	movabsq	$2361183241434822607, %r14      # imm = 0x20C49BA5E353F7CF
.LBB4_1:                                # =>This Inner Loop Header: Depth=1
	callq	_ZNSt6chrono3_V212steady_clock3nowEv
	movq	%rax, %r13
	movq	16(%rsp), %rdi
	movq	24(%rsp), %rsi
.Ltmp5:
	movl	%ebp, %edx
	callq	_Z13call_relu_gpuPfPKfi
.Ltmp6:
# %bb.2:                                #   in Loop: Header=BB4_1 Depth=1
.Ltmp7:
	callq	hipDeviceSynchronize
.Ltmp8:
# %bb.3:                                #   in Loop: Header=BB4_1 Depth=1
	callq	_ZNSt6chrono3_V212steady_clock3nowEv
	subq	%r13, %rax
	imulq	%r14
	movq	%rdx, %rax
	shrq	$63, %rax
	sarq	$7, %rdx
	addq	%rax, %rdx
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rdx, %xmm0
	mulsd	.LCPI4_0(%rip), %xmm0
	movsd	%xmm0, (%rbx,%r15,8)
	incq	%r15
	cmpq	$11, %r15
	jne	.LBB4_1
# %bb.4:
	movq	16(%rsp), %rsi
.Ltmp10:
	movq	32(%rsp), %rdi                  # 8-byte Reload
	movq	%r12, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
.Ltmp11:
	movq	8(%rsp), %r14                   # 8-byte Reload
# %bb.5:
	movq	24(%rsp), %rdi
.Ltmp12:
	callq	hipFree
.Ltmp13:
# %bb.6:
	movq	16(%rsp), %rdi
.Ltmp14:
	callq	hipFree
.Ltmp15:
# %bb.7:
.Ltmp16:
	movl	$6, %edx
	movq	%rbx, %rdi
	movq	%r14, %rsi
	callq	_ZSt16__introsort_loopIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEElNS0_5__ops15_Iter_less_iterEEvT_S9_T0_T1_
.Ltmp17:
# %bb.8:                                # %.noexc
.Ltmp18:
	movq	%rbx, %rdi
	movq	%r14, %rsi
	callq	_ZSt22__final_insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_
.Ltmp19:
# %bb.9:                                # %_ZNSt6vectorIdSaIdEED2Ev.exit
	movsd	40(%rbx), %xmm0                 # xmm0 = mem[0],zero
	movsd	%xmm0, 8(%rsp)                  # 8-byte Spill
	movq	%rbx, %rdi
	callq	_ZdlPv
	movsd	8(%rsp), %xmm0                  # 8-byte Reload
                                        # xmm0 = mem[0],zero
	addq	$40, %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
.LBB4_10:
	.cfi_def_cfa_offset 96
.Ltmp20:
	jmp	.LBB4_11
.LBB4_12:
.Ltmp9:
.LBB4_11:                               # %_ZNSt6vectorIdSaIdEED2Ev.exit21
	movq	%rax, %r14
	movq	%rbx, %rdi
	callq	_ZdlPv
	movq	%r14, %rdi
	callq	_Unwind_Resume@PLT
.Lfunc_end4:
	.size	_Z18relu_gpu_benchmarkPfPKfi, .Lfunc_end4-_Z18relu_gpu_benchmarkPfPKfi
	.cfi_endproc
	.section	.gcc_except_table,"a",@progbits
	.p2align	2, 0x0
GCC_except_table4:
.Lexception1:
	.byte	255                             # @LPStart Encoding = omit
	.byte	255                             # @TType Encoding = omit
	.byte	1                               # Call site Encoding = uleb128
	.uleb128 .Lcst_end1-.Lcst_begin1
.Lcst_begin1:
	.uleb128 .Lfunc_begin1-.Lfunc_begin1    # >> Call Site 1 <<
	.uleb128 .Ltmp5-.Lfunc_begin1           #   Call between .Lfunc_begin1 and .Ltmp5
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp5-.Lfunc_begin1           # >> Call Site 2 <<
	.uleb128 .Ltmp8-.Ltmp5                  #   Call between .Ltmp5 and .Ltmp8
	.uleb128 .Ltmp9-.Lfunc_begin1           #     jumps to .Ltmp9
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp10-.Lfunc_begin1          # >> Call Site 3 <<
	.uleb128 .Ltmp19-.Ltmp10                #   Call between .Ltmp10 and .Ltmp19
	.uleb128 .Ltmp20-.Lfunc_begin1          #     jumps to .Ltmp20
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp19-.Lfunc_begin1          # >> Call Site 4 <<
	.uleb128 .Lfunc_end4-.Ltmp19            #   Call between .Ltmp19 and .Lfunc_end4
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
.Lcst_end1:
	.p2align	2, 0x0
                                        # -- End function
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function _Z8validatePKfS0_i
.LCPI5_0:
	.long	0x7fffffff                      # float NaN
	.long	0x7fffffff                      # float NaN
	.long	0x7fffffff                      # float NaN
	.long	0x7fffffff                      # float NaN
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0
.LCPI5_1:
	.quad	0x3f1a36e2eb1c432d              # double 1.0E-4
	.text
	.globl	_Z8validatePKfS0_i
	.type	_Z8validatePKfS0_i,@function
_Z8validatePKfS0_i:                     # @_Z8validatePKfS0_i
	.cfi_startproc
# %bb.0:
	testl	%edx, %edx
	jle	.LBB5_1
# %bb.3:                                # %.lr.ph.preheader
	movl	%edx, %ecx
	movb	$1, %al
	xorl	%edx, %edx
	movaps	.LCPI5_0(%rip), %xmm0           # xmm0 = [NaN,NaN,NaN,NaN]
	movsd	.LCPI5_1(%rip), %xmm1           # xmm1 = mem[0],zero
.LBB5_4:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movss	(%rdi,%rdx,4), %xmm2            # xmm2 = mem[0],zero,zero,zero
	subss	(%rsi,%rdx,4), %xmm2
	andps	%xmm0, %xmm2
	cvtss2sd	%xmm2, %xmm2
	ucomisd	%xmm1, %xmm2
	movl	$0, %r8d
	ja	.LBB5_6
# %bb.5:                                # %.lr.ph
                                        #   in Loop: Header=BB5_4 Depth=1
	movzbl	%al, %r8d
.LBB5_6:                                # %.lr.ph
                                        #   in Loop: Header=BB5_4 Depth=1
	movl	%r8d, %eax
	incq	%rdx
	cmpq	%rdx, %rcx
	jne	.LBB5_4
# %bb.7:                                # %._crit_edge.loopexit
	andb	$1, %al
	retq
.LBB5_1:
	movb	$1, %al
	retq
.Lfunc_end5:
	.size	_Z8validatePKfS0_i, .Lfunc_end5-_Z8validatePKfS0_i
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function main
.LCPI6_0:
	.long	0xbf800000                      # float -1
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0
.LCPI6_1:
	.long	0x7fffffff                      # float NaN
	.long	0x7fffffff                      # float NaN
	.long	0x7fffffff                      # float NaN
	.long	0x7fffffff                      # float NaN
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0
.LCPI6_2:
	.quad	0x3f1a36e2eb1c432d              # double 1.0E-4
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
.Lfunc_begin2:
	.cfi_startproc
	.cfi_personality 3, __gxx_personality_v0
	.cfi_lsda 3, .Lexception2
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	subq	$104, %rsp
	.cfi_def_cfa_offset 144
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%rsi, %rbx
	cmpl	$1, %edi
	jg	.LBB6_6
# %bb.1:
	movl	$_ZSt4cerr, %edi
	movl	$.L.str, %esi
	movl	$7, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	(%rbx), %rbx
	testq	%rbx, %rbx
	je	.LBB6_2
# %bb.3:
	movq	%rbx, %rdi
	callq	strlen
	movl	$_ZSt4cerr, %edi
	movq	%rbx, %rsi
	movq	%rax, %rdx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	jmp	.LBB6_4
.LBB6_6:
	movq	8(%rbx), %rdi
	callq	atoi
	movl	%eax, %ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.2, %esi
	movl	$4, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movl	%ebx, %esi
	callq	_ZNSolsEi
	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	$1, 96(%rsp)
	movslq	%ebx, %r14
	leaq	72(%rsp), %rdi
	leaq	48(%rsp), %rdx
	movq	%r14, %rsi
	callq	_ZNSt6vectorIfSaIfEEC2EmRKS0_
.Ltmp21:
	leaq	48(%rsp), %rdi
	leaq	24(%rsp), %rdx
	movq	%r14, %rsi
	callq	_ZNSt6vectorIfSaIfEEC2EmRKS0_
.Ltmp22:
# %bb.7:
.Ltmp24:
	leaq	24(%rsp), %rdi
	leaq	15(%rsp), %rdx
	movq	%r14, %rsi
	callq	_ZNSt6vectorIfSaIfEEC2EmRKS0_
.Ltmp25:
# %bb.8:
	testl	%ebx, %ebx
	jle	.LBB6_9
# %bb.28:                               # %.lr.ph.preheader
	movl	%ebx, %r15d
	xorl	%r12d, %r12d
	leaq	96(%rsp), %r14
.LBB6_29:                               # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
.Ltmp27:
	movq	%r14, %rdi
	callq	_ZSt18generate_canonicalIfLm24ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_
.Ltmp28:
# %bb.30:                               #   in Loop: Header=BB6_29 Depth=1
	addss	%xmm0, %xmm0
	addss	.LCPI6_0(%rip), %xmm0
	movq	72(%rsp), %rsi
	movss	%xmm0, (%rsi,%r12,4)
	incq	%r12
	cmpq	%r12, %r15
	jne	.LBB6_29
	jmp	.LBB6_10
.LBB6_2:
	movq	_ZSt4cerr(%rip), %rax
	movq	-24(%rax), %rax
	leaq	_ZSt4cerr(%rax), %rdi
	movl	32(%rdi), %esi
	orl	$1, %esi
	callq	_ZNSt9basic_iosIcSt11char_traitsIcEE5clearESt12_Ios_Iostate
.LBB6_4:                                # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit
	movl	$_ZSt4cerr, %ebx
	movl	$_ZSt4cerr, %edi
	movl	$.L.str.1, %esi
	movl	$2, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cerr(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cerr, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	jmp	.LBB6_5
.LBB6_9:                                # %.._crit_edge_crit_edge
	movq	72(%rsp), %rsi
.LBB6_10:                               # %._crit_edge
	movq	48(%rsp), %rdi
.Ltmp30:
	movl	%ebx, %edx
	callq	_Z18relu_cpu_benchmarkPfPKfi
	movsd	%xmm0, 16(%rsp)                 # 8-byte Spill
.Ltmp31:
# %bb.11:
.Ltmp32:
	movl	$_ZSt4cout, %edi
	movl	$.L.str.3, %esi
	movl	$5, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp33:
# %bb.12:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit34
.Ltmp34:
	movl	$_ZSt4cout, %edi
	movsd	16(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
.Ltmp35:
# %bb.13:                               # %_ZNSolsEd.exit
.Ltmp36:
	movq	%rax, %r14
	movl	$.L.str.4, %esi
	movl	$5, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp37:
# %bb.14:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit35
	movq	(%r14), %rax
	movq	-24(%rax), %rdi
	addq	%r14, %rdi
.Ltmp38:
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
.Ltmp39:
# %bb.15:                               # %.noexc
.Ltmp40:
	movsbl	%al, %esi
	movq	%r14, %rdi
	callq	_ZNSo3putEc
.Ltmp41:
# %bb.16:                               # %.noexc54
.Ltmp42:
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.Ltmp43:
# %bb.17:                               # %_ZNSolsEPFRSoS_E.exit
	movq	24(%rsp), %rdi
	movq	72(%rsp), %rsi
.Ltmp45:
	movl	%ebx, %edx
	callq	_Z18relu_gpu_benchmarkPfPKfi
	movsd	%xmm0, 16(%rsp)                 # 8-byte Spill
.Ltmp46:
# %bb.18:
.Ltmp47:
	movl	$_ZSt4cout, %edi
	movl	$.L.str.5, %esi
	movl	$5, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp48:
# %bb.19:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit36
.Ltmp49:
	movl	$_ZSt4cout, %edi
	movsd	16(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
.Ltmp50:
# %bb.20:                               # %_ZNSolsEd.exit37
.Ltmp51:
	movq	%rax, %r14
	movl	$.L.str.4, %esi
	movl	$5, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp52:
# %bb.21:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit38
	movq	(%r14), %rax
	movq	-24(%rax), %rdi
	addq	%r14, %rdi
.Ltmp53:
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
.Ltmp54:
# %bb.22:                               # %.noexc56
.Ltmp55:
	movsbl	%al, %esi
	movq	%r14, %rdi
	callq	_ZNSo3putEc
.Ltmp56:
# %bb.23:                               # %.noexc57
.Ltmp57:
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.Ltmp58:
# %bb.24:                               # %_ZNSolsEPFRSoS_E.exit39
	testl	%ebx, %ebx
	jle	.LBB6_25
# %bb.32:                               # %.lr.ph.preheader.i
	movq	48(%rsp), %rax
	movq	24(%rsp), %rcx
	movl	%ebx, %edx
	movb	$1, %dil
	xorl	%esi, %esi
	movaps	.LCPI6_1(%rip), %xmm0           # xmm0 = [NaN,NaN,NaN,NaN]
	movsd	.LCPI6_2(%rip), %xmm1           # xmm1 = mem[0],zero
.LBB6_33:                               # %.lr.ph.i
                                        # =>This Inner Loop Header: Depth=1
	movss	(%rax,%rsi,4), %xmm2            # xmm2 = mem[0],zero,zero,zero
	subss	(%rcx,%rsi,4), %xmm2
	andps	%xmm0, %xmm2
	cvtss2sd	%xmm2, %xmm2
	ucomisd	%xmm1, %xmm2
	movl	$0, %r8d
	ja	.LBB6_35
# %bb.34:                               # %.lr.ph.i
                                        #   in Loop: Header=BB6_33 Depth=1
	movzbl	%dil, %r8d
.LBB6_35:                               # %.lr.ph.i
                                        #   in Loop: Header=BB6_33 Depth=1
	movl	%r8d, %edi
	incq	%rsi
	cmpq	%rsi, %rdx
	jne	.LBB6_33
# %bb.36:                               # %._crit_edge.loopexit.i
	testb	$1, %dil
	movl	$.L.str.8, %eax
	movl	$.L.str.7, %ebx
	cmoveq	%rax, %rbx
	jmp	.LBB6_37
.LBB6_25:
	movl	$.L.str.7, %ebx
.LBB6_37:                               # %_Z8validatePKfS0_i.exit
.Ltmp60:
	movl	$_ZSt4cout, %edi
	movl	$.L.str.6, %esi
	movl	$12, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp61:
# %bb.38:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit40
	movq	%rbx, %rdi
	callq	strlen
.Ltmp62:
	movl	$_ZSt4cout, %r14d
	movl	$_ZSt4cout, %edi
	movq	%rbx, %rsi
	movq	%rax, %rdx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp63:
# %bb.39:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit42
	movq	_ZSt4cout(%rip), %rax
	addq	-24(%rax), %r14
.Ltmp64:
	movq	%r14, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
.Ltmp65:
# %bb.40:                               # %.noexc60
.Ltmp66:
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
.Ltmp67:
# %bb.41:                               # %.noexc61
.Ltmp68:
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.Ltmp69:
# %bb.42:                               # %_ZNSolsEPFRSoS_E.exit43
	movq	24(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB6_44
# %bb.43:
	callq	_ZdlPv
.LBB6_44:                               # %_ZNSt6vectorIfSaIfEED2Ev.exit
	movq	48(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB6_46
# %bb.45:
	callq	_ZdlPv
.LBB6_46:                               # %_ZNSt6vectorIfSaIfEED2Ev.exit45
	movq	72(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB6_5
# %bb.47:
	callq	_ZdlPv
.LBB6_5:
	xorl	%eax, %eax
	addq	$104, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.LBB6_27:
	.cfi_def_cfa_offset 144
.Ltmp26:
	movq	%rax, %rbx
	jmp	.LBB6_53
.LBB6_26:
.Ltmp23:
	movq	%rax, %rbx
	jmp	.LBB6_55
.LBB6_50:
.Ltmp70:
	jmp	.LBB6_51
.LBB6_49:
.Ltmp59:
	jmp	.LBB6_51
.LBB6_48:
.Ltmp44:
	jmp	.LBB6_51
.LBB6_31:
.Ltmp29:
.LBB6_51:
	movq	%rax, %rbx
	movq	24(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB6_53
# %bb.52:
	callq	_ZdlPv
.LBB6_53:                               # %_ZNSt6vectorIfSaIfEED2Ev.exit49
	movq	48(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB6_55
# %bb.54:
	callq	_ZdlPv
.LBB6_55:                               # %_ZNSt6vectorIfSaIfEED2Ev.exit51
	movq	72(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB6_57
# %bb.56:
	callq	_ZdlPv
.LBB6_57:                               # %_ZNSt6vectorIfSaIfEED2Ev.exit53
	movq	%rbx, %rdi
	callq	_Unwind_Resume@PLT
.Lfunc_end6:
	.size	main, .Lfunc_end6-main
	.cfi_endproc
	.section	.gcc_except_table,"a",@progbits
	.p2align	2, 0x0
GCC_except_table6:
.Lexception2:
	.byte	255                             # @LPStart Encoding = omit
	.byte	255                             # @TType Encoding = omit
	.byte	1                               # Call site Encoding = uleb128
	.uleb128 .Lcst_end2-.Lcst_begin2
.Lcst_begin2:
	.uleb128 .Lfunc_begin2-.Lfunc_begin2    # >> Call Site 1 <<
	.uleb128 .Ltmp21-.Lfunc_begin2          #   Call between .Lfunc_begin2 and .Ltmp21
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp21-.Lfunc_begin2          # >> Call Site 2 <<
	.uleb128 .Ltmp22-.Ltmp21                #   Call between .Ltmp21 and .Ltmp22
	.uleb128 .Ltmp23-.Lfunc_begin2          #     jumps to .Ltmp23
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp24-.Lfunc_begin2          # >> Call Site 3 <<
	.uleb128 .Ltmp25-.Ltmp24                #   Call between .Ltmp24 and .Ltmp25
	.uleb128 .Ltmp26-.Lfunc_begin2          #     jumps to .Ltmp26
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp27-.Lfunc_begin2          # >> Call Site 4 <<
	.uleb128 .Ltmp28-.Ltmp27                #   Call between .Ltmp27 and .Ltmp28
	.uleb128 .Ltmp29-.Lfunc_begin2          #     jumps to .Ltmp29
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp28-.Lfunc_begin2          # >> Call Site 5 <<
	.uleb128 .Ltmp30-.Ltmp28                #   Call between .Ltmp28 and .Ltmp30
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp30-.Lfunc_begin2          # >> Call Site 6 <<
	.uleb128 .Ltmp43-.Ltmp30                #   Call between .Ltmp30 and .Ltmp43
	.uleb128 .Ltmp44-.Lfunc_begin2          #     jumps to .Ltmp44
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp45-.Lfunc_begin2          # >> Call Site 7 <<
	.uleb128 .Ltmp58-.Ltmp45                #   Call between .Ltmp45 and .Ltmp58
	.uleb128 .Ltmp59-.Lfunc_begin2          #     jumps to .Ltmp59
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp60-.Lfunc_begin2          # >> Call Site 8 <<
	.uleb128 .Ltmp69-.Ltmp60                #   Call between .Ltmp60 and .Ltmp69
	.uleb128 .Ltmp70-.Lfunc_begin2          #     jumps to .Ltmp70
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp69-.Lfunc_begin2          # >> Call Site 9 <<
	.uleb128 .Lfunc_end6-.Ltmp69            #   Call between .Ltmp69 and .Lfunc_end6
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
.Lcst_end2:
	.p2align	2, 0x0
                                        # -- End function
	.section	.text._ZNSt6vectorIfSaIfEEC2EmRKS0_,"axG",@progbits,_ZNSt6vectorIfSaIfEEC2EmRKS0_,comdat
	.weak	_ZNSt6vectorIfSaIfEEC2EmRKS0_   # -- Begin function _ZNSt6vectorIfSaIfEEC2EmRKS0_
	.p2align	1, 0x90
	.type	_ZNSt6vectorIfSaIfEEC2EmRKS0_,@function
_ZNSt6vectorIfSaIfEEC2EmRKS0_:          # @_ZNSt6vectorIfSaIfEEC2EmRKS0_
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%rsi, %rax
	shrq	$61, %rax
	jne	.LBB7_7
# %bb.1:                                # %_ZNSt6vectorIfSaIfEE17_S_check_init_lenEmRKS0_.exit
	movq	%rsi, %r14
	movq	%rdi, %rbx
	xorps	%xmm0, %xmm0
	movups	%xmm0, (%rdi)
	movq	$0, 16(%rdi)
	testq	%rsi, %rsi
	je	.LBB7_2
# %bb.3:
	movq	%rbx, %rdi
	movq	%r14, %rsi
	xorl	%edx, %edx
	callq	_ZNSt15__new_allocatorIfE8allocateEmPKv
	movq	%rax, (%rbx)
	leaq	(%rax,%r14,4), %r15
	movq	%r15, 16(%rbx)
	movl	$0, (%rax)
	addq	$4, %rax
	cmpq	$1, %r14
	jne	.LBB7_5
# %bb.4:
	movq	%rax, %r15
	jmp	.LBB7_6
.LBB7_2:                                # %_ZNSt12_Vector_baseIfSaIfEEC2EmRKS0_.exit.thread
	movups	%xmm0, (%rbx)
	movq	$0, 16(%rbx)
	xorl	%r15d, %r15d
	jmp	.LBB7_6
.LBB7_5:                                # %_ZSt6fill_nIPfmfET_S1_T0_RKT1_.exit.loopexit.i.i.i.i
	leaq	-4(,%r14,4), %rdx
	movq	%rax, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
.LBB7_6:
	movq	%r15, 8(%rbx)
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.LBB7_7:
	.cfi_def_cfa_offset 32
	movl	$.L.str.9, %edi
	callq	_ZSt20__throw_length_errorPKc
.Lfunc_end7:
	.size	_ZNSt6vectorIfSaIfEEC2EmRKS0_, .Lfunc_end7-_ZNSt6vectorIfSaIfEEC2EmRKS0_
	.cfi_endproc
                                        # -- End function
	.section	.text._ZSt16__introsort_loopIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEElNS0_5__ops15_Iter_less_iterEEvT_S9_T0_T1_,"axG",@progbits,_ZSt16__introsort_loopIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEElNS0_5__ops15_Iter_less_iterEEvT_S9_T0_T1_,comdat
	.weak	_ZSt16__introsort_loopIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEElNS0_5__ops15_Iter_less_iterEEvT_S9_T0_T1_ # -- Begin function _ZSt16__introsort_loopIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEElNS0_5__ops15_Iter_less_iterEEvT_S9_T0_T1_
	.type	_ZSt16__introsort_loopIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEElNS0_5__ops15_Iter_less_iterEEvT_S9_T0_T1_,@function
_ZSt16__introsort_loopIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEElNS0_5__ops15_Iter_less_iterEEvT_S9_T0_T1_: # @_ZSt16__introsort_loopIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEElNS0_5__ops15_Iter_less_iterEEvT_S9_T0_T1_
	.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	$24, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rsi, %r14
	subq	%rdi, %r14
	sarq	$3, %r14
	cmpq	$17, %r14
	jl	.LBB8_12
# %bb.1:                                # %.lr.ph
	movq	%rdx, %r15
	movq	%rsi, %r12
	movq	%rdi, %rbx
	leaq	8(%rdi), %rax
	movq	%rax, 16(%rsp)                  # 8-byte Spill
	movq	$-8, %rbp
	subq	%rdi, %rbp
.LBB8_2:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB8_5 Depth 2
                                        #       Child Loop BB8_6 Depth 3
                                        #       Child Loop BB8_8 Depth 3
	subq	$1, %r15
	jb	.LBB8_3
# %bb.4:                                #   in Loop: Header=BB8_2 Depth=1
	andq	$-2, %r14
	leaq	(%rbx,%r14,4), %rdx
	leaq	-8(%r12), %rcx
	movq	%rbx, %rdi
	movq	16(%rsp), %r14                  # 8-byte Reload
	movq	%r14, %rsi
	callq	_ZSt22__move_median_to_firstIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_S9_S9_T0_
	movq	%r14, %rax
	movq	%r12, %rcx
.LBB8_5:                                #   Parent Loop BB8_2 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB8_6 Depth 3
                                        #       Child Loop BB8_8 Depth 3
	movsd	(%rbx), %xmm0                   # xmm0 = mem[0],zero
	leaq	(%rax,%rbp), %r14
.LBB8_6:                                #   Parent Loop BB8_2 Depth=1
                                        #     Parent Loop BB8_5 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movsd	(%rax), %xmm1                   # xmm1 = mem[0],zero
	addq	$8, %rax
	addq	$8, %r14
	ucomisd	%xmm1, %xmm0
	ja	.LBB8_6
# %bb.7:                                # %.preheader.i.i.preheader
                                        #   in Loop: Header=BB8_5 Depth=2
	leaq	-8(%rax), %r13
.LBB8_8:                                # %.preheader.i.i
                                        #   Parent Loop BB8_2 Depth=1
                                        #     Parent Loop BB8_5 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movsd	-8(%rcx), %xmm2                 # xmm2 = mem[0],zero
	addq	$-8, %rcx
	ucomisd	%xmm0, %xmm2
	ja	.LBB8_8
# %bb.9:                                #   in Loop: Header=BB8_5 Depth=2
	cmpq	%rcx, %r13
	jae	.LBB8_11
# %bb.10:                               #   in Loop: Header=BB8_5 Depth=2
	movsd	%xmm2, (%r13)
	movsd	%xmm1, (%rcx)
	jmp	.LBB8_5
.LBB8_11:                               # %_ZSt27__unguarded_partition_pivotIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEET_S9_S9_T0_.exit
                                        #   in Loop: Header=BB8_2 Depth=1
	movq	%r13, %rdi
	movq	%r12, %rsi
	movq	%r15, %rdx
	callq	_ZSt16__introsort_loopIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEElNS0_5__ops15_Iter_less_iterEEvT_S9_T0_T1_
	sarq	$3, %r14
	movq	%r13, %r12
	cmpq	$16, %r14
	jg	.LBB8_2
	jmp	.LBB8_12
.LBB8_3:
	leaq	15(%rsp), %rdx
	movq	%rbx, %rdi
	movq	%r12, %rsi
	callq	_ZSt11__make_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_RT0_
	leaq	14(%rsp), %rdx
	movq	%rbx, %rdi
	movq	%r12, %rsi
	callq	_ZSt11__sort_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_RT0_
.LBB8_12:                               # %.loopexit
	addq	$24, %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_end8:
	.size	_ZSt16__introsort_loopIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEElNS0_5__ops15_Iter_less_iterEEvT_S9_T0_T1_, .Lfunc_end8-_ZSt16__introsort_loopIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEElNS0_5__ops15_Iter_less_iterEEvT_S9_T0_T1_
	.cfi_endproc
                                        # -- End function
	.section	.text._ZSt22__final_insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_,"axG",@progbits,_ZSt22__final_insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_,comdat
	.weak	_ZSt22__final_insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_ # -- Begin function _ZSt22__final_insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_
	.type	_ZSt22__final_insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_,@function
_ZSt22__final_insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_: # @_ZSt22__final_insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	pushq	%rax
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	movq	%rsi, %rbx
	movq	%rsi, %rax
	subq	%rdi, %rax
	cmpq	$129, %rax
	jl	.LBB9_8
# %bb.1:
	leaq	128(%rdi), %r14
	movq	%r14, %rsi
	callq	_ZSt16__insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_
.LBB9_2:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB9_5 Depth 2
	cmpq	%rbx, %r14
	je	.LBB9_7
# %bb.3:                                # %.lr.ph.i
                                        #   in Loop: Header=BB9_2 Depth=1
	movsd	-8(%r14), %xmm1                 # xmm1 = mem[0],zero
	movsd	(%r14), %xmm0                   # xmm0 = mem[0],zero
	ucomisd	%xmm0, %xmm1
	movq	%r14, %rax
	jbe	.LBB9_6
# %bb.4:                                # %.lr.ph.i.i.preheader
                                        #   in Loop: Header=BB9_2 Depth=1
	movq	%r14, %rax
.LBB9_5:                                # %.lr.ph.i.i
                                        #   Parent Loop BB9_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	%xmm1, (%rax)
	movsd	-16(%rax), %xmm1                # xmm1 = mem[0],zero
	addq	$-8, %rax
	ucomisd	%xmm0, %xmm1
	ja	.LBB9_5
.LBB9_6:                                # %_ZSt25__unguarded_linear_insertIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops14_Val_less_iterEEvT_T0_.exit.i
                                        #   in Loop: Header=BB9_2 Depth=1
	movsd	%xmm0, (%rax)
	addq	$8, %r14
	jmp	.LBB9_2
.LBB9_8:
	movq	%rbx, %rsi
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	jmp	_ZSt16__insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_ # TAILCALL
.LBB9_7:                                # %_ZSt26__unguarded_insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_.exit
	.cfi_def_cfa_offset 32
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end9:
	.size	_ZSt22__final_insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_, .Lfunc_end9-_ZSt22__final_insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_
	.cfi_endproc
                                        # -- End function
	.section	.text._ZSt11__sort_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_RT0_,"axG",@progbits,_ZSt11__sort_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_RT0_,comdat
	.weak	_ZSt11__sort_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_RT0_ # -- Begin function _ZSt11__sort_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_RT0_
	.type	_ZSt11__sort_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_RT0_,@function
_ZSt11__sort_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_RT0_: # @_ZSt11__sort_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_RT0_
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%rsi, %r15
	subq	%rdi, %r15
	cmpq	$9, %r15
	jl	.LBB10_3
# %bb.1:                                # %.lr.ph.preheader
	movq	%rsi, %rbx
	movq	%rdi, %r14
	addq	$-8, %rbx
.LBB10_2:                               # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movsd	(%rbx), %xmm0                   # xmm0 = mem[0],zero
	movsd	(%r14), %xmm1                   # xmm1 = mem[0],zero
	movsd	%xmm1, (%rbx)
	addq	$-8, %r15
	movq	%r15, %rdx
	sarq	$3, %rdx
	movq	%r14, %rdi
	xorl	%esi, %esi
	callq	_ZSt13__adjust_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEEldNS0_5__ops15_Iter_less_iterEEvT_T0_SA_T1_T2_
	addq	$-8, %rbx
	cmpq	$8, %r15
	jg	.LBB10_2
.LBB10_3:                               # %._crit_edge
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end10:
	.size	_ZSt11__sort_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_RT0_, .Lfunc_end10-_ZSt11__sort_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_RT0_
	.cfi_endproc
                                        # -- End function
	.section	.text._ZSt11__make_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_RT0_,"axG",@progbits,_ZSt11__make_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_RT0_,comdat
	.weak	_ZSt11__make_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_RT0_ # -- Begin function _ZSt11__make_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_RT0_
	.type	_ZSt11__make_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_RT0_,@function
_ZSt11__make_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_RT0_: # @_ZSt11__make_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_RT0_
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%rsi, %rbx
	subq	%rdi, %rbx
	sarq	$3, %rbx
	cmpq	$2, %rbx
	jl	.LBB11_3
# %bb.1:
	movq	%rdi, %r14
	leaq	-2(%rbx), %r15
	shrq	%r15
.LBB11_2:                               # =>This Inner Loop Header: Depth=1
	movsd	(%r14,%r15,8), %xmm0            # xmm0 = mem[0],zero
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%rbx, %rdx
	callq	_ZSt13__adjust_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEEldNS0_5__ops15_Iter_less_iterEEvT_T0_SA_T1_T2_
	addq	$-1, %r15
	jb	.LBB11_2
.LBB11_3:                               # %.loopexit
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end11:
	.size	_ZSt11__make_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_RT0_, .Lfunc_end11-_ZSt11__make_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_RT0_
	.cfi_endproc
                                        # -- End function
	.section	.text._ZSt13__adjust_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEEldNS0_5__ops15_Iter_less_iterEEvT_T0_SA_T1_T2_,"axG",@progbits,_ZSt13__adjust_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEEldNS0_5__ops15_Iter_less_iterEEvT_T0_SA_T1_T2_,comdat
	.weak	_ZSt13__adjust_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEEldNS0_5__ops15_Iter_less_iterEEvT_T0_SA_T1_T2_ # -- Begin function _ZSt13__adjust_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEEldNS0_5__ops15_Iter_less_iterEEvT_T0_SA_T1_T2_
	.type	_ZSt13__adjust_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEEldNS0_5__ops15_Iter_less_iterEEvT_T0_SA_T1_T2_,@function
_ZSt13__adjust_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEEldNS0_5__ops15_Iter_less_iterEEvT_T0_SA_T1_T2_: # @_ZSt13__adjust_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEEldNS0_5__ops15_Iter_less_iterEEvT_T0_SA_T1_T2_
	.cfi_startproc
# %bb.0:
	leaq	-1(%rdx), %rax
	shrq	$63, %rax
	leaq	(%rdx,%rax), %rcx
	decq	%rcx
	sarq	%rcx
	movq	%rsi, %rax
	cmpq	%rsi, %rcx
	jle	.LBB12_6
# %bb.1:                                # %.lr.ph.preheader
	movq	%rsi, %r8
.LBB12_2:                               # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	leaq	(%r8,%r8), %rax
	movsd	8(%rdi,%rax,8), %xmm1           # xmm1 = mem[0],zero
	ucomisd	16(%rdi,%rax,8), %xmm1
	ja	.LBB12_3
# %bb.4:                                # %.lr.ph
                                        #   in Loop: Header=BB12_2 Depth=1
	leaq	2(,%r8,2), %rax
	jmp	.LBB12_5
.LBB12_3:                               #   in Loop: Header=BB12_2 Depth=1
	leaq	1(,%r8,2), %rax
.LBB12_5:                               # %.lr.ph
                                        #   in Loop: Header=BB12_2 Depth=1
	movsd	(%rdi,%rax,8), %xmm1            # xmm1 = mem[0],zero
	movsd	%xmm1, (%rdi,%r8,8)
	movq	%rax, %r8
	cmpq	%rcx, %rax
	jl	.LBB12_2
.LBB12_6:                               # %._crit_edge
	testb	$1, %dl
	jne	.LBB12_9
# %bb.7:
	addq	$-2, %rdx
	sarq	%rdx
	cmpq	%rdx, %rax
	jne	.LBB12_9
# %bb.8:
	leaq	(%rax,%rax), %rcx
	movsd	8(%rdi,%rcx,8), %xmm1           # xmm1 = mem[0],zero
	movsd	%xmm1, (%rdi,%rax,8)
	leaq	1(,%rax,2), %rax
.LBB12_9:
	cmpq	%rsi, %rax
	jle	.LBB12_10
.LBB12_11:                              # %.lr.ph.i
                                        # =>This Inner Loop Header: Depth=1
	leaq	-1(%rax), %rcx
	shrq	$63, %rcx
	addq	%rax, %rcx
	decq	%rcx
	sarq	%rcx
	movsd	(%rdi,%rcx,8), %xmm1            # xmm1 = mem[0],zero
	ucomisd	%xmm1, %xmm0
	jbe	.LBB12_10
# %bb.12:                               #   in Loop: Header=BB12_11 Depth=1
	movsd	%xmm1, (%rdi,%rax,8)
	movq	%rcx, %rax
	cmpq	%rsi, %rcx
	jg	.LBB12_11
	jmp	.LBB12_13
.LBB12_10:
	movq	%rax, %rcx
.LBB12_13:                              # %_ZSt11__push_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEEldNS0_5__ops14_Iter_less_valEEvT_T0_SA_T1_RT2_.exit
	movsd	%xmm0, (%rdi,%rcx,8)
	retq
.Lfunc_end12:
	.size	_ZSt13__adjust_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEEldNS0_5__ops15_Iter_less_iterEEvT_T0_SA_T1_T2_, .Lfunc_end12-_ZSt13__adjust_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEEldNS0_5__ops15_Iter_less_iterEEvT_T0_SA_T1_T2_
	.cfi_endproc
                                        # -- End function
	.section	.text._ZSt22__move_median_to_firstIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_S9_S9_T0_,"axG",@progbits,_ZSt22__move_median_to_firstIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_S9_S9_T0_,comdat
	.weak	_ZSt22__move_median_to_firstIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_S9_S9_T0_ # -- Begin function _ZSt22__move_median_to_firstIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_S9_S9_T0_
	.type	_ZSt22__move_median_to_firstIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_S9_S9_T0_,@function
_ZSt22__move_median_to_firstIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_S9_S9_T0_: # @_ZSt22__move_median_to_firstIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_S9_S9_T0_
	.cfi_startproc
# %bb.0:
	movsd	(%rsi), %xmm1                   # xmm1 = mem[0],zero
	movsd	(%rdx), %xmm2                   # xmm2 = mem[0],zero
	ucomisd	%xmm1, %xmm2
	movsd	(%rcx), %xmm0                   # xmm0 = mem[0],zero
	jbe	.LBB13_6
# %bb.1:
	ucomisd	%xmm2, %xmm0
	jbe	.LBB13_3
# %bb.2:
	movsd	(%rdi), %xmm0                   # xmm0 = mem[0],zero
	movsd	%xmm2, (%rdi)
	movsd	%xmm0, (%rdx)
	retq
.LBB13_6:
	ucomisd	%xmm1, %xmm0
	jbe	.LBB13_8
# %bb.7:
	movsd	(%rdi), %xmm0                   # xmm0 = mem[0],zero
	movsd	%xmm1, (%rdi)
	movsd	%xmm0, (%rsi)
	retq
.LBB13_3:
	ucomisd	%xmm1, %xmm0
	movsd	(%rdi), %xmm2                   # xmm2 = mem[0],zero
	jbe	.LBB13_5
# %bb.4:
	movsd	%xmm0, (%rdi)
	movsd	%xmm2, (%rcx)
	retq
.LBB13_8:
	ucomisd	%xmm2, %xmm0
	movsd	(%rdi), %xmm1                   # xmm1 = mem[0],zero
	jbe	.LBB13_10
# %bb.9:
	movsd	%xmm0, (%rdi)
	movsd	%xmm1, (%rcx)
	retq
.LBB13_5:
	movsd	%xmm1, (%rdi)
	movsd	%xmm2, (%rsi)
	retq
.LBB13_10:
	movsd	%xmm2, (%rdi)
	movsd	%xmm1, (%rdx)
	retq
.Lfunc_end13:
	.size	_ZSt22__move_median_to_firstIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_S9_S9_T0_, .Lfunc_end13-_ZSt22__move_median_to_firstIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_S9_S9_T0_
	.cfi_endproc
                                        # -- End function
	.section	.text._ZSt16__insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_,"axG",@progbits,_ZSt16__insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_,comdat
	.weak	_ZSt16__insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_ # -- Begin function _ZSt16__insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_
	.type	_ZSt16__insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_,@function
_ZSt16__insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_: # @_ZSt16__insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_
	.cfi_startproc
# %bb.0:
	cmpq	%rsi, %rdi
	je	.LBB14_14
# %bb.1:                                # %.preheader
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	subq	$16, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%rsi, %rbx
	movq	%rdi, %r14
	leaq	8(%rdi), %rax
	cmpq	%rsi, %rax
	je	.LBB14_13
# %bb.2:                                # %.lr.ph
	movq	%r14, %r15
.LBB14_3:                               # =>This Loop Header: Depth=1
                                        #     Child Loop BB14_8 Depth 2
	movq	%r15, %rdi
	movq	%rax, %r15
	movsd	8(%rdi), %xmm1                  # xmm1 = mem[0],zero
	movsd	(%r14), %xmm0                   # xmm0 = mem[0],zero
	ucomisd	%xmm1, %xmm0
	jbe	.LBB14_6
# %bb.4:                                #   in Loop: Header=BB14_3 Depth=1
	movq	%r15, %rdx
	subq	%r14, %rdx
	movq	%rdx, %rax
	sarq	$3, %rax
	cmpq	$2, %rax
	jl	.LBB14_10
# %bb.5:                                #   in Loop: Header=BB14_3 Depth=1
	shlq	$3, %rax
	subq	%rax, %rdi
	addq	$16, %rdi
	movq	%r14, %rsi
	movsd	%xmm1, 8(%rsp)                  # 8-byte Spill
	callq	memmove@PLT
	movsd	8(%rsp), %xmm1                  # 8-byte Reload
                                        # xmm1 = mem[0],zero
.LBB14_12:                              # %_ZSt13move_backwardIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEES6_ET0_T_S8_S7_.exit
                                        #   in Loop: Header=BB14_3 Depth=1
	movq	%r14, %rax
	jmp	.LBB14_9
.LBB14_6:                               #   in Loop: Header=BB14_3 Depth=1
	movsd	(%rdi), %xmm0                   # xmm0 = mem[0],zero
	ucomisd	%xmm1, %xmm0
	movq	%r15, %rax
	jbe	.LBB14_9
# %bb.7:                                # %.lr.ph.i.preheader
                                        #   in Loop: Header=BB14_3 Depth=1
	movq	%r15, %rax
.LBB14_8:                               # %.lr.ph.i
                                        #   Parent Loop BB14_3 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	%xmm0, (%rax)
	movsd	-16(%rax), %xmm0                # xmm0 = mem[0],zero
	addq	$-8, %rax
	ucomisd	%xmm1, %xmm0
	ja	.LBB14_8
.LBB14_9:                               # %_ZSt13move_backwardIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEES6_ET0_T_S8_S7_.exit
                                        #   in Loop: Header=BB14_3 Depth=1
	movsd	%xmm1, (%rax)
	leaq	8(%r15), %rax
	cmpq	%rbx, %rax
	jne	.LBB14_3
	jmp	.LBB14_13
.LBB14_10:                              #   in Loop: Header=BB14_3 Depth=1
	movq	%r14, %rax
	cmpq	$8, %rdx
	jne	.LBB14_9
# %bb.11:                               #   in Loop: Header=BB14_3 Depth=1
	movsd	%xmm0, 8(%rdi)
	jmp	.LBB14_12
.LBB14_13:
	addq	$16, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	.cfi_restore %rbx
	.cfi_restore %r14
	.cfi_restore %r15
.LBB14_14:                              # %.loopexit
	retq
.Lfunc_end14:
	.size	_ZSt16__insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_, .Lfunc_end14-_ZSt16__insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_
	.cfi_endproc
                                        # -- End function
	.section	.text._ZNSt15__new_allocatorIfE8allocateEmPKv,"axG",@progbits,_ZNSt15__new_allocatorIfE8allocateEmPKv,comdat
	.weak	_ZNSt15__new_allocatorIfE8allocateEmPKv # -- Begin function _ZNSt15__new_allocatorIfE8allocateEmPKv
	.p2align	1, 0x90
	.type	_ZNSt15__new_allocatorIfE8allocateEmPKv,@function
_ZNSt15__new_allocatorIfE8allocateEmPKv: # @_ZNSt15__new_allocatorIfE8allocateEmPKv
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	movq	%rsi, %rax
	shrq	$61, %rax
	jne	.LBB15_1
# %bb.3:
	shlq	$2, %rsi
	movq	%rsi, %rdi
	popq	%rax
	.cfi_def_cfa_offset 8
	jmp	_Znwm                           # TAILCALL
.LBB15_1:
	.cfi_def_cfa_offset 16
	shrq	$62, %rsi
	je	.LBB15_2
# %bb.4:
	callq	_ZSt28__throw_bad_array_new_lengthv
.LBB15_2:
	callq	_ZSt17__throw_bad_allocv
.Lfunc_end15:
	.size	_ZNSt15__new_allocatorIfE8allocateEmPKv, .Lfunc_end15-_ZNSt15__new_allocatorIfE8allocateEmPKv
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _ZSt18generate_canonicalIfLm24ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_
.LCPI16_0:
	.quad	0x41dfffffff800000              # double 2147483646
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0
.LCPI16_1:
	.long	0x40000000                      # float 2
.LCPI16_2:
	.long	0x5f000000                      # float 9.22337203E+18
.LCPI16_3:
	.long	0x3f800000                      # float 1
	.section	.text._ZSt18generate_canonicalIfLm24ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_,"axG",@progbits,_ZSt18generate_canonicalIfLm24ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_,comdat
	.weak	_ZSt18generate_canonicalIfLm24ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_
	.type	_ZSt18generate_canonicalIfLm24ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_,@function
_ZSt18generate_canonicalIfLm24ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_: # @_ZSt18generate_canonicalIfLm24ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$64, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -16
	movq	%rdi, %rbx
	fldl	.LCPI16_0(%rip)
	fstpt	(%rsp)
	callq	logl
	fstpt	52(%rsp)                        # 10-byte Folded Spill
	flds	.LCPI16_1(%rip)
	fstpt	(%rsp)
	callq	logl
	fldt	52(%rsp)                        # 10-byte Folded Reload
	fdivp	%st, %st(1)
	flds	.LCPI16_2(%rip)
	xorl	%ecx, %ecx
	fxch	%st(1)
	fucomi	%st(1), %st
	fldz
	fcmovnb	%st(2), %st
	fstp	%st(2)
	fsubp	%st, %st(1)
	setae	%cl
	fnstcw	28(%rsp)
	movzwl	28(%rsp), %eax
	orl	$3072, %eax                     # imm = 0xC00
	movw	%ax, 30(%rsp)
	fldcw	30(%rsp)
	fistpll	40(%rsp)
	fldcw	28(%rsp)
	shlq	$63, %rcx
	xorq	40(%rsp), %rcx
	leaq	23(%rcx), %rax
	xorl	%edx, %edx
	divq	%rcx
	movq	%rax, %rcx
	cmpq	$1, %rax
	adcq	$0, %rcx
	movq	(%rbx), %rsi
	xorps	%xmm0, %xmm0
	movss	.LCPI16_3(%rip), %xmm1          # xmm1 = mem[0],zero,zero,zero
	movabsq	$8589934597, %rdi               # imm = 0x200000005
.LBB16_1:                               # %select.unfold
                                        # =>This Inner Loop Header: Depth=1
	imulq	$16807, %rsi, %rsi              # imm = 0x41A7
	movq	%rsi, %rax
	mulq	%rdi
	movq	%rsi, %rax
	subq	%rdx, %rax
	shrq	%rax
	addq	%rdx, %rax
	shrq	$30, %rax
	movq	%rax, %rdx
	shlq	$31, %rdx
	subq	%rdx, %rax
	leaq	(%rsi,%rax), %rdx
	decq	%rdx
	testq	%rdx, %rdx
	js	.LBB16_2
# %bb.3:                                # %select.unfold
                                        #   in Loop: Header=BB16_1 Depth=1
	xorps	%xmm2, %xmm2
	cvtsi2ss	%rdx, %xmm2
	jmp	.LBB16_4
.LBB16_2:                               #   in Loop: Header=BB16_1 Depth=1
	movq	%rdx, %r8
	shrq	%r8
	andl	$1, %edx
	orq	%r8, %rdx
	xorps	%xmm2, %xmm2
	cvtsi2ss	%rdx, %xmm2
	addss	%xmm2, %xmm2
.LBB16_4:                               # %select.unfold
                                        #   in Loop: Header=BB16_1 Depth=1
	addq	%rax, %rsi
	mulss	%xmm1, %xmm2
	movss	%xmm1, 36(%rsp)
	flds	36(%rsp)
	fmull	.LCPI16_0(%rip)
	addss	%xmm2, %xmm0
	fstps	32(%rsp)
	movss	32(%rsp), %xmm1                 # xmm1 = mem[0],zero,zero,zero
	decq	%rcx
	jne	.LBB16_1
# %bb.5:
	movq	%rsi, (%rbx)
	divss	%xmm1, %xmm0
	ucomiss	.LCPI16_3(%rip), %xmm0
	jae	.LBB16_7
# %bb.6:
	addq	$64, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.LBB16_7:
	.cfi_def_cfa_offset 80
	movss	.LCPI16_3(%rip), %xmm0          # xmm0 = mem[0],zero,zero,zero
	xorps	%xmm1, %xmm1
	addq	$64, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	nextafterf                      # TAILCALL
.Lfunc_end16:
	.size	_ZSt18generate_canonicalIfLm24ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_, .Lfunc_end16-_ZSt18generate_canonicalIfLm24ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_
	.cfi_endproc
                                        # -- End function
	.text
	.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	.LBB17_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB17_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z8relu_gpuPfPKfi, %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_end17:
	.size	__hip_module_ctor, .Lfunc_end17-__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	.LBB18_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
.LBB18_2:
	retq
.Lfunc_end18:
	.size	__hip_module_dtor, .Lfunc_end18-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z8relu_gpuPfPKfi,@object       # @_Z8relu_gpuPfPKfi
	.section	.rodata,"a",@progbits
	.globl	_Z8relu_gpuPfPKfi
	.p2align	3, 0x0
_Z8relu_gpuPfPKfi:
	.quad	_Z23__device_stub__relu_gpuPfPKfi
	.size	_Z8relu_gpuPfPKfi, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Usage: "
	.size	.L.str, 8

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	" n"
	.size	.L.str.1, 3

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"n = "
	.size	.L.str.2, 5

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"CPU: "
	.size	.L.str.3, 6

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	" [ms]"
	.size	.L.str.4, 6

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"GPU: "
	.size	.L.str.5, 6

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"Validation: "
	.size	.L.str.6, 13

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"Success"
	.size	.L.str.7, 8

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"Failed"
	.size	.L.str.8, 7

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"cannot create std::vector larger than max_size()"
	.size	.L.str.9, 49

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z8relu_gpuPfPKfi"
	.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

	.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 __gxx_personality_v0
	.addrsig_sym _Z23__device_stub__relu_gpuPfPKfi
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Unwind_Resume
	.addrsig_sym _Z8relu_gpuPfPKfi
	.addrsig_sym _ZSt4cerr
	.addrsig_sym _ZSt4cout
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z8relu_gpuPfPKfi       ; -- Begin function _Z8relu_gpuPfPKfi
	.globl	_Z8relu_gpuPfPKfi
	.p2align	8
	.type	_Z8relu_gpuPfPKfi,@function
_Z8relu_gpuPfPKfi:                      ; @_Z8relu_gpuPfPKfi
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x24
	s_load_b32 s3, s[0:1], 0x10
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_i32_e64 s3, v1
	s_cbranch_execz .LBB0_2
; %bb.1:
	s_load_b128 s[0:3], s[0:1], 0x0
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s2, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v3, vcc_lo, s3, v1, vcc_lo
	v_add_co_u32 v0, vcc_lo, s0, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_load_b32 v2, v[2:3], off
	s_waitcnt vmcnt(0)
	v_max_f32_e32 v2, v2, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_max_f32_e32 v2, 0, v2
	global_store_b32 v[0:1], v2, off
.LBB0_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z8relu_gpuPfPKfi
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 280
		.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 16
		.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	_Z8relu_gpuPfPKfi, .Lfunc_end0-_Z8relu_gpuPfPKfi
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 160
; NumSgprs: 18
; NumVgprs: 4
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; 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
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         28
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         36
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         38
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         40
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         42
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         44
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         46
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         88
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 280
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z8relu_gpuPfPKfi
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z8relu_gpuPfPKfi.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
