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

#include "cuda_preprocessing.cuh"
#include "util.cuh"


__global__ void boruvka_smallest_kernel(int n, float* x, float* y, float* pi, int* component, float* component_best, int* component_best_i, int* component_best_j, int* component_lock) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    int j = blockIdx.y * blockDim.x;
    __shared__ float shared_x[64];
    __shared__ float shared_y[64];
    __shared__ float shared_pi[64];
    __shared__ float shared_best[64];
    __shared__ float shared_best_j[64];
    __shared__ float shared_component[64];
    __shared__ float shared_component_i[64];
    if (j+threadIdx.x < n) {
        shared_x[threadIdx.x] = x[j+threadIdx.x];
        shared_y[threadIdx.x] = y[j+threadIdx.x];
        shared_component[threadIdx.x] = component[j+threadIdx.x];
        shared_pi[threadIdx.x] = pi[j+threadIdx.x];
    }
    if (i < n)
        shared_component_i[threadIdx.x] = component[i];
    shared_best[threadIdx.x] = FLT_MAX;
    shared_best_j[threadIdx.x] = 0;
    __syncthreads();
    if (i >= n)
        return;
    float best = FLT_MAX;
    int best_j = -1;
    int component_i = component[i];
    if (!(i >= n)) {
        float pi_i = pi[i], x_i = x[i], y_i = y[i];
        if (j + blockDim.x > n || (i >= j && i < j + blockDim.x)) {
            for (int k = 0; k < 64; ++k) {
                if (j+k >= n || shared_component[k] == component_i || j+k == i)
                    continue;
                float d_ij = pi_i + shared_pi[k];
                d_ij += (x_i - shared_x[k]) * (x_i - shared_x[k]);
                d_ij += (y_i - shared_y[k]) * (y_i - shared_y[k]);
                if (d_ij < best) {
                    best = d_ij;
                    best_j = j+k;
                }
            }
        } else {
            for (int k = 0; k < 64; ++k) {
                if (shared_component[k] == component_i)
                    continue;
                float d_ij = pi_i + shared_pi[k];
                d_ij += (x_i - shared_x[k]) * (x_i - shared_x[k]);
                d_ij += (y_i - shared_y[k]) * (y_i - shared_y[k]);
                if (d_ij < best) {
                    best = d_ij;
                    best_j = j+k;
                }
            }
        }
    }
    shared_best[threadIdx.x] = best;
    shared_best_j[threadIdx.x] = best_j;
    //printf("%d best %d %.3f\n", i, best_j, best);
    __syncthreads();
    if (threadIdx.x == 0) {
        for (int k = 0; k < 64; ++k) {
            int tmp_i = blockIdx.x * blockDim.x + k;
            if (tmp_i >= n)
                continue;
            component_i = shared_component_i[k];
            best = shared_best[k];
            best_j = shared_best_j[k];
            if (best < component_best[component_i] || best == component_best[component_i]) {
                while (atomicExch(&component_lock[component_i], 1) != 0);
                if (best < component_best[component_i]) {
                    component_best[component_i] = best;
                    component_best_i[component_i] = tmp_i;
                    component_best_j[component_i] = best_j;
                } else if (abs(best - component_best[component_i]) < 0.0000001) {
                    int mi_c = min(tmp_i, best_j), ma_c = max(tmp_i, best_j);
                    int mi_o = min(component_best_i[component_i], component_best_j[component_i]), ma_o = max(component_best_i[component_i], component_best_j[component_i]);
                    if (mi_c < mi_o) {
                        component_best_i[component_i] = tmp_i;
                        component_best_j[component_i] = best_j;
                    } else if (mi_c == mi_o && ma_c < ma_o) {
                        component_best_i[component_i] = tmp_i;
                        component_best_j[component_i] = best_j;
                    }
                }
                component_lock[component_i] = 0;
                __threadfence();
            }
        }
    }
}

__global__ void boruvka_update_components(int n, 
    int* component, int* successor, float* component_best, int* component_best_i, int* component_best_j, int* component_lock, 
    int* degrees, float* L_T, int* components) {
    int i = blockIdx.x;
    int component_i = component[i];
    int component_j = component[component_best_j[component_i]];
    
    component[i] = successor[i];
    int vertex_ii = component_best_i[component_i];
    int vertex_ij = component_best_j[component_i];
    int vertex_jj = component_best_j[component_j];
    if (i == vertex_ii) {
        if (vertex_jj == i) {
            if (i < vertex_ij) {
                atomicSub(components, 1);
                atomicAdd(L_T, component_best[component_i]);
            }
        } else {
            atomicSub(components, 1);
            atomicAdd(L_T, component_best[component_i]);
            atomicAdd(&degrees[i], 1);
            atomicAdd(&degrees[vertex_ij], 1);
            //printf("e %d %d %.3f\n", min(i, vertex_ij), max(i, vertex_ij), component_best[component_i]);
        }
        __threadfence();
    }
}

__global__ void boruvka_remove_cycles(int n, 
    int* component, float* component_best, int* component_best_i, int* component_best_j, int* component_lock, 
    int* degrees, float* L_T, int* components) {
    int i = blockIdx.x;
    int component_i = component[i];
    int component_j = component[component_best_j[component_i]];

    int vertex_ii = component_best_i[component_i];
    int vertex_ij = component_best_j[component_i];
    int vertex_jj = component_best_j[component_j];
    if (i == vertex_ii) {
        if (vertex_jj == i) {
            if (i < vertex_ij) {
                component_best_j[component_i] = i;
            }
        }
    }
}

__global__ void boruvka_pointer_doubling(int n, 
    int* component, float* component_best, int* component_best_i, int* component_best_j, int* component_lock, 
    int* degrees, float* L_T, int* components) {
    int i = blockIdx.x;
    component[i] = component[component_best_j[component[i]]];
    component[i] = component[component_best_j[component[i]]];
}

__global__ void second_closest_find(int n, 
    float* x, float* y, float* pi, int* degrees, float* closest, int* closest_i, int* lock) {
    int i = blockIdx.x;
    if (degrees[i] != 1)
        return;
    int c_j1 = 0, c_j2 = 0;
    float d_j1 = FLT_MAX, d_j2 = FLT_MAX;
    for (int j = 0; j < n; ++j) {
        if (i == j)
            continue;
        float d_ij = pi[j] + pi[i];
        
        d_ij += (x[j] - x[i]) * (x[j] - x[i]);
        d_ij += (y[j] - y[i]) * (y[j] - y[i]);
        if (d_ij < d_j1 && d_ij < d_j2) {
            c_j2 = c_j1;
            d_j2 = d_j1;
            c_j1 = j;
            d_j1 = d_ij;
        } else if (d_ij < d_j2) {
            c_j2 = j;
            d_j2 = d_ij;
        }
    }
    //printf("i sec %d %d %.3f %d %.3f %.3f\n", i, c_j1, d_j1, c_j2, d_j2, closest[0]);
    if (d_j2 > closest[0]) {
        while (atomicExch(&lock[0], 1) != 0);
        if (d_j2 > closest[0]) {
            //printf("lait i sec %d %d %.3f %d %.3f %.3f\n", i, c_j1, d_j1, c_j2, d_j2, closest[0]);
            closest[0] = d_j2;
            closest_i[0] = c_j2;
            closest_i[1] = i;
        }
        lock[0] = 0;
        __threadfence();
    }
}

__global__ void second_closest_set(float* closest, int* closest_i, int* degrees, float* L_T) {
    degrees[closest_i[0]]++;
    degrees[closest_i[1]]++;
    //printf("s %d %d %.3f\n", min(closest_i[0], closest_i[1]), max(closest_i[0], closest_i[1]), closest[0]);
    L_T[0] += closest[0];
}


std::pair<float, std::vector<int>> gpu_boruvka_onetree(int n, float* Gx, float* Gy, float* Gpi) {
    int* Gcomponent = NULL;
    cudaMalloc((void**)&Gcomponent, n * sizeof(int));
    std::vector<int> super_init;
    for (int i = 0; i < n; ++i)
        super_init.push_back(i);
    cudaMemcpy(Gcomponent, super_init.data(), n * sizeof(int), cudaMemcpyHostToDevice);

    int* Gsuccessor = NULL;
    cudaMalloc((void**)&Gsuccessor, n * sizeof(int));
    cudaMemcpy(Gsuccessor, super_init.data(), n * sizeof(int), cudaMemcpyHostToDevice);

    int* Gvertex_lock = NULL;
    cudaMalloc((void**)&Gvertex_lock, n * sizeof(int));
    cudaMemset(Gvertex_lock, 0, n*sizeof(int));

    float* Gsmallest_add = NULL;
    cudaMalloc((void**)&Gsmallest_add, n * sizeof(float));

    int* Gsmallest_i = NULL;
    cudaMalloc((void**)&Gsmallest_i, n * sizeof(int));
    int* Gsmallest_j = NULL;
    cudaMalloc((void**)&Gsmallest_j, n * sizeof(int));

    int* Gdegrees = NULL;
    cudaMalloc((void**)&Gdegrees, n * sizeof(int));
    cudaMemset(Gdegrees, 0, n*sizeof(int));

    float* GL_T = NULL;
    cudaMalloc((void**)&GL_T, 1 * sizeof(float));
    cudaMemset(GL_T, 0, 1*sizeof(float));

    int* Gcomponents = NULL;
    cudaMalloc((void**)&Gcomponents, 1 * sizeof(int));
    int tmp = n;
    cudaMemcpy(Gcomponents, &tmp, 1 * sizeof(int), cudaMemcpyHostToDevice);

    int components = n;
    dim3 dimBlock(64, 1);
    dim3 dimGrid(divup(n, 64), divup(n, 64));

    std::vector<float> inf(n, std::numeric_limits<float>::max());
    while (components > 1) {
        cudaMemcpy(Gsmallest_add, inf.data(), n * sizeof(float), cudaMemcpyHostToDevice);
        boruvka_smallest_kernel<<<dimGrid, dimBlock>>>(n, Gx, Gy, Gpi, Gcomponent, Gsmallest_add, Gsmallest_i, Gsmallest_j, Gvertex_lock);
        CHECK(cudaGetLastError());
        cudaDeviceSynchronize();
        boruvka_remove_cycles<<<dim3(n, 1), dim3(1, 1)>>>(n, Gcomponent, Gsmallest_add, Gsmallest_i, Gsmallest_j, Gvertex_lock, Gdegrees, GL_T, Gcomponents);
        CHECK(cudaGetLastError());
        cudaDeviceSynchronize();
        int n_pd = components;
        while (n_pd > 0) {
            boruvka_pointer_doubling<<<dim3(n, 1), dim3(1, 1)>>>(n, Gsuccessor, Gsmallest_add, Gsmallest_i, Gsmallest_j, Gvertex_lock, Gdegrees, GL_T, Gcomponents);
            CHECK(cudaGetLastError());
            cudaDeviceSynchronize();
            n_pd /= 2;
        }
        boruvka_update_components<<<dim3(n, 1), dim3(1, 1)>>>(n, Gcomponent, Gsuccessor, Gsmallest_add, Gsmallest_i, Gsmallest_j, Gvertex_lock, Gdegrees, GL_T, Gcomponents);
        CHECK(cudaGetLastError());
        cudaDeviceSynchronize();
        cudaMemcpy(&components, Gcomponents, 1 * sizeof(int), cudaMemcpyDeviceToHost);
        cudaDeviceSynchronize();
    }
    float* Gclosest = NULL;
    cudaMalloc((void**)&Gclosest, 2 * sizeof(float));
    std::vector<float> minf(2, -std::numeric_limits<float>::max());
    cudaMemcpy(Gclosest, minf.data(), 2 * sizeof(float), cudaMemcpyHostToDevice);
    int* Gclosest_idx = NULL;
    cudaMalloc((void**)&Gclosest_idx, 2 * sizeof(int));
    second_closest_find<<<dim3(n, 1), dim3(1, 1)>>>(n, Gx, Gy, Gpi, Gdegrees, Gclosest, Gclosest_idx, Gvertex_lock);
    CHECK(cudaGetLastError());
    cudaDeviceSynchronize();
    second_closest_set<<<dim3(1, 1), dim3(1, 1)>>>(Gclosest, Gclosest_idx, Gdegrees, GL_T);
    CHECK(cudaGetLastError());
    cudaDeviceSynchronize();

    float return_length;
    cudaMemcpy(&return_length, GL_T, 1 * sizeof(float), cudaMemcpyDeviceToHost);
    cudaDeviceSynchronize();
    std::vector<int> d(n);
    cudaMemcpy(d.data(), Gdegrees, n * sizeof(float), cudaMemcpyDeviceToHost);
    cudaDeviceSynchronize();

    cudaFree(Gcomponent);
    cudaFree(Gsuccessor);
    cudaFree(Gvertex_lock);
    cudaFree(Gsmallest_add);
    cudaFree(Gsmallest_i);
    cudaFree(Gsmallest_j);
    cudaFree(Gdegrees);
    cudaFree(GL_T);
    cudaFree(Gcomponents);
    cudaFree(Gclosest);
    cudaFree(Gclosest_idx);
    return {return_length, d};
}

std::vector<float> gpu_subgradient_opt_alpha(float* x, float* y, int n) {
    float* Gpi = NULL;
    cudaMalloc((void**)&Gpi, n * sizeof(float));
    cudaMemset(Gpi, 0, n*sizeof(float));
    float* Gx = NULL;
    cudaMalloc((void**)&Gx, n * sizeof(float));
    cudaMemcpy(Gx, x, n * sizeof(float), cudaMemcpyHostToDevice);
    float* Gy = NULL;
    cudaMalloc((void**)&Gy, n * sizeof(float));
    cudaMemcpy(Gy, y, n * sizeof(float), cudaMemcpyHostToDevice);

    std::vector<float> pi(n, 0), best_pi(n, 0);
    auto [init_w, init_d] = gpu_boruvka_onetree(n, Gx, Gy, Gpi);
    float best_w = init_w;
    std::vector<int> last_v(n), v(n);
    bool is_tour = true;
    for (int i = 0; i < n; ++i) {
        last_v[i] = init_d[i] - 2;
        v[i] = last_v[i];
        is_tour &= (last_v[i] == 0);
    }
    bool initial_phase = true;
    int initial_period = 1000;
    int period = initial_period;
    for (float t = 1.f; t > 0; t /= 2.f, period /= 2) {
        for (int p = 1; t > 0 && p <= period; ++p) {
            for (int i = 0; i < n; ++i) {
                pi[i] += t * ( 0.7f * v[i] + 0.3f * last_v[i]);
                //std::cout << pi[i] << " ";
            }
            //std::cout << std::endl;
            cudaMemcpy(Gpi, pi.data(), n * sizeof(float), cudaMemcpyHostToDevice);
            cudaDeviceSynchronize();
            last_v = v;
            auto [w, d] = gpu_boruvka_onetree(n, Gx, Gy, Gpi);
            for (int i = 0; i < n; ++i)
                w -= 2 * pi[i];
            is_tour = true;
            for (int i = 0; i < n; ++i) {
                v[i] = d[i] - 2;
                is_tour &= (v[i] == 0);
            }
            //printf("upt %.3f %.3f %.3f %d %d\n", w, best_w, t, p, period, initial_period);
            if (w > best_w) {
                best_w = w;
                best_pi = pi;
                if (initial_phase)
                    t *= 2.f;
                if (p == period)
                    period *= 2;
            } else if (initial_phase && p > initial_period / 2) {
                initial_phase = false;
                p = 0;
                t = 0.75f * t;
            }
        }
    }
    cudaFree(Gpi);
    cudaFree(Gx);
    cudaFree(Gy);
    return best_pi;
}

	.file	"tmpxft_0031d104_00000000-6_cuda_preprocessing.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4046:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4046:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.align 2
	.type	_ZNSt15__new_allocatorIiE8allocateEmPKv.isra.0, @function
_ZNSt15__new_allocatorIiE8allocateEmPKv.isra.0:
.LFB5003:
	.cfi_startproc
	movabsq	$2305843009213693951, %rax
	cmpq	%rdi, %rax
	jnb	.L3
	pushq	%rax
	.cfi_def_cfa_offset 16
	movabsq	$4611686018427387903, %rax
	cmpq	%rdi, %rax
	jnb	.L4
	call	_ZSt28__throw_bad_array_new_lengthv@PLT
.L4:
	call	_ZSt17__throw_bad_allocv@PLT
.L3:
	.cfi_def_cfa_offset 8
	salq	$2, %rdi
	jmp	_Znwm@PLT
	.cfi_endproc
.LFE5003:
	.size	_ZNSt15__new_allocatorIiE8allocateEmPKv.isra.0, .-_ZNSt15__new_allocatorIiE8allocateEmPKv.isra.0
	.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:
.LFB5004:
	.cfi_startproc
	testl	%edi, %edi
	je	.L9
	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
.L9:
	.cfi_def_cfa_offset 8
	.cfi_restore 3
	.cfi_restore 6
	ret
	.cfi_endproc
.LFE5004:
	.size	_ZL5check9cudaErrorPKc.constprop.0, .-_ZL5check9cudaErrorPKc.constprop.0
	.globl	_Z61__device_stub__Z23boruvka_smallest_kerneliPfS_S_PiS_S0_S0_S0_iPfS_S_PiS_S0_S0_S0_
	.type	_Z61__device_stub__Z23boruvka_smallest_kerneliPfS_S_PiS_S0_S0_S0_iPfS_S_PiS_S0_S0_S0_, @function
_Z61__device_stub__Z23boruvka_smallest_kerneliPfS_S_PiS_S0_S0_S0_iPfS_S_PiS_S0_S0_S0_:
.LFB4068:
	.cfi_startproc
	endbr64
	subq	$232, %rsp
	.cfi_def_cfa_offset 240
	movq	240(%rsp), %rax
	movl	%edi, 76(%rsp)
	leaq	96(%rsp), %rdi
	movq	%rsi, 64(%rsp)
	leaq	108(%rsp), %rsi
	movq	%rax, 24(%rsp)
	movq	248(%rsp), %rax
	movq	%rdx, 56(%rsp)
	leaq	80(%rsp), %rdx
	movq	%rax, 16(%rsp)
	movq	256(%rsp), %rax
	movq	%rcx, 48(%rsp)
	leaq	88(%rsp), %rcx
	movq	%r8, 40(%rsp)
	movq	%r9, 32(%rsp)
	movq	%rax, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 216(%rsp)
	xorl	%eax, %eax
	leaq	76(%rsp), %rax
	movl	$1, 104(%rsp)
	movq	%rax, 144(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 208(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 96(%rsp)
	movq	%rax, 108(%rsp)
	movl	$1, 116(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L14
	pushq	88(%rsp)
	.cfi_def_cfa_offset 248
	leaq	_Z23boruvka_smallest_kerneliPfS_S_PiS_S0_S0_S0_(%rip), %rdi
	pushq	88(%rsp)
	.cfi_def_cfa_offset 256
	movq	124(%rsp), %rcx
	movl	132(%rsp), %r8d
	movq	112(%rsp), %rsi
	movl	120(%rsp), %edx
	leaq	160(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 248
	popq	%rdx
	.cfi_def_cfa_offset 240
.L14:
	movq	216(%rsp), %rax
	subq	%fs:40, %rax
	je	.L16
	call	__stack_chk_fail@PLT
.L16:
	addq	$232, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4068:
	.size	_Z61__device_stub__Z23boruvka_smallest_kerneliPfS_S_PiS_S0_S0_S0_iPfS_S_PiS_S0_S0_S0_, .-_Z61__device_stub__Z23boruvka_smallest_kerneliPfS_S_PiS_S0_S0_S0_iPfS_S_PiS_S0_S0_S0_
	.globl	_Z23boruvka_smallest_kerneliPfS_S_PiS_S0_S0_S0_
	.type	_Z23boruvka_smallest_kerneliPfS_S_PiS_S0_S0_S0_, @function
_Z23boruvka_smallest_kerneliPfS_S_PiS_S0_S0_S0_:
.LFB4069:
	.cfi_startproc
	endbr64
	jmp	_Z61__device_stub__Z23boruvka_smallest_kerneliPfS_S_PiS_S0_S0_S0_iPfS_S_PiS_S0_S0_S0_
	.cfi_endproc
.LFE4069:
	.size	_Z23boruvka_smallest_kerneliPfS_S_PiS_S0_S0_S0_, .-_Z23boruvka_smallest_kerneliPfS_S_PiS_S0_S0_S0_
	.globl	_Z63__device_stub__Z25boruvka_update_componentsiPiS_PfS_S_S_S_S0_S_iPiS_PfS_S_S_S_S0_S_
	.type	_Z63__device_stub__Z25boruvka_update_componentsiPiS_PfS_S_S_S_S0_S_iPiS_PfS_S_S_S_S0_S_, @function
_Z63__device_stub__Z25boruvka_update_componentsiPiS_PfS_S_S_S_S0_S_iPiS_PfS_S_S_S_S0_S_:
.LFB4070:
	.cfi_startproc
	endbr64
	subq	$248, %rsp
	.cfi_def_cfa_offset 256
	movq	256(%rsp), %rax
	movl	%edi, 76(%rsp)
	leaq	104(%rsp), %rdi
	movq	%rsi, 64(%rsp)
	leaq	116(%rsp), %rsi
	movq	%rax, 24(%rsp)
	movq	264(%rsp), %rax
	movq	%rdx, 56(%rsp)
	leaq	88(%rsp), %rdx
	movq	%rax, 16(%rsp)
	movq	272(%rsp), %rax
	movq	%rcx, 48(%rsp)
	leaq	96(%rsp), %rcx
	movq	%rax, 8(%rsp)
	movq	280(%rsp), %rax
	movq	%r8, 40(%rsp)
	movq	%r9, 32(%rsp)
	movq	%rax, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 232(%rsp)
	xorl	%eax, %eax
	leaq	76(%rsp), %rax
	movl	$1, 112(%rsp)
	movq	%rax, 152(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 208(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 216(%rsp)
	movq	%rsp, %rax
	movq	%rax, 224(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 104(%rsp)
	movq	%rax, 116(%rsp)
	movl	$1, 124(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L19
	pushq	96(%rsp)
	.cfi_def_cfa_offset 264
	leaq	_Z25boruvka_update_componentsiPiS_PfS_S_S_S_S0_S_(%rip), %rdi
	pushq	96(%rsp)
	.cfi_def_cfa_offset 272
	movq	132(%rsp), %rcx
	movl	140(%rsp), %r8d
	movq	120(%rsp), %rsi
	movl	128(%rsp), %edx
	leaq	168(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 264
	popq	%rdx
	.cfi_def_cfa_offset 256
.L19:
	movq	232(%rsp), %rax
	subq	%fs:40, %rax
	je	.L21
	call	__stack_chk_fail@PLT
.L21:
	addq	$248, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4070:
	.size	_Z63__device_stub__Z25boruvka_update_componentsiPiS_PfS_S_S_S_S0_S_iPiS_PfS_S_S_S_S0_S_, .-_Z63__device_stub__Z25boruvka_update_componentsiPiS_PfS_S_S_S_S0_S_iPiS_PfS_S_S_S_S0_S_
	.globl	_Z25boruvka_update_componentsiPiS_PfS_S_S_S_S0_S_
	.type	_Z25boruvka_update_componentsiPiS_PfS_S_S_S_S0_S_, @function
_Z25boruvka_update_componentsiPiS_PfS_S_S_S_S0_S_:
.LFB4071:
	.cfi_startproc
	endbr64
	jmp	_Z63__device_stub__Z25boruvka_update_componentsiPiS_PfS_S_S_S_S0_S_iPiS_PfS_S_S_S_S0_S_
	.cfi_endproc
.LFE4071:
	.size	_Z25boruvka_update_componentsiPiS_PfS_S_S_S_S0_S_, .-_Z25boruvka_update_componentsiPiS_PfS_S_S_S_S0_S_
	.globl	_Z57__device_stub__Z21boruvka_remove_cyclesiPiPfS_S_S_S_S0_S_iPiPfS_S_S_S_S0_S_
	.type	_Z57__device_stub__Z21boruvka_remove_cyclesiPiPfS_S_S_S_S0_S_iPiPfS_S_S_S_S0_S_, @function
_Z57__device_stub__Z21boruvka_remove_cyclesiPiPfS_S_S_S_S0_S_iPiPfS_S_S_S_S0_S_:
.LFB4072:
	.cfi_startproc
	endbr64
	subq	$232, %rsp
	.cfi_def_cfa_offset 240
	movq	240(%rsp), %rax
	movl	%edi, 76(%rsp)
	leaq	96(%rsp), %rdi
	movq	%rsi, 64(%rsp)
	leaq	108(%rsp), %rsi
	movq	%rax, 24(%rsp)
	movq	248(%rsp), %rax
	movq	%rdx, 56(%rsp)
	leaq	80(%rsp), %rdx
	movq	%rax, 16(%rsp)
	movq	256(%rsp), %rax
	movq	%rcx, 48(%rsp)
	leaq	88(%rsp), %rcx
	movq	%r8, 40(%rsp)
	movq	%r9, 32(%rsp)
	movq	%rax, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 216(%rsp)
	xorl	%eax, %eax
	leaq	76(%rsp), %rax
	movl	$1, 104(%rsp)
	movq	%rax, 144(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 208(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 96(%rsp)
	movq	%rax, 108(%rsp)
	movl	$1, 116(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L24
	pushq	88(%rsp)
	.cfi_def_cfa_offset 248
	leaq	_Z21boruvka_remove_cyclesiPiPfS_S_S_S_S0_S_(%rip), %rdi
	pushq	88(%rsp)
	.cfi_def_cfa_offset 256
	movq	124(%rsp), %rcx
	movl	132(%rsp), %r8d
	movq	112(%rsp), %rsi
	movl	120(%rsp), %edx
	leaq	160(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 248
	popq	%rdx
	.cfi_def_cfa_offset 240
.L24:
	movq	216(%rsp), %rax
	subq	%fs:40, %rax
	je	.L26
	call	__stack_chk_fail@PLT
.L26:
	addq	$232, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4072:
	.size	_Z57__device_stub__Z21boruvka_remove_cyclesiPiPfS_S_S_S_S0_S_iPiPfS_S_S_S_S0_S_, .-_Z57__device_stub__Z21boruvka_remove_cyclesiPiPfS_S_S_S_S0_S_iPiPfS_S_S_S_S0_S_
	.globl	_Z21boruvka_remove_cyclesiPiPfS_S_S_S_S0_S_
	.type	_Z21boruvka_remove_cyclesiPiPfS_S_S_S_S0_S_, @function
_Z21boruvka_remove_cyclesiPiPfS_S_S_S_S0_S_:
.LFB4073:
	.cfi_startproc
	endbr64
	jmp	_Z57__device_stub__Z21boruvka_remove_cyclesiPiPfS_S_S_S_S0_S_iPiPfS_S_S_S_S0_S_
	.cfi_endproc
.LFE4073:
	.size	_Z21boruvka_remove_cyclesiPiPfS_S_S_S_S0_S_, .-_Z21boruvka_remove_cyclesiPiPfS_S_S_S_S0_S_
	.globl	_Z60__device_stub__Z24boruvka_pointer_doublingiPiPfS_S_S_S_S0_S_iPiPfS_S_S_S_S0_S_
	.type	_Z60__device_stub__Z24boruvka_pointer_doublingiPiPfS_S_S_S_S0_S_iPiPfS_S_S_S_S0_S_, @function
_Z60__device_stub__Z24boruvka_pointer_doublingiPiPfS_S_S_S_S0_S_iPiPfS_S_S_S_S0_S_:
.LFB4074:
	.cfi_startproc
	endbr64
	subq	$232, %rsp
	.cfi_def_cfa_offset 240
	movq	240(%rsp), %rax
	movl	%edi, 76(%rsp)
	leaq	96(%rsp), %rdi
	movq	%rsi, 64(%rsp)
	leaq	108(%rsp), %rsi
	movq	%rax, 24(%rsp)
	movq	248(%rsp), %rax
	movq	%rdx, 56(%rsp)
	leaq	80(%rsp), %rdx
	movq	%rax, 16(%rsp)
	movq	256(%rsp), %rax
	movq	%rcx, 48(%rsp)
	leaq	88(%rsp), %rcx
	movq	%r8, 40(%rsp)
	movq	%r9, 32(%rsp)
	movq	%rax, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 216(%rsp)
	xorl	%eax, %eax
	leaq	76(%rsp), %rax
	movl	$1, 104(%rsp)
	movq	%rax, 144(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 208(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 96(%rsp)
	movq	%rax, 108(%rsp)
	movl	$1, 116(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L29
	pushq	88(%rsp)
	.cfi_def_cfa_offset 248
	leaq	_Z24boruvka_pointer_doublingiPiPfS_S_S_S_S0_S_(%rip), %rdi
	pushq	88(%rsp)
	.cfi_def_cfa_offset 256
	movq	124(%rsp), %rcx
	movl	132(%rsp), %r8d
	movq	112(%rsp), %rsi
	movl	120(%rsp), %edx
	leaq	160(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 248
	popq	%rdx
	.cfi_def_cfa_offset 240
.L29:
	movq	216(%rsp), %rax
	subq	%fs:40, %rax
	je	.L31
	call	__stack_chk_fail@PLT
.L31:
	addq	$232, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4074:
	.size	_Z60__device_stub__Z24boruvka_pointer_doublingiPiPfS_S_S_S_S0_S_iPiPfS_S_S_S_S0_S_, .-_Z60__device_stub__Z24boruvka_pointer_doublingiPiPfS_S_S_S_S0_S_iPiPfS_S_S_S_S0_S_
	.globl	_Z24boruvka_pointer_doublingiPiPfS_S_S_S_S0_S_
	.type	_Z24boruvka_pointer_doublingiPiPfS_S_S_S_S0_S_, @function
_Z24boruvka_pointer_doublingiPiPfS_S_S_S_S0_S_:
.LFB4075:
	.cfi_startproc
	endbr64
	jmp	_Z60__device_stub__Z24boruvka_pointer_doublingiPiPfS_S_S_S_S0_S_iPiPfS_S_S_S_S0_S_
	.cfi_endproc
.LFE4075:
	.size	_Z24boruvka_pointer_doublingiPiPfS_S_S_S_S0_S_, .-_Z24boruvka_pointer_doublingiPiPfS_S_S_S_S0_S_
	.globl	_Z54__device_stub__Z19second_closest_findiPfS_S_PiS_S0_S0_iPfS_S_PiS_S0_S0_
	.type	_Z54__device_stub__Z19second_closest_findiPfS_S_PiS_S0_S0_iPfS_S_PiS_S0_S0_, @function
_Z54__device_stub__Z19second_closest_findiPfS_S_PiS_S0_S0_iPfS_S_PiS_S0_S0_:
.LFB4076:
	.cfi_startproc
	endbr64
	subq	$216, %rsp
	.cfi_def_cfa_offset 224
	movq	224(%rsp), %rax
	movl	%edi, 60(%rsp)
	leaq	88(%rsp), %rdi
	movq	%rsi, 48(%rsp)
	leaq	100(%rsp), %rsi
	movq	%rax, 8(%rsp)
	movq	232(%rsp), %rax
	movq	%rdx, 40(%rsp)
	leaq	72(%rsp), %rdx
	movq	%rcx, 32(%rsp)
	leaq	80(%rsp), %rcx
	movq	%r8, 24(%rsp)
	movq	%r9, 16(%rsp)
	movq	%rax, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 200(%rsp)
	xorl	%eax, %eax
	leaq	60(%rsp), %rax
	movl	$1, 96(%rsp)
	movq	%rax, 136(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 184(%rsp)
	movq	%rsp, %rax
	movq	%rax, 192(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 88(%rsp)
	movq	%rax, 100(%rsp)
	movl	$1, 108(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L34
	pushq	80(%rsp)
	.cfi_def_cfa_offset 232
	leaq	_Z19second_closest_findiPfS_S_PiS_S0_S0_(%rip), %rdi
	pushq	80(%rsp)
	.cfi_def_cfa_offset 240
	movq	116(%rsp), %rcx
	movl	124(%rsp), %r8d
	movq	104(%rsp), %rsi
	movl	112(%rsp), %edx
	leaq	152(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 232
	popq	%rdx
	.cfi_def_cfa_offset 224
.L34:
	movq	200(%rsp), %rax
	subq	%fs:40, %rax
	je	.L36
	call	__stack_chk_fail@PLT
.L36:
	addq	$216, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4076:
	.size	_Z54__device_stub__Z19second_closest_findiPfS_S_PiS_S0_S0_iPfS_S_PiS_S0_S0_, .-_Z54__device_stub__Z19second_closest_findiPfS_S_PiS_S0_S0_iPfS_S_PiS_S0_S0_
	.globl	_Z19second_closest_findiPfS_S_PiS_S0_S0_
	.type	_Z19second_closest_findiPfS_S_PiS_S0_S0_, @function
_Z19second_closest_findiPfS_S_PiS_S0_S0_:
.LFB4077:
	.cfi_startproc
	endbr64
	jmp	_Z54__device_stub__Z19second_closest_findiPfS_S_PiS_S0_S0_iPfS_S_PiS_S0_S0_
	.cfi_endproc
.LFE4077:
	.size	_Z19second_closest_findiPfS_S_PiS_S0_S0_, .-_Z19second_closest_findiPfS_S_PiS_S0_S0_
	.globl	_Z45__device_stub__Z18second_closest_setPfPiS0_S_PfPiS0_S_
	.type	_Z45__device_stub__Z18second_closest_setPfPiS0_S_PfPiS0_S_, @function
_Z45__device_stub__Z18second_closest_setPfPiS0_S_PfPiS0_S_:
.LFB4078:
	.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
	movq	%rcx, (%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)
	movq	%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	.L39
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z18second_closest_setPfPiS0_S_(%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
.L39:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L41
	call	__stack_chk_fail@PLT
.L41:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4078:
	.size	_Z45__device_stub__Z18second_closest_setPfPiS0_S_PfPiS0_S_, .-_Z45__device_stub__Z18second_closest_setPfPiS0_S_PfPiS0_S_
	.globl	_Z18second_closest_setPfPiS0_S_
	.type	_Z18second_closest_setPfPiS0_S_, @function
_Z18second_closest_setPfPiS0_S_:
.LFB4079:
	.cfi_startproc
	endbr64
	jmp	_Z45__device_stub__Z18second_closest_setPfPiS0_S_PfPiS0_S_
	.cfi_endproc
.LFE4079:
	.size	_Z18second_closest_setPfPiS0_S_, .-_Z18second_closest_setPfPiS0_S_
	.section	.rodata.str1.1
.LC3:
	.string	"_Z18second_closest_setPfPiS0_S_"
.LC4:
	.string	"_Z19second_closest_findiPfS_S_PiS_S0_S0_"
.LC5:
	.string	"_Z24boruvka_pointer_doublingiPiPfS_S_S_S_S0_S_"
.LC6:
	.string	"_Z21boruvka_remove_cyclesiPiPfS_S_S_S_S0_S_"
.LC7:
	.string	"_Z25boruvka_update_componentsiPiS_PfS_S_S_S_S0_S_"
.LC8:
	.string	"_Z23boruvka_smallest_kerneliPfS_S_PiS_S0_S0_S0_"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4081:
	.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	.LC3(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z18second_closest_setPfPiS0_S_(%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	.LC4(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z19second_closest_findiPfS_S_PiS_S0_S0_(%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	.LC5(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z24boruvka_pointer_doublingiPiPfS_S_S_S_S0_S_(%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	.LC6(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z21boruvka_remove_cyclesiPiPfS_S_S_S_S0_S_(%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	_Z25boruvka_update_componentsiPiS_PfS_S_S_S_S0_S_(%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	_Z23boruvka_smallest_kerneliPfS_S_PiS_S0_S0_S0_(%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
.LFE4081:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.text._ZNSt6vectorIiSaIiEED2Ev,"axG",@progbits,_ZNSt6vectorIiSaIiEED5Ev,comdat
	.align 2
	.weak	_ZNSt6vectorIiSaIiEED2Ev
	.type	_ZNSt6vectorIiSaIiEED2Ev, @function
_ZNSt6vectorIiSaIiEED2Ev:
.LFB4394:
	.cfi_startproc
	endbr64
	movq	(%rdi), %rax
	testq	%rax, %rax
	je	.L46
	movq	16(%rdi), %rsi
	movq	%rax, %rdi
	subq	%rax, %rsi
	jmp	_ZdlPvm@PLT
.L46:
	ret
	.cfi_endproc
.LFE4394:
	.size	_ZNSt6vectorIiSaIiEED2Ev, .-_ZNSt6vectorIiSaIiEED2Ev
	.weak	_ZNSt6vectorIiSaIiEED1Ev
	.set	_ZNSt6vectorIiSaIiEED1Ev,_ZNSt6vectorIiSaIiEED2Ev
	.section	.rodata._ZNSt6vectorIfSaIfEEC2EmRKfRKS0_.str1.1,"aMS",@progbits,1
.LC9:
	.string	"cannot create std::vector larger than max_size()"
	.section	.text._ZNSt6vectorIfSaIfEEC2EmRKfRKS0_,"axG",@progbits,_ZNSt6vectorIfSaIfEEC5EmRKfRKS0_,comdat
	.align 2
	.weak	_ZNSt6vectorIfSaIfEEC2EmRKfRKS0_
	.type	_ZNSt6vectorIfSaIfEEC2EmRKfRKS0_, @function
_ZNSt6vectorIfSaIfEEC2EmRKfRKS0_:
.LFB4406:
	.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	.L49
	leaq	.LC9(%rip), %rdi
	call	_ZSt20__throw_length_errorPKc@PLT
.L49:
	xorl	%eax, %eax
	movq	%rdi, %rbx
	movq	%rsi, %rbp
	movq	%rdx, %r12
	movq	%rax, (%rdi)
	movq	%rax, 8(%rdi)
	movq	%rax, 16(%rdi)
	xorl	%eax, %eax
	testq	%rsi, %rsi
	je	.L50
	movq	%rsi, %rdi
	call	_ZNSt15__new_allocatorIiE8allocateEmPKv.isra.0
.L50:
	leaq	(%rax,%rbp,4), %rdx
	movq	%rax, (%rbx)
	movq	%rdx, 16(%rbx)
	testq	%rbp, %rbp
	je	.L51
	movss	(%r12), %xmm0
.L52:
	cmpq	%rax, %rdx
	je	.L51
	movss	%xmm0, (%rax)
	addq	$4, %rax
	jmp	.L52
.L51:
	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
.LFE4406:
	.size	_ZNSt6vectorIfSaIfEEC2EmRKfRKS0_, .-_ZNSt6vectorIfSaIfEEC2EmRKfRKS0_
	.weak	_ZNSt6vectorIfSaIfEEC1EmRKfRKS0_
	.set	_ZNSt6vectorIfSaIfEEC1EmRKfRKS0_,_ZNSt6vectorIfSaIfEEC2EmRKfRKS0_
	.section	.text._ZNSt6vectorIfSaIfEED2Ev,"axG",@progbits,_ZNSt6vectorIfSaIfEED5Ev,comdat
	.align 2
	.weak	_ZNSt6vectorIfSaIfEED2Ev
	.type	_ZNSt6vectorIfSaIfEED2Ev, @function
_ZNSt6vectorIfSaIfEED2Ev:
.LFB4409:
	.cfi_startproc
	endbr64
	movq	(%rdi), %rax
	testq	%rax, %rax
	je	.L60
	movq	16(%rdi), %rsi
	movq	%rax, %rdi
	subq	%rax, %rsi
	jmp	_ZdlPvm@PLT
.L60:
	ret
	.cfi_endproc
.LFE4409:
	.size	_ZNSt6vectorIfSaIfEED2Ev, .-_ZNSt6vectorIfSaIfEED2Ev
	.weak	_ZNSt6vectorIfSaIfEED1Ev
	.set	_ZNSt6vectorIfSaIfEED1Ev,_ZNSt6vectorIfSaIfEED2Ev
	.section	.rodata._ZNSt6vectorIiSaIiEE17_M_realloc_insertIJRKiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_.str1.1,"aMS",@progbits,1
.LC10:
	.string	"vector::_M_realloc_insert"
	.section	.text._ZNSt6vectorIiSaIiEE17_M_realloc_insertIJRKiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_,"axG",@progbits,_ZNSt6vectorIiSaIiEE17_M_realloc_insertIJRKiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_,comdat
	.align 2
	.weak	_ZNSt6vectorIiSaIiEE17_M_realloc_insertIJRKiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_
	.type	_ZNSt6vectorIiSaIiEE17_M_realloc_insertIJRKiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_, @function
_ZNSt6vectorIiSaIiEE17_M_realloc_insertIJRKiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_:
.LFB4588:
	.cfi_startproc
	endbr64
	movabsq	$2305843009213693951, %rcx
	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
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	movq	8(%rdi), %rax
	movq	(%rdi), %r14
	movq	%rax, %rbx
	movq	%rax, (%rsp)
	subq	%r14, %rbx
	sarq	$2, %rbx
	cmpq	%rcx, %rbx
	jne	.L63
	leaq	.LC10(%rip), %rdi
	call	_ZSt20__throw_length_errorPKc@PLT
.L63:
	cmpq	%r14, (%rsp)
	movl	$1, %eax
	movq	%rsi, %r15
	movq	%rdi, %r13
	cmovne	%rbx, %rax
	movq	%rsi, %rbp
	addq	%rax, %rbx
	setc	%al
	subq	%r14, %r15
	movzbl	%al, %eax
	testq	%rax, %rax
	jne	.L73
	xorl	%r12d, %r12d
	testq	%rbx, %rbx
	je	.L68
	movabsq	$2305843009213693951, %rax
	cmpq	%rax, %rbx
	cmova	%rax, %rbx
	jmp	.L67
.L73:
	movq	%rcx, %rbx
.L67:
	movq	%rbx, %rdi
	movq	%rdx, 8(%rsp)
	call	_ZNSt15__new_allocatorIiE8allocateEmPKv.isra.0
	movq	8(%rsp), %rdx
	movq	%rax, %r12
.L68:
	movl	(%rdx), %eax
	movl	%eax, (%r12,%r15)
	testq	%r15, %r15
	jle	.L69
	movq	%r15, %rdx
	movq	%r14, %rsi
	movq	%r12, %rdi
	call	memmove@PLT
.L69:
	movq	(%rsp), %rax
	leaq	4(%r12,%r15), %rdx
	subq	%rbp, %rax
	testq	%rax, %rax
	jle	.L70
	movq	%rdx, %rdi
	movq	%rbp, %rsi
	movq	%rax, %rcx
	rep movsb
.L70:
	leaq	(%rdx,%rax), %rbp
	testq	%r14, %r14
	je	.L71
	movq	16(%r13), %rsi
	movq	%r14, %rdi
	subq	%r14, %rsi
	call	_ZdlPvm@PLT
.L71:
	leaq	(%r12,%rbx,4), %rax
	movq	%r12, 0(%r13)
	movq	%rbp, 8(%r13)
	movq	%rax, 16(%r13)
	addq	$24, %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
.LFE4588:
	.size	_ZNSt6vectorIiSaIiEE17_M_realloc_insertIJRKiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_, .-_ZNSt6vectorIiSaIiEE17_M_realloc_insertIJRKiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_
	.section	.text._ZNSt12_Vector_baseIiSaIiEE17_M_create_storageEm,"axG",@progbits,_ZNSt12_Vector_baseIiSaIiEE17_M_create_storageEm,comdat
	.align 2
	.weak	_ZNSt12_Vector_baseIiSaIiEE17_M_create_storageEm
	.type	_ZNSt12_Vector_baseIiSaIiEE17_M_create_storageEm, @function
_ZNSt12_Vector_baseIiSaIiEE17_M_create_storageEm:
.LFB4777:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rsi, %rbp
	xorl	%eax, %eax
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rdi, %rbx
	pushq	%rdx
	.cfi_def_cfa_offset 32
	testq	%rsi, %rsi
	je	.L81
	movq	%rsi, %rdi
	call	_ZNSt15__new_allocatorIiE8allocateEmPKv.isra.0
.L81:
	movq	%rax, (%rbx)
	movq	%rax, 8(%rbx)
	leaq	(%rax,%rbp,4), %rax
	movq	%rax, 16(%rbx)
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4777:
	.size	_ZNSt12_Vector_baseIiSaIiEE17_M_create_storageEm, .-_ZNSt12_Vector_baseIiSaIiEE17_M_create_storageEm
	.section	.text._ZNSt6vectorIiSaIiEEC2EmRKS0_,"axG",@progbits,_ZNSt6vectorIiSaIiEEC5EmRKS0_,comdat
	.align 2
	.weak	_ZNSt6vectorIiSaIiEEC2EmRKS0_
	.type	_ZNSt6vectorIiSaIiEEC2EmRKS0_, @function
_ZNSt6vectorIiSaIiEEC2EmRKS0_:
.LFB4416:
	.cfi_startproc
	endbr64
	movabsq	$2305843009213693951, %rax
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	pushq	%r10
	.cfi_def_cfa_offset 32
	cmpq	%rsi, %rax
	jnb	.L87
	leaq	.LC9(%rip), %rdi
	call	_ZSt20__throw_length_errorPKc@PLT
.L87:
	xorl	%r9d, %r9d
	movq	%rdi, %rbp
	movq	%rsi, %rbx
	movq	%r9, (%rdi)
	movq	%r9, 8(%rdi)
	movq	%r9, 16(%rdi)
	call	_ZNSt12_Vector_baseIiSaIiEE17_M_create_storageEm
	movq	0(%rbp), %rax
	testq	%rbx, %rbx
	je	.L88
	xorl	%r8d, %r8d
	leaq	4(%rax), %rdx
	movl	%r8d, (%rax)
	decq	%rbx
	je	.L91
	leaq	(%rdx,%rbx,4), %rcx
	movq	%rdx, %rsi
.L89:
	cmpq	%rsi, %rcx
	je	.L96
	xorl	%edi, %edi
	addq	$4, %rsi
	movl	%edi, -4(%rsi)
	jmp	.L89
.L96:
	subq	%rax, %rcx
	leaq	-4(%rdx,%rcx), %rax
	jmp	.L88
.L91:
	movq	%rdx, %rax
.L88:
	movq	%rax, 8(%rbp)
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4416:
	.size	_ZNSt6vectorIiSaIiEEC2EmRKS0_, .-_ZNSt6vectorIiSaIiEEC2EmRKS0_
	.weak	_ZNSt6vectorIiSaIiEEC1EmRKS0_
	.set	_ZNSt6vectorIiSaIiEEC1EmRKS0_,_ZNSt6vectorIiSaIiEEC2EmRKS0_
	.section	.text._ZNSt11__copy_moveILb0ELb1ESt26random_access_iterator_tagE8__copy_mIKiiEEPT0_PT_S7_S5_,"axG",@progbits,_ZNSt11__copy_moveILb0ELb1ESt26random_access_iterator_tagE8__copy_mIKiiEEPT0_PT_S7_S5_,comdat
	.weak	_ZNSt11__copy_moveILb0ELb1ESt26random_access_iterator_tagE8__copy_mIKiiEEPT0_PT_S7_S5_
	.type	_ZNSt11__copy_moveILb0ELb1ESt26random_access_iterator_tagE8__copy_mIKiiEEPT0_PT_S7_S5_, @function
_ZNSt11__copy_moveILb0ELb1ESt26random_access_iterator_tagE8__copy_mIKiiEEPT0_PT_S7_S5_:
.LFB4971:
	.cfi_startproc
	endbr64
	subq	%rdi, %rsi
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rdx, %rcx
	movq	%rsi, %rbx
	cmpq	$4, %rsi
	jle	.L98
	movq	%rsi, %rdx
	movq	%rdi, %rsi
	movq	%rcx, %rdi
	call	memmove@PLT
	movq	%rax, %rcx
	jmp	.L99
.L98:
	jne	.L99
	movl	(%rdi), %eax
	movl	%eax, (%rdx)
.L99:
	leaq	(%rcx,%rbx), %rax
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4971:
	.size	_ZNSt11__copy_moveILb0ELb1ESt26random_access_iterator_tagE8__copy_mIKiiEEPT0_PT_S7_S5_, .-_ZNSt11__copy_moveILb0ELb1ESt26random_access_iterator_tagE8__copy_mIKiiEEPT0_PT_S7_S5_
	.text
	.globl	_Z19gpu_boruvka_onetreeiPfS_S_
	.type	_Z19gpu_boruvka_onetreeiPfS_S_, @function
_Z19gpu_boruvka_onetreeiPfS_S_:
.LFB4009:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4009
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register 6
	pushq	%r15
	pushq	%r14
	.cfi_offset 15, -24
	.cfi_offset 14, -32
	movslq	%esi, %r14
	leaq	-80(%rbp), %r15
	pushq	%r13
	.cfi_offset 13, -40
	movq	%r8, %r13
	pushq	%r12
	.cfi_offset 12, -48
	xorl	%r12d, %r12d
	pushq	%rbx
	.cfi_offset 3, -56
	movq	%r14, %rbx
	subq	$312, %rsp
	movq	%rdi, -320(%rbp)
	leaq	-280(%rbp), %rdi
	movq	%rcx, -344(%rbp)
	xorl	%ecx, %ecx
	movq	%rdx, -336(%rbp)
	movq	%fs:40, %rax
	movq	%rax, -56(%rbp)
	xorl	%eax, %eax
	leaq	0(,%r14,4), %rax
	movq	%rcx, -280(%rbp)
	movq	%rax, %rsi
	movq	%rax, -312(%rbp)
.LEHB0:
	call	cudaMalloc@PLT
.LEHE0:
	xorl	%esi, %esi
	movq	%rsi, -152(%rbp)
	movq	%rsi, -144(%rbp)
	movq	%rsi, -136(%rbp)
.L102:
	movl	%r12d, -80(%rbp)
	cmpl	%r12d, %ebx
	jle	.L146
	movq	-144(%rbp), %rsi
	cmpq	-136(%rbp), %rsi
	je	.L103
	movl	%r12d, (%rsi)
	addq	$4, %rsi
	movq	%rsi, -144(%rbp)
	jmp	.L104
.L103:
	leaq	-152(%rbp), %rdi
	movq	%r15, %rdx
.LEHB1:
	call	_ZNSt6vectorIiSaIiEE17_M_realloc_insertIJRKiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_
.L104:
	incl	%r12d
	jmp	.L102
.L146:
	movq	-152(%rbp), %r12
	movq	-312(%rbp), %rdx
	movl	$1, %ecx
	movq	-280(%rbp), %rdi
	movq	%r12, %rsi
	call	cudaMemcpy@PLT
	movq	-312(%rbp), %rsi
	xorl	%edx, %edx
	leaq	-272(%rbp), %rdi
	movq	%rdx, -272(%rbp)
	call	cudaMalloc@PLT
	movq	-312(%rbp), %rdx
	movq	-272(%rbp), %rdi
	movl	$1, %ecx
	movq	%r12, %rsi
	call	cudaMemcpy@PLT
	movq	-312(%rbp), %rsi
	xorl	%eax, %eax
	leaq	-264(%rbp), %rdi
	movq	%rax, -264(%rbp)
	call	cudaMalloc@PLT
	movq	-312(%rbp), %rdx
	movq	-264(%rbp), %rdi
	xorl	%esi, %esi
	call	cudaMemset@PLT
	movq	-312(%rbp), %rsi
	xorl	%eax, %eax
	leaq	-256(%rbp), %rdi
	movq	%rax, -256(%rbp)
	call	cudaMalloc@PLT
	movq	-312(%rbp), %rsi
	xorl	%eax, %eax
	leaq	-248(%rbp), %rdi
	movq	%rax, -248(%rbp)
	call	cudaMalloc@PLT
	movq	-312(%rbp), %rsi
	xorl	%eax, %eax
	leaq	-240(%rbp), %rdi
	movq	%rax, -240(%rbp)
	call	cudaMalloc@PLT
	movq	-312(%rbp), %rsi
	xorl	%eax, %eax
	leaq	-232(%rbp), %rdi
	movq	%rax, -232(%rbp)
	call	cudaMalloc@PLT
	movq	-312(%rbp), %rdx
	movq	-232(%rbp), %rdi
	xorl	%esi, %esi
	call	cudaMemset@PLT
	xorl	%r15d, %r15d
	leaq	-224(%rbp), %rdi
	movl	$4, %esi
	movq	%r15, -224(%rbp)
	call	cudaMalloc@PLT
	movq	-224(%rbp), %rdi
	movl	$4, %edx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	xorl	%r12d, %r12d
	leaq	-216(%rbp), %rdi
	movl	$4, %esi
	movq	%r12, -216(%rbp)
	call	cudaMalloc@PLT
	movl	$1, %ecx
	movl	$4, %edx
	movq	-216(%rbp), %rdi
	leaq	-288(%rbp), %rsi
	movl	%ebx, -288(%rbp)
	call	cudaMemcpy@PLT
	movl	$67108865, %eax
	movl	$64, %ecx
	leaq	-128(%rbp), %rdi
	movq	%r14, %rsi
	salq	$6, %rax
	movl	%ebx, -284(%rbp)
	movq	%rax, -188(%rbp)
	leal	63(%rbx), %eax
	cltd
	movl	$0x7f7fffff, -80(%rbp)
	idivl	%ecx
	movl	$1, -180(%rbp)
	movl	$1, -168(%rbp)
	leaq	-104(%rbp), %rcx
	movl	%eax, -176(%rbp)
	movl	%eax, -172(%rbp)
	leaq	-80(%rbp), %rax
	movq	%rax, %rdx
	movq	%rax, -328(%rbp)
	call	_ZNSt6vectorIfSaIfEEC1EmRKfRKS0_
.LEHE1:
	movabsq	$4294967297, %r12
.L106:
	cmpl	$1, -284(%rbp)
	jle	.L147
	movq	-312(%rbp), %rdx
	movq	-128(%rbp), %rsi
	movl	$1, %ecx
	movq	-256(%rbp), %rdi
.LEHB2:
	call	cudaMemcpy@PLT
	movl	-180(%rbp), %ecx
	movq	-188(%rbp), %rdx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	-176(%rbp), %rdi
	movl	-168(%rbp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	je	.L107
.L109:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	call	_ZL5check9cudaErrorPKc.constprop.0
	jmp	.L148
.L107:
	pushq	%r11
	movq	-256(%rbp), %r9
	movq	%r13, %rcx
	movl	%ebx, %edi
	pushq	-264(%rbp)
	movq	-344(%rbp), %rdx
	movq	-336(%rbp), %rsi
	pushq	-240(%rbp)
	pushq	-248(%rbp)
	movq	-280(%rbp), %r8
	.cfi_escape 0x2e,0x20
	call	_Z61__device_stub__Z23boruvka_smallest_kerneliPfS_S_PiS_S0_S0_S0_iPfS_S_PiS_S0_S0_S0_
	addq	$32, %rsp
	jmp	.L109
.L148:
	.cfi_escape 0x2e,0
	call	cudaDeviceSynchronize@PLT
	movq	%r12, -100(%rbp)
	movl	-96(%rbp), %esi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	%ebx, -104(%rbp)
	movq	-104(%rbp), %rdi
	movq	%r12, %rdx
	movl	$1, %ecx
	movq	%r12, -80(%rbp)
	movl	$1, -72(%rbp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	je	.L110
.L112:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	call	_ZL5check9cudaErrorPKc.constprop.0
	jmp	.L149
.L110:
	pushq	%r10
	movl	%ebx, %edi
	pushq	-216(%rbp)
	pushq	-224(%rbp)
	movq	-256(%rbp), %rdx
	pushq	-232(%rbp)
	movq	-264(%rbp), %r9
	movq	-240(%rbp), %r8
	movq	-248(%rbp), %rcx
	movq	-280(%rbp), %rsi
	.cfi_escape 0x2e,0x20
	call	_Z57__device_stub__Z21boruvka_remove_cyclesiPiPfS_S_S_S_S0_S_iPiPfS_S_S_S_S0_S_
	addq	$32, %rsp
	jmp	.L112
.L149:
	.cfi_escape 0x2e,0
	call	cudaDeviceSynchronize@PLT
	movl	-284(%rbp), %r15d
.L113:
	movq	%r12, -100(%rbp)
	xorl	%r9d, %r9d
	movl	-96(%rbp), %esi
	xorl	%r8d, %r8d
	movl	%ebx, -104(%rbp)
	movq	%r12, %rdx
	movq	-104(%rbp), %rdi
	movl	$1, %ecx
	movq	%r12, -80(%rbp)
	movl	$1, -72(%rbp)
	testl	%r15d, %r15d
	jle	.L150
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	je	.L114
.L116:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	call	_ZL5check9cudaErrorPKc.constprop.0
	jmp	.L151
.L114:
	pushq	%r9
	movl	%ebx, %edi
	pushq	-216(%rbp)
	pushq	-224(%rbp)
	movq	-256(%rbp), %rdx
	pushq	-232(%rbp)
	movq	-264(%rbp), %r9
	movq	-240(%rbp), %r8
	movq	-248(%rbp), %rcx
	movq	-272(%rbp), %rsi
	.cfi_escape 0x2e,0x20
	call	_Z60__device_stub__Z24boruvka_pointer_doublingiPiPfS_S_S_S_S0_S_iPiPfS_S_S_S_S0_S_
	addq	$32, %rsp
	jmp	.L116
.L151:
	.cfi_escape 0x2e,0
	call	cudaDeviceSynchronize@PLT
	sarl	%r15d
	jmp	.L113
.L150:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	je	.L118
.L120:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	call	_ZL5check9cudaErrorPKc.constprop.0
	jmp	.L152
.L118:
	pushq	-216(%rbp)
	movl	%ebx, %edi
	pushq	-224(%rbp)
	movq	-256(%rbp), %rcx
	pushq	-232(%rbp)
	pushq	-264(%rbp)
	movq	-240(%rbp), %r9
	movq	-248(%rbp), %r8
	movq	-272(%rbp), %rdx
	movq	-280(%rbp), %rsi
	.cfi_escape 0x2e,0x20
	call	_Z63__device_stub__Z25boruvka_update_componentsiPiS_PfS_S_S_S_S0_S_iPiS_PfS_S_S_S_S0_S_
	addq	$32, %rsp
	jmp	.L120
.L152:
	.cfi_escape 0x2e,0
	call	cudaDeviceSynchronize@PLT
	movq	-216(%rbp), %rsi
	movl	$2, %ecx
	movl	$4, %edx
	leaq	-284(%rbp), %rdi
	call	cudaMemcpy@PLT
	call	cudaDeviceSynchronize@PLT
	jmp	.L106
.L147:
	xorl	%r8d, %r8d
	leaq	-208(%rbp), %rdi
	movl	$8, %esi
	movq	%r8, -208(%rbp)
	call	cudaMalloc@PLT
	leaq	-104(%rbp), %r12
	movq	-328(%rbp), %rdx
	leaq	-164(%rbp), %rcx
	movl	$2, %esi
	movq	%r12, %rdi
	movl	$0xff7fffff, -80(%rbp)
	call	_ZNSt6vectorIfSaIfEEC1EmRKfRKS0_
.LEHE2:
	movq	-104(%rbp), %rsi
	movl	$1, %ecx
	movl	$8, %edx
	movq	-208(%rbp), %rdi
.LEHB3:
	call	cudaMemcpy@PLT
	xorl	%edi, %edi
	movl	$8, %esi
	movq	%rdi, -200(%rbp)
	leaq	-200(%rbp), %rdi
	call	cudaMalloc@PLT
	movabsq	$4294967297, %rdx
	movl	%ebx, -164(%rbp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%rdx, -160(%rbp)
	movq	-164(%rbp), %rdi
	movl	$1, %ecx
	movl	-156(%rbp), %esi
	movq	%rdx, -80(%rbp)
	movl	$1, -72(%rbp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	je	.L122
.L124:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	call	_ZL5check9cudaErrorPKc.constprop.0
	jmp	.L153
.L122:
	pushq	-264(%rbp)
	movq	-208(%rbp), %r9
	movq	%r13, %rcx
	movl	%ebx, %edi
	movq	-344(%rbp), %rdx
	pushq	-200(%rbp)
	movq	-336(%rbp), %rsi
	movq	-232(%rbp), %r8
	.cfi_escape 0x2e,0x10
	call	_Z54__device_stub__Z19second_closest_findiPfS_S_PiS_S0_S0_iPfS_S_PiS_S0_S0_
	popq	%rcx
	popq	%rsi
	jmp	.L124
.L153:
	.cfi_escape 0x2e,0
	call	cudaDeviceSynchronize@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movl	$1, %esi
	movabsq	$4294967297, %rdi
	movl	$1, -72(%rbp)
	movq	%rdi, %rdx
	movq	%rdi, -80(%rbp)
	movq	%rdi, -164(%rbp)
	movl	$1, -156(%rbp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	je	.L125
.L127:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	call	_ZL5check9cudaErrorPKc.constprop.0
	jmp	.L154
.L125:
	movq	-224(%rbp), %rcx
	movq	-232(%rbp), %rdx
	movq	-200(%rbp), %rsi
	movq	-208(%rbp), %rdi
	call	_Z45__device_stub__Z18second_closest_setPfPiS0_S_PfPiS0_S_
	jmp	.L127
.L154:
	call	cudaDeviceSynchronize@PLT
	movq	-224(%rbp), %rsi
	movl	$2, %ecx
	movl	$4, %edx
	leaq	-164(%rbp), %rdi
	call	cudaMemcpy@PLT
	call	cudaDeviceSynchronize@PLT
	movq	-328(%rbp), %rdi
	leaq	-289(%rbp), %rdx
	movq	%r14, %rsi
	call	_ZNSt6vectorIiSaIiEEC1EmRKS0_
.LEHE3:
	movq	-80(%rbp), %rbx
	movq	-312(%rbp), %rdx
	movl	$2, %ecx
	movq	-232(%rbp), %rsi
	movq	%rbx, %rdi
.LEHB4:
	call	cudaMemcpy@PLT
	call	cudaDeviceSynchronize@PLT
	movq	-280(%rbp), %rdi
	call	cudaFree@PLT
	movq	-272(%rbp), %rdi
	call	cudaFree@PLT
	movq	-264(%rbp), %rdi
	call	cudaFree@PLT
	movq	-256(%rbp), %rdi
	call	cudaFree@PLT
	movq	-248(%rbp), %rdi
	call	cudaFree@PLT
	movq	-240(%rbp), %rdi
	call	cudaFree@PLT
	movq	-232(%rbp), %rdi
	call	cudaFree@PLT
	movq	-224(%rbp), %rdi
	call	cudaFree@PLT
	movq	-216(%rbp), %rdi
	call	cudaFree@PLT
	movq	-208(%rbp), %rdi
	call	cudaFree@PLT
	movq	-200(%rbp), %rdi
	call	cudaFree@PLT
	movq	-320(%rbp), %rax
	movq	-72(%rbp), %r13
	xorl	%edx, %edx
	movss	-164(%rbp), %xmm0
	movq	%r13, %rsi
	movq	%rdx, 8(%rax)
	leaq	8(%rax), %rdi
	movq	%rdx, 16(%rax)
	subq	%rbx, %rsi
	movq	%rdx, 24(%rax)
	sarq	$2, %rsi
	movss	%xmm0, (%rax)
	call	_ZNSt12_Vector_baseIiSaIiEE17_M_create_storageEm
.LEHE4:
	movq	-320(%rbp), %rax
	movq	%r13, %rsi
	movq	%rbx, %rdi
	movq	8(%rax), %rdx
	call	_ZNSt11__copy_moveILb0ELb1ESt26random_access_iterator_tagE8__copy_mIKiiEEPT0_PT_S7_S5_
	movq	-320(%rbp), %rbx
	leaq	-80(%rbp), %rdi
	movq	%rax, 16(%rbx)
	call	_ZNSt6vectorIiSaIiEED1Ev
	movq	%r12, %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	leaq	-128(%rbp), %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	leaq	-152(%rbp), %rdi
	call	_ZNSt6vectorIiSaIiEED1Ev
	movq	-56(%rbp), %rax
	subq	%fs:40, %rax
	je	.L133
	jmp	.L144
.L137:
	endbr64
	movq	%rax, %rbx
.L128:
	movq	-328(%rbp), %rdi
	call	_ZNSt6vectorIiSaIiEED1Ev
	jmp	.L129
.L136:
	endbr64
	movq	%rax, %rbx
.L129:
	movq	%r12, %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	jmp	.L130
.L135:
	endbr64
	movq	%rax, %rbx
.L130:
	leaq	-128(%rbp), %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	jmp	.L131
.L134:
	endbr64
	movq	%rax, %rbx
.L131:
	leaq	-152(%rbp), %rdi
	call	_ZNSt6vectorIiSaIiEED1Ev
	movq	-56(%rbp), %rax
	subq	%fs:40, %rax
	jne	.L144
	movq	%rbx, %rdi
.LEHB5:
	call	_Unwind_Resume@PLT
.LEHE5:
.L144:
	call	__stack_chk_fail@PLT
.L133:
	movq	-320(%rbp), %rax
	leaq	-40(%rbp), %rsp
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa 7, 8
	ret
	.cfi_endproc
.LFE4009:
	.globl	__gxx_personality_v0
	.section	.gcc_except_table,"a",@progbits
.LLSDA4009:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4009-.LLSDACSB4009
.LLSDACSB4009:
	.uleb128 .LEHB0-.LFB4009
	.uleb128 .LEHE0-.LEHB0
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB1-.LFB4009
	.uleb128 .LEHE1-.LEHB1
	.uleb128 .L134-.LFB4009
	.uleb128 0
	.uleb128 .LEHB2-.LFB4009
	.uleb128 .LEHE2-.LEHB2
	.uleb128 .L135-.LFB4009
	.uleb128 0
	.uleb128 .LEHB3-.LFB4009
	.uleb128 .LEHE3-.LEHB3
	.uleb128 .L136-.LFB4009
	.uleb128 0
	.uleb128 .LEHB4-.LFB4009
	.uleb128 .LEHE4-.LEHB4
	.uleb128 .L137-.LFB4009
	.uleb128 0
	.uleb128 .LEHB5-.LFB4009
	.uleb128 .LEHE5-.LEHB5
	.uleb128 0
	.uleb128 0
.LLSDACSE4009:
	.text
	.size	_Z19gpu_boruvka_onetreeiPfS_S_, .-_Z19gpu_boruvka_onetreeiPfS_S_
	.section	.text._ZNSt11__copy_moveILb0ELb1ESt26random_access_iterator_tagE8__copy_mIKffEEPT0_PT_S7_S5_,"axG",@progbits,_ZNSt11__copy_moveILb0ELb1ESt26random_access_iterator_tagE8__copy_mIKffEEPT0_PT_S7_S5_,comdat
	.weak	_ZNSt11__copy_moveILb0ELb1ESt26random_access_iterator_tagE8__copy_mIKffEEPT0_PT_S7_S5_
	.type	_ZNSt11__copy_moveILb0ELb1ESt26random_access_iterator_tagE8__copy_mIKffEEPT0_PT_S7_S5_, @function
_ZNSt11__copy_moveILb0ELb1ESt26random_access_iterator_tagE8__copy_mIKffEEPT0_PT_S7_S5_:
.LFB4974:
	.cfi_startproc
	endbr64
	subq	%rdi, %rsi
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rdx, %rcx
	movq	%rsi, %rbx
	cmpq	$4, %rsi
	jle	.L156
	movq	%rsi, %rdx
	movq	%rdi, %rsi
	movq	%rcx, %rdi
	call	memmove@PLT
	movq	%rax, %rcx
	jmp	.L157
.L156:
	jne	.L157
	movss	(%rdi), %xmm0
	movss	%xmm0, (%rdx)
.L157:
	leaq	(%rcx,%rbx), %rax
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4974:
	.size	_ZNSt11__copy_moveILb0ELb1ESt26random_access_iterator_tagE8__copy_mIKffEEPT0_PT_S7_S5_, .-_ZNSt11__copy_moveILb0ELb1ESt26random_access_iterator_tagE8__copy_mIKffEEPT0_PT_S7_S5_
	.text
	.globl	_Z25gpu_subgradient_opt_alphaPfS_i
	.type	_Z25gpu_subgradient_opt_alphaPfS_i, @function
_Z25gpu_subgradient_opt_alphaPfS_i:
.LFB4027:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4027
	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
	movq	%rsi, %r13
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdx, %r12
	xorl	%edx, %edx
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movslq	%ecx, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rdi, %rbx
	subq	$248, %rsp
	.cfi_def_cfa_offset 304
	movl	%ecx, 12(%rsp)
	leaq	72(%rsp), %rdi
	movq	%fs:40, %rax
	movq	%rax, 232(%rsp)
	xorl	%eax, %eax
	leaq	0(,%rbp,4), %rax
	movq	%rdx, 72(%rsp)
	movq	%rax, %rsi
	movq	%rax, (%rsp)
.LEHB6:
	call	cudaMalloc@PLT
	movq	(%rsp), %rdx
	movq	72(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaMemset@PLT
	movq	(%rsp), %rsi
	xorl	%ecx, %ecx
	leaq	80(%rsp), %rdi
	movq	%rcx, 80(%rsp)
	call	cudaMalloc@PLT
	movq	(%rsp), %rdx
	movq	80(%rsp), %rdi
	movq	%r13, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	xorl	%esi, %esi
	leaq	88(%rsp), %rdi
	movq	%rsi, 88(%rsp)
	movq	(%rsp), %rsi
	call	cudaMalloc@PLT
	movq	(%rsp), %rdx
	movq	88(%rsp), %rdi
	movq	%r12, %rsi
	movl	$1, %ecx
	leaq	168(%rsp), %r12
	call	cudaMemcpy@PLT
	leaq	200(%rsp), %rax
	movq	%r12, %rcx
	movq	%rbp, %rsi
	movq	%rax, 24(%rsp)
	leaq	96(%rsp), %rax
	movq	24(%rsp), %rdx
	movq	%rax, %rdi
	movq	%rax, 40(%rsp)
	movl	$0x00000000, 200(%rsp)
	call	_ZNSt6vectorIfSaIfEEC1EmRKfRKS0_
.LEHE6:
	movq	24(%rsp), %rdx
	movq	%r12, %rcx
	movq	%rbp, %rsi
	movl	$0x00000000, 200(%rsp)
	movq	%rbx, %rdi
.LEHB7:
	call	_ZNSt6vectorIfSaIfEEC1EmRKfRKS0_
.LEHE7:
	movq	72(%rsp), %r8
	movq	88(%rsp), %rcx
	movq	%r12, %rdi
	movq	80(%rsp), %rdx
	movl	12(%rsp), %esi
.LEHB8:
	call	_Z19gpu_boruvka_onetreeiPfS_S_
.LEHE8:
	movq	24(%rsp), %rdx
	leaq	120(%rsp), %rdi
	movq	%rbp, %rsi
	movss	168(%rsp), %xmm6
	movss	%xmm6, 52(%rsp)
.LEHB9:
	call	_ZNSt6vectorIiSaIiEEC1EmRKS0_
.LEHE9:
	movq	24(%rsp), %rdx
	leaq	144(%rsp), %rdi
	movq	%rbp, %rsi
.LEHB10:
	call	_ZNSt6vectorIiSaIiEEC1EmRKS0_
.LEHE10:
	movq	176(%rsp), %rcx
	movq	120(%rsp), %rsi
	xorl	%eax, %eax
	movq	144(%rsp), %r12
.L160:
	cmpl	%eax, 12(%rsp)
	jle	.L237
	movl	(%rcx,%rax,4), %edx
	subl	$2, %edx
	movl	%edx, (%rsi,%rax,4)
	movl	%edx, (%r12,%rax,4)
	incq	%rax
	jmp	.L160
.L237:
	movss	.LC14(%rip), %xmm7
	movl	$1000, 16(%rsp)
	movb	$1, 51(%rsp)
	movss	%xmm7, 8(%rsp)
.L162:
	cmpl	$0, 16(%rsp)
	movl	$1, 20(%rsp)
	jle	.L189
.L194:
	movq	120(%rsp), %r13
	movq	96(%rsp), %rbp
	xorl	%eax, %eax
	movss	.LC16(%rip), %xmm2
.L191:
	cmpl	%eax, 12(%rsp)
	jle	.L238
	cvtsi2ssl	0(%r13,%rax,4), %xmm1
	cvtsi2ssl	(%r12,%rax,4), %xmm0
	mulss	.LC15(%rip), %xmm0
	mulss	%xmm2, %xmm1
	addss	%xmm1, %xmm0
	mulss	8(%rsp), %xmm0
	addss	0(%rbp,%rax,4), %xmm0
	movss	%xmm0, 0(%rbp,%rax,4)
	incq	%rax
	jmp	.L191
.L238:
	movq	(%rsp), %rdx
	movq	72(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbp, %rsi
.LEHB11:
	call	cudaMemcpy@PLT
	call	cudaDeviceSynchronize@PLT
	movq	152(%rsp), %rax
	movq	%rax, %r14
	movq	%rax, 32(%rsp)
	movq	136(%rsp), %rax
	subq	%r12, %r14
	subq	%r13, %rax
	movq	%rax, 56(%rsp)
	cmpq	%r14, %rax
	jnb	.L164
	movq	%r14, %rdi
	sarq	$2, %rdi
	je	.L203
	call	_ZNSt15__new_allocatorIiE8allocateEmPKv.isra.0
	movq	%rax, %r15
	jmp	.L165
.L203:
	xorl	%r15d, %r15d
.L165:
	movq	32(%rsp), %rsi
	movq	%r15, %rdx
	movq	%r12, %rdi
	call	_ZNSt11__copy_moveILb0ELb1ESt26random_access_iterator_tagE8__copy_mIKiiEEPT0_PT_S7_S5_
	testq	%r13, %r13
	je	.L166
	movq	56(%rsp), %rsi
	movq	%r13, %rdi
	call	_ZdlPvm@PLT
.L166:
	movq	%r15, 120(%rsp)
	addq	%r14, %r15
	movq	%r15, 136(%rsp)
	jmp	.L167
.L164:
	movq	128(%rsp), %rcx
	movq	%rcx, %r15
	subq	%r13, %r15
	cmpq	%r14, %r15
	jb	.L168
	movq	32(%rsp), %rsi
	movq	%r13, %rdx
	movq	%r12, %rdi
	call	_ZNSt11__copy_moveILb0ELb1ESt26random_access_iterator_tagE8__copy_mIKiiEEPT0_PT_S7_S5_
	jmp	.L167
.L168:
	cmpq	$4, %r15
	jle	.L169
	movq	%r15, %rdx
	movq	%r12, %rsi
	movq	%r13, %rdi
	movq	%rcx, 56(%rsp)
	call	memmove@PLT
	movq	56(%rsp), %rcx
	jmp	.L170
.L169:
	jne	.L170
	movl	(%r12), %eax
	movl	%eax, 0(%r13)
.L170:
	movq	32(%rsp), %rdx
	leaq	(%r12,%r15), %rsi
	subq	%rsi, %rdx
	cmpq	$4, %rdx
	jle	.L171
	movq	%rcx, %rdi
	call	memmove@PLT
	jmp	.L167
.L171:
	jne	.L167
	movl	(%rsi), %eax
	movl	%eax, (%rcx)
.L167:
	movq	72(%rsp), %r8
	movq	88(%rsp), %rcx
	movq	80(%rsp), %rdx
	movl	12(%rsp), %esi
	movq	24(%rsp), %rdi
	addq	120(%rsp), %r14
	movq	%r14, 128(%rsp)
	call	_Z19gpu_boruvka_onetreeiPfS_S_
.LEHE11:
	movss	200(%rsp), %xmm1
	xorl	%eax, %eax
	xorl	%edx, %edx
.L172:
	cmpl	%eax, 12(%rsp)
	jle	.L239
	movss	0(%rbp,%rax,4), %xmm0
	movb	$1, %dl
	incq	%rax
	addss	%xmm0, %xmm0
	subss	%xmm0, %xmm1
	jmp	.L172
.L239:
	testb	%dl, %dl
	je	.L174
	movss	%xmm1, 200(%rsp)
.L174:
	movq	208(%rsp), %rcx
	xorl	%eax, %eax
.L175:
	cmpl	%eax, 12(%rsp)
	jle	.L240
	movl	(%rcx,%rax,4), %esi
	leal	-2(%rsi), %edx
	movl	%edx, (%r12,%rax,4)
	incq	%rax
	jmp	.L175
.L240:
	movss	200(%rsp), %xmm2
	comiss	52(%rsp), %xmm2
	movss	%xmm2, 32(%rsp)
	jbe	.L232
	movq	104(%rsp), %r15
	movq	(%rbx), %rdi
	movq	16(%rbx), %rax
	movq	%r15, %r13
	subq	%rbp, %r13
	subq	%rdi, %rax
	cmpq	%r13, %rax
	jnb	.L179
	movq	%r13, %rdi
	sarq	$2, %rdi
	je	.L204
.LEHB12:
	call	_ZNSt15__new_allocatorIiE8allocateEmPKv.isra.0
.LEHE12:
	movq	%rax, %r14
	jmp	.L180
.L204:
	xorl	%r14d, %r14d
.L180:
	movq	%rbp, %rdi
	movq	%r14, %rdx
	movq	%r15, %rsi
	call	_ZNSt11__copy_moveILb0ELb1ESt26random_access_iterator_tagE8__copy_mIKffEEPT0_PT_S7_S5_
	movq	(%rbx), %rdi
	testq	%rdi, %rdi
	je	.L181
	movq	16(%rbx), %rsi
	subq	%rdi, %rsi
	call	_ZdlPvm@PLT
.L181:
	movq	%r14, (%rbx)
	addq	%r13, %r14
	movq	%r14, 16(%rbx)
	jmp	.L182
.L179:
	movq	8(%rbx), %rdx
	subq	%rdi, %rdx
	cmpq	%r13, %rdx
	jb	.L183
	movq	%rdi, %rdx
	movq	%r15, %rsi
	movq	%rbp, %rdi
	call	_ZNSt11__copy_moveILb0ELb1ESt26random_access_iterator_tagE8__copy_mIKffEEPT0_PT_S7_S5_
	jmp	.L182
.L183:
	cmpq	$4, %rdx
	jle	.L184
	movq	%rbp, %rsi
	call	memmove@PLT
	jmp	.L185
.L184:
	jne	.L185
	movss	0(%rbp), %xmm0
	movss	%xmm0, (%rdi)
.L185:
	movq	8(%rbx), %rdi
	movq	%r15, %rdx
	movq	%rdi, %rsi
	subq	(%rbx), %rsi
	addq	%rbp, %rsi
	subq	%rsi, %rdx
	cmpq	$4, %rdx
	jle	.L186
	call	memmove@PLT
	jmp	.L182
.L186:
	jne	.L182
	movss	(%rsi), %xmm0
	movss	%xmm0, (%rdi)
.L182:
	addq	(%rbx), %r13
	cmpb	$0, 51(%rsp)
	movq	%r13, 8(%rbx)
	je	.L187
	movss	8(%rsp), %xmm3
	addss	%xmm3, %xmm3
	movss	%xmm3, 8(%rsp)
.L187:
	movss	32(%rsp), %xmm5
	movl	16(%rsp), %ecx
	movss	%xmm5, 52(%rsp)
	cmpl	%ecx, 20(%rsp)
	jne	.L188
	movl	20(%rsp), %eax
	addl	%eax, %eax
	movl	%eax, 16(%rsp)
	jmp	.L188
.L232:
	cmpb	$0, 51(%rsp)
	je	.L188
	cmpl	$500, 20(%rsp)
	jle	.L188
	movss	8(%rsp), %xmm4
	mulss	.LC17(%rip), %xmm4
	xorl	%eax, %eax
	movb	$0, 51(%rsp)
	movl	%eax, 20(%rsp)
	movss	%xmm4, 8(%rsp)
.L188:
	leaq	208(%rsp), %rdi
	call	_ZNSt6vectorIiSaIiEED1Ev
	movss	8(%rsp), %xmm2
	xorps	%xmm0, %xmm0
	incl	20(%rsp)
	comiss	%xmm0, %xmm2
	jbe	.L233
	movl	20(%rsp), %ecx
	cmpl	%ecx, 16(%rsp)
	jge	.L194
	jmp	.L189
.L233:
	movl	$0x00000000, 8(%rsp)
.L189:
	movl	16(%rsp), %eax
	movl	$2, %ecx
	movss	8(%rsp), %xmm5
	xorps	%xmm0, %xmm0
	mulss	.LC18(%rip), %xmm5
	cltd
	idivl	%ecx
	movss	%xmm5, 8(%rsp)
	comiss	%xmm0, %xmm5
	movl	%eax, 16(%rsp)
	ja	.L162
	movq	72(%rsp), %rdi
.LEHB13:
	call	cudaFree@PLT
	movq	80(%rsp), %rdi
	call	cudaFree@PLT
	movq	88(%rsp), %rdi
	call	cudaFree@PLT
.LEHE13:
	leaq	144(%rsp), %rdi
	call	_ZNSt6vectorIiSaIiEED1Ev
	leaq	120(%rsp), %rdi
	call	_ZNSt6vectorIiSaIiEED1Ev
	leaq	176(%rsp), %rdi
	call	_ZNSt6vectorIiSaIiEED1Ev
	movq	40(%rsp), %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	movq	232(%rsp), %rax
	subq	%fs:40, %rax
	je	.L202
	jmp	.L235
.L212:
	endbr64
	movq	%rax, %rbp
.L195:
	leaq	208(%rsp), %rdi
	call	_ZNSt6vectorIiSaIiEED1Ev
	jmp	.L196
.L213:
	endbr64
	movq	%rax, %rbp
.L196:
	leaq	144(%rsp), %rdi
	call	_ZNSt6vectorIiSaIiEED1Ev
	jmp	.L197
.L211:
	endbr64
	movq	%rax, %rbp
.L197:
	leaq	120(%rsp), %rdi
	call	_ZNSt6vectorIiSaIiEED1Ev
	jmp	.L198
.L210:
	endbr64
	movq	%rax, %rbp
.L198:
	leaq	176(%rsp), %rdi
	call	_ZNSt6vectorIiSaIiEED1Ev
	jmp	.L199
.L209:
	endbr64
	movq	%rax, %rbp
.L199:
	movq	%rbx, %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	jmp	.L200
.L208:
	endbr64
	movq	%rax, %rbp
.L200:
	movq	40(%rsp), %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	movq	232(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L235
	movq	%rbp, %rdi
.LEHB14:
	call	_Unwind_Resume@PLT
.LEHE14:
.L235:
	call	__stack_chk_fail@PLT
.L202:
	addq	$248, %rsp
	.cfi_def_cfa_offset 56
	movq	%rbx, %rax
	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
.LFE4027:
	.section	.gcc_except_table
.LLSDA4027:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4027-.LLSDACSB4027
.LLSDACSB4027:
	.uleb128 .LEHB6-.LFB4027
	.uleb128 .LEHE6-.LEHB6
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB7-.LFB4027
	.uleb128 .LEHE7-.LEHB7
	.uleb128 .L208-.LFB4027
	.uleb128 0
	.uleb128 .LEHB8-.LFB4027
	.uleb128 .LEHE8-.LEHB8
	.uleb128 .L209-.LFB4027
	.uleb128 0
	.uleb128 .LEHB9-.LFB4027
	.uleb128 .LEHE9-.LEHB9
	.uleb128 .L210-.LFB4027
	.uleb128 0
	.uleb128 .LEHB10-.LFB4027
	.uleb128 .LEHE10-.LEHB10
	.uleb128 .L211-.LFB4027
	.uleb128 0
	.uleb128 .LEHB11-.LFB4027
	.uleb128 .LEHE11-.LEHB11
	.uleb128 .L213-.LFB4027
	.uleb128 0
	.uleb128 .LEHB12-.LFB4027
	.uleb128 .LEHE12-.LEHB12
	.uleb128 .L212-.LFB4027
	.uleb128 0
	.uleb128 .LEHB13-.LFB4027
	.uleb128 .LEHE13-.LEHB13
	.uleb128 .L213-.LFB4027
	.uleb128 0
	.uleb128 .LEHB14-.LFB4027
	.uleb128 .LEHE14-.LEHB14
	.uleb128 0
	.uleb128 0
.LLSDACSE4027:
	.text
	.size	_Z25gpu_subgradient_opt_alphaPfS_i, .-_Z25gpu_subgradient_opt_alphaPfS_i
	.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.cst4,"aM",@progbits,4
	.align 4
.LC14:
	.long	1065353216
	.align 4
.LC15:
	.long	1060320051
	.align 4
.LC16:
	.long	1050253722
	.align 4
.LC17:
	.long	1061158912
	.align 4
.LC18:
	.long	1056964608
	.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>

std::vector<float> gpu_subgradient_opt_alpha(float* x, float* y, int n);
	.text
	.file	"cuda_preprocessing.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
