/*
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 <fstream>
#include <algorithm>
#include <random>
#include <chrono>
#include <cmath>


static constexpr int NUM_TRIALS = 11;


// CPU版行列積カーネル
void matmul_cpu(float *C, const float *A, const float *B, int n){
	for(int i = 0; i < n; ++i){
		for(int j = 0; j < n; ++j){
			C[i * n + j] = 0.0f;
			for(int k = 0; k < n; ++k){
				C[i * n + j] += A[i * n + k] * B[k * n + j];
			}
		}
	}
}


// GPU版行列積カーネル
__global__ void matmul_gpu(float *C, const float *A, const float *B, int n){
	const int i = blockIdx.y * blockDim.y + threadIdx.y;
	const int j = blockIdx.x * blockDim.x + threadIdx.x;
	float sum = 0.0f;
	for(int k = 0; k < n; ++k){
		sum += A[i * n + k] * B[k * n + j];
	}
	C[i * n + j] = sum;
}

void call_matmul_gpu(float *C, const float *A, const float *B, int n){
	const dim3 bdim(16, 16, 1), gdim(n / 16, n / 16, 1);
	matmul_gpu<<<gdim, bdim>>>(C, A, B, n);
}

// GPU版処理時間計測
// NUM_TRIALS 回計測して中央値を求める
double matmul_gpu_benchmark(float *h_C, const float *h_A, const float *h_B, int n){
	// デバイスメモリの確保
	float *d_C = nullptr, *d_A = nullptr, *d_B = nullptr;
	cudaMalloc(&d_A, sizeof(float) * n * n);
	cudaMalloc(&d_B, sizeof(float) * n * n);
	cudaMalloc(&d_C, sizeof(float) * n * n);
	// 入力データの転送
	cudaMemcpy(d_A, h_A, sizeof(float) * n * n, cudaMemcpyHostToDevice);
	cudaMemcpy(d_B, h_B, sizeof(float) * n * 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_matmul_gpu(d_C, d_A, d_B, 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_C, d_C, sizeof(float) * n * n, cudaMemcpyDeviceToHost);
	// デバイスメモリの開放
	cudaFree(d_A);
	cudaFree(d_B);
	cudaFree(d_C);

	// 中央値を求める
	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 * n; ++i){
		if(std::fabs(expect[i] - actual[i]) > 1e-4){
			std::cerr << "(" << i / n << ", " << i % n << "): " << expect[i] << " != " << actual[i] << std::endl;
			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> A(n * n), B(n * n), cpu_C(n * n), gpu_C(n * n);
	for(int i = 0; i < n * n; ++i){
		A[i] = dist(engine);
		B[i] = dist(engine);
	}

	// CPU側の計算が遅いのでキャッシュする
	const std::string cache_name = "cache_" + std::to_string(n);
	std::ifstream cache_ifs(cache_name, std::ios::binary);
	if(cache_ifs){
		cache_ifs.read(reinterpret_cast<char*>(cpu_C.data()), sizeof(float) * n * n);
	}else{
		std::ofstream cache_ofs(cache_name);
		matmul_cpu(cpu_C.data(), A.data(), B.data(), n);
		cache_ofs.write(reinterpret_cast<const char*>(cpu_C.data()), sizeof(float) * n * n);
	}

	const auto gpu_duration =
		matmul_gpu_benchmark(gpu_C.data(), A.data(), B.data(), n);
	std::cout << "GPU: " << gpu_duration << " [ms]" << std::endl;
	std::cout << "     " << 2.0 * n * n * n / gpu_duration * 1e-9 << " [TFLOPS]" << std::endl;

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

	return 0;
}

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

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB5176:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE5176:
	.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:
.LFB6219:
	.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
.LFE6219:
	.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:
.LFB6221:
	.cfi_startproc
	jmp	*%rsi
	.cfi_endproc
.LFE6221:
	.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:
.LFB6222:
	.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
.LFE6222:
	.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:
.LFB6224:
	.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
.LFE6224:
	.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
	.globl	_Z10matmul_cpuPfPKfS1_i
	.type	_Z10matmul_cpuPfPKfS1_i, @function
_Z10matmul_cpuPfPKfS1_i:
.LFB5166:
	.cfi_startproc
	endbr64
	movq	%rsi, %r11
	movslq	%ecx, %rsi
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movq	%rdi, %r10
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	salq	$2, %rsi
	xorl	%r9d, %r9d
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	%rdx, %rbx
	xorl	%edx, %edx
.L75:
	cmpl	%ecx, %r9d
	jge	.L82
	leaq	(%r10,%rdx), %rax
	xorl	%edi, %edi
	leaq	(%r11,%rdx), %r12
.L77:
	movl	$0x00000000, (%rax)
	leaq	(%rbx,%rdi,4), %rbp
	xorl	%r8d, %r8d
.L76:
	movss	(%r12,%r8,4), %xmm0
	mulss	0(%rbp), %xmm0
	incq	%r8
	addq	%rsi, %rbp
	addss	(%rax), %xmm0
	movss	%xmm0, (%rax)
	cmpl	%r8d, %ecx
	jg	.L76
	incq	%rdi
	addq	$4, %rax
	cmpl	%edi, %ecx
	jg	.L77
	incl	%r9d
	addq	%rsi, %rdx
	jmp	.L75
.L82:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE5166:
	.size	_Z10matmul_cpuPfPKfS1_i, .-_Z10matmul_cpuPfPKfS1_i
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC3:
	.string	"("
.LC4:
	.string	", "
.LC5:
	.string	"): "
.LC6:
	.string	" != "
	.text
	.globl	_Z8validatePKfS0_i
	.type	_Z8validatePKfS0_i, @function
_Z8validatePKfS0_i:
.LFB5172:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movb	$1, %al
	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
	movl	%edx, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	imull	%edx, %r13d
	movq	%rsi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movl	%edx, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	xorl	%ebx, %ebx
	pushq	%rcx
	.cfi_def_cfa_offset 64
.L84:
	cmpl	%ebx, %r13d
	jle	.L90
	movss	(%r14,%rbx,4), %xmm0
	subss	(%r12,%rbx,4), %xmm0
	andps	.LC1(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	comisd	.LC2(%rip), %xmm0
	jbe	.L85
	leaq	.LC3(%rip), %rsi
	leaq	_ZSt4cerr(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	movl	%ebx, %eax
	cltd
	idivl	%ebp
	movl	%eax, %esi
	movl	%edx, %r15d
	call	_ZNSolsEi@PLT
	leaq	.LC4(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%r15d, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC5(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cvtss2sd	(%r14,%rbx,4), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	leaq	.LC6(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cvtss2sd	(%r12,%rbx,4), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	xorl	%eax, %eax
.L85:
	incq	%rbx
	jmp	.L84
.L90:
	popq	%rdx
	.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
.LFE5172:
	.size	_Z8validatePKfS0_i, .-_Z8validatePKfS0_i
	.globl	_Z37__device_stub__Z10matmul_gpuPfPKfS1_iPfPKfS1_i
	.type	_Z37__device_stub__Z10matmul_gpuPfPKfS1_iPfPKfS1_i, @function
_Z37__device_stub__Z10matmul_gpuPfPKfS1_iPfPKfS1_i:
.LFB5198:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 4(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L91
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z10matmul_gpuPfPKfS1_i(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 176
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L91:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L93
	call	__stack_chk_fail@PLT
.L93:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE5198:
	.size	_Z37__device_stub__Z10matmul_gpuPfPKfS1_iPfPKfS1_i, .-_Z37__device_stub__Z10matmul_gpuPfPKfS1_iPfPKfS1_i
	.globl	_Z10matmul_gpuPfPKfS1_i
	.type	_Z10matmul_gpuPfPKfS1_i, @function
_Z10matmul_gpuPfPKfS1_i:
.LFB5199:
	.cfi_startproc
	endbr64
	jmp	_Z37__device_stub__Z10matmul_gpuPfPKfS1_iPfPKfS1_i
	.cfi_endproc
.LFE5199:
	.size	_Z10matmul_gpuPfPKfS1_i, .-_Z10matmul_gpuPfPKfS1_i
	.globl	_Z15call_matmul_gpuPfPKfS1_i
	.type	_Z15call_matmul_gpuPfPKfS1_i, @function
_Z15call_matmul_gpuPfPKfS1_i:
.LFB5167:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movq	%rdx, %r13
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movq	%rsi, %r12
	movl	$1, %esi
	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
	movl	%ecx, %ebx
	movl	$16, %ecx
	movl	%ebx, %eax
	cltd
	idivl	%ecx
	subq	$40, %rsp
	.cfi_def_cfa_offset 80
	movl	$1, %ecx
	movabsq	$68719476752, %rdx
	movl	%eax, 20(%rsp)
	movl	%eax, 24(%rsp)
	movq	20(%rsp), %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L96
	addq	$40, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 40
	movl	%ebx, %ecx
	movq	%r13, %rdx
	movq	%r12, %rsi
	popq	%rbx
	.cfi_def_cfa_offset 32
	movq	%rbp, %rdi
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	jmp	_Z37__device_stub__Z10matmul_gpuPfPKfS1_iPfPKfS1_i
.L96:
	.cfi_restore_state
	addq	$40, %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
.LFE5167:
	.size	_Z15call_matmul_gpuPfPKfS1_i, .-_Z15call_matmul_gpuPfPKfS1_i
	.section	.rodata.str1.1
.LC7:
	.string	"_Z10matmul_gpuPfPKfS1_i"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB5201:
	.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	.LC7(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z10matmul_gpuPfPKfS1_i(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE5201:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.rodata._ZNSt6vectorIfSaIfEEC2EmRKS0_.str1.1,"aMS",@progbits,1
.LC8:
	.string	"cannot create std::vector larger than max_size()"
	.section	.text._ZNSt6vectorIfSaIfEEC2EmRKS0_,"axG",@progbits,_ZNSt6vectorIfSaIfEEC5EmRKS0_,comdat
	.align 2
	.weak	_ZNSt6vectorIfSaIfEEC2EmRKS0_
	.type	_ZNSt6vectorIfSaIfEEC2EmRKS0_, @function
_ZNSt6vectorIfSaIfEEC2EmRKS0_:
.LFB5598:
	.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	.L102
	leaq	.LC8(%rip), %rdi
	call	_ZSt20__throw_length_errorPKc@PLT
.L102:
	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	.L103
	movq	%r12, %rdi
	call	_Znwm@PLT
.L103:
	addq	%rax, %r12
	movq	%rax, (%rbx)
	movq	%rax, 8(%rbx)
	movq	%r12, 16(%rbx)
	testq	%rbp, %rbp
	je	.L104
	movl	$0x00000000, (%rax)
	addq	$4, %rax
	decq	%rbp
	je	.L104
.L105:
	cmpq	%rax, %r12
	je	.L104
	movl	$0x00000000, (%rax)
	addq	$4, %rax
	jmp	.L105
.L104:
	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
.LFE5598:
	.size	_ZNSt6vectorIfSaIfEEC2EmRKS0_, .-_ZNSt6vectorIfSaIfEEC2EmRKS0_
	.weak	_ZNSt6vectorIfSaIfEEC1EmRKS0_
	.set	_ZNSt6vectorIfSaIfEEC1EmRKS0_,_ZNSt6vectorIfSaIfEEC2EmRKS0_
	.section	.text._ZNSt6vectorIfSaIfEED2Ev,"axG",@progbits,_ZNSt6vectorIfSaIfEED5Ev,comdat
	.align 2
	.weak	_ZNSt6vectorIfSaIfEED2Ev
	.type	_ZNSt6vectorIfSaIfEED2Ev, @function
_ZNSt6vectorIfSaIfEED2Ev:
.LFB5601:
	.cfi_startproc
	endbr64
	movq	(%rdi), %rax
	testq	%rax, %rax
	je	.L116
	movq	16(%rdi), %rsi
	movq	%rax, %rdi
	subq	%rax, %rsi
	jmp	_ZdlPvm@PLT
.L116:
	ret
	.cfi_endproc
.LFE5601:
	.size	_ZNSt6vectorIfSaIfEED2Ev, .-_ZNSt6vectorIfSaIfEED2Ev
	.weak	_ZNSt6vectorIfSaIfEED1Ev
	.set	_ZNSt6vectorIfSaIfEED1Ev,_ZNSt6vectorIfSaIfEED2Ev
	.section	.text._ZNSt12_Vector_baseIdSaIdEED2Ev,"axG",@progbits,_ZNSt12_Vector_baseIdSaIdEED5Ev,comdat
	.align 2
	.weak	_ZNSt12_Vector_baseIdSaIdEED2Ev
	.type	_ZNSt12_Vector_baseIdSaIdEED2Ev, @function
_ZNSt12_Vector_baseIdSaIdEED2Ev:
.LFB5801:
	.cfi_startproc
	endbr64
	movq	(%rdi), %rax
	testq	%rax, %rax
	je	.L118
	movq	16(%rdi), %rsi
	movq	%rax, %rdi
	subq	%rax, %rsi
	jmp	_ZdlPvm@PLT
.L118:
	ret
	.cfi_endproc
.LFE5801:
	.size	_ZNSt12_Vector_baseIdSaIdEED2Ev, .-_ZNSt12_Vector_baseIdSaIdEED2Ev
	.weak	_ZNSt12_Vector_baseIdSaIdEED1Ev
	.set	_ZNSt12_Vector_baseIdSaIdEED1Ev,_ZNSt12_Vector_baseIdSaIdEED2Ev
	.text
	.globl	_Z20matmul_gpu_benchmarkPfPKfS1_i
	.type	_Z20matmul_gpu_benchmarkPfPKfS1_i, @function
_Z20matmul_gpu_benchmarkPfPKfS1_i:
.LFB5168:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA5168
	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
	movl	%ecx, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rsi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rdx, %rbx
	subq	$88, %rsp
	.cfi_def_cfa_offset 144
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	xorl	%eax, %eax
	leaq	32(%rsp), %rdi
	movq	%rax, 24(%rsp)
	movq	%rax, 32(%rsp)
	movq	%rax, 40(%rsp)
	movslq	%ecx, %rax
	imulq	%rax, %rax
	salq	$2, %rax
	movq	%rax, %rsi
	movq	%rax, 8(%rsp)
.LEHB0:
	call	cudaMalloc@PLT
	movq	8(%rsp), %rsi
	leaq	40(%rsp), %rdi
	call	cudaMalloc@PLT
	movq	8(%rsp), %rsi
	leaq	24(%rsp), %rdi
	call	cudaMalloc@PLT
	movq	8(%rsp), %rdx
	movq	32(%rsp), %rdi
	movq	%rbp, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rdx
	movq	40(%rsp), %rdi
	movq	%rbx, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movl	$88, %edi
	call	_Znwm@PLT
.LEHE0:
	leaq	88(%rax), %rbp
	movq	%rax, 48(%rsp)
	movq	%rax, %rbx
	movq	%rbp, 64(%rsp)
	movq	$0x000000000, (%rax)
	xorl	%eax, %eax
.L121:
	addq	$8, %rax
	cmpq	$88, %rax
	je	.L131
	movq	$0x000000000, (%rbx,%rax)
	jmp	.L121
.L131:
	movq	%rbp, 56(%rsp)
	xorl	%r12d, %r12d
.L123:
	call	_ZNSt6chrono3_V212steady_clock3nowEv@PLT
	movq	40(%rsp), %rdx
	movq	32(%rsp), %rsi
	movl	%r13d, %ecx
	movq	24(%rsp), %rdi
	movq	%rax, %r15
.LEHB1:
	call	_Z15call_matmul_gpuPfPKfS1_i
	call	cudaDeviceSynchronize@PLT
	call	_ZNSt6chrono3_V212steady_clock3nowEv@PLT
	movl	$1000, %ecx
	subq	%r15, %rax
	cqto
	idivq	%rcx
	cvtsi2sdq	%rax, %xmm0
	mulsd	.LC10(%rip), %xmm0
	movsd	%xmm0, (%rbx,%r12,8)
	incq	%r12
	cmpq	$11, %r12
	jne	.L123
	movq	8(%rsp), %rdx
	movq	24(%rsp), %rsi
	movl	$2, %ecx
	movq	%r14, %rdi
	call	cudaMemcpy@PLT
	movq	32(%rsp), %rdi
	call	cudaFree@PLT
	movq	40(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
.LEHE1:
	movl	$6, %edx
	movq	%rbp, %rsi
	movq	%rbx, %rdi
	call	_ZSt16__introsort_loopIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEElNS0_5__ops15_Iter_less_iterEEvT_S9_T0_T1_.isra.0
	movq	%rbp, %rsi
	movq	%rbx, %rdi
	call	_ZSt16__insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_.isra.0
	movsd	40(%rbx), %xmm0
	leaq	48(%rsp), %rdi
	movsd	%xmm0, 8(%rsp)
	call	_ZNSt12_Vector_baseIdSaIdEED2Ev
	movsd	8(%rsp), %xmm0
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L126
	jmp	.L129
.L127:
	endbr64
	movq	%rax, %rbx
.L124:
	leaq	48(%rsp), %rdi
	call	_ZNSt12_Vector_baseIdSaIdEED2Ev
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L129
	movq	%rbx, %rdi
.LEHB2:
	call	_Unwind_Resume@PLT
.LEHE2:
.L129:
	call	__stack_chk_fail@PLT
.L126:
	addq	$88, %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
.LFE5168:
	.globl	__gxx_personality_v0
	.section	.gcc_except_table,"a",@progbits
.LLSDA5168:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE5168-.LLSDACSB5168
.LLSDACSB5168:
	.uleb128 .LEHB0-.LFB5168
	.uleb128 .LEHE0-.LEHB0
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB1-.LFB5168
	.uleb128 .LEHE1-.LEHB1
	.uleb128 .L127-.LFB5168
	.uleb128 0
	.uleb128 .LEHB2-.LFB5168
	.uleb128 .LEHE2-.LEHB2
	.uleb128 0
	.uleb128 0
.LLSDACSE5168:
	.text
	.size	_Z20matmul_gpu_benchmarkPfPKfS1_i, .-_Z20matmul_gpu_benchmarkPfPKfS1_i
	.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_:
.LFB6106:
	.cfi_startproc
	endbr64
	imulq	$16807, (%rdi), %rax
	movl	$2147483647, %ecx
	xorl	%edx, %edx
	divq	%rcx
	movq	%rdx, (%rdi)
	decq	%rdx
	js	.L133
	cvtsi2ssq	%rdx, %xmm0
	jmp	.L134
.L133:
	movq	%rdx, %rax
	andl	$1, %edx
	shrq	%rax
	orq	%rdx, %rax
	cvtsi2ssq	%rax, %xmm0
	addss	%xmm0, %xmm0
.L134:
	xorps	%xmm1, %xmm1
	addss	%xmm1, %xmm0
	mulss	.LC12(%rip), %xmm0
	comiss	.LC13(%rip), %xmm0
	jb	.L132
	movss	.LC11(%rip), %xmm0
.L132:
	ret
	.cfi_endproc
.LFE6106:
	.size	_ZSt18generate_canonicalIfLm24ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_, .-_ZSt18generate_canonicalIfLm24ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_
	.section	.rodata.str1.1
.LC14:
	.string	"Success"
.LC15:
	.string	"Failed"
.LC16:
	.string	"Usage: "
.LC17:
	.string	" n"
.LC18:
	.string	"n = "
.LC20:
	.string	"cache_"
.LC21:
	.string	"GPU: "
.LC22:
	.string	" [ms]"
.LC23:
	.string	"     "
.LC25:
	.string	" [TFLOPS]"
.LC26:
	.string	"Validation: "
.LC19:
	.string	"00010203040506070809101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899"
	.section	.text.startup
	.globl	main
	.type	main, @function
main:
.LFB5173:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA5173
	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	$1448, %rsp
	.cfi_def_cfa_offset 1504
	movq	%fs:40, %rax
	movq	%rax, 1432(%rsp)
	xorl	%eax, %eax
	decl	%edi
	jg	.L138
	leaq	.LC16(%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	.LC17(%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	.L139
.L138:
	movq	8(%rsi), %rdi
	leaq	_ZSt4cout(%rip), %r14
	call	atoi@PLT
	leaq	.LC18(%rip), %rsi
	movq	%r14, %rdi
	movl	%eax, %ebx
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %r12d
	movl	%ebx, %esi
	movq	%rax, %rdi
	imull	%ebx, %r12d
	call	_ZNSolsEi@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	136(%rsp), %rax
	movslq	%r12d, %rbp
	movq	$1, 56(%rsp)
	movq	%rax, 8(%rsp)
	leaq	64(%rsp), %rax
	movq	8(%rsp), %rdx
	movq	%rbp, %rsi
	movq	%rax, %rdi
	movq	%rax, 16(%rsp)
	call	_ZNSt6vectorIfSaIfEEC1EmRKS0_
.LEHE3:
	leaq	88(%rsp), %rax
	movq	8(%rsp), %rdx
	movq	%rbp, %rsi
	movq	%rax, %rdi
	movq	%rax, 24(%rsp)
.LEHB4:
	call	_ZNSt6vectorIfSaIfEEC1EmRKS0_
.LEHE4:
	leaq	112(%rsp), %rax
	movq	8(%rsp), %rdx
	movq	%rbp, %rsi
	movq	%rax, %rdi
	movq	%rax, 32(%rsp)
.LEHB5:
	call	_ZNSt6vectorIfSaIfEEC1EmRKS0_
.LEHE5:
	movq	8(%rsp), %rdi
	leaq	55(%rsp), %rdx
	movq	%rbp, %rsi
.LEHB6:
	call	_ZNSt6vectorIfSaIfEEC1EmRKS0_
.LEHE6:
	xorl	%r13d, %r13d
.L140:
	cmpl	%r13d, %r12d
	jle	.L176
	leaq	56(%rsp), %rbp
	movq	%rbp, %rdi
	call	_ZSt18generate_canonicalIfLm24ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_
	movq	64(%rsp), %rax
	movq	%rbp, %rdi
	addss	%xmm0, %xmm0
	subss	.LC13(%rip), %xmm0
	movss	%xmm0, (%rax,%r13,4)
	call	_ZSt18generate_canonicalIfLm24ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_
	movq	88(%rsp), %rax
	addss	%xmm0, %xmm0
	subss	.LC13(%rip), %xmm0
	movss	%xmm0, (%rax,%r13,4)
	incq	%r13
	jmp	.L140
.L176:
	movl	%ebx, %r15d
	movl	%ebx, %ebp
	movl	$1, %r13d
	movl	$10000, %ecx
	shrl	$31, %r15d
	negl	%ebp
	cmovs	%ebx, %ebp
	movl	%ebp, %eax
.L147:
	cmpl	$9, %eax
	jbe	.L143
	cmpl	$99, %eax
	ja	.L144
	incl	%r13d
	jmp	.L143
.L144:
	cmpl	$999, %eax
	ja	.L145
	addl	$2, %r13d
	jmp	.L143
.L145:
	cmpl	$9999, %eax
	ja	.L146
	addl	$3, %r13d
	jmp	.L143
.L146:
	movl	%eax, %eax
	xorl	%edx, %edx
	addl	$4, %r13d
	divq	%rcx
	jmp	.L147
.L143:
	leaq	704(%rsp), %r12
	leal	(%r15,%r13), %esi
	movl	$45, %edx
	decl	%r13d
	leaq	720(%rsp), %rax
	movq	%r12, %rdi
	movq	%rax, 704(%rsp)
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructEmc@PLT
	movzbl	%r15b, %r8d
	movl	$201, %ecx
	leaq	1231(%rsp), %rdi
	leaq	.LC19(%rip), %rsi
	addq	704(%rsp), %r8
	rep movsb
	movl	$100, %ecx
.L148:
	cmpl	$99, %ebp
	jbe	.L177
	movl	%ebp, %eax
	xorl	%edx, %edx
	movl	%r13d, %esi
	divl	%ecx
	addl	%edx, %edx
	movl	%eax, %ebp
	leal	1(%rdx), %eax
	movb	1231(%rsp,%rdx), %dl
	movb	1231(%rsp,%rax), %al
	movb	%al, (%r8,%rsi)
	leal	-1(%r13), %eax
	subl	$2, %r13d
	movb	%dl, (%r8,%rax)
	jmp	.L148
.L177:
	leal	48(%rbp), %eax
	cmpl	$9, %ebp
	jbe	.L151
	addl	%ebp, %ebp
	leal	1(%rbp), %eax
	movb	1231(%rsp,%rax), %al
	movb	%al, 1(%r8)
	movb	1231(%rsp,%rbp), %al
.L151:
	movb	%al, (%r8)
	leaq	.LC20(%rip), %rdx
	xorl	%esi, %esi
	movq	%r12, %rdi
.LEHB7:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE6insertEmPKc@PLT
.LEHE7:
	leaq	160(%rsp), %r13
	movq	%rax, %rsi
	movq	%r13, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC1EOS4_@PLT
	movq	%r12, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	movl	$4, %edx
	movq	%r13, %rsi
	movq	%r12, %rdi
.LEHB8:
	call	_ZNSt14basic_ifstreamIcSt11char_traitsIcEEC1ERKNSt7__cxx1112basic_stringIcS1_SaIcEEESt13_Ios_Openmode@PLT
.LEHE8:
	testb	$5, 992(%rsp)
	movslq	%ebx, %rbp
	jne	.L152
	imulq	%rbp, %rbp
	movq	112(%rsp), %rsi
	movq	%r12, %rdi
	leaq	0(,%rbp,4), %rdx
.LEHB9:
	call	_ZNSi4readEPcl@PLT
	jmp	.L153
.L152:
	leaq	192(%rsp), %rdi
	movl	$16, %edx
	movq	%r13, %rsi
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEEC1ERKNSt7__cxx1112basic_stringIcS1_SaIcEEESt13_Ios_Openmode@PLT
.LEHE9:
	movq	112(%rsp), %r15
	imulq	%rbp, %rbp
	movq	88(%rsp), %rdx
	movl	%ebx, %ecx
	movq	64(%rsp), %rsi
	movq	%r15, %rdi
	call	_Z10matmul_cpuPfPKfS1_i
	leaq	0(,%rbp,4), %rdx
	movq	%r15, %rsi
	leaq	192(%rsp), %rbp
	movq	%rbp, %rdi
.LEHB10:
	call	_ZNSo5writeEPKcl@PLT
.LEHE10:
	movq	%rbp, %rdi
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEED1Ev@PLT
.L153:
	movq	136(%rsp), %rbp
	movq	88(%rsp), %rdx
	movl	%ebx, %ecx
	movq	64(%rsp), %rsi
	movq	%rbp, %rdi
.LEHB11:
	call	_Z20matmul_gpu_benchmarkPfPKfS1_i
	leaq	.LC21(%rip), %rsi
	movq	%r14, %rdi
	movsd	%xmm0, 40(%rsp)
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	40(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	%rax, %rdi
	leaq	.LC22(%rip), %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
	leaq	.LC23(%rip), %rsi
	movq	%r14, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cvtsi2sdl	%ebx, %xmm1
	movq	%rax, %rdi
	movaps	%xmm1, %xmm0
	addsd	%xmm1, %xmm0
	mulsd	%xmm1, %xmm0
	mulsd	%xmm1, %xmm0
	divsd	40(%rsp), %xmm0
	mulsd	.LC24(%rip), %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	%rax, %rdi
	leaq	.LC25(%rip), %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	112(%rsp), %rdi
	movl	%ebx, %edx
	movq	%rbp, %rsi
	call	_Z8validatePKfS0_i
	leaq	.LC26(%rip), %rsi
	movq	%r14, %rdi
	movl	%eax, %ebx
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	testb	%bl, %bl
	leaq	.LC15(%rip), %rsi
	leaq	.LC14(%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
.LEHE11:
	movq	%r12, %rdi
	call	_ZNSt14basic_ifstreamIcSt11char_traitsIcEED1Ev@PLT
	movq	%r13, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	movq	8(%rsp), %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	movq	32(%rsp), %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	movq	24(%rsp), %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	movq	16(%rsp), %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
.L139:
	movq	1432(%rsp), %rax
	subq	%fs:40, %rax
	je	.L164
	jmp	.L173
.L169:
	endbr64
	movq	%rax, %rbx
.L155:
	movq	%r12, %rdi
	jmp	.L175
.L172:
	endbr64
	movq	%rax, %rbx
.L157:
	movq	%rbp, %rdi
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEED1Ev@PLT
	jmp	.L158
.L171:
	endbr64
	movq	%rax, %rbx
.L158:
	movq	%r12, %rdi
	call	_ZNSt14basic_ifstreamIcSt11char_traitsIcEED1Ev@PLT
	jmp	.L159
.L170:
	endbr64
	movq	%rax, %rbx
.L159:
	movq	%r13, %rdi
.L175:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	movq	8(%rsp), %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	jmp	.L160
.L168:
	endbr64
	movq	%rax, %rbx
.L160:
	movq	32(%rsp), %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	jmp	.L161
.L167:
	endbr64
	movq	%rax, %rbx
.L161:
	movq	24(%rsp), %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	jmp	.L162
.L166:
	endbr64
	movq	%rax, %rbx
.L162:
	movq	16(%rsp), %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	movq	1432(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L173
	movq	%rbx, %rdi
.LEHB12:
	call	_Unwind_Resume@PLT
.LEHE12:
.L173:
	call	__stack_chk_fail@PLT
.L164:
	addq	$1448, %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
.LFE5173:
	.section	.gcc_except_table
.LLSDA5173:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE5173-.LLSDACSB5173
.LLSDACSB5173:
	.uleb128 .LEHB3-.LFB5173
	.uleb128 .LEHE3-.LEHB3
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB4-.LFB5173
	.uleb128 .LEHE4-.LEHB4
	.uleb128 .L166-.LFB5173
	.uleb128 0
	.uleb128 .LEHB5-.LFB5173
	.uleb128 .LEHE5-.LEHB5
	.uleb128 .L167-.LFB5173
	.uleb128 0
	.uleb128 .LEHB6-.LFB5173
	.uleb128 .LEHE6-.LEHB6
	.uleb128 .L168-.LFB5173
	.uleb128 0
	.uleb128 .LEHB7-.LFB5173
	.uleb128 .LEHE7-.LEHB7
	.uleb128 .L169-.LFB5173
	.uleb128 0
	.uleb128 .LEHB8-.LFB5173
	.uleb128 .LEHE8-.LEHB8
	.uleb128 .L170-.LFB5173
	.uleb128 0
	.uleb128 .LEHB9-.LFB5173
	.uleb128 .LEHE9-.LEHB9
	.uleb128 .L171-.LFB5173
	.uleb128 0
	.uleb128 .LEHB10-.LFB5173
	.uleb128 .LEHE10-.LEHB10
	.uleb128 .L172-.LFB5173
	.uleb128 0
	.uleb128 .LEHB11-.LFB5173
	.uleb128 .LEHE11-.LEHB11
	.uleb128 .L171-.LFB5173
	.uleb128 0
	.uleb128 .LEHB12-.LFB5173
	.uleb128 .LEHE12-.LEHB12
	.uleb128 0
	.uleb128 0
.LLSDACSE5173:
	.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
.LC10:
	.long	-755914244
	.long	1062232653
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC11:
	.long	1065353215
	.align 4
.LC12:
	.long	805306368
	.align 4
.LC13:
	.long	1065353216
	.section	.rodata.cst8
	.align 8
.LC24:
	.long	-400107883
	.long	1041313291
	.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 <fstream>
#include <algorithm>
#include <random>
#include <chrono>
#include <cmath>


static constexpr int NUM_TRIALS = 11;


// CPU版行列積カーネル
void matmul_cpu(float *C, const float *A, const float *B, int n){
	for(int i = 0; i < n; ++i){
		for(int j = 0; j < n; ++j){
			C[i * n + j] = 0.0f;
			for(int k = 0; k < n; ++k){
				C[i * n + j] += A[i * n + k] * B[k * n + j];
			}
		}
	}
}


// GPU版行列積カーネル
__global__ void matmul_gpu(float *C, const float *A, const float *B, int n){
	const int i = blockIdx.y * blockDim.y + threadIdx.y;
	const int j = blockIdx.x * blockDim.x + threadIdx.x;
	float sum = 0.0f;
	for(int k = 0; k < n; ++k){
		sum += A[i * n + k] * B[k * n + j];
	}
	C[i * n + j] = sum;
}

void call_matmul_gpu(float *C, const float *A, const float *B, int n){
	const dim3 bdim(16, 16, 1), gdim(n / 16, n / 16, 1);
	matmul_gpu<<<gdim, bdim>>>(C, A, B, n);
}

// GPU版処理時間計測
// NUM_TRIALS 回計測して中央値を求める
double matmul_gpu_benchmark(float *h_C, const float *h_A, const float *h_B, int n){
	// デバイスメモリの確保
	float *d_C = nullptr, *d_A = nullptr, *d_B = nullptr;
	hipMalloc(&d_A, sizeof(float) * n * n);
	hipMalloc(&d_B, sizeof(float) * n * n);
	hipMalloc(&d_C, sizeof(float) * n * n);
	// 入力データの転送
	hipMemcpy(d_A, h_A, sizeof(float) * n * n, hipMemcpyHostToDevice);
	hipMemcpy(d_B, h_B, sizeof(float) * n * 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_matmul_gpu(d_C, d_A, d_B, 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_C, d_C, sizeof(float) * n * n, hipMemcpyDeviceToHost);
	// デバイスメモリの開放
	hipFree(d_A);
	hipFree(d_B);
	hipFree(d_C);

	// 中央値を求める
	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 * n; ++i){
		if(std::fabs(expect[i] - actual[i]) > 1e-4){
			std::cerr << "(" << i / n << ", " << i % n << "): " << expect[i] << " != " << actual[i] << std::endl;
			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> A(n * n), B(n * n), cpu_C(n * n), gpu_C(n * n);
	for(int i = 0; i < n * n; ++i){
		A[i] = dist(engine);
		B[i] = dist(engine);
	}

	// CPU側の計算が遅いのでキャッシュする
	const std::string cache_name = "cache_" + std::to_string(n);
	std::ifstream cache_ifs(cache_name, std::ios::binary);
	if(cache_ifs){
		cache_ifs.read(reinterpret_cast<char*>(cpu_C.data()), sizeof(float) * n * n);
	}else{
		std::ofstream cache_ofs(cache_name);
		matmul_cpu(cpu_C.data(), A.data(), B.data(), n);
		cache_ofs.write(reinterpret_cast<const char*>(cpu_C.data()), sizeof(float) * n * n);
	}

	const auto gpu_duration =
		matmul_gpu_benchmark(gpu_C.data(), A.data(), B.data(), n);
	std::cout << "GPU: " << gpu_duration << " [ms]" << std::endl;
	std::cout << "     " << 2.0 * n * n * n / gpu_duration * 1e-9 << " [TFLOPS]" << std::endl;

	const auto valid = validate(cpu_C.data(), gpu_C.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	_Z10matmul_cpuPfPKfS1_i         # -- Begin function _Z10matmul_cpuPfPKfS1_i
	.type	_Z10matmul_cpuPfPKfS1_i,@function
_Z10matmul_cpuPfPKfS1_i:                # @_Z10matmul_cpuPfPKfS1_i
	.cfi_startproc
# %bb.0:
	testl	%ecx, %ecx
	jle	.LBB0_8
# %bb.1:                                # %.preheader.lr.ph
	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
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movl	%ecx, %eax
	leaq	(,%rax,4), %r8
	xorl	%r9d, %r9d
	xorl	%r10d, %r10d
.LBB0_2:                                # %.lr.ph29
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB0_3 Depth 2
                                        #       Child Loop BB0_4 Depth 3
	movl	%r9d, %r11d
	leaq	(%rsi,%r11,4), %r11
	movq	%r10, %rbx
	imulq	%rax, %rbx
	leaq	(%rdi,%rbx,4), %rbx
	movq	%rdx, %r14
	xorl	%r15d, %r15d
.LBB0_3:                                # %.lr.ph
                                        #   Parent Loop BB0_2 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB0_4 Depth 3
	movl	$0, (%rbx,%r15,4)
	xorps	%xmm0, %xmm0
	movq	%r14, %r12
	xorl	%r13d, %r13d
.LBB0_4:                                #   Parent Loop BB0_2 Depth=1
                                        #     Parent Loop BB0_3 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movss	(%r11,%r13,4), %xmm1            # xmm1 = mem[0],zero,zero,zero
	mulss	(%r12), %xmm1
	addss	%xmm1, %xmm0
	movss	%xmm0, (%rbx,%r15,4)
	incq	%r13
	addq	%r8, %r12
	cmpq	%r13, %rax
	jne	.LBB0_4
# %bb.5:                                # %._crit_edge
                                        #   in Loop: Header=BB0_3 Depth=2
	incq	%r15
	addq	$4, %r14
	cmpq	%rax, %r15
	jne	.LBB0_3
# %bb.6:                                # %._crit_edge30
                                        #   in Loop: Header=BB0_2 Depth=1
	incq	%r10
	addl	%ecx, %r9d
	cmpq	%rax, %r10
	jne	.LBB0_2
# %bb.7:
	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
	.cfi_restore %rbx
	.cfi_restore %r12
	.cfi_restore %r13
	.cfi_restore %r14
	.cfi_restore %r15
.LBB0_8:                                # %._crit_edge32
	retq
.Lfunc_end0:
	.size	_Z10matmul_cpuPfPKfS1_i, .Lfunc_end0-_Z10matmul_cpuPfPKfS1_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z25__device_stub__matmul_gpuPfPKfS1_i # -- Begin function _Z25__device_stub__matmul_gpuPfPKfS1_i
	.type	_Z25__device_stub__matmul_gpuPfPKfS1_i,@function
_Z25__device_stub__matmul_gpuPfPKfS1_i: # @_Z25__device_stub__matmul_gpuPfPKfS1_i
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	4(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z10matmul_gpuPfPKfS1_i, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z25__device_stub__matmul_gpuPfPKfS1_i, .Lfunc_end1-_Z25__device_stub__matmul_gpuPfPKfS1_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z15call_matmul_gpuPfPKfS1_i    # -- Begin function _Z15call_matmul_gpuPfPKfS1_i
	.type	_Z15call_matmul_gpuPfPKfS1_i,@function
_Z15call_matmul_gpuPfPKfS1_i:           # @_Z15call_matmul_gpuPfPKfS1_i
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movl	%ecx, %ebx
	movq	%rdx, %r14
	movq	%rsi, %r15
	movq	%rdi, %r12
	leal	15(%rbx), %eax
	testl	%ecx, %ecx
	cmovnsl	%ecx, %eax
	sarl	$4, %eax
	movq	%rax, %rdi
	shlq	$32, %rdi
	orq	%rax, %rdi
	movabsq	$68719476752, %rdx              # imm = 0x1000000010
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	je	.LBB2_2
# %bb.1:
	addq	$8, %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
.LBB2_2:
	.cfi_def_cfa_offset 48
	movq	%r12, %rdi
	movq	%r15, %rsi
	movq	%r14, %rdx
	movl	%ebx, %ecx
	addq	$8, %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
	jmp	_Z25__device_stub__matmul_gpuPfPKfS1_i # TAILCALL
.Lfunc_end2:
	.size	_Z15call_matmul_gpuPfPKfS1_i, .Lfunc_end2-_Z15call_matmul_gpuPfPKfS1_i
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _Z20matmul_gpu_benchmarkPfPKfS1_i
.LCPI3_0:
	.quad	0x3f50624dd2f1a9fc              # double 0.001
	.text
	.globl	_Z20matmul_gpu_benchmarkPfPKfS1_i
	.type	_Z20matmul_gpu_benchmarkPfPKfS1_i,@function
_Z20matmul_gpu_benchmarkPfPKfS1_i:      # @_Z20matmul_gpu_benchmarkPfPKfS1_i
.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	$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	%ecx, %ebp
	movq	%rdx, (%rsp)                    # 8-byte Spill
	movq	%rsi, %r15
	movq	%rdi, 32(%rsp)                  # 8-byte Spill
	xorl	%r14d, %r14d
	leaq	8(%rsp), %rax
	movq	%r14, (%rax)
	leaq	24(%rsp), %rbx
	movq	%r14, (%rbx)
	leaq	16(%rsp), %r13
	movq	%r14, (%r13)
	movslq	%ecx, %r12
	imulq	%r12, %r12
	shlq	$2, %r12
	movq	%rbx, %rdi
	movq	%r12, %rsi
	callq	hipMalloc
	movq	%r13, %rdi
	movq	%r12, %rsi
	callq	hipMalloc
	leaq	8(%rsp), %rdi
	movq	%r12, %rsi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r13), %rdi
	movq	(%rsp), %rsi                    # 8-byte Reload
	movq	%r12, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$88, %edi
	callq	_Znwm
	movq	%rax, %rbx
	addq	$88, %rax
	movq	%rax, (%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	%r14, 80(%rbx)
	movabsq	$2361183241434822607, %r15      # imm = 0x20C49BA5E353F7CF
.LBB3_1:                                # =>This Inner Loop Header: Depth=1
	callq	_ZNSt6chrono3_V212steady_clock3nowEv
	movq	%rax, %r13
	movq	8(%rsp), %rdi
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdx
.Ltmp0:
	movl	%ebp, %ecx
	callq	_Z15call_matmul_gpuPfPKfS1_i
.Ltmp1:
# %bb.2:                                #   in Loop: Header=BB3_1 Depth=1
.Ltmp2:
	callq	hipDeviceSynchronize
.Ltmp3:
# %bb.3:                                #   in Loop: Header=BB3_1 Depth=1
	callq	_ZNSt6chrono3_V212steady_clock3nowEv
	subq	%r13, %rax
	imulq	%r15
	movq	%rdx, %rax
	shrq	$63, %rax
	sarq	$7, %rdx
	addq	%rax, %rdx
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rdx, %xmm0
	mulsd	.LCPI3_0(%rip), %xmm0
	movsd	%xmm0, (%rbx,%r14,8)
	incq	%r14
	cmpq	$11, %r14
	jne	.LBB3_1
# %bb.4:
	movq	8(%rsp), %rsi
.Ltmp5:
	movq	32(%rsp), %rdi                  # 8-byte Reload
	movq	%r12, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
.Ltmp6:
	movq	(%rsp), %r14                    # 8-byte Reload
# %bb.5:
	movq	24(%rsp), %rdi
.Ltmp7:
	callq	hipFree
.Ltmp8:
# %bb.6:
	movq	16(%rsp), %rdi
.Ltmp9:
	callq	hipFree
.Ltmp10:
# %bb.7:
	movq	8(%rsp), %rdi
.Ltmp11:
	callq	hipFree
.Ltmp12:
# %bb.8:
.Ltmp13:
	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_
.Ltmp14:
# %bb.9:                                # %.noexc
.Ltmp15:
	movq	%rbx, %rdi
	movq	%r14, %rsi
	callq	_ZSt22__final_insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_
.Ltmp16:
# %bb.10:                               # %_ZNSt6vectorIdSaIdEED2Ev.exit
	movsd	40(%rbx), %xmm0                 # xmm0 = mem[0],zero
	movsd	%xmm0, (%rsp)                   # 8-byte Spill
	movq	%rbx, %rdi
	callq	_ZdlPv
	movsd	(%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
.LBB3_11:
	.cfi_def_cfa_offset 96
.Ltmp17:
	jmp	.LBB3_12
.LBB3_13:
.Ltmp4:
.LBB3_12:                               # %_ZNSt6vectorIdSaIdEED2Ev.exit30
	movq	%rax, %r14
	movq	%rbx, %rdi
	callq	_ZdlPv
	movq	%r14, %rdi
	callq	_Unwind_Resume@PLT
.Lfunc_end3:
	.size	_Z20matmul_gpu_benchmarkPfPKfS1_i, .Lfunc_end3-_Z20matmul_gpu_benchmarkPfPKfS1_i
	.cfi_endproc
	.section	.gcc_except_table,"a",@progbits
	.p2align	2, 0x0
GCC_except_table3:
.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 .Ltmp5-.Lfunc_begin0           # >> Call Site 3 <<
	.uleb128 .Ltmp16-.Ltmp5                 #   Call between .Ltmp5 and .Ltmp16
	.uleb128 .Ltmp17-.Lfunc_begin0          #     jumps to .Ltmp17
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp16-.Lfunc_begin0          # >> Call Site 4 <<
	.uleb128 .Lfunc_end3-.Ltmp16            #   Call between .Ltmp16 and .Lfunc_end3
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
.Lcst_end0:
	.p2align	2, 0x0
                                        # -- End function
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function _Z8validatePKfS0_i
.LCPI4_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
.LCPI4_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
	je	.LBB4_6
# %bb.1:                                # %.lr.ph.preheader
	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
	pushq	%rax
	.cfi_def_cfa_offset 64
	.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, %r15
	movl	%edx, %r14d
	imull	%r14d, %r14d
	cmpl	$1, %r14d
	adcl	$0, %r14d
	movb	$1, %al
	xorl	%r12d, %r12d
	movaps	.LCPI4_0(%rip), %xmm1           # xmm1 = [NaN,NaN,NaN,NaN]
	movsd	.LCPI4_1(%rip), %xmm2           # xmm2 = mem[0],zero
	movq	%r14, (%rsp)                    # 8-byte Spill
.LBB4_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movss	(%r15,%r12,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
	subss	(%rbx,%r12,4), %xmm0
	andps	%xmm1, %xmm0
	cvtss2sd	%xmm0, %xmm0
	ucomisd	%xmm2, %xmm0
	jbe	.LBB4_4
# %bb.3:                                #   in Loop: Header=BB4_2 Depth=1
	movl	$_ZSt4cerr, %edi
	movl	$.L.str, %esi
	movl	$1, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	%r12d, %eax
	cltd
	idivl	%ebp
	movq	%r15, %r14
	movq	%rbx, %r15
	movl	%ebp, %ebx
	movl	%edx, %ebp
	movl	$_ZSt4cerr, %edi
	movl	%eax, %esi
	callq	_ZNSolsEi
	movq	%rax, %r13
	movl	$.L.str.1, %esi
	movl	$2, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%r13, %rdi
	movl	%ebp, %esi
	movl	%ebx, %ebp
	movq	%r15, %rbx
	movq	%r14, %r15
	callq	_ZNSolsEi
	movq	%rax, %r13
	movl	$.L.str.2, %esi
	movl	$3, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	xorps	%xmm0, %xmm0
	cvtss2sd	(%r14,%r12,4), %xmm0
	movq	%r13, %rdi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r13
	movl	$.L.str.3, %esi
	movl	$4, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	xorps	%xmm0, %xmm0
	cvtss2sd	(%rbx,%r12,4), %xmm0
	movq	%r13, %rdi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r13
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r13, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r13, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movsd	.LCPI4_1(%rip), %xmm2           # xmm2 = mem[0],zero
	movaps	.LCPI4_0(%rip), %xmm1           # xmm1 = [NaN,NaN,NaN,NaN]
	movq	(%rsp), %r14                    # 8-byte Reload
	xorl	%eax, %eax
.LBB4_4:                                #   in Loop: Header=BB4_2 Depth=1
	incq	%r12
	cmpq	%r12, %r14
	jne	.LBB4_2
# %bb.5:                                # %._crit_edge.loopexit
	andb	$1, %al
	addq	$8, %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
	.cfi_restore %rbx
	.cfi_restore %r12
	.cfi_restore %r13
	.cfi_restore %r14
	.cfi_restore %r15
	.cfi_restore %rbp
                                        # kill: def $al killed $al killed $eax
	retq
.LBB4_6:
	movb	$1, %al
                                        # kill: def $al killed $al killed $eax
	retq
.Lfunc_end4:
	.size	_Z8validatePKfS0_i, .Lfunc_end4-_Z8validatePKfS0_i
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function main
.LCPI5_0:
	.long	0xbf800000                      # float -1
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0
.LCPI5_1:
	.quad	0x3e112e0be826d695              # double 1.0000000000000001E-9
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
.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	$1176, %rsp                     # imm = 0x498
	.cfi_def_cfa_offset 1232
	.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, %rbx
	cmpl	$1, %edi
	jg	.LBB5_6
# %bb.1:
	movl	$_ZSt4cerr, %edi
	movl	$.L.str.4, %esi
	movl	$7, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	(%rbx), %rbx
	testq	%rbx, %rbx
	je	.LBB5_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	.LBB5_4
.LBB5_6:
	movq	8(%rbx), %rdi
	callq	atoi
	movl	%eax, %ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.6, %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, 112(%rsp)
	movl	%ebx, %r14d
	imull	%r14d, %r14d
	leaq	88(%rsp), %rdi
	leaq	144(%rsp), %rdx
	movq	%r14, %rsi
	callq	_ZNSt6vectorIfSaIfEEC2EmRKS0_
.Ltmp18:
	leaq	64(%rsp), %rdi
	leaq	144(%rsp), %rdx
	movq	%r14, %rsi
	callq	_ZNSt6vectorIfSaIfEEC2EmRKS0_
.Ltmp19:
# %bb.7:
.Ltmp21:
	leaq	8(%rsp), %rdi
	leaq	144(%rsp), %rdx
	movq	%r14, %rsi
	callq	_ZNSt6vectorIfSaIfEEC2EmRKS0_
.Ltmp22:
# %bb.8:
.Ltmp24:
	leaq	120(%rsp), %rdi
	leaq	144(%rsp), %rdx
	movq	%r14, %rsi
	callq	_ZNSt6vectorIfSaIfEEC2EmRKS0_
.Ltmp25:
# %bb.9:
	testl	%ebx, %ebx
	je	.LBB5_14
# %bb.10:                               # %.lr.ph.preheader
	cmpl	$2, %r14d
	movl	$1, %r15d
	cmovael	%r14d, %r15d
	xorl	%r12d, %r12d
	leaq	112(%rsp), %r14
.LBB5_11:                               # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
.Ltmp27:
	movq	%r14, %rdi
	callq	_ZSt18generate_canonicalIfLm24ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_
.Ltmp28:
	movss	.LCPI5_0(%rip), %xmm1           # xmm1 = mem[0],zero,zero,zero
# %bb.12:                               #   in Loop: Header=BB5_11 Depth=1
	addss	%xmm0, %xmm0
	addss	%xmm1, %xmm0
	movq	88(%rsp), %rax
	movss	%xmm0, (%rax,%r12,4)
.Ltmp29:
	movq	%r14, %rdi
	callq	_ZSt18generate_canonicalIfLm24ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_
	movss	.LCPI5_0(%rip), %xmm1           # xmm1 = mem[0],zero,zero,zero
.Ltmp30:
# %bb.13:                               #   in Loop: Header=BB5_11 Depth=1
	addss	%xmm0, %xmm0
	addss	%xmm1, %xmm0
	movq	64(%rsp), %rax
	movss	%xmm0, (%rax,%r12,4)
	incq	%r12
	cmpq	%r12, %r15
	jne	.LBB5_11
.LBB5_14:                               # %._crit_edge
	movl	%ebx, %ebp
	negl	%ebp
	cmovsl	%ebx, %ebp
	movl	$1, %r14d
	cmpl	$10, %ebp
	jb	.LBB5_23
# %bb.15:                               # %.lr.ph.i.i.preheader
	movl	$4, %r14d
	movl	$3518437209, %eax               # imm = 0xD1B71759
	movl	%ebp, %ecx
.LBB5_16:                               # %.lr.ph.i.i
                                        # =>This Inner Loop Header: Depth=1
	cmpl	$99, %ecx
	jbe	.LBB5_17
# %bb.18:                               #   in Loop: Header=BB5_16 Depth=1
	cmpl	$999, %ecx                      # imm = 0x3E7
	jbe	.LBB5_19
# %bb.20:                               #   in Loop: Header=BB5_16 Depth=1
	cmpl	$10000, %ecx                    # imm = 0x2710
	jb	.LBB5_23
# %bb.21:                               #   in Loop: Header=BB5_16 Depth=1
	movl	%ecx, %edx
	imulq	%rax, %rdx
	shrq	$45, %rdx
	addl	$4, %r14d
	cmpl	$99999, %ecx                    # imm = 0x1869F
	movl	%edx, %ecx
	ja	.LBB5_16
# %bb.22:                               # %_ZNSt8__detail14__to_chars_lenIjEEjT_i.exit.i.loopexit
	addl	$-3, %r14d
	jmp	.LBB5_23
.LBB5_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
.LBB5_4:                                # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit
	movl	$_ZSt4cerr, %ebx
	movl	$_ZSt4cerr, %edi
	movl	$.L.str.5, %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	.LBB5_5
.LBB5_17:
	addl	$-2, %r14d
	jmp	.LBB5_23
.LBB5_19:
	decl	%r14d
.LBB5_23:                               # %_ZNSt8__detail14__to_chars_lenIjEEjT_i.exit.i
	movl	%ebx, %r15d
	shrl	$31, %r15d
	leal	(%r14,%r15), %esi
	leaq	160(%rsp), %r12
	movq	%r12, -16(%r12)
.Ltmp32:
	leaq	144(%rsp), %rdi
	movl	$45, %edx
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructEmc
.Ltmp33:
# %bb.24:                               # %_ZNSt7__cxx119to_stringEi.exit
	movl	%r15d, %edi
	leaq	144(%rsp), %r15
	addq	(%r15), %rdi
	movl	%r14d, %esi
	movl	%ebp, %edx
	callq	_ZNSt8__detail18__to_chars_10_implIjEEvPcjT_
.Ltmp35:
	movl	$.L.str.7, %ecx
	movl	$6, %r8d
	movq	%r15, %rdi
	xorl	%esi, %esi
	xorl	%edx, %edx
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE7replaceEmmPKcm
.Ltmp36:
# %bb.25:                               # %.noexc
	movq	%rax, %r15
	leaq	48(%rsp), %r14
	movq	%r14, -16(%r14)
	movq	(%rax), %rsi
	movq	%rax, %r13
	addq	$16, %r13
	cmpq	%r13, %rsi
	je	.LBB5_26
# %bb.31:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i
	movq	%rsi, 32(%rsp)
	movq	16(%r15), %rax
	movq	%rax, 48(%rsp)
	movq	8(%r15), %rbp
	jmp	.LBB5_32
.LBB5_26:
	movq	8(%r15), %rbp
	leaq	1(%rbp), %rdx
	movq	%r14, %rdi
	callq	memcpy@PLT
.LBB5_32:
	movq	%rbp, 40(%rsp)
	movq	%r13, (%r15)
	movq	$0, 8(%r15)
	movb	$0, 16(%r15)
	movq	144(%rsp), %rdi
	cmpq	%r12, %rdi
	je	.LBB5_34
# %bb.33:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i61
	callq	_ZdlPv
.LBB5_34:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit
.Ltmp38:
	leaq	144(%rsp), %rdi
	leaq	32(%rsp), %rsi
	movl	$4, %edx
	callq	_ZNSt14basic_ifstreamIcSt11char_traitsIcEEC1ERKNSt7__cxx1112basic_stringIcS1_SaIcEEESt13_Ios_Openmode
.Ltmp39:
# %bb.35:
	movq	144(%rsp), %rax
	movq	-24(%rax), %rax
	testb	$5, 176(%rsp,%rax)
	je	.LBB5_36
# %bb.40:
.Ltmp41:
	leaq	664(%rsp), %rdi
	leaq	32(%rsp), %rsi
	movl	$16, %edx
	callq	_ZNSt14basic_ofstreamIcSt11char_traitsIcEEC1ERKNSt7__cxx1112basic_stringIcS1_SaIcEEESt13_Ios_Openmode
.Ltmp42:
# %bb.41:
	movq	8(%rsp), %rdi
	movq	88(%rsp), %rsi
	movq	64(%rsp), %rdx
	movl	%ebx, %ecx
	callq	_Z10matmul_cpuPfPKfS1_i
	movq	8(%rsp), %rsi
	movslq	%ebx, %rdx
	imulq	%rdx, %rdx
	shlq	$2, %rdx
.Ltmp44:
	leaq	664(%rsp), %rdi
	callq	_ZNSo5writeEPKcl
.Ltmp45:
# %bb.42:
	movq	_ZTTSt14basic_ofstreamIcSt11char_traitsIcEE(%rip), %rax
	leaq	672(%rsp), %rdi
	movq	%rax, -8(%rdi)
	movq	_ZTTSt14basic_ofstreamIcSt11char_traitsIcEE+24(%rip), %rcx
	movq	-24(%rax), %rax
	movq	%rcx, 664(%rsp,%rax)
	callq	_ZNSt13basic_filebufIcSt11char_traitsIcEED2Ev
	leaq	912(%rsp), %rdi
	callq	_ZNSt8ios_baseD2Ev
	jmp	.LBB5_43
.LBB5_36:
	movq	8(%rsp), %rsi
	movslq	%ebx, %rdx
	imulq	%rdx, %rdx
	shlq	$2, %rdx
.Ltmp47:
	leaq	144(%rsp), %rdi
	callq	_ZNSi4readEPcl
.Ltmp48:
.LBB5_43:
	movq	120(%rsp), %rdi
	movq	88(%rsp), %rsi
	movq	64(%rsp), %rdx
.Ltmp50:
	movl	%ebx, %ecx
	callq	_Z20matmul_gpu_benchmarkPfPKfS1_i
	movsd	%xmm0, (%rsp)                   # 8-byte Spill
.Ltmp51:
# %bb.44:
.Ltmp52:
	movl	$_ZSt4cout, %edi
	movl	$.L.str.8, %esi
	movl	$5, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp53:
# %bb.45:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit67
.Ltmp54:
	movl	$_ZSt4cout, %edi
	movsd	(%rsp), %xmm0                   # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
.Ltmp55:
# %bb.46:                               # %_ZNSolsEd.exit
.Ltmp56:
	movq	%rax, %r15
	movl	$.L.str.9, %esi
	movl	$5, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp57:
# %bb.47:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit70
	movq	(%r15), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
.Ltmp58:
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
.Ltmp59:
# %bb.48:                               # %.noexc107
.Ltmp60:
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
.Ltmp61:
# %bb.49:                               # %.noexc108
.Ltmp62:
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.Ltmp63:
# %bb.50:                               # %_ZNSolsEPFRSoS_E.exit
.Ltmp64:
	movl	$_ZSt4cout, %edi
	movl	$.L.str.10, %esi
	movl	$5, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp65:
# %bb.51:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit73
	cvtsi2sd	%ebx, %xmm1
	movapd	%xmm1, %xmm0
	addsd	%xmm1, %xmm0
	mulsd	%xmm1, %xmm0
	mulsd	%xmm1, %xmm0
	divsd	(%rsp), %xmm0                   # 8-byte Folded Reload
	mulsd	.LCPI5_1(%rip), %xmm0
.Ltmp66:
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
.Ltmp67:
# %bb.52:                               # %_ZNSolsEd.exit75
.Ltmp68:
	movq	%rax, %r15
	movl	$.L.str.11, %esi
	movl	$9, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp69:
# %bb.53:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit77
	movq	(%r15), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
.Ltmp70:
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
.Ltmp71:
# %bb.54:                               # %.noexc110
.Ltmp72:
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
.Ltmp73:
# %bb.55:                               # %.noexc111
.Ltmp74:
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.Ltmp75:
# %bb.56:                               # %_ZNSolsEPFRSoS_E.exit79
	movq	8(%rsp), %rdi
	movq	120(%rsp), %rsi
.Ltmp77:
	movl	%ebx, %edx
	callq	_Z8validatePKfS0_i
.Ltmp78:
# %bb.57:
.Ltmp79:
	movl	%eax, %ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.12, %esi
	movl	$12, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp80:
# %bb.58:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit81
	movl	$.L.str.13, %eax
	movl	$.L.str.14, %esi
	testb	%bl, %bl
	cmovneq	%rax, %rsi
	movzbl	%bl, %edx
	orq	$6, %rdx
.Ltmp81:
	movl	$_ZSt4cout, %ebx
	movl	$_ZSt4cout, %edi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp82:
# %bb.59:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit84
	movq	_ZSt4cout(%rip), %rax
	addq	-24(%rax), %rbx
.Ltmp83:
	movq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
.Ltmp84:
# %bb.60:                               # %.noexc114
.Ltmp85:
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
.Ltmp86:
# %bb.61:                               # %.noexc115
.Ltmp87:
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.Ltmp88:
# %bb.62:                               # %_ZNSolsEPFRSoS_E.exit86
	leaq	144(%rsp), %rdi
	callq	_ZNSt14basic_ifstreamIcSt11char_traitsIcEED1Ev
	movq	32(%rsp), %rdi
	cmpq	%r14, %rdi
	je	.LBB5_64
# %bb.63:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i87
	callq	_ZdlPv
.LBB5_64:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit89
	movq	120(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB5_66
# %bb.65:
	callq	_ZdlPv
.LBB5_66:                               # %_ZNSt6vectorIfSaIfEED2Ev.exit
	movq	8(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB5_68
# %bb.67:
	callq	_ZdlPv
.LBB5_68:                               # %_ZNSt6vectorIfSaIfEED2Ev.exit91
	movq	64(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB5_70
# %bb.69:
	callq	_ZdlPv
.LBB5_70:                               # %_ZNSt6vectorIfSaIfEED2Ev.exit93
	movq	88(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB5_5
# %bb.71:
	callq	_ZdlPv
.LBB5_5:
	xorl	%eax, %eax
	addq	$1176, %rsp                     # imm = 0x498
	.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
.LBB5_39:
	.cfi_def_cfa_offset 1232
.Ltmp49:
	jmp	.LBB5_76
.LBB5_73:
.Ltmp46:
	movq	%rax, %rbx
	leaq	664(%rsp), %rdi
	callq	_ZNSt14basic_ofstreamIcSt11char_traitsIcEED1Ev
	jmp	.LBB5_77
.LBB5_72:
.Ltmp43:
	jmp	.LBB5_76
.LBB5_38:
.Ltmp40:
	movq	%rax, %rbx
	jmp	.LBB5_78
.LBB5_37:
.Ltmp37:
	movq	%rax, %rbx
	movq	144(%rsp), %rdi
	cmpq	%r12, %rdi
	jne	.LBB5_79
	jmp	.LBB5_80
.LBB5_89:
.Ltmp34:
	movq	%rax, %rdi
	callq	__clang_call_terminate
.LBB5_29:
.Ltmp26:
	movq	%rax, %rbx
	jmp	.LBB5_82
.LBB5_28:
.Ltmp23:
	movq	%rax, %rbx
	jmp	.LBB5_84
.LBB5_27:
.Ltmp20:
	movq	%rax, %rbx
	jmp	.LBB5_86
.LBB5_75:
.Ltmp89:
	jmp	.LBB5_76
.LBB5_74:
.Ltmp76:
.LBB5_76:
	movq	%rax, %rbx
.LBB5_77:
	leaq	144(%rsp), %rdi
	callq	_ZNSt14basic_ifstreamIcSt11char_traitsIcEED1Ev
.LBB5_78:
	movq	32(%rsp), %rdi
	cmpq	%r14, %rdi
	je	.LBB5_80
.LBB5_79:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i96
	callq	_ZdlPv
	jmp	.LBB5_80
.LBB5_30:
.Ltmp31:
	movq	%rax, %rbx
.LBB5_80:
	movq	120(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB5_82
# %bb.81:
	callq	_ZdlPv
.LBB5_82:                               # %_ZNSt6vectorIfSaIfEED2Ev.exit100
	movq	8(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB5_84
# %bb.83:
	callq	_ZdlPv
.LBB5_84:                               # %_ZNSt6vectorIfSaIfEED2Ev.exit102
	movq	64(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB5_86
# %bb.85:
	callq	_ZdlPv
.LBB5_86:                               # %_ZNSt6vectorIfSaIfEED2Ev.exit104
	movq	88(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB5_88
# %bb.87:
	callq	_ZdlPv
.LBB5_88:                               # %_ZNSt6vectorIfSaIfEED2Ev.exit106
	movq	%rbx, %rdi
	callq	_Unwind_Resume@PLT
.Lfunc_end5:
	.size	main, .Lfunc_end5-main
	.cfi_endproc
	.section	.gcc_except_table,"a",@progbits
	.p2align	2, 0x0
GCC_except_table5:
.Lexception1:
	.byte	255                             # @LPStart Encoding = omit
	.byte	3                               # @TType Encoding = udata4
	.uleb128 .Lttbase0-.Lttbaseref0
.Lttbaseref0:
	.byte	1                               # Call site Encoding = uleb128
	.uleb128 .Lcst_end1-.Lcst_begin1
.Lcst_begin1:
	.uleb128 .Lfunc_begin1-.Lfunc_begin1    # >> Call Site 1 <<
	.uleb128 .Ltmp18-.Lfunc_begin1          #   Call between .Lfunc_begin1 and .Ltmp18
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp18-.Lfunc_begin1          # >> Call Site 2 <<
	.uleb128 .Ltmp19-.Ltmp18                #   Call between .Ltmp18 and .Ltmp19
	.uleb128 .Ltmp20-.Lfunc_begin1          #     jumps to .Ltmp20
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp21-.Lfunc_begin1          # >> Call Site 3 <<
	.uleb128 .Ltmp22-.Ltmp21                #   Call between .Ltmp21 and .Ltmp22
	.uleb128 .Ltmp23-.Lfunc_begin1          #     jumps to .Ltmp23
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp24-.Lfunc_begin1          # >> Call Site 4 <<
	.uleb128 .Ltmp25-.Ltmp24                #   Call between .Ltmp24 and .Ltmp25
	.uleb128 .Ltmp26-.Lfunc_begin1          #     jumps to .Ltmp26
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp27-.Lfunc_begin1          # >> Call Site 5 <<
	.uleb128 .Ltmp30-.Ltmp27                #   Call between .Ltmp27 and .Ltmp30
	.uleb128 .Ltmp31-.Lfunc_begin1          #     jumps to .Ltmp31
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp30-.Lfunc_begin1          # >> Call Site 6 <<
	.uleb128 .Ltmp32-.Ltmp30                #   Call between .Ltmp30 and .Ltmp32
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp32-.Lfunc_begin1          # >> Call Site 7 <<
	.uleb128 .Ltmp33-.Ltmp32                #   Call between .Ltmp32 and .Ltmp33
	.uleb128 .Ltmp34-.Lfunc_begin1          #     jumps to .Ltmp34
	.byte	1                               #   On action: 1
	.uleb128 .Ltmp35-.Lfunc_begin1          # >> Call Site 8 <<
	.uleb128 .Ltmp36-.Ltmp35                #   Call between .Ltmp35 and .Ltmp36
	.uleb128 .Ltmp37-.Lfunc_begin1          #     jumps to .Ltmp37
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp36-.Lfunc_begin1          # >> Call Site 9 <<
	.uleb128 .Ltmp38-.Ltmp36                #   Call between .Ltmp36 and .Ltmp38
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp38-.Lfunc_begin1          # >> Call Site 10 <<
	.uleb128 .Ltmp39-.Ltmp38                #   Call between .Ltmp38 and .Ltmp39
	.uleb128 .Ltmp40-.Lfunc_begin1          #     jumps to .Ltmp40
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp41-.Lfunc_begin1          # >> Call Site 11 <<
	.uleb128 .Ltmp42-.Ltmp41                #   Call between .Ltmp41 and .Ltmp42
	.uleb128 .Ltmp43-.Lfunc_begin1          #     jumps to .Ltmp43
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp44-.Lfunc_begin1          # >> Call Site 12 <<
	.uleb128 .Ltmp45-.Ltmp44                #   Call between .Ltmp44 and .Ltmp45
	.uleb128 .Ltmp46-.Lfunc_begin1          #     jumps to .Ltmp46
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp47-.Lfunc_begin1          # >> Call Site 13 <<
	.uleb128 .Ltmp48-.Ltmp47                #   Call between .Ltmp47 and .Ltmp48
	.uleb128 .Ltmp49-.Lfunc_begin1          #     jumps to .Ltmp49
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp50-.Lfunc_begin1          # >> Call Site 14 <<
	.uleb128 .Ltmp75-.Ltmp50                #   Call between .Ltmp50 and .Ltmp75
	.uleb128 .Ltmp76-.Lfunc_begin1          #     jumps to .Ltmp76
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp77-.Lfunc_begin1          # >> Call Site 15 <<
	.uleb128 .Ltmp88-.Ltmp77                #   Call between .Ltmp77 and .Ltmp88
	.uleb128 .Ltmp89-.Lfunc_begin1          #     jumps to .Ltmp89
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp88-.Lfunc_begin1          # >> Call Site 16 <<
	.uleb128 .Lfunc_end5-.Ltmp88            #   Call between .Ltmp88 and .Lfunc_end5
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
.Lcst_end1:
	.byte	1                               # >> Action Record 1 <<
                                        #   Catch TypeInfo 1
	.byte	0                               #   No further actions
	.p2align	2, 0x0
                                        # >> Catch TypeInfos <<
	.long	0                               # TypeInfo 1
.Lttbase0:
	.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	.LBB6_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	.LBB6_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	.LBB6_5
# %bb.4:
	movq	%rax, %r15
	jmp	.LBB6_6
.LBB6_2:                                # %_ZNSt12_Vector_baseIfSaIfEEC2EmRKS0_.exit.thread
	movups	%xmm0, (%rbx)
	movq	$0, 16(%rbx)
	xorl	%r15d, %r15d
	jmp	.LBB6_6
.LBB6_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
.LBB6_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
.LBB6_7:
	.cfi_def_cfa_offset 32
	movl	$.L.str.17, %edi
	callq	_ZSt20__throw_length_errorPKc
.Lfunc_end6:
	.size	_ZNSt6vectorIfSaIfEEC2EmRKS0_, .Lfunc_end6-_ZNSt6vectorIfSaIfEEC2EmRKS0_
	.cfi_endproc
                                        # -- End function
	.section	.text.__clang_call_terminate,"axG",@progbits,__clang_call_terminate,comdat
	.hidden	__clang_call_terminate          # -- Begin function __clang_call_terminate
	.weak	__clang_call_terminate
	.type	__clang_call_terminate,@function
__clang_call_terminate:                 # @__clang_call_terminate
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__cxa_begin_catch
	callq	_ZSt9terminatev
.Lfunc_end7:
	.size	__clang_call_terminate, .Lfunc_end7-__clang_call_terminate
	.cfi_endproc
                                        # -- End function
	.section	.text._ZNSt8__detail18__to_chars_10_implIjEEvPcjT_,"axG",@progbits,_ZNSt8__detail18__to_chars_10_implIjEEvPcjT_,comdat
	.weak	_ZNSt8__detail18__to_chars_10_implIjEEvPcjT_ # -- Begin function _ZNSt8__detail18__to_chars_10_implIjEEvPcjT_
	.type	_ZNSt8__detail18__to_chars_10_implIjEEvPcjT_,@function
_ZNSt8__detail18__to_chars_10_implIjEEvPcjT_: # @_ZNSt8__detail18__to_chars_10_implIjEEvPcjT_
	.cfi_startproc
# %bb.0:
                                        # kill: def $esi killed $esi def $rsi
	cmpl	$100, %edx
	jb	.LBB8_1
# %bb.2:                                # %.lr.ph.preheader
	decl	%esi
.LBB8_3:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	%edx, %eax
	imulq	$1374389535, %rax, %rax         # imm = 0x51EB851F
	shrq	$37, %rax
	imull	$100, %eax, %ecx
	movl	%edx, %r8d
	subl	%ecx, %r8d
	movb	.L__const._ZNSt8__detail18__to_chars_10_implIjEEvPcjT_.__digits+1(%r8,%r8), %cl
	movl	%esi, %r9d
	movb	%cl, (%rdi,%r9)
	movb	.L__const._ZNSt8__detail18__to_chars_10_implIjEEvPcjT_.__digits(%r8,%r8), %cl
	leal	-1(%rsi), %r8d
	movb	%cl, (%rdi,%r8)
	addl	$-2, %esi
	cmpl	$9999, %edx                     # imm = 0x270F
	movl	%eax, %edx
	ja	.LBB8_3
	jmp	.LBB8_4
.LBB8_1:
	movl	%edx, %eax
.LBB8_4:                                # %._crit_edge
	cmpl	$10, %eax
	jb	.LBB8_6
# %bb.5:
	movl	%eax, %eax
	leaq	(%rax,%rax), %rcx
	movl	%ecx, %ecx
	movb	.L__const._ZNSt8__detail18__to_chars_10_implIjEEvPcjT_.__digits+1(%rcx), %cl
	movb	%cl, 1(%rdi)
	movb	.L__const._ZNSt8__detail18__to_chars_10_implIjEEvPcjT_.__digits(%rax,%rax), %al
	jmp	.LBB8_7
.LBB8_6:
	orb	$48, %al
.LBB8_7:
	movb	%al, (%rdi)
	retq
.Lfunc_end8:
	.size	_ZNSt8__detail18__to_chars_10_implIjEEvPcjT_, .Lfunc_end8-_ZNSt8__detail18__to_chars_10_implIjEEvPcjT_
	.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	.LBB9_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
.LBB9_2:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB9_5 Depth 2
                                        #       Child Loop BB9_6 Depth 3
                                        #       Child Loop BB9_8 Depth 3
	subq	$1, %r15
	jb	.LBB9_3
# %bb.4:                                #   in Loop: Header=BB9_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
.LBB9_5:                                #   Parent Loop BB9_2 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB9_6 Depth 3
                                        #       Child Loop BB9_8 Depth 3
	movsd	(%rbx), %xmm0                   # xmm0 = mem[0],zero
	leaq	(%rax,%rbp), %r14
.LBB9_6:                                #   Parent Loop BB9_2 Depth=1
                                        #     Parent Loop BB9_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	.LBB9_6
# %bb.7:                                # %.preheader.i.i.preheader
                                        #   in Loop: Header=BB9_5 Depth=2
	leaq	-8(%rax), %r13
.LBB9_8:                                # %.preheader.i.i
                                        #   Parent Loop BB9_2 Depth=1
                                        #     Parent Loop BB9_5 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movsd	-8(%rcx), %xmm2                 # xmm2 = mem[0],zero
	addq	$-8, %rcx
	ucomisd	%xmm0, %xmm2
	ja	.LBB9_8
# %bb.9:                                #   in Loop: Header=BB9_5 Depth=2
	cmpq	%rcx, %r13
	jae	.LBB9_11
# %bb.10:                               #   in Loop: Header=BB9_5 Depth=2
	movsd	%xmm2, (%r13)
	movsd	%xmm1, (%rcx)
	jmp	.LBB9_5
.LBB9_11:                               # %_ZSt27__unguarded_partition_pivotIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEET_S9_S9_T0_.exit
                                        #   in Loop: Header=BB9_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	.LBB9_2
	jmp	.LBB9_12
.LBB9_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_
.LBB9_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_end9:
	.size	_ZSt16__introsort_loopIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEElNS0_5__ops15_Iter_less_iterEEvT_S9_T0_T1_, .Lfunc_end9-_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	.LBB10_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_
.LBB10_2:                               # =>This Loop Header: Depth=1
                                        #     Child Loop BB10_5 Depth 2
	cmpq	%rbx, %r14
	je	.LBB10_7
# %bb.3:                                # %.lr.ph.i
                                        #   in Loop: Header=BB10_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	.LBB10_6
# %bb.4:                                # %.lr.ph.i.i.preheader
                                        #   in Loop: Header=BB10_2 Depth=1
	movq	%r14, %rax
.LBB10_5:                               # %.lr.ph.i.i
                                        #   Parent Loop BB10_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	.LBB10_5
.LBB10_6:                               # %_ZSt25__unguarded_linear_insertIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops14_Val_less_iterEEvT_T0_.exit.i
                                        #   in Loop: Header=BB10_2 Depth=1
	movsd	%xmm0, (%rax)
	addq	$8, %r14
	jmp	.LBB10_2
.LBB10_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
.LBB10_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_end10:
	.size	_ZSt22__final_insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_, .Lfunc_end10-_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	.LBB11_3
# %bb.1:                                # %.lr.ph.preheader
	movq	%rsi, %rbx
	movq	%rdi, %r14
	addq	$-8, %rbx
.LBB11_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	.LBB11_2
.LBB11_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_end11:
	.size	_ZSt11__sort_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_RT0_, .Lfunc_end11-_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	.LBB12_3
# %bb.1:
	movq	%rdi, %r14
	leaq	-2(%rbx), %r15
	shrq	%r15
.LBB12_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	.LBB12_2
.LBB12_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_end12:
	.size	_ZSt11__make_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_RT0_, .Lfunc_end12-_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	.LBB13_6
# %bb.1:                                # %.lr.ph.preheader
	movq	%rsi, %r8
.LBB13_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	.LBB13_3
# %bb.4:                                # %.lr.ph
                                        #   in Loop: Header=BB13_2 Depth=1
	leaq	2(,%r8,2), %rax
	jmp	.LBB13_5
.LBB13_3:                               #   in Loop: Header=BB13_2 Depth=1
	leaq	1(,%r8,2), %rax
.LBB13_5:                               # %.lr.ph
                                        #   in Loop: Header=BB13_2 Depth=1
	movsd	(%rdi,%rax,8), %xmm1            # xmm1 = mem[0],zero
	movsd	%xmm1, (%rdi,%r8,8)
	movq	%rax, %r8
	cmpq	%rcx, %rax
	jl	.LBB13_2
.LBB13_6:                               # %._crit_edge
	testb	$1, %dl
	jne	.LBB13_9
# %bb.7:
	addq	$-2, %rdx
	sarq	%rdx
	cmpq	%rdx, %rax
	jne	.LBB13_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
.LBB13_9:
	cmpq	%rsi, %rax
	jle	.LBB13_10
.LBB13_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	.LBB13_10
# %bb.12:                               #   in Loop: Header=BB13_11 Depth=1
	movsd	%xmm1, (%rdi,%rax,8)
	movq	%rcx, %rax
	cmpq	%rsi, %rcx
	jg	.LBB13_11
	jmp	.LBB13_13
.LBB13_10:
	movq	%rax, %rcx
.LBB13_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_end13:
	.size	_ZSt13__adjust_heapIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEEldNS0_5__ops15_Iter_less_iterEEvT_T0_SA_T1_T2_, .Lfunc_end13-_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	.LBB14_6
# %bb.1:
	ucomisd	%xmm2, %xmm0
	jbe	.LBB14_3
# %bb.2:
	movsd	(%rdi), %xmm0                   # xmm0 = mem[0],zero
	movsd	%xmm2, (%rdi)
	movsd	%xmm0, (%rdx)
	retq
.LBB14_6:
	ucomisd	%xmm1, %xmm0
	jbe	.LBB14_8
# %bb.7:
	movsd	(%rdi), %xmm0                   # xmm0 = mem[0],zero
	movsd	%xmm1, (%rdi)
	movsd	%xmm0, (%rsi)
	retq
.LBB14_3:
	ucomisd	%xmm1, %xmm0
	movsd	(%rdi), %xmm2                   # xmm2 = mem[0],zero
	jbe	.LBB14_5
# %bb.4:
	movsd	%xmm0, (%rdi)
	movsd	%xmm2, (%rcx)
	retq
.LBB14_8:
	ucomisd	%xmm2, %xmm0
	movsd	(%rdi), %xmm1                   # xmm1 = mem[0],zero
	jbe	.LBB14_10
# %bb.9:
	movsd	%xmm0, (%rdi)
	movsd	%xmm1, (%rcx)
	retq
.LBB14_5:
	movsd	%xmm1, (%rdi)
	movsd	%xmm2, (%rsi)
	retq
.LBB14_10:
	movsd	%xmm2, (%rdi)
	movsd	%xmm1, (%rdx)
	retq
.Lfunc_end14:
	.size	_ZSt22__move_median_to_firstIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_S9_S9_T0_, .Lfunc_end14-_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	.LBB15_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	.LBB15_13
# %bb.2:                                # %.lr.ph
	movq	%r14, %r15
.LBB15_3:                               # =>This Loop Header: Depth=1
                                        #     Child Loop BB15_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	.LBB15_6
# %bb.4:                                #   in Loop: Header=BB15_3 Depth=1
	movq	%r15, %rdx
	subq	%r14, %rdx
	movq	%rdx, %rax
	sarq	$3, %rax
	cmpq	$2, %rax
	jl	.LBB15_10
# %bb.5:                                #   in Loop: Header=BB15_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
.LBB15_12:                              # %_ZSt13move_backwardIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEES6_ET0_T_S8_S7_.exit
                                        #   in Loop: Header=BB15_3 Depth=1
	movq	%r14, %rax
	jmp	.LBB15_9
.LBB15_6:                               #   in Loop: Header=BB15_3 Depth=1
	movsd	(%rdi), %xmm0                   # xmm0 = mem[0],zero
	ucomisd	%xmm1, %xmm0
	movq	%r15, %rax
	jbe	.LBB15_9
# %bb.7:                                # %.lr.ph.i.preheader
                                        #   in Loop: Header=BB15_3 Depth=1
	movq	%r15, %rax
.LBB15_8:                               # %.lr.ph.i
                                        #   Parent Loop BB15_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	.LBB15_8
.LBB15_9:                               # %_ZSt13move_backwardIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEES6_ET0_T_S8_S7_.exit
                                        #   in Loop: Header=BB15_3 Depth=1
	movsd	%xmm1, (%rax)
	leaq	8(%r15), %rax
	cmpq	%rbx, %rax
	jne	.LBB15_3
	jmp	.LBB15_13
.LBB15_10:                              #   in Loop: Header=BB15_3 Depth=1
	movq	%r14, %rax
	cmpq	$8, %rdx
	jne	.LBB15_9
# %bb.11:                               #   in Loop: Header=BB15_3 Depth=1
	movsd	%xmm0, 8(%rdi)
	jmp	.LBB15_12
.LBB15_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
.LBB15_14:                              # %.loopexit
	retq
.Lfunc_end15:
	.size	_ZSt16__insertion_sortIN9__gnu_cxx17__normal_iteratorIPdSt6vectorIdSaIdEEEENS0_5__ops15_Iter_less_iterEEvT_S9_T0_, .Lfunc_end15-_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	.LBB16_1
# %bb.3:
	shlq	$2, %rsi
	movq	%rsi, %rdi
	popq	%rax
	.cfi_def_cfa_offset 8
	jmp	_Znwm                           # TAILCALL
.LBB16_1:
	.cfi_def_cfa_offset 16
	shrq	$62, %rsi
	je	.LBB16_2
# %bb.4:
	callq	_ZSt28__throw_bad_array_new_lengthv
.LBB16_2:
	callq	_ZSt17__throw_bad_allocv
.Lfunc_end16:
	.size	_ZNSt15__new_allocatorIfE8allocateEmPKv, .Lfunc_end16-_ZNSt15__new_allocatorIfE8allocateEmPKv
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _ZSt18generate_canonicalIfLm24ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_
.LCPI17_0:
	.quad	0x41dfffffff800000              # double 2147483646
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0
.LCPI17_1:
	.long	0x40000000                      # float 2
.LCPI17_2:
	.long	0x5f000000                      # float 9.22337203E+18
.LCPI17_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	.LCPI17_0(%rip)
	fstpt	(%rsp)
	callq	logl
	fstpt	52(%rsp)                        # 10-byte Folded Spill
	flds	.LCPI17_1(%rip)
	fstpt	(%rsp)
	callq	logl
	fldt	52(%rsp)                        # 10-byte Folded Reload
	fdivp	%st, %st(1)
	flds	.LCPI17_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	.LCPI17_3(%rip), %xmm1          # xmm1 = mem[0],zero,zero,zero
	movabsq	$8589934597, %rdi               # imm = 0x200000005
.LBB17_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	.LBB17_2
# %bb.3:                                # %select.unfold
                                        #   in Loop: Header=BB17_1 Depth=1
	xorps	%xmm2, %xmm2
	cvtsi2ss	%rdx, %xmm2
	jmp	.LBB17_4
.LBB17_2:                               #   in Loop: Header=BB17_1 Depth=1
	movq	%rdx, %r8
	shrq	%r8
	andl	$1, %edx
	orq	%r8, %rdx
	xorps	%xmm2, %xmm2
	cvtsi2ss	%rdx, %xmm2
	addss	%xmm2, %xmm2
.LBB17_4:                               # %select.unfold
                                        #   in Loop: Header=BB17_1 Depth=1
	addq	%rax, %rsi
	mulss	%xmm1, %xmm2
	movss	%xmm1, 36(%rsp)
	flds	36(%rsp)
	fmull	.LCPI17_0(%rip)
	addss	%xmm2, %xmm0
	fstps	32(%rsp)
	movss	32(%rsp), %xmm1                 # xmm1 = mem[0],zero,zero,zero
	decq	%rcx
	jne	.LBB17_1
# %bb.5:
	movq	%rsi, (%rbx)
	divss	%xmm1, %xmm0
	ucomiss	.LCPI17_3(%rip), %xmm0
	jae	.LBB17_7
# %bb.6:
	addq	$64, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.LBB17_7:
	.cfi_def_cfa_offset 80
	movss	.LCPI17_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_end17:
	.size	_ZSt18generate_canonicalIfLm24ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_, .Lfunc_end17-_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	.LBB18_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB18_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z10matmul_gpuPfPKfS1_i, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end18:
	.size	__hip_module_ctor, .Lfunc_end18-__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	.LBB19_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
.LBB19_2:
	retq
.Lfunc_end19:
	.size	__hip_module_dtor, .Lfunc_end19-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z10matmul_gpuPfPKfS1_i,@object # @_Z10matmul_gpuPfPKfS1_i
	.section	.rodata,"a",@progbits
	.globl	_Z10matmul_gpuPfPKfS1_i
	.p2align	3, 0x0
_Z10matmul_gpuPfPKfS1_i:
	.quad	_Z25__device_stub__matmul_gpuPfPKfS1_i
	.size	_Z10matmul_gpuPfPKfS1_i, 8

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

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

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"): "
	.size	.L.str.2, 4

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	" != "
	.size	.L.str.3, 5

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"Usage: "
	.size	.L.str.4, 8

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

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

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

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

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

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

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	" [TFLOPS]"
	.size	.L.str.11, 10

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

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

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

	.type	.L__const._ZNSt8__detail18__to_chars_10_implIjEEvPcjT_.__digits,@object # @__const._ZNSt8__detail18__to_chars_10_implIjEEvPcjT_.__digits
	.section	.rodata.str1.16,"aMS",@progbits,1
	.p2align	4, 0x0
.L__const._ZNSt8__detail18__to_chars_10_implIjEEvPcjT_.__digits:
	.asciz	"00010203040506070809101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899"
	.size	.L__const._ZNSt8__detail18__to_chars_10_implIjEEvPcjT_.__digits, 201

	.type	.L.str.17,@object               # @.str.17
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.17:
	.asciz	"cannot create std::vector larger than max_size()"
	.size	.L.str.17, 49

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z10matmul_gpuPfPKfS1_i"
	.size	.L__unnamed_1, 24

	.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 _Z25__device_stub__matmul_gpuPfPKfS1_i
	.addrsig_sym __gxx_personality_v0
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Unwind_Resume
	.addrsig_sym _Z10matmul_gpuPfPKfS1_i
	.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	_Z10matmul_gpuPfPKfS1_i ; -- Begin function _Z10matmul_gpuPfPKfS1_i
	.globl	_Z10matmul_gpuPfPKfS1_i
	.p2align	8
	.type	_Z10matmul_gpuPfPKfS1_i,@function
_Z10matmul_gpuPfPKfS1_i:                ; @_Z10matmul_gpuPfPKfS1_i
; %bb.0:
	s_clause 0x1
	s_load_b32 s3, s[0:1], 0x2c
	s_load_b32 s2, s[0:1], 0x18
	v_bfe_u32 v1, v0, 10, 10
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[0:1], s[0:1], 0x10
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s8, s3, 16
	s_and_b32 s3, s3, 0xffff
	v_mad_u64_u32 v[2:3], null, s15, s8, v[1:2]
	v_and_b32_e32 v3, 0x3ff, v0
	s_cmp_lt_i32 s2, 1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mad_u64_u32 v[0:1], null, s14, s3, v[3:4]
	v_mul_lo_u32 v1, v2, s2
	s_cbranch_scc1 .LBB0_3
; %bb.1:                                ; %.lr.ph.preheader
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_ashrrev_i32_e32 v2, 31, v1
	v_mov_b32_e32 v6, 0
	v_mov_b32_e32 v4, v0
	s_mov_b32 s3, s2
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[2:3], 2, v[1:2]
	v_add_co_u32 v2, vcc_lo, s6, v2
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v3, vcc_lo, s7, v3, vcc_lo
.LBB0_2:                                ; %.lr.ph
                                        ; =>This Inner Loop Header: Depth=1
	v_ashrrev_i32_e32 v5, 31, v4
	s_add_i32 s3, s3, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_cmp_eq_u32 s3, 0
	v_lshlrev_b64 v[7:8], 2, v[4:5]
	v_add_nc_u32_e32 v4, s2, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v7, vcc_lo, s0, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s1, v8, vcc_lo
	global_load_b32 v5, v[2:3], off
	global_load_b32 v7, v[7:8], off
	v_add_co_u32 v2, vcc_lo, v2, 4
	v_add_co_ci_u32_e32 v3, vcc_lo, 0, v3, vcc_lo
	s_waitcnt vmcnt(0)
	v_fmac_f32_e32 v6, v5, v7
	s_cbranch_scc0 .LBB0_2
	s_branch .LBB0_4
.LBB0_3:
	v_mov_b32_e32 v6, 0
.LBB0_4:                                ; %Flow40
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v0, v1, v0
	v_ashrrev_i32_e32 v1, 31, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[0:1]
	v_add_co_u32 v0, vcc_lo, s4, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s5, v1, vcc_lo
	global_store_b32 v[0:1], v6, off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z10matmul_gpuPfPKfS1_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.amdhsa_user_sgpr_count 14
		.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 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 9
		.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	_Z10matmul_gpuPfPKfS1_i, .Lfunc_end0-_Z10matmul_gpuPfPKfS1_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 308
; NumSgprs: 18
; NumVgprs: 9
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 9
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.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
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z10matmul_gpuPfPKfS1_i
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z10matmul_gpuPfPKfS1_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     9
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
