#include <stdio.h>
#include <cstdlib>
#include <iostream>
#include <cuda_runtime.h>
#include <float.h>

#include "cuda_2opt.cuh"
#include "util.cuh"


struct best_struct {
    float best;
    int i;
    int j;
};

__global__ void two_opt_kernel(const float* x, const float* y, best_struct* return_best, int n, int* lock) {
    int i = threadIdx.x + blockIdx.x * blockDim.x+1;
    int j = blockIdx.y * blockDim.x;
    __shared__ float shared_x[65], shared_y[65], shared_best[64];
    __shared__ int shared_j[64];
    shared_best[threadIdx.x] = 0;
    shared_j[threadIdx.x] = 0;
    if (j+threadIdx.x < n) {
        shared_x[threadIdx.x] = x[j+threadIdx.x];
        shared_y[threadIdx.x] = y[j+threadIdx.x];
    }
    if (threadIdx.x == blockDim.x-1) {
        shared_x[blockDim.x] = x[j+threadIdx.x+1];
        shared_y[blockDim.x] = y[j+threadIdx.x+1];
    }
    if (j+blockDim.x+1 <= i || i >= n-2)
        return;
    __syncthreads();
    float best = 0;
    int best_j = 0;
    if (i > 0) {
        float xi = x[i], xim = x[i-1], yi = y[i], yim = y[i-1];
        float i_dist = dist(xi, yi, xim, yim);
        if (j <= i || j + blockDim.x >= n) {
            for (int k = 0; k < blockDim.x; ++k) {
                if (j+k > i && j+k < n-1) {
                    float k_dist = i_dist + dist(shared_x[k], shared_y[k], shared_x[k+1], shared_y[k+1]) - 
                    (dist(xi, yi, shared_x[k+1], shared_y[k+1]) + dist(xim, yim, shared_x[k], shared_y[k]));
                    if (k_dist > best) {
                        best = k_dist;
                        best_j = j+k;
                    }
                }
            }
        } else {
            for (int k = 0; k < blockDim.x; ++k) {
                float k_dist = i_dist + dist(shared_x[k], shared_y[k], shared_x[k+1], shared_y[k+1]) - 
                (dist(xi, yi, shared_x[k+1], shared_y[k+1]) + dist(xim, yim, shared_x[k], shared_y[k]));
                if (k_dist > best) {
                    best = k_dist;
                    best_j = j+k;
                }
            }
        }
    }

    int out_w = divupg(n, 64);
    shared_best[threadIdx.x] = best;
    shared_j[threadIdx.x] = best_j;
    int best_i = i;
    __syncthreads();
    if (threadIdx.x == 0) {
        for (int i2 = 0; i2 < blockDim.x; ++i2) {
            if (i + i2 >= n-2)
                continue;
            if (shared_best[i2] > best) {
                best_i = i+i2;
                best = shared_best[i2];
                best_j = shared_j[i2];
            }
        }
        if (best > return_best[0].best) {
            while (atomicExch(&lock[0], 1) != 0);
            if (best > return_best[0].best) {
                return_best[0].best = best;
                return_best[0].i = best_i;
                return_best[0].j = best_j;
            }
            lock[0] = 0;
            __threadfence();
        }
    }
}

__global__ void two_opt_swap_kernel(float* x, float* y, int i, int j) {
    int k = blockIdx.x;
    float tmp = x[i+k];
    x[i+k] = x[j-k];
    x[j-k] = tmp;
    tmp = y[i+k];
    y[i+k] = y[j-k];
    y[j-k] = tmp;
}

void run_gpu_2opt(float* x, float* y, int n) {
    float* xGPU = NULL;
    cudaMalloc((void**)&xGPU, n * sizeof(float));
    cudaMemcpy(xGPU, x, n * sizeof(float), cudaMemcpyHostToDevice);
    float* yGPU = NULL;
    cudaMalloc((void**)&yGPU, n * sizeof(float));
    cudaMemcpy(yGPU, y, n * sizeof(float), cudaMemcpyHostToDevice);
    best_struct* bestGPU = NULL;
    cudaMalloc((void**)&bestGPU, sizeof(best_struct));
    int* lock = NULL;
    cudaMalloc((void**)&lock, 1 * sizeof(int));

    dim3 dimBlock(64, 1);
    dim3 dimGrid(divup(n, 64), divup(n, 64));
    do {
        CHECK(cudaGetLastError());
        cudaDeviceSynchronize();
        cudaMemset(bestGPU, 0, sizeof(best_struct));
        CHECK(cudaGetLastError());
        cudaDeviceSynchronize();
        cudaMemset(lock, 0, 1*sizeof(int));
        CHECK(cudaGetLastError());
        cudaDeviceSynchronize();
        two_opt_kernel<<<dimGrid, dimBlock>>>(xGPU, yGPU, bestGPU, n, lock);
        CHECK(cudaGetLastError());
        cudaDeviceSynchronize();
        best_struct* best = (best_struct*)malloc(sizeof(best_struct));
        cudaMemcpy(best, bestGPU, sizeof(best_struct), cudaMemcpyDeviceToHost);
        cudaDeviceSynchronize();
        if (abs(best[0].best) < 0.000001)
            break;
        //printf("Improvement %f %d %d\n", best[0].best, best[0].i, best[0].j);
        two_opt_swap_kernel<<<dim3((best[0].j-best[0].i+1)/2, 1), dim3(1, 1)>>>(xGPU, yGPU, best[0].i, best[0].j);
        CHECK(cudaGetLastError());
        cudaDeviceSynchronize();
    } while(true);

    cudaMemcpy(x, xGPU, n * sizeof(float), cudaMemcpyDeviceToHost);
    cudaMemcpy(y, yGPU, n * sizeof(float), cudaMemcpyDeviceToHost);
    cudaFree(xGPU);
    cudaFree(yGPU);
    cudaFree(bestGPU);
    cudaFree(lock);
}

__global__ void two_opt_swap_kernel_id(float* x, float* y, int* id, int i, int j) {
    int k = blockIdx.x;
    float tmp1 = x[i+k];
    x[i+k] = x[j-k];
    x[j-k] = tmp1;
    float tmp2 = y[i+k];
    y[i+k] = y[j-k];
    y[j-k] = tmp2;
    int tmp3 = id[i+k];
    id[i+k] = id[j-k];
    id[j-k] = tmp3;
}

__global__ void id_map_calculate(int* id, int* id_map) {
    int k = blockIdx.x;
    id_map[id[k]] = k;
}

__global__ void two_opt_kernel_restricted(const float* x, const float* y, const int* moves, const int* id, const int* id_map, best_struct* return_best, int n, const int allowed_moves, int* lock) {
    int i = blockIdx.x;
    float best = 0;
    int best_j = 0;
    if (i > 0) {
        int id_i = id[i];
        int id_i_prev = id[i-1];
        const int* i_prev_moves = moves + id_i_prev * allowed_moves;
        const int* i_moves = moves + id_i * allowed_moves;
        
        for (int k = 0; k < allowed_moves; ++k) {
            int id_j = i_prev_moves[k];
            int j = id_map[id_j];
            if (j != 0 && j != n-1 && j != i-1) {
                int jp = id[j+1];
                for (int k2 = 0; k2 < allowed_moves; ++k2) {
                    int id_jp = i_moves[k2];
                    if (jp == id_jp) {
                        int i2 = min(i, j);
                        int j2 = max(i, j);
                        float k_dist = dist(x[i2], y[i2], x[i2-1], y[i2-1]) + dist(x[j2], y[j2], x[j2+1], y[j2+1]) - 
                        (dist(x[i2], y[i2], x[j2+1], y[j2+1]) + dist(x[i2-1], y[i2-1], x[j2], y[j2]));
                        //printf("sis %d %d %d %d %f\n", i, j, id_i, jp, k_dist);
                        if (k_dist > best) {
                            best = k_dist;
                            best_j = j;
                        }
                    }
                }
            }
        }
    }
    //printf("%d parasa %.3f\n", i, best);
    if (best > return_best[0].best) {
        while (atomicExch(&lock[0], 1) != 0);
        if (best > return_best[0].best) {
            return_best[0].best = best;
            return_best[0].i = min(i, best_j);
            return_best[0].j = max(i, best_j);
        }
        lock[0] = 0;
    }
}

void run_gpu_2opt_restricted(float* x, float* y, int* id, int* moves, int n, int allowed_moves) {
    float* xGPU = NULL;
    cudaMalloc((void**)&xGPU, n * sizeof(float));
    cudaMemcpy(xGPU, x, n * sizeof(float), cudaMemcpyHostToDevice);
    float* yGPU = NULL;
    cudaMalloc((void**)&yGPU, n * sizeof(float));
    cudaMemcpy(yGPU, y, n * sizeof(float), cudaMemcpyHostToDevice);
    best_struct* bestGPU = NULL;
    cudaMalloc((void**)&bestGPU, sizeof(best_struct));
    int* lock = NULL;
    cudaMalloc((void**)&lock, 1 * sizeof(int));

    int* Gmoves = NULL;
    cudaMalloc((void**)&Gmoves, n * allowed_moves * sizeof(int));
    cudaMemcpy(Gmoves, moves, n * allowed_moves * sizeof(float), cudaMemcpyHostToDevice);

    int* Gid = NULL;
    cudaMalloc((void**)&Gid, n * sizeof(int));
    cudaMemcpy(Gid, id, n * sizeof(float), cudaMemcpyHostToDevice);

    std::vector<int> id_map(n);
    for (int i = 0; i < n; ++i)
        id_map[id[i]] = i;

    int* Gid_map = NULL;
    cudaMalloc((void**)&Gid_map, n * sizeof(int));
    cudaMemcpy(Gid_map, id_map.data(), n * sizeof(float), cudaMemcpyHostToDevice);

    dim3 dimBlock(1, 1);
    dim3 dimGrid(n, 1);
    do {
        CHECK(cudaGetLastError());
        cudaDeviceSynchronize();
        cudaMemset(bestGPU, 0, sizeof(best_struct));
        CHECK(cudaGetLastError());
        cudaDeviceSynchronize();
        cudaMemset(lock, 0, 1*sizeof(int));
        CHECK(cudaGetLastError());
        cudaDeviceSynchronize();
        two_opt_kernel_restricted<<<dimGrid, dimBlock>>>(xGPU, yGPU, Gmoves, Gid, Gid_map, bestGPU, n, allowed_moves, lock);
        CHECK(cudaGetLastError());
        cudaDeviceSynchronize();
        best_struct* best = (best_struct*)malloc(sizeof(best_struct));
        cudaMemcpy(best, bestGPU, sizeof(best_struct), cudaMemcpyDeviceToHost);
        cudaDeviceSynchronize();
        //printf("Improvement %f %d %d\n", best[0].best, best[0].i, best[0].j);
        //fflush(stdout);
        if (abs(best[0].best) < 0.000001)
            break;
        
        two_opt_swap_kernel_id<<<dim3((best[0].j-best[0].i+1)/2, 1), dim3(1, 1)>>>(xGPU, yGPU, Gid, best[0].i, best[0].j);
        CHECK(cudaGetLastError());
        cudaDeviceSynchronize();
        id_map_calculate<<<dim3(n, 1), dim3(1, 1)>>>(Gid, Gid_map);
        CHECK(cudaGetLastError());
        cudaDeviceSynchronize();
    } while(true);

    cudaMemcpy(x, xGPU, n * sizeof(float), cudaMemcpyDeviceToHost);
    cudaMemcpy(y, yGPU, n * sizeof(float), cudaMemcpyDeviceToHost);
    cudaMemcpy(id, Gid, n * sizeof(float), cudaMemcpyDeviceToHost);
    cudaDeviceSynchronize();
    cudaFree(xGPU);
    cudaFree(yGPU);
    cudaFree(bestGPU);
    cudaFree(lock);
    cudaFree(Gmoves);
    cudaFree(Gid);
}
	.file	"tmpxft_00260529_00000000-6_cuda_2opt.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4008:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4008:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"CUDA error: "
.LC1:
	.string	"cudaGetLastError()"
.LC2:
	.string	": "
	.text
	.type	_ZL5check9cudaErrorPKc.constprop.0, @function
_ZL5check9cudaErrorPKc.constprop.0:
.LFB4746:
	.cfi_startproc
	testl	%edi, %edi
	je	.L2
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	leaq	.LC0(%rip), %rsi
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movl	%edi, %ebx
	leaq	_ZSt4cerr(%rip), %rdi
	pushq	%rax
	.cfi_def_cfa_offset 32
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	leaq	.LC1(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	leaq	.LC2(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %edi
	movq	%rax, %rbp
	call	cudaGetErrorString@PLT
	movq	%rbp, %rdi
	movq	%rax, %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	movl	$1, %edi
	call	exit@PLT
.L2:
	.cfi_def_cfa_offset 8
	.cfi_restore 3
	.cfi_restore 6
	ret
	.cfi_endproc
.LFE4746:
	.size	_ZL5check9cudaErrorPKc.constprop.0, .-_ZL5check9cudaErrorPKc.constprop.0
	.globl	_Z55__device_stub__Z14two_opt_kernelPKfS0_P11best_structiPiPKfS0_P11best_structiPi
	.type	_Z55__device_stub__Z14two_opt_kernelPKfS0_P11best_structiPiPKfS0_P11best_structiPi, @function
_Z55__device_stub__Z14two_opt_kernelPKfS0_P11best_structiPiPKfS0_P11best_structiPi:
.LFB4030:
	.cfi_startproc
	endbr64
	subq	$168, %rsp
	.cfi_def_cfa_offset 176
	movq	%rdi, 40(%rsp)
	leaq	64(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	48(%rsp), %rdx
	movl	%ecx, 20(%rsp)
	leaq	56(%rsp), %rcx
	movq	%r8, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 112(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 144(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L8
	pushq	56(%rsp)
	.cfi_def_cfa_offset 184
	leaq	_Z14two_opt_kernelPKfS0_P11best_structiPi(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 192
	movq	92(%rsp), %rcx
	movl	100(%rsp), %r8d
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	leaq	128(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 184
	popq	%rdx
	.cfi_def_cfa_offset 176
.L8:
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$168, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4030:
	.size	_Z55__device_stub__Z14two_opt_kernelPKfS0_P11best_structiPiPKfS0_P11best_structiPi, .-_Z55__device_stub__Z14two_opt_kernelPKfS0_P11best_structiPiPKfS0_P11best_structiPi
	.globl	_Z14two_opt_kernelPKfS0_P11best_structiPi
	.type	_Z14two_opt_kernelPKfS0_P11best_structiPi, @function
_Z14two_opt_kernelPKfS0_P11best_structiPi:
.LFB4031:
	.cfi_startproc
	endbr64
	jmp	_Z55__device_stub__Z14two_opt_kernelPKfS0_P11best_structiPiPKfS0_P11best_structiPi
	.cfi_endproc
.LFE4031:
	.size	_Z14two_opt_kernelPKfS0_P11best_structiPi, .-_Z14two_opt_kernelPKfS0_P11best_structiPi
	.globl	_Z43__device_stub__Z19two_opt_swap_kernelPfS_iiPfS_ii
	.type	_Z43__device_stub__Z19two_opt_swap_kernelPfS_iiPfS_ii, @function
_Z43__device_stub__Z19two_opt_swap_kernelPfS_iiPfS_ii:
.LFB4032:
	.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
	movl	%edx, 12(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 8(%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	12(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	8(%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	.L13
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z19two_opt_swap_kernelPfS_ii(%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
.L13:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L15
	call	__stack_chk_fail@PLT
.L15:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4032:
	.size	_Z43__device_stub__Z19two_opt_swap_kernelPfS_iiPfS_ii, .-_Z43__device_stub__Z19two_opt_swap_kernelPfS_iiPfS_ii
	.globl	_Z19two_opt_swap_kernelPfS_ii
	.type	_Z19two_opt_swap_kernelPfS_ii, @function
_Z19two_opt_swap_kernelPfS_ii:
.LFB4033:
	.cfi_startproc
	endbr64
	jmp	_Z43__device_stub__Z19two_opt_swap_kernelPfS_iiPfS_ii
	.cfi_endproc
.LFE4033:
	.size	_Z19two_opt_swap_kernelPfS_ii, .-_Z19two_opt_swap_kernelPfS_ii
	.globl	_Z12run_gpu_2optPfS_i
	.type	_Z12run_gpu_2optPfS_i, @function
_Z12run_gpu_2optPfS_i:
.LFB4004:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movq	%rdi, %r14
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	movq	%rsi, %r13
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	movslq	%edx, %rbx
	movq	%rbx, %r12
	salq	$2, %rbx
	movq	%rbx, %rsi
	subq	$96, %rsp
	.cfi_def_cfa_offset 144
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rdi
	xorl	%eax, %eax
	movq	%rax, 8(%rsp)
	call	cudaMalloc@PLT
	movq	8(%rsp), %rdi
	movq	%rbx, %rdx
	movq	%r14, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	xorl	%edx, %edx
	leaq	16(%rsp), %rdi
	movq	%rbx, %rsi
	movq	%rdx, 16(%rsp)
	call	cudaMalloc@PLT
	movq	16(%rsp), %rdi
	movq	%rbx, %rdx
	movq	%r13, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	xorl	%ecx, %ecx
	leaq	24(%rsp), %rdi
	movl	$12, %esi
	movq	%rcx, 24(%rsp)
	call	cudaMalloc@PLT
	xorl	%esi, %esi
	leaq	32(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	movl	$4, %esi
	call	cudaMalloc@PLT
	movl	$67108865, %eax
	movl	$64, %ecx
	movl	$1, 48(%rsp)
	salq	$6, %rax
	movl	$1, 60(%rsp)
	movq	%rax, 40(%rsp)
	leal	63(%r12), %eax
	cltd
	idivl	%ecx
	movl	%eax, 52(%rsp)
	movl	%eax, 56(%rsp)
.L22:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	call	_ZL5check9cudaErrorPKc.constprop.0
	call	cudaDeviceSynchronize@PLT
	movq	24(%rsp), %rdi
	xorl	%esi, %esi
	movl	$12, %edx
	call	cudaMemset@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	call	_ZL5check9cudaErrorPKc.constprop.0
	call	cudaDeviceSynchronize@PLT
	movq	32(%rsp), %rdi
	xorl	%esi, %esi
	movl	$4, %edx
	call	cudaMemset@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	call	_ZL5check9cudaErrorPKc.constprop.0
	call	cudaDeviceSynchronize@PLT
	movl	48(%rsp), %ecx
	movl	60(%rsp), %esi
	xorl	%r9d, %r9d
	movq	40(%rsp), %rdx
	movq	52(%rsp), %rdi
	xorl	%r8d, %r8d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L19
	movq	32(%rsp), %r8
	movq	24(%rsp), %rdx
	movl	%r12d, %ecx
	movq	16(%rsp), %rsi
	movq	8(%rsp), %rdi
	call	_Z55__device_stub__Z14two_opt_kernelPKfS0_P11best_structiPiPKfS0_P11best_structiPi
.L19:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	call	_ZL5check9cudaErrorPKc.constprop.0
	call	cudaDeviceSynchronize@PLT
	movl	$12, %edi
	call	malloc@PLT
	movq	24(%rsp), %rsi
	movl	$2, %ecx
	movl	$12, %edx
	movq	%rax, %rbp
	movq	%rax, %rdi
	call	cudaMemcpy@PLT
	call	cudaDeviceSynchronize@PLT
	movss	0(%rbp), %xmm0
	movsd	.LC4(%rip), %xmm1
	andps	.LC3(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	comisd	%xmm0, %xmm1
	ja	.L20
	movl	8(%rbp), %eax
	subl	4(%rbp), %eax
	movl	$2, %esi
	xorl	%r9d, %r9d
	movabsq	$4294967297, %rcx
	incl	%eax
	xorl	%r8d, %r8d
	cltd
	movq	%rcx, 68(%rsp)
	idivl	%esi
	movl	72(%rsp), %esi
	movq	%rcx, %rdx
	movl	$1, %ecx
	movl	%eax, 64(%rsp)
	movq	64(%rsp), %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L21
	movl	8(%rbp), %ecx
	movl	4(%rbp), %edx
	movq	16(%rsp), %rsi
	movq	8(%rsp), %rdi
	call	_Z43__device_stub__Z19two_opt_swap_kernelPfS_iiPfS_ii
.L21:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	call	_ZL5check9cudaErrorPKc.constprop.0
	call	cudaDeviceSynchronize@PLT
	jmp	.L22
.L20:
	movq	8(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	%r14, %rdi
	call	cudaMemcpy@PLT
	movq	16(%rsp), %rsi
	movq	%rbx, %rdx
	movq	%r13, %rdi
	movl	$2, %ecx
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	32(%rsp), %rdi
	call	cudaFree@PLT
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L23
	call	__stack_chk_fail@PLT
.L23:
	addq	$96, %rsp
	.cfi_def_cfa_offset 48
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4004:
	.size	_Z12run_gpu_2optPfS_i, .-_Z12run_gpu_2optPfS_i
	.globl	_Z48__device_stub__Z22two_opt_swap_kernel_idPfS_PiiiPfS_Piii
	.type	_Z48__device_stub__Z22two_opt_swap_kernel_idPfS_PiiiPfS_Piii, @function
_Z48__device_stub__Z22two_opt_swap_kernel_idPfS_PiiiPfS_Piii:
.LFB4034:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movl	%ecx, 4(%rsp)
	leaq	40(%rsp), %rcx
	movl	%r8d, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	%rsp, %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L25
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z22two_opt_swap_kernel_idPfS_Piii(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 176
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L25:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L27
	call	__stack_chk_fail@PLT
.L27:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4034:
	.size	_Z48__device_stub__Z22two_opt_swap_kernel_idPfS_PiiiPfS_Piii, .-_Z48__device_stub__Z22two_opt_swap_kernel_idPfS_PiiiPfS_Piii
	.globl	_Z22two_opt_swap_kernel_idPfS_Piii
	.type	_Z22two_opt_swap_kernel_idPfS_Piii, @function
_Z22two_opt_swap_kernel_idPfS_Piii:
.LFB4035:
	.cfi_startproc
	endbr64
	jmp	_Z48__device_stub__Z22two_opt_swap_kernel_idPfS_PiiiPfS_Piii
	.cfi_endproc
.LFE4035:
	.size	_Z22two_opt_swap_kernel_idPfS_Piii, .-_Z22two_opt_swap_kernel_idPfS_Piii
	.globl	_Z38__device_stub__Z16id_map_calculatePiS_PiS_
	.type	_Z38__device_stub__Z16id_map_calculatePiS_PiS_, @function
_Z38__device_stub__Z16id_map_calculatePiS_PiS_:
.LFB4036:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L30
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z16id_map_calculatePiS_(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L30:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L32
	call	__stack_chk_fail@PLT
.L32:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4036:
	.size	_Z38__device_stub__Z16id_map_calculatePiS_PiS_, .-_Z38__device_stub__Z16id_map_calculatePiS_PiS_
	.globl	_Z16id_map_calculatePiS_
	.type	_Z16id_map_calculatePiS_, @function
_Z16id_map_calculatePiS_:
.LFB4037:
	.cfi_startproc
	endbr64
	jmp	_Z38__device_stub__Z16id_map_calculatePiS_PiS_
	.cfi_endproc
.LFE4037:
	.size	_Z16id_map_calculatePiS_, .-_Z16id_map_calculatePiS_
	.globl	_Z76__device_stub__Z25two_opt_kernel_restrictedPKfS0_PKiS2_S2_P11best_structiiPiPKfS0_PKiS2_S2_P11best_structiiPi
	.type	_Z76__device_stub__Z25two_opt_kernel_restrictedPKfS0_PKiS2_S2_P11best_structiiPiPKfS0_PKiS2_S2_P11best_structiiPi, @function
_Z76__device_stub__Z25two_opt_kernel_restrictedPKfS0_PKiS2_S2_P11best_structiiPiPKfS0_PKiS2_S2_P11best_structiiPi:
.LFB4038:
	.cfi_startproc
	endbr64
	subq	$216, %rsp
	.cfi_def_cfa_offset 224
	movq	240(%rsp), %rax
	movq	%rdi, 56(%rsp)
	leaq	80(%rsp), %rdi
	movq	%rsi, 48(%rsp)
	leaq	92(%rsp), %rsi
	movq	%rdx, 40(%rsp)
	leaq	64(%rsp), %rdx
	movq	%rcx, 32(%rsp)
	leaq	72(%rsp), %rcx
	movq	%r8, 24(%rsp)
	movq	%r9, 16(%rsp)
	movq	%rax, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 200(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rax
	movl	$1, 88(%rsp)
	movq	%rax, 128(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	224(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	232(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 192(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 80(%rsp)
	movq	%rax, 92(%rsp)
	movl	$1, 100(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L35
	pushq	72(%rsp)
	.cfi_def_cfa_offset 232
	leaq	_Z25two_opt_kernel_restrictedPKfS0_PKiS2_S2_P11best_structiiPi(%rip), %rdi
	pushq	72(%rsp)
	.cfi_def_cfa_offset 240
	movq	108(%rsp), %rcx
	movl	116(%rsp), %r8d
	movq	96(%rsp), %rsi
	movl	104(%rsp), %edx
	leaq	144(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 232
	popq	%rdx
	.cfi_def_cfa_offset 224
.L35:
	movq	200(%rsp), %rax
	subq	%fs:40, %rax
	je	.L37
	call	__stack_chk_fail@PLT
.L37:
	addq	$216, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4038:
	.size	_Z76__device_stub__Z25two_opt_kernel_restrictedPKfS0_PKiS2_S2_P11best_structiiPiPKfS0_PKiS2_S2_P11best_structiiPi, .-_Z76__device_stub__Z25two_opt_kernel_restrictedPKfS0_PKiS2_S2_P11best_structiiPiPKfS0_PKiS2_S2_P11best_structiiPi
	.globl	_Z25two_opt_kernel_restrictedPKfS0_PKiS2_S2_P11best_structiiPi
	.type	_Z25two_opt_kernel_restrictedPKfS0_PKiS2_S2_P11best_structiiPi, @function
_Z25two_opt_kernel_restrictedPKfS0_PKiS2_S2_P11best_structiiPi:
.LFB4039:
	.cfi_startproc
	endbr64
	jmp	_Z76__device_stub__Z25two_opt_kernel_restrictedPKfS0_PKiS2_S2_P11best_structiiPiPKfS0_PKiS2_S2_P11best_structiiPi
	.cfi_endproc
.LFE4039:
	.size	_Z25two_opt_kernel_restrictedPKfS0_PKiS2_S2_P11best_structiiPi, .-_Z25two_opt_kernel_restrictedPKfS0_PKiS2_S2_P11best_structiiPi
	.section	.rodata.str1.1
.LC5:
	.string	"_Z25two_opt_kernel_restrictedPKfS0_PKiS2_S2_P11best_structiiPi"
.LC6:
	.string	"_Z16id_map_calculatePiS_"
.LC7:
	.string	"_Z22two_opt_swap_kernel_idPfS_Piii"
.LC8:
	.string	"_Z19two_opt_swap_kernelPfS_ii"
.LC9:
	.string	"_Z14two_opt_kernelPKfS0_P11best_structiPi"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4041:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC5(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z25two_opt_kernel_restrictedPKfS0_PKiS2_S2_P11best_structiiPi(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC6(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z16id_map_calculatePiS_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC7(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z22two_opt_swap_kernel_idPfS_Piii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC8(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z19two_opt_swap_kernelPfS_ii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC9(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z14two_opt_kernelPKfS0_P11best_structiPi(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE4041:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.rodata._ZNSt6vectorIiSaIiEEC2EmRKS0_.str1.1,"aMS",@progbits,1
.LC10:
	.string	"cannot create std::vector larger than max_size()"
	.section	.text._ZNSt6vectorIiSaIiEEC2EmRKS0_,"axG",@progbits,_ZNSt6vectorIiSaIiEEC5EmRKS0_,comdat
	.align 2
	.weak	_ZNSt6vectorIiSaIiEEC2EmRKS0_
	.type	_ZNSt6vectorIiSaIiEEC2EmRKS0_, @function
_ZNSt6vectorIiSaIiEEC2EmRKS0_:
.LFB4351:
	.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	.L43
	leaq	.LC10(%rip), %rdi
	call	_ZSt20__throw_length_errorPKc@PLT
.L43:
	movq	%rsi, %rbp
	xorl	%esi, %esi
	movq	%rdi, %rbx
	xorl	%eax, %eax
	movq	%rsi, (%rdi)
	leaq	0(,%rbp,4), %r12
	movq	%rsi, 8(%rdi)
	movq	%rsi, 16(%rdi)
	testq	%rbp, %rbp
	je	.L44
	movq	%r12, %rdi
	call	_Znwm@PLT
.L44:
	addq	%rax, %r12
	movq	%rax, (%rbx)
	movq	%rax, 8(%rbx)
	movq	%r12, 16(%rbx)
	testq	%rbp, %rbp
	je	.L45
	xorl	%ecx, %ecx
	addq	$4, %rax
	movl	%ecx, -4(%rax)
	decq	%rbp
	je	.L45
.L46:
	cmpq	%rax, %r12
	je	.L45
	xorl	%edx, %edx
	addq	$4, %rax
	movl	%edx, -4(%rax)
	jmp	.L46
.L45:
	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
.LFE4351:
	.size	_ZNSt6vectorIiSaIiEEC2EmRKS0_, .-_ZNSt6vectorIiSaIiEEC2EmRKS0_
	.weak	_ZNSt6vectorIiSaIiEEC1EmRKS0_
	.set	_ZNSt6vectorIiSaIiEEC1EmRKS0_,_ZNSt6vectorIiSaIiEEC2EmRKS0_
	.section	.text._ZNSt6vectorIiSaIiEED2Ev,"axG",@progbits,_ZNSt6vectorIiSaIiEED5Ev,comdat
	.align 2
	.weak	_ZNSt6vectorIiSaIiEED2Ev
	.type	_ZNSt6vectorIiSaIiEED2Ev, @function
_ZNSt6vectorIiSaIiEED2Ev:
.LFB4354:
	.cfi_startproc
	endbr64
	movq	(%rdi), %rax
	testq	%rax, %rax
	je	.L57
	movq	16(%rdi), %rsi
	movq	%rax, %rdi
	subq	%rax, %rsi
	jmp	_ZdlPvm@PLT
.L57:
	ret
	.cfi_endproc
.LFE4354:
	.size	_ZNSt6vectorIiSaIiEED2Ev, .-_ZNSt6vectorIiSaIiEED2Ev
	.weak	_ZNSt6vectorIiSaIiEED1Ev
	.set	_ZNSt6vectorIiSaIiEED1Ev,_ZNSt6vectorIiSaIiEED2Ev
	.text
	.globl	_Z23run_gpu_2opt_restrictedPfS_PiS0_ii
	.type	_Z23run_gpu_2opt_restrictedPfS_PiS0_ii, @function
_Z23run_gpu_2opt_restrictedPfS_PiS0_ii:
.LFB4005:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4005
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register 6
	pushq	%r15
	.cfi_offset 15, -24
	movl	%r9d, %r15d
	pushq	%r14
	.cfi_offset 14, -32
	movq	%rdi, %r14
	leaq	-184(%rbp), %rdi
	pushq	%r13
	.cfi_offset 13, -40
	movq	%rsi, %r13
	pushq	%r12
	.cfi_offset 12, -48
	movq	%rdx, %r12
	pushq	%rbx
	.cfi_offset 3, -56
	movl	%r8d, %ebx
	subq	$184, %rsp
	movq	%rcx, -208(%rbp)
	xorl	%ecx, %ecx
	movq	%fs:40, %rax
	movq	%rax, -56(%rbp)
	xorl	%eax, %eax
	movslq	%r8d, %rax
	movq	%rcx, -184(%rbp)
	salq	$2, %rax
	movq	%rax, %rsi
	movq	%rax, -200(%rbp)
.LEHB0:
	call	cudaMalloc@PLT
	movq	-200(%rbp), %rdx
	movl	$1, %ecx
	movq	%r14, %rsi
	movq	-184(%rbp), %rdi
	call	cudaMemcpy@PLT
	xorl	%esi, %esi
	leaq	-176(%rbp), %rdi
	movq	%rsi, -176(%rbp)
	movq	-200(%rbp), %rsi
	call	cudaMalloc@PLT
	movq	-200(%rbp), %rdx
	movl	$1, %ecx
	movq	%r13, %rsi
	movq	-176(%rbp), %rdi
	call	cudaMemcpy@PLT
	xorl	%edi, %edi
	movl	$12, %esi
	movq	%rdi, -168(%rbp)
	leaq	-168(%rbp), %rdi
	call	cudaMalloc@PLT
	xorl	%r8d, %r8d
	movl	$4, %esi
	leaq	-160(%rbp), %rdi
	movq	%r8, -160(%rbp)
	call	cudaMalloc@PLT
	movl	%ebx, %edx
	xorl	%r9d, %r9d
	leaq	-152(%rbp), %rdi
	imull	%r15d, %edx
	movq	%r9, -152(%rbp)
	movslq	%edx, %rdx
	salq	$2, %rdx
	movq	%rdx, %rsi
	movq	%rdx, -216(%rbp)
	call	cudaMalloc@PLT
	movq	-216(%rbp), %rdx
	movl	$1, %ecx
	movq	-208(%rbp), %rsi
	movq	-152(%rbp), %rdi
	call	cudaMemcpy@PLT
	movq	-200(%rbp), %rsi
	xorl	%r10d, %r10d
	leaq	-144(%rbp), %rdi
	movq	%r10, -144(%rbp)
	call	cudaMalloc@PLT
	movq	-200(%rbp), %rdx
	movl	$1, %ecx
	movq	%r12, %rsi
	movq	-144(%rbp), %rdi
	call	cudaMemcpy@PLT
	leaq	-80(%rbp), %rax
	leaq	-92(%rbp), %rdx
	movslq	%ebx, %rsi
	movq	%rax, %rdi
	movq	%rax, -216(%rbp)
	call	_ZNSt6vectorIiSaIiEEC1EmRKS0_
.LEHE0:
	movq	-80(%rbp), %rax
	movq	%rax, -208(%rbp)
	xorl	%eax, %eax
.L60:
	cmpl	%eax, %ebx
	jle	.L73
	movslq	(%r12,%rax,4), %rdx
	movq	-208(%rbp), %rcx
	movl	%eax, (%rcx,%rdx,4)
	incq	%rax
	jmp	.L60
.L73:
	movq	-200(%rbp), %rsi
	xorl	%edx, %edx
	leaq	-136(%rbp), %rdi
	movq	%rdx, -136(%rbp)
.LEHB1:
	call	cudaMalloc@PLT
	movq	-200(%rbp), %rdx
	movq	-208(%rbp), %rsi
	movl	$1, %ecx
	movq	-136(%rbp), %rdi
	call	cudaMemcpy@PLT
	movabsq	$4294967297, %rax
	movl	$1, -120(%rbp)
	movq	%rax, -128(%rbp)
	movl	%ebx, -116(%rbp)
	movq	%rax, -112(%rbp)
.L66:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	call	_ZL5check9cudaErrorPKc.constprop.0
	call	cudaDeviceSynchronize@PLT
	movq	-168(%rbp), %rdi
	movl	$12, %edx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	call	_ZL5check9cudaErrorPKc.constprop.0
	call	cudaDeviceSynchronize@PLT
	movq	-160(%rbp), %rdi
	movl	$4, %edx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	call	_ZL5check9cudaErrorPKc.constprop.0
	call	cudaDeviceSynchronize@PLT
	movl	-120(%rbp), %ecx
	movq	-128(%rbp), %rdx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	-116(%rbp), %rdi
	movl	-108(%rbp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L62
	pushq	%rax
	movq	-168(%rbp), %r9
	pushq	-160(%rbp)
	movq	-176(%rbp), %rsi
	movq	-136(%rbp), %r8
	pushq	%r15
	movq	-144(%rbp), %rcx
	movq	-152(%rbp), %rdx
	pushq	%rbx
	movq	-184(%rbp), %rdi
	.cfi_escape 0x2e,0x20
	call	_Z76__device_stub__Z25two_opt_kernel_restrictedPKfS0_PKiS2_S2_P11best_structiiPiPKfS0_PKiS2_S2_P11best_structiiPi
	addq	$32, %rsp
.L62:
	.cfi_escape 0x2e,0
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	call	_ZL5check9cudaErrorPKc.constprop.0
	call	cudaDeviceSynchronize@PLT
	movl	$12, %edi
	call	malloc@PLT
	movl	$2, %ecx
	movl	$12, %edx
	movq	-168(%rbp), %rsi
	movq	%rax, %rdi
	movq	%rax, -208(%rbp)
	call	cudaMemcpy@PLT
	call	cudaDeviceSynchronize@PLT
	movq	-208(%rbp), %rax
	movsd	.LC4(%rip), %xmm1
	movss	(%rax), %xmm0
	andps	.LC3(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	comisd	%xmm0, %xmm1
	ja	.L63
	movq	-208(%rbp), %rsi
	movl	8(%rax), %eax
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movabsq	$4294967297, %rcx
	movl	$1, -84(%rbp)
	subl	4(%rsi), %eax
	movl	$2, %esi
	movq	%rcx, -100(%rbp)
	incl	%eax
	movq	%rcx, -92(%rbp)
	cltd
	idivl	%esi
	movl	-96(%rbp), %esi
	movq	%rcx, %rdx
	movl	$1, %ecx
	movl	%eax, -104(%rbp)
	movq	-104(%rbp), %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L64
	movq	-208(%rbp), %rax
	movq	-144(%rbp), %rdx
	movq	-176(%rbp), %rsi
	movq	-184(%rbp), %rdi
	movl	4(%rax), %ecx
	movl	8(%rax), %r8d
	call	_Z48__device_stub__Z22two_opt_swap_kernel_idPfS_PiiiPfS_Piii
.L64:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	call	_ZL5check9cudaErrorPKc.constprop.0
	call	cudaDeviceSynchronize@PLT
	movl	%ebx, -104(%rbp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movabsq	$4294967297, %rdx
	movl	$1, -84(%rbp)
	movq	%rdx, -100(%rbp)
	movq	-104(%rbp), %rdi
	movl	-96(%rbp), %esi
	movq	%rdx, -92(%rbp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L65
	movq	-136(%rbp), %rsi
	movq	-144(%rbp), %rdi
	call	_Z38__device_stub__Z16id_map_calculatePiS_PiS_
.L65:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	call	_ZL5check9cudaErrorPKc.constprop.0
	call	cudaDeviceSynchronize@PLT
	jmp	.L66
.L63:
	movq	-200(%rbp), %rdx
	movq	-184(%rbp), %rsi
	movl	$2, %ecx
	movq	%r14, %rdi
	call	cudaMemcpy@PLT
	movq	-200(%rbp), %rdx
	movq	-176(%rbp), %rsi
	movl	$2, %ecx
	movq	%r13, %rdi
	call	cudaMemcpy@PLT
	movq	-200(%rbp), %rdx
	movq	-144(%rbp), %rsi
	movl	$2, %ecx
	movq	%r12, %rdi
	call	cudaMemcpy@PLT
	call	cudaDeviceSynchronize@PLT
	movq	-184(%rbp), %rdi
	call	cudaFree@PLT
	movq	-176(%rbp), %rdi
	call	cudaFree@PLT
	movq	-168(%rbp), %rdi
	call	cudaFree@PLT
	movq	-160(%rbp), %rdi
	call	cudaFree@PLT
	movq	-152(%rbp), %rdi
	call	cudaFree@PLT
	movq	-144(%rbp), %rdi
	call	cudaFree@PLT
.LEHE1:
	leaq	-80(%rbp), %rdi
	call	_ZNSt6vectorIiSaIiEED1Ev
	movq	-56(%rbp), %rax
	subq	%fs:40, %rax
	je	.L69
	jmp	.L71
.L70:
	endbr64
	movq	%rax, %rbx
.L67:
	movq	-216(%rbp), %rdi
	call	_ZNSt6vectorIiSaIiEED1Ev
	movq	-56(%rbp), %rax
	subq	%fs:40, %rax
	jne	.L71
	movq	%rbx, %rdi
.LEHB2:
	call	_Unwind_Resume@PLT
.LEHE2:
.L71:
	call	__stack_chk_fail@PLT
.L69:
	leaq	-40(%rbp), %rsp
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa 7, 8
	ret
	.cfi_endproc
.LFE4005:
	.globl	__gxx_personality_v0
	.section	.gcc_except_table,"a",@progbits
.LLSDA4005:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4005-.LLSDACSB4005
.LLSDACSB4005:
	.uleb128 .LEHB0-.LFB4005
	.uleb128 .LEHE0-.LEHB0
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB1-.LFB4005
	.uleb128 .LEHE1-.LEHB1
	.uleb128 .L70-.LFB4005
	.uleb128 0
	.uleb128 .LEHB2-.LFB4005
	.uleb128 .LEHE2-.LEHB2
	.uleb128 0
	.uleb128 0
.LLSDACSE4005:
	.text
	.size	_Z23run_gpu_2opt_restrictedPfS_PiS0_ii, .-_Z23run_gpu_2opt_restrictedPfS_PiS0_ii
	.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
.LC3:
	.long	2147483647
	.long	0
	.long	0
	.long	0
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC4:
	.long	-1598689907
	.long	1051772663
	.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:

#pragma once
#include <hip/hip_runtime.h>

#include <vector>

void run_gpu_2opt(float* x, float* y, int n);

void run_gpu_2opt_restricted(float* x, float* y, int* id, int* moves, int n, int allowed_moves);
	.text
	.file	"cuda_2opt.hip"
	.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 __hip_cuid_

	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
