/*
    CUDA good to knows:
        Basics:
            Per thread:
                registers (fast)
                local memory (off-chip [still on the GPU though], slow)

            Per block:
                multiple threads
                shared memory (semi-fast)

            Per GPU:
                Multiple kernels that each run multiple blocks
                Global memory (off-chip [still on the GPU though], slow)


            Threads are executed by thread processors
            
            Threads reside in thread blocks
            
            Thread blocks are executed by multiprocessors

            Several concurrent thread blocks can reside on one multiprocessor
                Limited by multiprocessor resources (shared memory and registers)

            A kernel is launched as a grid of thread blocks. Only one kernel can execute on a device at a time.
        Advanced:
            cudaMemcpy(dst, src, size, direction)
                blocks CPU thread.
            

    Compiler tips:
        nvcc <filename>.cu [-o <executable>]
            Builds release mode

        nvcc -g <filename>.cu
            Builds debug mode
            Can debug host code but not device code

        nvcc -deviceemu <filename>.cu
            Builds device emulation mode
            All code runs on CPU, no debug symbols

        nvcc -deviceemu -g <filename>.cu
            Builds debug device emulation mode
            All code runs on CPU, with debug symbols

    Tips and tricks:
        If our arrays A,B,C are shorter than 1024 elements, N < 1024, then
            – one thread block is enough
            – N threads in the thread block
        If our arrays are longer than 1024, then
            – Choose the number of threads in the thread blocks to be
            integer*32
            – Calculate how many thread blocks you need
            – There will be some threads that should do nothing
        Why multiples of 32?
            – Threads are executed synchronously in bunches of 32 =
            warp
            – All threads must have their data ready before the warp runs
            – Cache lines are 4 B x warp size = 128 B
            – GPU resources can be fully utilized when these parameters
            are used
        # of blocks = ceil(N/threadsInBlock)
                    = (N+threadsInBlock-1)/threadsInBlock



Compile:
nvcc -o galaxy galaxy_program.cu -res-usage

Run:
time ./galaxy

*/
#include <stdio.h>
#include <iostream>
#include <fstream>

using namespace std;


// Declare functions and classes that are below main.
class GalaxyFile{
public:
    int number_of_galaxies;
    float *alphas, *deltas;

    GalaxyFile(){}

    GalaxyFile(int num, float *as, float *ds)
    {
        number_of_galaxies = num;
        alphas = as;
        deltas = ds;
    }
};
void print_omegas(float*, int);
void write_omegas_to_file(string, float*);
void write_histogram_to_file(string, int*);
void print_histogram(string, int*, int);
GalaxyFile readFile(string);


// Define some useful macros
#define BIN_WIDTH 0.25f
#define BIN_MIN 0.0f
#define BIN_MAX 180.0f
#define NUMBER_OF_BINS (int)(BIN_MAX*(1.0f/BIN_WIDTH))

// Google is your friend.
#define ARCMINS_TO_RADIANS 0.000290888209f
#define RADIANS_TO_DEGREES 57.295779513f


__global__
void angle_between_galaxies(float *alphas1, float *deltas1, float *alphas2, float *deltas2, int *gpu_hist){
    int idx = blockDim.x*blockIdx.x + threadIdx.x;
    
    if(idx < 100000){
        for(int i=0; i<100000; i++){
                float angle = 0.0f;
                // Don't do duplicates
                if( alphas1[i] != alphas2[idx] && deltas1[i] != deltas2[idx] ) {
                    float x = sin(deltas1[i]) * sin(deltas2[idx]) + cos(deltas1[i]) * cos(deltas2[idx]) * cos(alphas1[i] - alphas2[idx]);
                    angle = acosf(fmaxf(-1.0f, fminf(x, 1.0f))) * RADIANS_TO_DEGREES;
                }

                int ix = (int)(floor(angle * (1.0f/BIN_WIDTH))) % NUMBER_OF_BINS;
                __syncthreads();
                atomicAdd(&gpu_hist[ix], 1);
        }
    }
}

int* calculate_histogram(GalaxyFile galaxies1, GalaxyFile galaxies2){
    // Declare and allocate memory for histogram arrays that will be accessible on CPU
    float galaxy_array_size = galaxies1.number_of_galaxies * sizeof(float);
    float histogram_size = NUMBER_OF_BINS * sizeof(int);

    int *histogram;
    int *total_histogram;
    histogram = (int *) malloc(NUMBER_OF_BINS*sizeof(int));
    total_histogram = (int *) malloc(NUMBER_OF_BINS*sizeof(int));

    memset(total_histogram, 0, NUMBER_OF_BINS*sizeof(int));

    // Declare angle arrays that will be accessible on GPU
    float *gpu_alphas1;
    float *gpu_deltas1;
    float *gpu_alphas2;
    float *gpu_deltas2;
    int *gpu_histogram;

    // Allocate memory on GPU for angle arrays
    cudaMalloc((void**) &gpu_alphas1, galaxy_array_size);
    cudaMalloc((void**) &gpu_deltas1, galaxy_array_size);
    cudaMalloc((void**) &gpu_alphas2, galaxy_array_size);
    cudaMalloc((void**) &gpu_deltas2, galaxy_array_size);
    cudaMalloc((void**) &gpu_histogram, NUMBER_OF_BINS*sizeof(int));

	// Copy angles from CPU onto GPU
	cudaMemcpy(gpu_alphas1, galaxies1.alphas, galaxy_array_size, cudaMemcpyHostToDevice);
	cudaMemcpy(gpu_deltas1, galaxies1.deltas, galaxy_array_size, cudaMemcpyHostToDevice);
	cudaMemcpy(gpu_alphas2, galaxies2.alphas, galaxy_array_size, cudaMemcpyHostToDevice);
	cudaMemcpy(gpu_deltas2, galaxies2.deltas, galaxy_array_size, cudaMemcpyHostToDevice);
	cudaMemcpy(gpu_histogram, histogram, galaxy_array_size, cudaMemcpyHostToDevice);
    
    int warp_size = 32;
    int threadsInBlock = 11 * warp_size;
    int blocksInGrid = ceil((galaxies1.number_of_galaxies + threadsInBlock) / threadsInBlock);

	// Define the grid size (blocks per grid)
    dim3 dimGrid(blocksInGrid);
    
	// Define block size (threads per block)
	dim3 dimBlock(threadsInBlock);

    // Write histogram full of zeros
    cudaMemset(gpu_histogram, 0, histogram_size);

    // Calculate angles between galaxies1[i] and every galaxy in galaxies2
    angle_between_galaxies<<<dimGrid, dimBlock>>>(gpu_alphas1, gpu_deltas1, gpu_alphas2, gpu_deltas2, gpu_histogram);
    
    // Copy result histogram into CPU histogram
    cudaMemcpy(histogram, gpu_histogram, histogram_size, cudaMemcpyDeviceToHost);
    
	// Free all the memory we allocated on the GPU
	cudaFree( gpu_alphas1 );
	cudaFree( gpu_deltas1 );
	cudaFree( gpu_alphas2 );
	cudaFree( gpu_deltas2 );
    cudaFree( gpu_histogram );

    return histogram;
}

float* calculate_omegas(int* DD, int* DR, int* RR){
    float* omegas;
    omegas = (float *) malloc(NUMBER_OF_BINS*sizeof(float));

    for(int i=0; i<NUMBER_OF_BINS; i++){
        if(RR[i] != 0.0f){
            omegas[i] = (DD[i] - 2.0f*DR[i] + RR[i]) / RR[i];
        }else{
            omegas[i] = 0.0f;
        }
    }
    return omegas;
}

// CUDA program that calculates distribution of galaxies
int main()
{
    // Read files and store data in GalaxyFile classes.
    GalaxyFile galaxies1;
    GalaxyFile galaxies2;
    galaxies1 = readFile("test_data/flat_100k_arcmin.txt");
    galaxies2 = readFile("test_data/data_100k_arcmin.txt");

    int* DD_hist = calculate_histogram(galaxies1, galaxies1);
    int* DR_hist = calculate_histogram(galaxies1, galaxies2);
    int* RR_hist = calculate_histogram(galaxies2, galaxies2);

    print_histogram("DD", DD_hist, 20);
    print_histogram("DR", DR_hist, 20);
    print_histogram("RR", RR_hist, 20);

    write_histogram_to_file("dd_histogram.txt", DD_hist);
    write_histogram_to_file("dr_histogram.txt", DR_hist);
    write_histogram_to_file("rr_histogram.txt", RR_hist);

    float* omegas = calculate_omegas(DD_hist, DR_hist, RR_hist);
    print_omegas(omegas, 15);
    write_omegas_to_file("omegas.txt", omegas);

	return EXIT_SUCCESS;
}


/* UTILITY FUNCTIONS/CLASSES BELOW */
GalaxyFile readFile(string filename)
{
    ifstream infile(filename);
    int number_of_galaxies;

    // Read first line which is the number of galaxies that's stored in the file.
    infile >> number_of_galaxies;

    float galaxy_array_size = number_of_galaxies * sizeof(float);

    float *alphas, *deltas;
    alphas = (float*) malloc(galaxy_array_size);
    deltas = (float*) malloc(galaxy_array_size);

    float alpha;
    float delta;

    // Read arc minute angles for each galaxy
    // Then convert those angles to radians and store those in alphas and deltas arrays
    for(int i=0; i<number_of_galaxies; i++) {
        infile >> alpha >> delta;

        alphas[i] = alpha * ARCMINS_TO_RADIANS;
        deltas[i] = delta * ARCMINS_TO_RADIANS;
    }
    infile.close();

    GalaxyFile galaxyFile(number_of_galaxies, alphas, deltas);
    return galaxyFile;
}

void print_omegas(float* omegas, int bins_to_print){
    for (int i=0; i<NUMBER_OF_BINS; i++){
        if(omegas[i] != 0.0f && i < bins_to_print){
            printf("omegas[%d]: %f\n", i, omegas[i]);
        }
    }
}

void print_histogram(string label, int *histogram, int bins_to_print){
    long long galaxies_counted = 0;
    // Print each bucket bin that has 1 or more galaxy-pair-angle in it.
    for (int i=0; i<NUMBER_OF_BINS; i++) {
        float bucket_min = (float)i / (1.0f/BIN_WIDTH);
        float bucket_max = (float)i / (1.0f/BIN_WIDTH) + BIN_WIDTH;
        int bucket_value = histogram[i];

        galaxies_counted += histogram[i];
        if(bucket_value > 0 && i < bins_to_print){
            printf("[%f, %f]: %d\n", bucket_min, bucket_max, bucket_value);
        }
    }

    cout << "Galaxy pairs counted in " << label << ": " << galaxies_counted << endl;
}

void write_omegas_to_file(string filename, float* omegas){
    ofstream file;
    file.open("output/"+filename);

    for (int i=0; i<NUMBER_OF_BINS; i++){
        file << omegas[i];
        if(i<NUMBER_OF_BINS-1) file << "\n";
    }
    file.close();
}

void write_histogram_to_file(string filename, int* histogram){
    ofstream file;
    file.open("output/"+filename);
    
    for (int i=0; i<NUMBER_OF_BINS; i++){
        file << histogram[i]; 
        if(i<NUMBER_OF_BINS-1) file << "\n";
    }
    file.close();
}
	.file	"tmpxft_0036b38b_00000000-6_galaxy_program.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3782:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3782:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.rodata._ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE6appendEPKcm.isra.0.str1.1,"aMS",@progbits,1
.LC0:
	.string	"basic_string::append"
	.section	.text._ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE6appendEPKcm.isra.0,"axG",@progbits,_ZStplIcSt11char_traitsIcESaIcEENSt7__cxx1112basic_stringIT_T0_T1_EEPKS5_RKS8_,comdat
	.align 2
	.type	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE6appendEPKcm.isra.0, @function
_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE6appendEPKcm.isra.0:
.LFB4531:
	.cfi_startproc
	movabsq	$4611686018427387903, %rax
	subq	8(%rdi), %rax
	cmpq	%rdx, %rax
	jnb	.L3
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	.LC0(%rip), %rdi
	call	_ZSt20__throw_length_errorPKc@PLT
.L3:
	.cfi_def_cfa_offset 8
	jmp	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE9_M_appendEPKcm@PLT
	.cfi_endproc
.LFE4531:
	.size	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE6appendEPKcm.isra.0, .-_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE6appendEPKcm.isra.0
	.text
	.align 2
	.type	_ZNSt14basic_ofstreamIcSt11char_traitsIcEE4openERKNSt7__cxx1112basic_stringIcS1_SaIcEEESt13_Ios_Openmode.isra.0, @function
_ZNSt14basic_ofstreamIcSt11char_traitsIcEE4openERKNSt7__cxx1112basic_stringIcS1_SaIcEEESt13_Ios_Openmode.isra.0:
.LFB4533:
	.cfi_startproc
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	orl	$16, %edx
	movq	%rdi, %rbx
	addq	$8, %rdi
	call	_ZNSt13basic_filebufIcSt11char_traitsIcEE4openEPKcSt13_Ios_Openmode@PLT
	testq	%rax, %rax
	movq	(%rbx), %rax
	jne	.L9
	addq	-24(%rax), %rbx
	movl	$4, %esi
	movq	%rbx, %rdi
	popq	%rbx
	.cfi_remember_state
	.cfi_def_cfa_offset 8
	jmp	_ZNSt9basic_iosIcSt11char_traitsIcEE8setstateESt12_Ios_Iostate@PLT
.L9:
	.cfi_restore_state
	addq	-24(%rax), %rbx
	xorl	%esi, %esi
	movq	%rbx, %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	_ZNSt9basic_iosIcSt11char_traitsIcEE5clearESt12_Ios_Iostate@PLT
	.cfi_endproc
.LFE4533:
	.size	_ZNSt14basic_ofstreamIcSt11char_traitsIcEE4openERKNSt7__cxx1112basic_stringIcS1_SaIcEEESt13_Ios_Openmode.isra.0, .-_ZNSt14basic_ofstreamIcSt11char_traitsIcEE4openERKNSt7__cxx1112basic_stringIcS1_SaIcEEESt13_Ios_Openmode.isra.0
	.globl	_Z16calculate_omegasPiS_S_
	.type	_Z16calculate_omegasPiS_S_, @function
_Z16calculate_omegasPiS_S_:
.LFB3773:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movq	%rdx, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movq	%rsi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	%rdi, %rbx
	movl	$2880, %edi
	call	malloc@PLT
	xorl	%ecx, %ecx
	xorps	%xmm3, %xmm3
.L14:
	cvtsi2ssl	(%r12,%rcx), %xmm2
	ucomiss	%xmm3, %xmm2
	jp	.L16
	xorps	%xmm0, %xmm0
	je	.L12
.L16:
	cvtsi2ssl	0(%rbp,%rcx), %xmm1
	cvtsi2ssl	(%rbx,%rcx), %xmm0
	addss	%xmm1, %xmm1
	subss	%xmm1, %xmm0
	addss	%xmm2, %xmm0
	divss	%xmm2, %xmm0
.L12:
	movss	%xmm0, (%rax,%rcx)
	addq	$4, %rcx
	cmpq	$2880, %rcx
	jne	.L14
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3773:
	.size	_Z16calculate_omegasPiS_S_, .-_Z16calculate_omegasPiS_S_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC2:
	.string	"omegas[%d]: %f\n"
	.text
	.globl	_Z12print_omegasPfi
	.type	_Z12print_omegasPfi, @function
_Z12print_omegasPfi:
.LFB3776:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	leaq	.LC2(%rip), %r13
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movl	%esi, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	xorl	%ebx, %ebx
	pushq	%rdx
	.cfi_def_cfa_offset 48
.L21:
	movss	0(%rbp,%rbx,4), %xmm0
	cmpl	%ebx, %r12d
	jle	.L20
	xorps	%xmm1, %xmm1
	movb	$1, %dl
	ucomiss	%xmm1, %xmm0
	setp	%al
	cmovne	%edx, %eax
	testb	%al, %al
	je	.L20
	movl	%ebx, %edx
	movq	%r13, %rsi
	movl	$2, %edi
	movb	$1, %al
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
.L20:
	incq	%rbx
	cmpq	$720, %rbx
	jne	.L21
	popq	%rax
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3776:
	.size	_Z12print_omegasPfi, .-_Z12print_omegasPfi
	.section	.rodata.str1.1
.LC4:
	.string	"[%f, %f]: %d\n"
.LC5:
	.string	"Galaxy pairs counted in "
.LC6:
	.string	": "
	.text
	.globl	_Z15print_histogramNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPii
	.type	_Z15print_histogramNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPii, @function
_Z15print_histogramNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPii:
.LFB3777:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movq	%rsi, %r14
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	movl	%edx, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	movq	%rdi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	xorl	%ebp, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	xorl	%ebx, %ebx
.L32:
	movslq	(%r14,%rbx,4), %rax
	movq	%rax, %rdx
	addq	%rax, %rbp
	testl	%eax, %eax
	jle	.L31
	cmpl	%ebx, %r13d
	jle	.L31
	cvtsi2ssl	%ebx, %xmm0
	mulss	.LC3(%rip), %xmm0
	movl	$2, %edi
	movb	$2, %al
	movss	.LC3(%rip), %xmm1
	leaq	.LC4(%rip), %rsi
	addss	%xmm0, %xmm1
	cvtss2sd	%xmm0, %xmm0
	cvtss2sd	%xmm1, %xmm1
	call	__printf_chk@PLT
.L31:
	incq	%rbx
	cmpq	$720, %rbx
	jne	.L32
	leaq	.LC5(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%r12, %rsi
	movq	%rax, %rdi
	call	_ZStlsIcSt11char_traitsIcESaIcEERSt13basic_ostreamIT_T0_ES7_RKNSt7__cxx1112basic_stringIS4_S5_T1_EE@PLT
	leaq	.LC6(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rbp, %rsi
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIxEERSoT_@PLT
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	movq	%rax, %rdi
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	jmp	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	.cfi_endproc
.LFE3777:
	.size	_Z15print_histogramNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPii, .-_Z15print_histogramNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPii
	.globl	_Z50__device_stub__Z22angle_between_galaxiesPfS_S_S_PiPfS_S_S_Pi
	.type	_Z50__device_stub__Z22angle_between_galaxiesPfS_S_S_PiPfS_S_S_Pi, @function
_Z50__device_stub__Z22angle_between_galaxiesPfS_S_S_PiPfS_S_S_Pi:
.LFB3804:
	.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
	movq	%rcx, 16(%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	16(%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	.L41
	pushq	56(%rsp)
	.cfi_def_cfa_offset 184
	leaq	_Z22angle_between_galaxiesPfS_S_S_Pi(%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
.L41:
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L43
	call	__stack_chk_fail@PLT
.L43:
	addq	$168, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3804:
	.size	_Z50__device_stub__Z22angle_between_galaxiesPfS_S_S_PiPfS_S_S_Pi, .-_Z50__device_stub__Z22angle_between_galaxiesPfS_S_S_PiPfS_S_S_Pi
	.globl	_Z22angle_between_galaxiesPfS_S_S_Pi
	.type	_Z22angle_between_galaxiesPfS_S_S_Pi, @function
_Z22angle_between_galaxiesPfS_S_S_Pi:
.LFB3805:
	.cfi_startproc
	endbr64
	jmp	_Z50__device_stub__Z22angle_between_galaxiesPfS_S_S_PiPfS_S_S_Pi
	.cfi_endproc
.LFE3805:
	.size	_Z22angle_between_galaxiesPfS_S_S_Pi, .-_Z22angle_between_galaxiesPfS_S_S_Pi
	.globl	_Z19calculate_histogram10GalaxyFileS_
	.type	_Z19calculate_histogram10GalaxyFileS_, @function
_Z19calculate_histogram10GalaxyFileS_:
.LFB3772:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movl	$2880, %edi
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	subq	$80, %rsp
	.cfi_def_cfa_offset 112
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	movl	112(%rsp), %ebp
	call	malloc@PLT
	movq	%rax, %r12
	movslq	%ebp, %rax
	salq	$2, %rax
	js	.L47
	cvtsi2ssq	%rax, %xmm0
	jmp	.L48
.L47:
	shrq	%rax
	cvtsi2ssq	%rax, %xmm0
	addss	%xmm0, %xmm0
.L48:
	comiss	.LC7(%rip), %xmm0
	jnb	.L49
	cvttss2siq	%xmm0, %rbx
	jmp	.L50
.L49:
	subss	.LC7(%rip), %xmm0
	cvttss2siq	%xmm0, %rbx
	btcq	$63, %rbx
.L50:
	leaq	8(%rsp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	leaq	16(%rsp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	leaq	24(%rsp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	leaq	32(%rsp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	leaq	40(%rsp), %rdi
	movl	$2880, %esi
	call	cudaMalloc@PLT
	movq	120(%rsp), %rsi
	movq	8(%rsp), %rdi
	movq	%rbx, %rdx
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movq	128(%rsp), %rsi
	movq	16(%rsp), %rdi
	movq	%rbx, %rdx
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movq	144(%rsp), %rsi
	movq	24(%rsp), %rdi
	movq	%rbx, %rdx
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movq	152(%rsp), %rsi
	movq	32(%rsp), %rdi
	movq	%rbx, %rdx
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movq	40(%rsp), %rdi
	movq	%rbx, %rdx
	movq	%r12, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	leal	352(%rbp), %eax
	movl	$352, %ecx
	xorl	%esi, %esi
	cltd
	movq	40(%rsp), %rdi
	movl	$1, 68(%rsp)
	idivl	%ecx
	movl	$2880, %edx
	movl	%eax, 48(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 52(%rsp)
	addq	$351, %rax
	movq	%rax, 60(%rsp)
	call	cudaMemset@PLT
	movl	68(%rsp), %ecx
	movl	56(%rsp), %esi
	xorl	%r9d, %r9d
	movq	60(%rsp), %rdx
	movq	48(%rsp), %rdi
	xorl	%r8d, %r8d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L51
	movq	40(%rsp), %r8
	movq	32(%rsp), %rcx
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rsi
	movq	8(%rsp), %rdi
	call	_Z50__device_stub__Z22angle_between_galaxiesPfS_S_S_PiPfS_S_S_Pi
.L51:
	movq	40(%rsp), %rsi
	movl	$2, %ecx
	movl	$2880, %edx
	movq	%r12, %rdi
	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	40(%rsp), %rdi
	call	cudaFree@PLT
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L52
	call	__stack_chk_fail@PLT
.L52:
	addq	$80, %rsp
	.cfi_def_cfa_offset 32
	movq	%r12, %rax
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3772:
	.size	_Z19calculate_histogram10GalaxyFileS_, .-_Z19calculate_histogram10GalaxyFileS_
	.section	.rodata.str1.1
.LC8:
	.string	"_Z22angle_between_galaxiesPfS_S_S_Pi"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3807:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC8(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z22angle_between_galaxiesPfS_S_S_Pi(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE3807:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.rodata._ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.str1.1,"aMS",@progbits,1
.LC9:
	.string	"basic_string: construction from null is not valid"
	.section	.text._ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_,"axG",@progbits,_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC5IS3_EEPKcRKS3_,comdat
	.align 2
	.weak	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_
	.type	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_, @function
_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_:
.LFB4125:
	.cfi_startproc
	endbr64
	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
	subq	$16, %rsp
	.cfi_def_cfa_offset 48
	movq	%fs:40, %rax
	movq	%rax, 8(%rsp)
	xorl	%eax, %eax
	leaq	16(%rdi), %rax
	movq	%rax, (%rdi)
	testq	%rsi, %rsi
	jne	.L57
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L62
	leaq	.LC9(%rip), %rdi
	call	_ZSt19__throw_logic_errorPKc@PLT
.L57:
	movq	%rdi, %rbx
	movq	%rsi, %rdi
	movq	%rsi, %rbp
	call	strlen@PLT
	movq	%rax, (%rsp)
	leaq	0(%rbp,%rax), %r12
	cmpq	$15, %rax
	jbe	.L59
	movq	%rsp, %rsi
	xorl	%edx, %edx
	movq	%rbx, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE9_M_createERmm@PLT
	movq	%rax, (%rbx)
	movq	(%rsp), %rax
	movq	%rax, 16(%rbx)
.L59:
	movq	(%rbx), %rdi
	movq	%r12, %rdx
	movq	%rbp, %rsi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE13_S_copy_charsEPcPKcS7_@PLT
	movq	(%rsp), %rax
	movq	(%rbx), %rdx
	movq	%rax, 8(%rbx)
	movb	$0, (%rdx,%rax)
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	je	.L60
.L62:
	call	__stack_chk_fail@PLT
.L60:
	addq	$16, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4125:
	.size	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_, .-_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_
	.weak	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC1IS3_EEPKcRKS3_
	.set	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC1IS3_EEPKcRKS3_,_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_
	.text
	.globl	_Z8readFileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE
	.type	_Z8readFileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE, @function
_Z8readFileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE:
.LFB3775:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA3775
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movl	$8, %edx
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	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
	movq	%rdi, %rbx
	subq	$544, %rsp
	.cfi_def_cfa_offset 592
	movq	%fs:40, %rax
	movq	%rax, 536(%rsp)
	xorl	%eax, %eax
	leaq	16(%rsp), %rbp
	movq	%rbp, %rdi
.LEHB0:
	call	_ZNSt14basic_ifstreamIcSt11char_traitsIcEEC1ERKNSt7__cxx1112basic_stringIcS1_SaIcEEESt13_Ios_Openmode@PLT
.LEHE0:
	leaq	4(%rsp), %rsi
	movq	%rbp, %rdi
.LEHB1:
	call	_ZNSirsERi@PLT
	movslq	4(%rsp), %rax
	salq	$2, %rax
	js	.L64
	cvtsi2ssq	%rax, %xmm0
	jmp	.L65
.L64:
	shrq	%rax
	cvtsi2ssq	%rax, %xmm0
	addss	%xmm0, %xmm0
.L65:
	comiss	.LC7(%rip), %xmm0
	jnb	.L66
	cvttss2siq	%xmm0, %r12
	jmp	.L67
.L66:
	subss	.LC7(%rip), %xmm0
	cvttss2siq	%xmm0, %r12
	btcq	$63, %r12
.L67:
	movq	%r12, %rdi
	call	malloc@PLT
	movq	%r12, %rdi
	xorl	%r12d, %r12d
	movq	%rax, %r14
	call	malloc@PLT
	movq	%rax, %r13
.L68:
	cmpl	%r12d, 4(%rsp)
	jle	.L76
	leaq	8(%rsp), %rsi
	movq	%rbp, %rdi
	call	_ZNSi10_M_extractIfEERSiRT_@PLT
	movq	%rax, %rdi
	leaq	12(%rsp), %rsi
	call	_ZNSi10_M_extractIfEERSiRT_@PLT
	movss	.LC10(%rip), %xmm0
	mulss	8(%rsp), %xmm0
	movss	%xmm0, (%r14,%r12,4)
	movss	.LC10(%rip), %xmm0
	mulss	12(%rsp), %xmm0
	movss	%xmm0, 0(%r13,%r12,4)
	incq	%r12
	jmp	.L68
.L76:
	movq	%rbp, %rdi
	call	_ZNSt14basic_ifstreamIcSt11char_traitsIcEE5closeEv@PLT
.LEHE1:
	movl	4(%rsp), %eax
	movq	%rbp, %rdi
	movq	%r14, 8(%rbx)
	movq	%r13, 16(%rbx)
	movl	%eax, (%rbx)
	call	_ZNSt14basic_ifstreamIcSt11char_traitsIcEED1Ev@PLT
	movq	536(%rsp), %rax
	subq	%fs:40, %rax
	je	.L72
	jmp	.L74
.L73:
	endbr64
	movq	%rax, %rbx
.L70:
	movq	%rbp, %rdi
	call	_ZNSt14basic_ifstreamIcSt11char_traitsIcEED1Ev@PLT
	movq	536(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L74
	movq	%rbx, %rdi
.LEHB2:
	call	_Unwind_Resume@PLT
.LEHE2:
.L74:
	call	__stack_chk_fail@PLT
.L72:
	addq	$544, %rsp
	.cfi_def_cfa_offset 48
	movq	%rbx, %rax
	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
.LFE3775:
	.globl	__gxx_personality_v0
	.section	.gcc_except_table,"a",@progbits
.LLSDA3775:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE3775-.LLSDACSB3775
.LLSDACSB3775:
	.uleb128 .LEHB0-.LFB3775
	.uleb128 .LEHE0-.LEHB0
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB1-.LFB3775
	.uleb128 .LEHE1-.LEHB1
	.uleb128 .L73-.LFB3775
	.uleb128 0
	.uleb128 .LEHB2-.LFB3775
	.uleb128 .LEHE2-.LEHB2
	.uleb128 0
	.uleb128 0
.LLSDACSE3775:
	.text
	.size	_Z8readFileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE, .-_Z8readFileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE
	.section	.text._ZStplIcSt11char_traitsIcESaIcEENSt7__cxx1112basic_stringIT_T0_T1_EEPKS5_RKS8_,"axG",@progbits,_ZStplIcSt11char_traitsIcESaIcEENSt7__cxx1112basic_stringIT_T0_T1_EEPKS5_RKS8_,comdat
	.weak	_ZStplIcSt11char_traitsIcESaIcEENSt7__cxx1112basic_stringIT_T0_T1_EEPKS5_RKS8_
	.type	_ZStplIcSt11char_traitsIcESaIcEENSt7__cxx1112basic_stringIT_T0_T1_EEPKS5_RKS8_, @function
_ZStplIcSt11char_traitsIcESaIcEENSt7__cxx1112basic_stringIT_T0_T1_EEPKS5_RKS8_:
.LFB4148:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4148
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	movq	%rsi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	movq	%rdi, %rbx
	movq	%rsi, %rdi
	movq	8(%rdx), %r13
	movq	(%rdx), %r14
	call	strlen@PLT
	movb	$0, 16(%rbx)
	movq	%rbx, %rdi
	movq	%rax, %r12
	leaq	16(%rbx), %rax
	movq	%rax, (%rbx)
	xorl	%eax, %eax
	leaq	0(%r13,%r12), %rsi
	movq	%rax, 8(%rbx)
.LEHB3:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE7reserveEm@PLT
	movq	%r12, %rdx
	movq	%rbp, %rsi
	movq	%rbx, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE6appendEPKcm.isra.0
	movq	%r13, %rdx
	movq	%r14, %rsi
	movq	%rbx, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE6appendEPKcm.isra.0
.LEHE3:
	movq	%rbx, %rax
	popq	%rbx
	.cfi_remember_state
	.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
.L80:
	.cfi_restore_state
	endbr64
	movq	%rax, %rbp
.L79:
	movq	%rbx, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	movq	%rbp, %rdi
.LEHB4:
	call	_Unwind_Resume@PLT
.LEHE4:
	.cfi_endproc
.LFE4148:
	.section	.gcc_except_table
.LLSDA4148:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4148-.LLSDACSB4148
.LLSDACSB4148:
	.uleb128 .LEHB3-.LFB4148
	.uleb128 .LEHE3-.LEHB3
	.uleb128 .L80-.LFB4148
	.uleb128 0
	.uleb128 .LEHB4-.LFB4148
	.uleb128 .LEHE4-.LEHB4
	.uleb128 0
	.uleb128 0
.LLSDACSE4148:
	.section	.text._ZStplIcSt11char_traitsIcESaIcEENSt7__cxx1112basic_stringIT_T0_T1_EEPKS5_RKS8_,"axG",@progbits,_ZStplIcSt11char_traitsIcESaIcEENSt7__cxx1112basic_stringIT_T0_T1_EEPKS5_RKS8_,comdat
	.size	_ZStplIcSt11char_traitsIcESaIcEENSt7__cxx1112basic_stringIT_T0_T1_EEPKS5_RKS8_, .-_ZStplIcSt11char_traitsIcESaIcEENSt7__cxx1112basic_stringIT_T0_T1_EEPKS5_RKS8_
	.section	.rodata.str1.1
.LC11:
	.string	"output/"
.LC12:
	.string	"\n"
	.text
	.globl	_Z20write_omegas_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPf
	.type	_Z20write_omegas_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPf, @function
_Z20write_omegas_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPf:
.LFB3778:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA3778
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movq	%rsi, %r13
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movq	%rdi, %rbx
	subq	$568, %rsp
	.cfi_def_cfa_offset 608
	movq	%fs:40, %rax
	movq	%rax, 552(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rbp
	leaq	8(%rsp), %r12
	movq	%rbp, %rdi
.LEHB5:
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEEC1Ev@PLT
.LEHE5:
	movq	%rbx, %rdx
	leaq	.LC11(%rip), %rsi
	movq	%r12, %rdi
.LEHB6:
	call	_ZStplIcSt11char_traitsIcESaIcEENSt7__cxx1112basic_stringIT_T0_T1_EEPKS5_RKS8_
.LEHE6:
	movq	8(%rsp), %rsi
	movl	$16, %edx
	movq	%rbp, %rdi
.LEHB7:
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEE4openERKNSt7__cxx1112basic_stringIcS1_SaIcEEESt13_Ios_Openmode.isra.0
.LEHE7:
	movq	%r12, %rdi
	xorl	%ebx, %ebx
	leaq	.LC12(%rip), %r12
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
.L86:
	movq	%rbp, %rdi
	cvtss2sd	0(%r13,%rbx,4), %xmm0
.LEHB8:
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	cmpq	$719, %rbx
	jne	.L84
	movq	%rbp, %rdi
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEE5closeEv@PLT
	jmp	.L96
.L84:
	movq	%r12, %rsi
	movq	%rbp, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
.LEHE8:
	incq	%rbx
	jmp	.L86
.L96:
	movq	%rbp, %rdi
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEED1Ev@PLT
	movq	552(%rsp), %rax
	subq	%fs:40, %rax
	je	.L90
	jmp	.L94
.L91:
	endbr64
	movq	%rax, %rbx
.L87:
	movq	%r12, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	jmp	.L88
.L92:
	endbr64
	movq	%rax, %rbx
.L88:
	movq	%rbp, %rdi
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEED1Ev@PLT
	movq	552(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L94
	movq	%rbx, %rdi
.LEHB9:
	call	_Unwind_Resume@PLT
.LEHE9:
.L94:
	call	__stack_chk_fail@PLT
.L90:
	addq	$568, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3778:
	.section	.gcc_except_table
.LLSDA3778:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE3778-.LLSDACSB3778
.LLSDACSB3778:
	.uleb128 .LEHB5-.LFB3778
	.uleb128 .LEHE5-.LEHB5
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB6-.LFB3778
	.uleb128 .LEHE6-.LEHB6
	.uleb128 .L92-.LFB3778
	.uleb128 0
	.uleb128 .LEHB7-.LFB3778
	.uleb128 .LEHE7-.LEHB7
	.uleb128 .L91-.LFB3778
	.uleb128 0
	.uleb128 .LEHB8-.LFB3778
	.uleb128 .LEHE8-.LEHB8
	.uleb128 .L92-.LFB3778
	.uleb128 0
	.uleb128 .LEHB9-.LFB3778
	.uleb128 .LEHE9-.LEHB9
	.uleb128 0
	.uleb128 0
.LLSDACSE3778:
	.text
	.size	_Z20write_omegas_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPf, .-_Z20write_omegas_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPf
	.globl	_Z23write_histogram_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPi
	.type	_Z23write_histogram_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPi, @function
_Z23write_histogram_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPi:
.LFB3779:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA3779
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movq	%rsi, %r13
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movq	%rdi, %rbx
	subq	$568, %rsp
	.cfi_def_cfa_offset 608
	movq	%fs:40, %rax
	movq	%rax, 552(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rbp
	leaq	8(%rsp), %r12
	movq	%rbp, %rdi
.LEHB10:
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEEC1Ev@PLT
.LEHE10:
	movq	%rbx, %rdx
	leaq	.LC11(%rip), %rsi
	movq	%r12, %rdi
.LEHB11:
	call	_ZStplIcSt11char_traitsIcESaIcEENSt7__cxx1112basic_stringIT_T0_T1_EEPKS5_RKS8_
.LEHE11:
	movq	8(%rsp), %rsi
	movl	$16, %edx
	movq	%rbp, %rdi
.LEHB12:
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEE4openERKNSt7__cxx1112basic_stringIcS1_SaIcEEESt13_Ios_Openmode.isra.0
.LEHE12:
	movq	%r12, %rdi
	xorl	%ebx, %ebx
	leaq	.LC12(%rip), %r12
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
.L100:
	movl	0(%r13,%rbx,4), %esi
	movq	%rbp, %rdi
.LEHB13:
	call	_ZNSolsEi@PLT
	cmpq	$719, %rbx
	jne	.L98
	movq	%rbp, %rdi
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEE5closeEv@PLT
	jmp	.L110
.L98:
	movq	%r12, %rsi
	movq	%rbp, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
.LEHE13:
	incq	%rbx
	jmp	.L100
.L110:
	movq	%rbp, %rdi
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEED1Ev@PLT
	movq	552(%rsp), %rax
	subq	%fs:40, %rax
	je	.L104
	jmp	.L108
.L105:
	endbr64
	movq	%rax, %rbx
.L101:
	movq	%r12, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	jmp	.L102
.L106:
	endbr64
	movq	%rax, %rbx
.L102:
	movq	%rbp, %rdi
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEED1Ev@PLT
	movq	552(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L108
	movq	%rbx, %rdi
.LEHB14:
	call	_Unwind_Resume@PLT
.LEHE14:
.L108:
	call	__stack_chk_fail@PLT
.L104:
	addq	$568, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3779:
	.section	.gcc_except_table
.LLSDA3779:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE3779-.LLSDACSB3779
.LLSDACSB3779:
	.uleb128 .LEHB10-.LFB3779
	.uleb128 .LEHE10-.LEHB10
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB11-.LFB3779
	.uleb128 .LEHE11-.LEHB11
	.uleb128 .L106-.LFB3779
	.uleb128 0
	.uleb128 .LEHB12-.LFB3779
	.uleb128 .LEHE12-.LEHB12
	.uleb128 .L105-.LFB3779
	.uleb128 0
	.uleb128 .LEHB13-.LFB3779
	.uleb128 .LEHE13-.LEHB13
	.uleb128 .L106-.LFB3779
	.uleb128 0
	.uleb128 .LEHB14-.LFB3779
	.uleb128 .LEHE14-.LEHB14
	.uleb128 0
	.uleb128 0
.LLSDACSE3779:
	.text
	.size	_Z23write_histogram_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPi, .-_Z23write_histogram_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPi
	.section	.rodata.str1.1
.LC13:
	.string	"test_data/flat_100k_arcmin.txt"
.LC14:
	.string	"test_data/data_100k_arcmin.txt"
.LC15:
	.string	"DD"
.LC16:
	.string	"DR"
.LC17:
	.string	"RR"
.LC18:
	.string	"dd_histogram.txt"
.LC19:
	.string	"dr_histogram.txt"
.LC20:
	.string	"rr_histogram.txt"
.LC21:
	.string	"omegas.txt"
	.section	.text.startup
	.globl	main
	.type	main, @function
main:
.LFB3774:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA3774
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	leaq	.LC13(%rip), %rsi
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	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
	subq	$96, %rsp
	.cfi_def_cfa_offset 144
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rbx
	leaq	32(%rsp), %r12
	movq	%rbx, %rdi
	movq	%r12, %rdx
.LEHB15:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC1IS3_EEPKcRKS3_
.LEHE15:
	leaq	8(%rsp), %rdi
	movq	%rbx, %rsi
.LEHB16:
	call	_Z8readFileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE
.LEHE16:
	movq	%rbx, %rdi
	leaq	7(%rsp), %rbp
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	leaq	.LC14(%rip), %rsi
	movq	%rbx, %rdi
	movq	%rbp, %rdx
.LEHB17:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC1IS3_EEPKcRKS3_
.LEHE17:
	movq	%rbx, %rsi
	movq	%r12, %rdi
.LEHB18:
	call	_Z8readFileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE
.LEHE18:
	movq	%rbx, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	subq	$48, %rsp
	.cfi_def_cfa_offset 192
	movups	56(%rsp), %xmm0
	movq	72(%rsp), %rax
	movups	%xmm0, 24(%rsp)
	movups	%xmm0, (%rsp)
	movq	%rax, 40(%rsp)
	movq	%rax, 16(%rsp)
.LEHB19:
	call	_Z19calculate_histogram10GalaxyFileS_
	movups	80(%rsp), %xmm1
	movups	56(%rsp), %xmm2
	movq	%rax, %r12
	movq	96(%rsp), %rax
	movups	%xmm1, 24(%rsp)
	movq	%rax, 40(%rsp)
	movq	72(%rsp), %rax
	movups	%xmm2, (%rsp)
	movq	%rax, 16(%rsp)
	call	_Z19calculate_histogram10GalaxyFileS_
	movups	80(%rsp), %xmm0
	movq	%rax, %r13
	movq	96(%rsp), %rax
	movups	%xmm0, 24(%rsp)
	movups	%xmm0, (%rsp)
	movq	%rax, 40(%rsp)
	movq	%rax, 16(%rsp)
	call	_Z19calculate_histogram10GalaxyFileS_
	movq	%rbp, %rdx
	movq	%rbx, %rdi
	addq	$48, %rsp
	.cfi_def_cfa_offset 144
	leaq	.LC15(%rip), %rsi
	movq	%rax, %r14
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC1IS3_EEPKcRKS3_
.LEHE19:
	movl	$20, %edx
	movq	%r12, %rsi
	movq	%rbx, %rdi
.LEHB20:
	call	_Z15print_histogramNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPii
.LEHE20:
	movq	%rbx, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	movq	%rbp, %rdx
	leaq	.LC16(%rip), %rsi
	movq	%rbx, %rdi
.LEHB21:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC1IS3_EEPKcRKS3_
.LEHE21:
	movl	$20, %edx
	movq	%r13, %rsi
	movq	%rbx, %rdi
.LEHB22:
	call	_Z15print_histogramNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPii
.LEHE22:
	movq	%rbx, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	movq	%rbp, %rdx
	leaq	.LC17(%rip), %rsi
	movq	%rbx, %rdi
.LEHB23:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC1IS3_EEPKcRKS3_
.LEHE23:
	movl	$20, %edx
	movq	%r14, %rsi
	movq	%rbx, %rdi
.LEHB24:
	call	_Z15print_histogramNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPii
.LEHE24:
	movq	%rbx, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	leaq	.LC18(%rip), %rsi
	movq	%rbx, %rdi
	movq	%rbp, %rdx
.LEHB25:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC1IS3_EEPKcRKS3_
.LEHE25:
	movq	%r12, %rsi
	movq	%rbx, %rdi
.LEHB26:
	call	_Z23write_histogram_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPi
.LEHE26:
	movq	%rbx, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	leaq	.LC19(%rip), %rsi
	movq	%rbx, %rdi
	movq	%rbp, %rdx
.LEHB27:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC1IS3_EEPKcRKS3_
.LEHE27:
	movq	%r13, %rsi
	movq	%rbx, %rdi
.LEHB28:
	call	_Z23write_histogram_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPi
.LEHE28:
	movq	%rbx, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	leaq	.LC20(%rip), %rsi
	movq	%rbx, %rdi
	movq	%rbp, %rdx
.LEHB29:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC1IS3_EEPKcRKS3_
.LEHE29:
	movq	%r14, %rsi
	movq	%rbx, %rdi
.LEHB30:
	call	_Z23write_histogram_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPi
.LEHE30:
	movq	%rbx, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	movq	%r14, %rdx
	movq	%r12, %rdi
	movq	%r13, %rsi
	call	_Z16calculate_omegasPiS_S_
	movl	$15, %esi
	movq	%rax, %rdi
	movq	%rax, %r12
.LEHB31:
	call	_Z12print_omegasPfi
	leaq	.LC21(%rip), %rsi
	movq	%rbx, %rdi
	movq	%rbp, %rdx
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC1IS3_EEPKcRKS3_
.LEHE31:
	movq	%r12, %rsi
	movq	%rbx, %rdi
.LEHB32:
	call	_Z20write_omegas_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPf
.LEHE32:
	movq	%rbx, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L130
	jmp	.L140
.L131:
	endbr64
	jmp	.L143
.L132:
	endbr64
	jmp	.L143
.L133:
	endbr64
	jmp	.L143
.L134:
	endbr64
	jmp	.L143
.L135:
	endbr64
	jmp	.L143
.L136:
	endbr64
	jmp	.L143
.L137:
	endbr64
	jmp	.L143
.L138:
	endbr64
	jmp	.L143
.L139:
	endbr64
.L143:
	movq	%rax, %rbp
.L128:
	movq	%rbx, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L140
	movq	%rbp, %rdi
.LEHB33:
	call	_Unwind_Resume@PLT
.LEHE33:
.L140:
	call	__stack_chk_fail@PLT
.L130:
	addq	$96, %rsp
	.cfi_def_cfa_offset 48
	xorl	%eax, %eax
	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
.LFE3774:
	.section	.gcc_except_table
.LLSDA3774:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE3774-.LLSDACSB3774
.LLSDACSB3774:
	.uleb128 .LEHB15-.LFB3774
	.uleb128 .LEHE15-.LEHB15
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB16-.LFB3774
	.uleb128 .LEHE16-.LEHB16
	.uleb128 .L131-.LFB3774
	.uleb128 0
	.uleb128 .LEHB17-.LFB3774
	.uleb128 .LEHE17-.LEHB17
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB18-.LFB3774
	.uleb128 .LEHE18-.LEHB18
	.uleb128 .L132-.LFB3774
	.uleb128 0
	.uleb128 .LEHB19-.LFB3774
	.uleb128 .LEHE19-.LEHB19
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB20-.LFB3774
	.uleb128 .LEHE20-.LEHB20
	.uleb128 .L133-.LFB3774
	.uleb128 0
	.uleb128 .LEHB21-.LFB3774
	.uleb128 .LEHE21-.LEHB21
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB22-.LFB3774
	.uleb128 .LEHE22-.LEHB22
	.uleb128 .L134-.LFB3774
	.uleb128 0
	.uleb128 .LEHB23-.LFB3774
	.uleb128 .LEHE23-.LEHB23
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB24-.LFB3774
	.uleb128 .LEHE24-.LEHB24
	.uleb128 .L135-.LFB3774
	.uleb128 0
	.uleb128 .LEHB25-.LFB3774
	.uleb128 .LEHE25-.LEHB25
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB26-.LFB3774
	.uleb128 .LEHE26-.LEHB26
	.uleb128 .L136-.LFB3774
	.uleb128 0
	.uleb128 .LEHB27-.LFB3774
	.uleb128 .LEHE27-.LEHB27
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB28-.LFB3774
	.uleb128 .LEHE28-.LEHB28
	.uleb128 .L137-.LFB3774
	.uleb128 0
	.uleb128 .LEHB29-.LFB3774
	.uleb128 .LEHE29-.LEHB29
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB30-.LFB3774
	.uleb128 .LEHE30-.LEHB30
	.uleb128 .L138-.LFB3774
	.uleb128 0
	.uleb128 .LEHB31-.LFB3774
	.uleb128 .LEHE31-.LEHB31
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB32-.LFB3774
	.uleb128 .LEHE32-.LEHB32
	.uleb128 .L139-.LFB3774
	.uleb128 0
	.uleb128 .LEHB33-.LFB3774
	.uleb128 .LEHE33-.LEHB33
	.uleb128 0
	.uleb128 0
.LLSDACSE3774:
	.section	.text.startup
	.size	main, .-main
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC3:
	.long	1048576000
	.align 4
.LC7:
	.long	1593835520
	.align 4
.LC10:
	.long	966296155
	.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:

/*
    CUDA good to knows:
        Basics:
            Per thread:
                registers (fast)
                local memory (off-chip [still on the GPU though], slow)

            Per block:
                multiple threads
                shared memory (semi-fast)

            Per GPU:
                Multiple kernels that each run multiple blocks
                Global memory (off-chip [still on the GPU though], slow)


            Threads are executed by thread processors
            
            Threads reside in thread blocks
            
            Thread blocks are executed by multiprocessors

            Several concurrent thread blocks can reside on one multiprocessor
                Limited by multiprocessor resources (shared memory and registers)

            A kernel is launched as a grid of thread blocks. Only one kernel can execute on a device at a time.
        Advanced:
            cudaMemcpy(dst, src, size, direction)
                blocks CPU thread.
            

    Compiler tips:
        nvcc <filename>.cu [-o <executable>]
            Builds release mode

        nvcc -g <filename>.cu
            Builds debug mode
            Can debug host code but not device code

        nvcc -deviceemu <filename>.cu
            Builds device emulation mode
            All code runs on CPU, no debug symbols

        nvcc -deviceemu -g <filename>.cu
            Builds debug device emulation mode
            All code runs on CPU, with debug symbols

    Tips and tricks:
        If our arrays A,B,C are shorter than 1024 elements, N < 1024, then
            – one thread block is enough
            – N threads in the thread block
        If our arrays are longer than 1024, then
            – Choose the number of threads in the thread blocks to be
            integer*32
            – Calculate how many thread blocks you need
            – There will be some threads that should do nothing
        Why multiples of 32?
            – Threads are executed synchronously in bunches of 32 =
            warp
            – All threads must have their data ready before the warp runs
            – Cache lines are 4 B x warp size = 128 B
            – GPU resources can be fully utilized when these parameters
            are used
        # of blocks = ceil(N/threadsInBlock)
                    = (N+threadsInBlock-1)/threadsInBlock



Compile:
nvcc -o galaxy galaxy_program.cu -res-usage

Run:
time ./galaxy

*/

#include <hip/hip_runtime.h>
#include <stdio.h>
#include <iostream>
#include <fstream>

using namespace std;


// Declare functions and classes that are below main.
class GalaxyFile{
public:
    int number_of_galaxies;
    float *alphas, *deltas;

    GalaxyFile(){}

    GalaxyFile(int num, float *as, float *ds)
    {
        number_of_galaxies = num;
        alphas = as;
        deltas = ds;
    }
};
void print_omegas(float*, int);
void write_omegas_to_file(string, float*);
void write_histogram_to_file(string, int*);
void print_histogram(string, int*, int);
GalaxyFile readFile(string);


// Define some useful macros
#define BIN_WIDTH 0.25f
#define BIN_MIN 0.0f
#define BIN_MAX 180.0f
#define NUMBER_OF_BINS (int)(BIN_MAX*(1.0f/BIN_WIDTH))

// Google is your friend.
#define ARCMINS_TO_RADIANS 0.000290888209f
#define RADIANS_TO_DEGREES 57.295779513f


__global__
void angle_between_galaxies(float *alphas1, float *deltas1, float *alphas2, float *deltas2, int *gpu_hist){
    int idx = blockDim.x*blockIdx.x + threadIdx.x;
    
    if(idx < 100000){
        for(int i=0; i<100000; i++){
                float angle = 0.0f;
                // Don't do duplicates
                if( alphas1[i] != alphas2[idx] && deltas1[i] != deltas2[idx] ) {
                    float x = sin(deltas1[i]) * sin(deltas2[idx]) + cos(deltas1[i]) * cos(deltas2[idx]) * cos(alphas1[i] - alphas2[idx]);
                    angle = acosf(fmaxf(-1.0f, fminf(x, 1.0f))) * RADIANS_TO_DEGREES;
                }

                int ix = (int)(floor(angle * (1.0f/BIN_WIDTH))) % NUMBER_OF_BINS;
                __syncthreads();
                atomicAdd(&gpu_hist[ix], 1);
        }
    }
}

int* calculate_histogram(GalaxyFile galaxies1, GalaxyFile galaxies2){
    // Declare and allocate memory for histogram arrays that will be accessible on CPU
    float galaxy_array_size = galaxies1.number_of_galaxies * sizeof(float);
    float histogram_size = NUMBER_OF_BINS * sizeof(int);

    int *histogram;
    int *total_histogram;
    histogram = (int *) malloc(NUMBER_OF_BINS*sizeof(int));
    total_histogram = (int *) malloc(NUMBER_OF_BINS*sizeof(int));

    memset(total_histogram, 0, NUMBER_OF_BINS*sizeof(int));

    // Declare angle arrays that will be accessible on GPU
    float *gpu_alphas1;
    float *gpu_deltas1;
    float *gpu_alphas2;
    float *gpu_deltas2;
    int *gpu_histogram;

    // Allocate memory on GPU for angle arrays
    hipMalloc((void**) &gpu_alphas1, galaxy_array_size);
    hipMalloc((void**) &gpu_deltas1, galaxy_array_size);
    hipMalloc((void**) &gpu_alphas2, galaxy_array_size);
    hipMalloc((void**) &gpu_deltas2, galaxy_array_size);
    hipMalloc((void**) &gpu_histogram, NUMBER_OF_BINS*sizeof(int));

	// Copy angles from CPU onto GPU
	hipMemcpy(gpu_alphas1, galaxies1.alphas, galaxy_array_size, hipMemcpyHostToDevice);
	hipMemcpy(gpu_deltas1, galaxies1.deltas, galaxy_array_size, hipMemcpyHostToDevice);
	hipMemcpy(gpu_alphas2, galaxies2.alphas, galaxy_array_size, hipMemcpyHostToDevice);
	hipMemcpy(gpu_deltas2, galaxies2.deltas, galaxy_array_size, hipMemcpyHostToDevice);
	hipMemcpy(gpu_histogram, histogram, galaxy_array_size, hipMemcpyHostToDevice);
    
    int warp_size = 32;
    int threadsInBlock = 11 * warp_size;
    int blocksInGrid = ceil((galaxies1.number_of_galaxies + threadsInBlock) / threadsInBlock);

	// Define the grid size (blocks per grid)
    dim3 dimGrid(blocksInGrid);
    
	// Define block size (threads per block)
	dim3 dimBlock(threadsInBlock);

    // Write histogram full of zeros
    hipMemset(gpu_histogram, 0, histogram_size);

    // Calculate angles between galaxies1[i] and every galaxy in galaxies2
    angle_between_galaxies<<<dimGrid, dimBlock>>>(gpu_alphas1, gpu_deltas1, gpu_alphas2, gpu_deltas2, gpu_histogram);
    
    // Copy result histogram into CPU histogram
    hipMemcpy(histogram, gpu_histogram, histogram_size, hipMemcpyDeviceToHost);
    
	// Free all the memory we allocated on the GPU
	hipFree( gpu_alphas1 );
	hipFree( gpu_deltas1 );
	hipFree( gpu_alphas2 );
	hipFree( gpu_deltas2 );
    hipFree( gpu_histogram );

    return histogram;
}

float* calculate_omegas(int* DD, int* DR, int* RR){
    float* omegas;
    omegas = (float *) malloc(NUMBER_OF_BINS*sizeof(float));

    for(int i=0; i<NUMBER_OF_BINS; i++){
        if(RR[i] != 0.0f){
            omegas[i] = (DD[i] - 2.0f*DR[i] + RR[i]) / RR[i];
        }else{
            omegas[i] = 0.0f;
        }
    }
    return omegas;
}

// CUDA program that calculates distribution of galaxies
int main()
{
    // Read files and store data in GalaxyFile classes.
    GalaxyFile galaxies1;
    GalaxyFile galaxies2;
    galaxies1 = readFile("test_data/flat_100k_arcmin.txt");
    galaxies2 = readFile("test_data/data_100k_arcmin.txt");

    int* DD_hist = calculate_histogram(galaxies1, galaxies1);
    int* DR_hist = calculate_histogram(galaxies1, galaxies2);
    int* RR_hist = calculate_histogram(galaxies2, galaxies2);

    print_histogram("DD", DD_hist, 20);
    print_histogram("DR", DR_hist, 20);
    print_histogram("RR", RR_hist, 20);

    write_histogram_to_file("dd_histogram.txt", DD_hist);
    write_histogram_to_file("dr_histogram.txt", DR_hist);
    write_histogram_to_file("rr_histogram.txt", RR_hist);

    float* omegas = calculate_omegas(DD_hist, DR_hist, RR_hist);
    print_omegas(omegas, 15);
    write_omegas_to_file("omegas.txt", omegas);

	return EXIT_SUCCESS;
}


/* UTILITY FUNCTIONS/CLASSES BELOW */
GalaxyFile readFile(string filename)
{
    ifstream infile(filename);
    int number_of_galaxies;

    // Read first line which is the number of galaxies that's stored in the file.
    infile >> number_of_galaxies;

    float galaxy_array_size = number_of_galaxies * sizeof(float);

    float *alphas, *deltas;
    alphas = (float*) malloc(galaxy_array_size);
    deltas = (float*) malloc(galaxy_array_size);

    float alpha;
    float delta;

    // Read arc minute angles for each galaxy
    // Then convert those angles to radians and store those in alphas and deltas arrays
    for(int i=0; i<number_of_galaxies; i++) {
        infile >> alpha >> delta;

        alphas[i] = alpha * ARCMINS_TO_RADIANS;
        deltas[i] = delta * ARCMINS_TO_RADIANS;
    }
    infile.close();

    GalaxyFile galaxyFile(number_of_galaxies, alphas, deltas);
    return galaxyFile;
}

void print_omegas(float* omegas, int bins_to_print){
    for (int i=0; i<NUMBER_OF_BINS; i++){
        if(omegas[i] != 0.0f && i < bins_to_print){
            printf("omegas[%d]: %f\n", i, omegas[i]);
        }
    }
}

void print_histogram(string label, int *histogram, int bins_to_print){
    long long galaxies_counted = 0;
    // Print each bucket bin that has 1 or more galaxy-pair-angle in it.
    for (int i=0; i<NUMBER_OF_BINS; i++) {
        float bucket_min = (float)i / (1.0f/BIN_WIDTH);
        float bucket_max = (float)i / (1.0f/BIN_WIDTH) + BIN_WIDTH;
        int bucket_value = histogram[i];

        galaxies_counted += histogram[i];
        if(bucket_value > 0 && i < bins_to_print){
            printf("[%f, %f]: %d\n", bucket_min, bucket_max, bucket_value);
        }
    }

    cout << "Galaxy pairs counted in " << label << ": " << galaxies_counted << endl;
}

void write_omegas_to_file(string filename, float* omegas){
    ofstream file;
    file.open("output/"+filename);

    for (int i=0; i<NUMBER_OF_BINS; i++){
        file << omegas[i];
        if(i<NUMBER_OF_BINS-1) file << "\n";
    }
    file.close();
}

void write_histogram_to_file(string filename, int* histogram){
    ofstream file;
    file.open("output/"+filename);
    
    for (int i=0; i<NUMBER_OF_BINS; i++){
        file << histogram[i]; 
        if(i<NUMBER_OF_BINS-1) file << "\n";
    }
    file.close();
}
	.text
	.file	"galaxy_program.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.globl	_Z37__device_stub__angle_between_galaxiesPfS_S_S_Pi # -- Begin function _Z37__device_stub__angle_between_galaxiesPfS_S_S_Pi
	.type	_Z37__device_stub__angle_between_galaxiesPfS_S_S_Pi,@function
_Z37__device_stub__angle_between_galaxiesPfS_S_S_Pi: # @_Z37__device_stub__angle_between_galaxiesPfS_S_S_Pi
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$144, %rsp
	.cfi_def_cfa_offset 192
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	56(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	48(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	24(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z22angle_between_galaxiesPfS_S_S_Pi, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$160, %rsp
	.cfi_adjust_cfa_offset -160
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z37__device_stub__angle_between_galaxiesPfS_S_S_Pi, .Lfunc_end0-_Z37__device_stub__angle_between_galaxiesPfS_S_S_Pi
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _Z19calculate_histogram10GalaxyFileS_
.LCPI1_0:
	.long	0x5f000000                      # float 9.22337203E+18
	.text
	.globl	_Z19calculate_histogram10GalaxyFileS_
	.type	_Z19calculate_histogram10GalaxyFileS_,@function
_Z19calculate_histogram10GalaxyFileS_:  # @_Z19calculate_histogram10GalaxyFileS_
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$88, %rsp
	.cfi_def_cfa_offset 144
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movslq	144(%rsp), %rax
	movq	%rax, 56(%rsp)                  # 8-byte Spill
	shlq	$2, %rax
	testq	%rax, %rax
	js	.LBB1_1
# %bb.2:
	cvtsi2ss	%rax, %xmm0
	jmp	.LBB1_3
.LBB1_1:
	shrq	%rax
	cvtsi2ss	%rax, %xmm0
	addss	%xmm0, %xmm0
.LBB1_3:
	movaps	%xmm0, 64(%rsp)                 # 16-byte Spill
	movl	$2880, %edi                     # imm = 0xB40
	callq	malloc
	movq	%rax, 48(%rsp)                  # 8-byte Spill
	movaps	64(%rsp), %xmm0                 # 16-byte Reload
	cvttss2si	%xmm0, %rax
	movq	%rax, %rcx
	sarq	$63, %rcx
	subss	.LCPI1_0(%rip), %xmm0
	cvttss2si	%xmm0, %r14
	andq	%rcx, %r14
	orq	%rax, %r14
	leaq	40(%rsp), %rbx
	movq	%rbx, %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	leaq	32(%rsp), %rbp
	movq	%rbp, %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	leaq	24(%rsp), %r13
	movq	%r13, %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	leaq	16(%rsp), %r12
	movq	%r12, %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	leaq	8(%rsp), %r15
	movl	$2880, %esi                     # imm = 0xB40
	movq	%r15, %rdi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movq	152(%rsp), %rsi
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%rbp), %rdi
	movq	160(%rsp), %rsi
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r13), %rdi
	movq	176(%rsp), %rsi
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r12), %rdi
	movq	48(%rsp), %r12                  # 8-byte Reload
	movq	184(%rsp), %rsi
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r15), %rdi
	movq	%r12, %rsi
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	56(%rsp), %rax                  # 8-byte Reload
	addl	$352, %eax                      # imm = 0x160
	cltq
	imulq	$780903145, %rax, %rbx          # imm = 0x2E8BA2E9
	movq	%rbx, %rax
	shrq	$63, %rax
	sarq	$38, %rbx
	addl	%eax, %ebx
	btsq	$32, %rbx
	movq	(%r15), %rdi
	movl	$2880, %edx                     # imm = 0xB40
	xorl	%esi, %esi
	callq	hipMemset
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$352, %rdx                      # imm = 0x160
	movq	%rbx, %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_5
# %bb.4:
	movq	40(%rsp), %rdi
	movq	32(%rsp), %rsi
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rcx
	movq	8(%rsp), %r8
	callq	_Z37__device_stub__angle_between_galaxiesPfS_S_S_Pi
.LBB1_5:
	movq	8(%rsp), %rsi
	movl	$2880, %edx                     # imm = 0xB40
	movq	%r12, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	40(%rsp), %rdi
	callq	hipFree
	movq	32(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	%r12, %rax
	addq	$88, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z19calculate_histogram10GalaxyFileS_, .Lfunc_end1-_Z19calculate_histogram10GalaxyFileS_
	.cfi_endproc
                                        # -- End function
	.globl	_Z16calculate_omegasPiS_S_      # -- Begin function _Z16calculate_omegasPiS_S_
	.type	_Z16calculate_omegasPiS_S_,@function
_Z16calculate_omegasPiS_S_:             # @_Z16calculate_omegasPiS_S_
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%rdx, %rbx
	movq	%rsi, %r14
	movq	%rdi, %r15
	movl	$2880, %edi                     # imm = 0xB40
	callq	malloc
	xorl	%ecx, %ecx
.LBB2_1:                                # =>This Inner Loop Header: Depth=1
	movl	(%rbx,%rcx,4), %edx
	testl	%edx, %edx
	je	.LBB2_2
# %bb.3:                                #   in Loop: Header=BB2_1 Depth=1
	xorps	%xmm1, %xmm1
	cvtsi2ss	%edx, %xmm1
	xorps	%xmm0, %xmm0
	cvtsi2ssl	(%r15,%rcx,4), %xmm0
	xorps	%xmm2, %xmm2
	cvtsi2ssl	(%r14,%rcx,4), %xmm2
	addss	%xmm2, %xmm2
	subss	%xmm2, %xmm0
	addss	%xmm1, %xmm0
	divss	%xmm1, %xmm0
	jmp	.LBB2_4
.LBB2_2:                                #   in Loop: Header=BB2_1 Depth=1
	xorps	%xmm0, %xmm0
.LBB2_4:                                #   in Loop: Header=BB2_1 Depth=1
	movss	%xmm0, (%rax,%rcx,4)
	incq	%rcx
	cmpq	$720, %rcx                      # imm = 0x2D0
	jne	.LBB2_1
# %bb.5:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	_Z16calculate_omegasPiS_S_, .Lfunc_end2-_Z16calculate_omegasPiS_S_
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
.Lfunc_begin0:
	.cfi_startproc
	.cfi_personality 3, __gxx_personality_v0
	.cfi_lsda 3, .Lexception0
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$416, %rsp                      # imm = 0x1A0
	.cfi_def_cfa_offset 464
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	176(%rsp), %r14
	movq	%r14, -16(%r14)
.Ltmp0:
	leaq	160(%rsp), %rdi
	movl	$.L.str, %esi
	movl	$.L.str+30, %edx
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag
.Ltmp1:
# %bb.1:                                # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.exit
.Ltmp3:
	leaq	104(%rsp), %rdi
	leaq	160(%rsp), %rsi
	callq	_Z8readFileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE
.Ltmp4:
# %bb.2:
	movq	120(%rsp), %rax
	movq	%rax, 96(%rsp)
	movups	104(%rsp), %xmm0
	movaps	%xmm0, 80(%rsp)
	movq	160(%rsp), %rdi
	cmpq	%r14, %rdi
	je	.LBB3_4
# %bb.3:                                # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i
	callq	_ZdlPv
.LBB3_4:                                # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit
	leaq	144(%rsp), %r14
	movq	%r14, -16(%r14)
.Ltmp6:
	leaq	128(%rsp), %rdi
	movl	$.L.str.1, %esi
	movl	$.L.str.1+30, %edx
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag
.Ltmp7:
# %bb.5:                                # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.exit42
.Ltmp9:
	leaq	104(%rsp), %rdi
	leaq	128(%rsp), %rsi
	callq	_Z8readFileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE
.Ltmp10:
# %bb.6:
	movq	120(%rsp), %rax
	movq	%rax, 64(%rsp)
	movups	104(%rsp), %xmm0
	movaps	%xmm0, 48(%rsp)
	movq	128(%rsp), %rdi
	cmpq	%r14, %rdi
	je	.LBB3_8
# %bb.7:                                # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i43
	callq	_ZdlPv
.LBB3_8:                                # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit45
	movq	96(%rsp), %rax
	movq	%rax, 40(%rsp)
	movaps	80(%rsp), %xmm0
	movups	%xmm0, 24(%rsp)
	movq	%rax, 16(%rsp)
	movups	%xmm0, (%rsp)
	callq	_Z19calculate_histogram10GalaxyFileS_
	movq	%rax, %rbx
	movq	64(%rsp), %rax
	movq	%rax, 40(%rsp)
	movaps	48(%rsp), %xmm0
	movups	%xmm0, 24(%rsp)
	movq	96(%rsp), %rax
	movq	%rax, 16(%rsp)
	movaps	80(%rsp), %xmm0
	movups	%xmm0, (%rsp)
	callq	_Z19calculate_histogram10GalaxyFileS_
	movq	%rax, %r14
	movq	64(%rsp), %rax
	movq	%rax, 40(%rsp)
	movaps	48(%rsp), %xmm0
	movups	%xmm0, 24(%rsp)
	movq	%rax, 16(%rsp)
	movups	%xmm0, (%rsp)
	callq	_Z19calculate_histogram10GalaxyFileS_
	movq	%rax, %r15
	leaq	400(%rsp), %r13
	movq	%r13, -16(%r13)
	leaq	384(%rsp), %r12
	movl	$.L.str.2, %esi
	movl	$.L.str.2+2, %edx
	movq	%r12, %rdi
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag
.Ltmp12:
	movq	%r12, %rdi
	movq	%rbx, %rsi
	movl	$20, %edx
	callq	_Z15print_histogramNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPii
.Ltmp13:
# %bb.9:
	movq	384(%rsp), %rdi
	cmpq	%r13, %rdi
	je	.LBB3_11
# %bb.10:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i47
	callq	_ZdlPv
.LBB3_11:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit49
	leaq	368(%rsp), %r13
	movq	%r13, -16(%r13)
	leaq	352(%rsp), %r12
	movl	$.L.str.3, %esi
	movl	$.L.str.3+2, %edx
	movq	%r12, %rdi
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag
.Ltmp15:
	movq	%r12, %rdi
	movq	%r14, %rsi
	movl	$20, %edx
	callq	_Z15print_histogramNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPii
.Ltmp16:
# %bb.12:
	movq	352(%rsp), %rdi
	cmpq	%r13, %rdi
	je	.LBB3_14
# %bb.13:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i51
	callq	_ZdlPv
.LBB3_14:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit53
	leaq	336(%rsp), %r13
	movq	%r13, -16(%r13)
	leaq	320(%rsp), %r12
	movl	$.L.str.4, %esi
	movl	$.L.str.4+2, %edx
	movq	%r12, %rdi
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag
.Ltmp18:
	movq	%r12, %rdi
	movq	%r15, %rsi
	movl	$20, %edx
	callq	_Z15print_histogramNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPii
.Ltmp19:
# %bb.15:
	movq	320(%rsp), %rdi
	cmpq	%r13, %rdi
	je	.LBB3_17
# %bb.16:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i55
	callq	_ZdlPv
.LBB3_17:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit57
	leaq	304(%rsp), %r13
	movq	%r13, -16(%r13)
	leaq	288(%rsp), %r12
	movl	$.L.str.5, %esi
	movl	$.L.str.5+16, %edx
	movq	%r12, %rdi
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag
.Ltmp21:
	movq	%r12, %rdi
	movq	%rbx, %rsi
	callq	_Z23write_histogram_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPi
.Ltmp22:
# %bb.18:
	movq	288(%rsp), %rdi
	cmpq	%r13, %rdi
	je	.LBB3_20
# %bb.19:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i59
	callq	_ZdlPv
.LBB3_20:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit61
	leaq	272(%rsp), %r13
	movq	%r13, -16(%r13)
	leaq	256(%rsp), %r12
	movl	$.L.str.6, %esi
	movl	$.L.str.6+16, %edx
	movq	%r12, %rdi
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag
.Ltmp24:
	movq	%r12, %rdi
	movq	%r14, %rsi
	callq	_Z23write_histogram_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPi
.Ltmp25:
# %bb.21:
	movq	256(%rsp), %rdi
	cmpq	%r13, %rdi
	je	.LBB3_23
# %bb.22:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i63
	callq	_ZdlPv
.LBB3_23:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit65
	leaq	240(%rsp), %r13
	movq	%r13, -16(%r13)
	leaq	224(%rsp), %r12
	movl	$.L.str.7, %esi
	movl	$.L.str.7+16, %edx
	movq	%r12, %rdi
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag
.Ltmp27:
	movq	%r12, %rdi
	movq	%r15, %rsi
	callq	_Z23write_histogram_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPi
.Ltmp28:
# %bb.24:
	movq	224(%rsp), %rdi
	cmpq	%r13, %rdi
	je	.LBB3_26
# %bb.25:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i67
	callq	_ZdlPv
.LBB3_26:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit69
	movq	%rbx, %rdi
	movq	%r14, %rsi
	movq	%r15, %rdx
	callq	_Z16calculate_omegasPiS_S_
	movq	%rax, %rbx
	movq	%rax, %rdi
	movl	$15, %esi
	callq	_Z12print_omegasPfi
	leaq	208(%rsp), %r15
	movq	%r15, -16(%r15)
	leaq	192(%rsp), %r14
	movl	$.L.str.8, %esi
	movl	$.L.str.8+10, %edx
	movq	%r14, %rdi
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag
.Ltmp30:
	movq	%r14, %rdi
	movq	%rbx, %rsi
	callq	_Z20write_omegas_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPf
.Ltmp31:
# %bb.27:
	movq	192(%rsp), %rdi
	cmpq	%r15, %rdi
	je	.LBB3_29
# %bb.28:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i71
	callq	_ZdlPv
.LBB3_29:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit73
	xorl	%eax, %eax
	addq	$416, %rsp                      # imm = 0x1A0
	.cfi_def_cfa_offset 48
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.LBB3_30:
	.cfi_def_cfa_offset 464
.Ltmp32:
	movq	%rax, %rbx
	movq	192(%rsp), %rdi
	cmpq	%r15, %rdi
	jne	.LBB3_42
	jmp	.LBB3_43
.LBB3_31:
.Ltmp29:
	movq	%rax, %rbx
	movq	224(%rsp), %rdi
	jmp	.LBB3_37
.LBB3_32:
.Ltmp26:
	movq	%rax, %rbx
	movq	256(%rsp), %rdi
	jmp	.LBB3_37
.LBB3_33:
.Ltmp23:
	movq	%rax, %rbx
	movq	288(%rsp), %rdi
	jmp	.LBB3_37
.LBB3_34:
.Ltmp20:
	movq	%rax, %rbx
	movq	320(%rsp), %rdi
	jmp	.LBB3_37
.LBB3_35:
.Ltmp17:
	movq	%rax, %rbx
	movq	352(%rsp), %rdi
	jmp	.LBB3_37
.LBB3_36:
.Ltmp14:
	movq	%rax, %rbx
	movq	384(%rsp), %rdi
.LBB3_37:
	cmpq	%r13, %rdi
	jne	.LBB3_42
	jmp	.LBB3_43
.LBB3_38:
.Ltmp11:
	movq	%rax, %rbx
	movq	128(%rsp), %rdi
	jmp	.LBB3_41
.LBB3_39:
.Ltmp8:
	jmp	.LBB3_45
.LBB3_40:
.Ltmp5:
	movq	%rax, %rbx
	movq	160(%rsp), %rdi
.LBB3_41:
	cmpq	%r14, %rdi
	je	.LBB3_43
.LBB3_42:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i98
	callq	_ZdlPv
	jmp	.LBB3_43
.LBB3_44:
.Ltmp2:
.LBB3_45:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit82
	movq	%rax, %rbx
.LBB3_43:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit82
	movq	%rbx, %rdi
	callq	_Unwind_Resume@PLT
.Lfunc_end3:
	.size	main, .Lfunc_end3-main
	.cfi_endproc
	.section	.gcc_except_table,"a",@progbits
	.p2align	2, 0x0
GCC_except_table3:
.Lexception0:
	.byte	255                             # @LPStart Encoding = omit
	.byte	255                             # @TType Encoding = omit
	.byte	1                               # Call site Encoding = uleb128
	.uleb128 .Lcst_end0-.Lcst_begin0
.Lcst_begin0:
	.uleb128 .Ltmp0-.Lfunc_begin0           # >> Call Site 1 <<
	.uleb128 .Ltmp1-.Ltmp0                  #   Call between .Ltmp0 and .Ltmp1
	.uleb128 .Ltmp2-.Lfunc_begin0           #     jumps to .Ltmp2
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp3-.Lfunc_begin0           # >> Call Site 2 <<
	.uleb128 .Ltmp4-.Ltmp3                  #   Call between .Ltmp3 and .Ltmp4
	.uleb128 .Ltmp5-.Lfunc_begin0           #     jumps to .Ltmp5
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp6-.Lfunc_begin0           # >> Call Site 3 <<
	.uleb128 .Ltmp7-.Ltmp6                  #   Call between .Ltmp6 and .Ltmp7
	.uleb128 .Ltmp8-.Lfunc_begin0           #     jumps to .Ltmp8
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp9-.Lfunc_begin0           # >> Call Site 4 <<
	.uleb128 .Ltmp10-.Ltmp9                 #   Call between .Ltmp9 and .Ltmp10
	.uleb128 .Ltmp11-.Lfunc_begin0          #     jumps to .Ltmp11
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp10-.Lfunc_begin0          # >> Call Site 5 <<
	.uleb128 .Ltmp12-.Ltmp10                #   Call between .Ltmp10 and .Ltmp12
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp12-.Lfunc_begin0          # >> Call Site 6 <<
	.uleb128 .Ltmp13-.Ltmp12                #   Call between .Ltmp12 and .Ltmp13
	.uleb128 .Ltmp14-.Lfunc_begin0          #     jumps to .Ltmp14
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp13-.Lfunc_begin0          # >> Call Site 7 <<
	.uleb128 .Ltmp15-.Ltmp13                #   Call between .Ltmp13 and .Ltmp15
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp15-.Lfunc_begin0          # >> Call Site 8 <<
	.uleb128 .Ltmp16-.Ltmp15                #   Call between .Ltmp15 and .Ltmp16
	.uleb128 .Ltmp17-.Lfunc_begin0          #     jumps to .Ltmp17
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp16-.Lfunc_begin0          # >> Call Site 9 <<
	.uleb128 .Ltmp18-.Ltmp16                #   Call between .Ltmp16 and .Ltmp18
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp18-.Lfunc_begin0          # >> Call Site 10 <<
	.uleb128 .Ltmp19-.Ltmp18                #   Call between .Ltmp18 and .Ltmp19
	.uleb128 .Ltmp20-.Lfunc_begin0          #     jumps to .Ltmp20
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp19-.Lfunc_begin0          # >> Call Site 11 <<
	.uleb128 .Ltmp21-.Ltmp19                #   Call between .Ltmp19 and .Ltmp21
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp21-.Lfunc_begin0          # >> Call Site 12 <<
	.uleb128 .Ltmp22-.Ltmp21                #   Call between .Ltmp21 and .Ltmp22
	.uleb128 .Ltmp23-.Lfunc_begin0          #     jumps to .Ltmp23
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp22-.Lfunc_begin0          # >> Call Site 13 <<
	.uleb128 .Ltmp24-.Ltmp22                #   Call between .Ltmp22 and .Ltmp24
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp24-.Lfunc_begin0          # >> Call Site 14 <<
	.uleb128 .Ltmp25-.Ltmp24                #   Call between .Ltmp24 and .Ltmp25
	.uleb128 .Ltmp26-.Lfunc_begin0          #     jumps to .Ltmp26
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp25-.Lfunc_begin0          # >> Call Site 15 <<
	.uleb128 .Ltmp27-.Ltmp25                #   Call between .Ltmp25 and .Ltmp27
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp27-.Lfunc_begin0          # >> Call Site 16 <<
	.uleb128 .Ltmp28-.Ltmp27                #   Call between .Ltmp27 and .Ltmp28
	.uleb128 .Ltmp29-.Lfunc_begin0          #     jumps to .Ltmp29
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp28-.Lfunc_begin0          # >> Call Site 17 <<
	.uleb128 .Ltmp30-.Ltmp28                #   Call between .Ltmp28 and .Ltmp30
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp30-.Lfunc_begin0          # >> Call Site 18 <<
	.uleb128 .Ltmp31-.Ltmp30                #   Call between .Ltmp30 and .Ltmp31
	.uleb128 .Ltmp32-.Lfunc_begin0          #     jumps to .Ltmp32
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp31-.Lfunc_begin0          # >> Call Site 19 <<
	.uleb128 .Lfunc_end3-.Ltmp31            #   Call between .Ltmp31 and .Lfunc_end3
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
.Lcst_end0:
	.p2align	2, 0x0
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _Z8readFileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE
.LCPI4_0:
	.long	0x5f000000                      # float 9.22337203E+18
.LCPI4_1:
	.long	0x3998825b                      # float 2.90888216E-4
	.text
	.globl	_Z8readFileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE
	.type	_Z8readFileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE,@function
_Z8readFileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE: # @_Z8readFileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE
.Lfunc_begin1:
	.cfi_startproc
	.cfi_personality 3, __gxx_personality_v0
	.cfi_lsda 3, .Lexception1
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$552, %rsp                      # imm = 0x228
	.cfi_def_cfa_offset 608
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rdi, %r15
	leaq	32(%rsp), %r14
	movq	%r14, %rdi
	movl	$8, %edx
	callq	_ZNSt14basic_ifstreamIcSt11char_traitsIcEEC1ERKNSt7__cxx1112basic_stringIcS1_SaIcEEESt13_Ios_Openmode
.Ltmp33:
	leaq	12(%rsp), %rsi
	movq	%r14, %rdi
	callq	_ZNSirsERi
.Ltmp34:
# %bb.1:
	movslq	12(%rsp), %rbx
	leaq	(,%rbx,4), %rax
	testq	%rax, %rax
	movq	%r15, 24(%rsp)                  # 8-byte Spill
	js	.LBB4_2
# %bb.3:
	cvtsi2ss	%rax, %xmm0
	jmp	.LBB4_4
.LBB4_2:
	shrq	%rax
	cvtsi2ss	%rax, %xmm0
	addss	%xmm0, %xmm0
.LBB4_4:
	cvttss2si	%xmm0, %rax
	movq	%rax, %rcx
	sarq	$63, %rcx
	subss	.LCPI4_0(%rip), %xmm0
	cvttss2si	%xmm0, %r15
	andq	%rcx, %r15
	orq	%rax, %r15
	movq	%r15, %rdi
	callq	malloc
	movq	%rax, %r14
	movq	%r15, %rdi
	callq	malloc
	movq	%rax, %r15
	testq	%rbx, %rbx
	jle	.LBB4_9
# %bb.5:                                # %.lr.ph.preheader
	xorl	%ebx, %ebx
	leaq	32(%rsp), %r12
	leaq	20(%rsp), %r13
	leaq	16(%rsp), %rbp
.LBB4_6:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
.Ltmp36:
	movq	%r12, %rdi
	movq	%r13, %rsi
	callq	_ZNSi10_M_extractIfEERSiRT_
.Ltmp37:
# %bb.7:                                # %_ZNSirsERf.exit
                                        #   in Loop: Header=BB4_6 Depth=1
.Ltmp38:
	movq	%rax, %rdi
	movq	%rbp, %rsi
	callq	_ZNSi10_M_extractIfEERSiRT_
.Ltmp39:
# %bb.8:                                # %_ZNSirsERf.exit16
                                        #   in Loop: Header=BB4_6 Depth=1
	movss	20(%rsp), %xmm0                 # xmm0 = mem[0],zero,zero,zero
	movss	.LCPI4_1(%rip), %xmm1           # xmm1 = mem[0],zero,zero,zero
	mulss	%xmm1, %xmm0
	movss	%xmm0, (%r14,%rbx,4)
	movss	16(%rsp), %xmm0                 # xmm0 = mem[0],zero,zero,zero
	mulss	%xmm1, %xmm0
	movss	%xmm0, (%r15,%rbx,4)
	incq	%rbx
	movslq	12(%rsp), %rax
	cmpq	%rax, %rbx
	jl	.LBB4_6
.LBB4_9:                                # %._crit_edge
.Ltmp41:
	leaq	32(%rsp), %rdi
	callq	_ZNSt14basic_ifstreamIcSt11char_traitsIcEE5closeEv
.Ltmp42:
# %bb.10:
	movl	12(%rsp), %eax
	movq	24(%rsp), %rbx                  # 8-byte Reload
	movl	%eax, (%rbx)
	movq	%r14, 8(%rbx)
	movq	%r15, 16(%rbx)
	leaq	32(%rsp), %rdi
	callq	_ZNSt14basic_ifstreamIcSt11char_traitsIcEED1Ev
	movq	%rbx, %rax
	addq	$552, %rsp                      # imm = 0x228
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.LBB4_11:
	.cfi_def_cfa_offset 608
.Ltmp43:
	jmp	.LBB4_12
.LBB4_13:
.Ltmp35:
	jmp	.LBB4_12
.LBB4_14:
.Ltmp40:
.LBB4_12:
	movq	%rax, %rbx
	leaq	32(%rsp), %rdi
	callq	_ZNSt14basic_ifstreamIcSt11char_traitsIcEED1Ev
	movq	%rbx, %rdi
	callq	_Unwind_Resume@PLT
.Lfunc_end4:
	.size	_Z8readFileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE, .Lfunc_end4-_Z8readFileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE
	.cfi_endproc
	.section	.gcc_except_table,"a",@progbits
	.p2align	2, 0x0
GCC_except_table4:
.Lexception1:
	.byte	255                             # @LPStart Encoding = omit
	.byte	255                             # @TType Encoding = omit
	.byte	1                               # Call site Encoding = uleb128
	.uleb128 .Lcst_end1-.Lcst_begin1
.Lcst_begin1:
	.uleb128 .Lfunc_begin1-.Lfunc_begin1    # >> Call Site 1 <<
	.uleb128 .Ltmp33-.Lfunc_begin1          #   Call between .Lfunc_begin1 and .Ltmp33
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp33-.Lfunc_begin1          # >> Call Site 2 <<
	.uleb128 .Ltmp34-.Ltmp33                #   Call between .Ltmp33 and .Ltmp34
	.uleb128 .Ltmp35-.Lfunc_begin1          #     jumps to .Ltmp35
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp36-.Lfunc_begin1          # >> Call Site 3 <<
	.uleb128 .Ltmp39-.Ltmp36                #   Call between .Ltmp36 and .Ltmp39
	.uleb128 .Ltmp40-.Lfunc_begin1          #     jumps to .Ltmp40
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp41-.Lfunc_begin1          # >> Call Site 4 <<
	.uleb128 .Ltmp42-.Ltmp41                #   Call between .Ltmp41 and .Ltmp42
	.uleb128 .Ltmp43-.Lfunc_begin1          #     jumps to .Ltmp43
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp42-.Lfunc_begin1          # >> Call Site 5 <<
	.uleb128 .Lfunc_end4-.Ltmp42            #   Call between .Ltmp42 and .Lfunc_end4
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
.Lcst_end1:
	.p2align	2, 0x0
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _Z15print_histogramNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPii
.LCPI5_0:
	.long	0x3e800000                      # float 0.25
	.text
	.globl	_Z15print_histogramNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPii
	.type	_Z15print_histogramNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPii,@function
_Z15print_histogramNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPii: # @_Z15print_histogramNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPii
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	pushq	%rax
	.cfi_def_cfa_offset 64
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rsi, %r15
	movq	%rdi, %r14
	movslq	%edx, %r13
	xorl	%ebp, %ebp
	xorl	%ebx, %ebx
.LBB5_1:                                # =>This Inner Loop Header: Depth=1
	movslq	(%r15,%rbp,4), %r12
	cmpq	%r13, %rbp
	jge	.LBB5_4
# %bb.2:                                #   in Loop: Header=BB5_1 Depth=1
	testl	%r12d, %r12d
	jle	.LBB5_4
# %bb.3:                                #   in Loop: Header=BB5_1 Depth=1
	xorps	%xmm1, %xmm1
	cvtsi2ss	%ebp, %xmm1
	movss	.LCPI5_0(%rip), %xmm2           # xmm2 = mem[0],zero,zero,zero
	mulss	%xmm2, %xmm1
	xorps	%xmm0, %xmm0
	cvtss2sd	%xmm1, %xmm0
	addss	%xmm2, %xmm1
	cvtss2sd	%xmm1, %xmm1
	movl	$.L.str.10, %edi
	movl	%r12d, %esi
	movb	$2, %al
	callq	printf
.LBB5_4:                                #   in Loop: Header=BB5_1 Depth=1
	addq	%r12, %rbx
	incq	%rbp
	cmpq	$720, %rbp                      # imm = 0x2D0
	jne	.LBB5_1
# %bb.5:
	movl	$_ZSt4cout, %edi
	movl	$.L.str.11, %esi
	movl	$24, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	(%r14), %rsi
	movq	8(%r14), %rdx
	movl	$_ZSt4cout, %edi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%rax, %r14
	movl	$.L.str.12, %esi
	movl	$2, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%r14, %rdi
	movq	%rbx, %rsi
	callq	_ZNSo9_M_insertIxEERSoT_
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	addq	$8, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	_ZNSo5flushEv                   # TAILCALL
.Lfunc_end5:
	.size	_Z15print_histogramNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPii, .Lfunc_end5-_Z15print_histogramNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPii
	.cfi_endproc
                                        # -- End function
	.globl	_Z23write_histogram_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPi # -- Begin function _Z23write_histogram_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPi
	.type	_Z23write_histogram_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPi,@function
_Z23write_histogram_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPi: # @_Z23write_histogram_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPi
.Lfunc_begin2:
	.cfi_startproc
	.cfi_personality 3, __gxx_personality_v0
	.cfi_lsda 3, .Lexception2
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	subq	$552, %rsp                      # imm = 0x228
	.cfi_def_cfa_offset 592
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%rsi, %rbx
	movq	%rdi, %r15
	leaq	40(%rsp), %rdi
	callq	_ZNSt14basic_ofstreamIcSt11char_traitsIcEEC1Ev
	movq	(%r15), %r14
	movq	8(%r15), %r15
	leaq	24(%rsp), %r12
	movq	%r12, -16(%r12)
	movq	$0, -8(%r12)
	movb	$0, (%r12)
	leaq	7(%r15), %rsi
.Ltmp44:
	leaq	8(%rsp), %rdi
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE7reserveEm
.Ltmp45:
# %bb.1:
.Ltmp46:
	leaq	8(%rsp), %rdi
	movl	$.L.str.13, %esi
	movl	$7, %edx
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE6appendEPKcm
.Ltmp47:
# %bb.2:
.Ltmp48:
	leaq	8(%rsp), %rdi
	movq	%r14, %rsi
	movq	%r15, %rdx
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE6appendEPKcm
.Ltmp49:
# %bb.3:                                # %_ZStplIcSt11char_traitsIcESaIcEENSt7__cxx1112basic_stringIT_T0_T1_EEPKS5_RKS8_.exit
.Ltmp51:
	leaq	40(%rsp), %rdi
	leaq	8(%rsp), %rsi
	movl	$16, %edx
	callq	_ZNSt14basic_ofstreamIcSt11char_traitsIcEE4openERKNSt7__cxx1112basic_stringIcS1_SaIcEEESt13_Ios_Openmode
.Ltmp52:
# %bb.4:
	movq	8(%rsp), %rdi
	cmpq	%r12, %rdi
	je	.LBB6_6
# %bb.5:                                # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i
	callq	_ZdlPv
.LBB6_6:                                # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit
	xorl	%r15d, %r15d
	leaq	40(%rsp), %r14
.LBB6_7:                                # =>This Inner Loop Header: Depth=1
	movl	(%rbx,%r15,4), %esi
.Ltmp54:
	movq	%r14, %rdi
	callq	_ZNSolsEi
.Ltmp55:
# %bb.8:                                #   in Loop: Header=BB6_7 Depth=1
	cmpq	$719, %r15                      # imm = 0x2CF
	je	.LBB6_16
# %bb.9:                                #   in Loop: Header=BB6_7 Depth=1
.Ltmp56:
	movl	$.L.str.14, %esi
	movl	$1, %edx
	movq	%r14, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp57:
# %bb.10:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit
                                        #   in Loop: Header=BB6_7 Depth=1
	incq	%r15
	jmp	.LBB6_7
.LBB6_16:
.Ltmp59:
	leaq	40(%rsp), %rdi
	callq	_ZNSt14basic_ofstreamIcSt11char_traitsIcEE5closeEv
.Ltmp60:
# %bb.17:
	movq	_ZTTSt14basic_ofstreamIcSt11char_traitsIcEE(%rip), %rax
	leaq	48(%rsp), %rdi
	movq	%rax, -8(%rdi)
	movq	_ZTTSt14basic_ofstreamIcSt11char_traitsIcEE+24(%rip), %rcx
	movq	-24(%rax), %rax
	movq	%rcx, 40(%rsp,%rax)
	callq	_ZNSt13basic_filebufIcSt11char_traitsIcEED2Ev
	leaq	288(%rsp), %rdi
	callq	_ZNSt8ios_baseD2Ev
	addq	$552, %rsp                      # imm = 0x228
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.LBB6_18:
	.cfi_def_cfa_offset 592
.Ltmp61:
	jmp	.LBB6_19
.LBB6_12:
.Ltmp53:
	jmp	.LBB6_13
.LBB6_11:
.Ltmp50:
.LBB6_13:
	movq	%rax, %rbx
	movq	8(%rsp), %rdi
	cmpq	%r12, %rdi
	je	.LBB6_20
# %bb.14:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i12
	callq	_ZdlPv
	jmp	.LBB6_20
.LBB6_15:
.Ltmp58:
.LBB6_19:
	movq	%rax, %rbx
.LBB6_20:
	leaq	40(%rsp), %rdi
	callq	_ZNSt14basic_ofstreamIcSt11char_traitsIcEED1Ev
	movq	%rbx, %rdi
	callq	_Unwind_Resume@PLT
.Lfunc_end6:
	.size	_Z23write_histogram_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPi, .Lfunc_end6-_Z23write_histogram_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPi
	.cfi_endproc
	.section	.gcc_except_table,"a",@progbits
	.p2align	2, 0x0
GCC_except_table6:
.Lexception2:
	.byte	255                             # @LPStart Encoding = omit
	.byte	255                             # @TType Encoding = omit
	.byte	1                               # Call site Encoding = uleb128
	.uleb128 .Lcst_end2-.Lcst_begin2
.Lcst_begin2:
	.uleb128 .Lfunc_begin2-.Lfunc_begin2    # >> Call Site 1 <<
	.uleb128 .Ltmp44-.Lfunc_begin2          #   Call between .Lfunc_begin2 and .Ltmp44
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp44-.Lfunc_begin2          # >> Call Site 2 <<
	.uleb128 .Ltmp49-.Ltmp44                #   Call between .Ltmp44 and .Ltmp49
	.uleb128 .Ltmp50-.Lfunc_begin2          #     jumps to .Ltmp50
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp51-.Lfunc_begin2          # >> Call Site 3 <<
	.uleb128 .Ltmp52-.Ltmp51                #   Call between .Ltmp51 and .Ltmp52
	.uleb128 .Ltmp53-.Lfunc_begin2          #     jumps to .Ltmp53
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp54-.Lfunc_begin2          # >> Call Site 4 <<
	.uleb128 .Ltmp57-.Ltmp54                #   Call between .Ltmp54 and .Ltmp57
	.uleb128 .Ltmp58-.Lfunc_begin2          #     jumps to .Ltmp58
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp59-.Lfunc_begin2          # >> Call Site 5 <<
	.uleb128 .Ltmp60-.Ltmp59                #   Call between .Ltmp59 and .Ltmp60
	.uleb128 .Ltmp61-.Lfunc_begin2          #     jumps to .Ltmp61
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp60-.Lfunc_begin2          # >> Call Site 6 <<
	.uleb128 .Lfunc_end6-.Ltmp60            #   Call between .Ltmp60 and .Lfunc_end6
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
.Lcst_end2:
	.p2align	2, 0x0
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _Z12print_omegasPfi
.LCPI7_0:
	.long	0x00000000                      # float 0
	.text
	.globl	_Z12print_omegasPfi
	.type	_Z12print_omegasPfi,@function
_Z12print_omegasPfi:                    # @_Z12print_omegasPfi
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%rdi, %rbx
	movslq	%esi, %r15
	xorl	%r14d, %r14d
.LBB7_1:                                # =>This Inner Loop Header: Depth=1
	cmpq	%r15, %r14
	jge	.LBB7_4
# %bb.2:                                #   in Loop: Header=BB7_1 Depth=1
	movss	(%rbx,%r14,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
	ucomiss	.LCPI7_0(%rip), %xmm0
	jne	.LBB7_3
	jnp	.LBB7_4
.LBB7_3:                                #   in Loop: Header=BB7_1 Depth=1
	cvtss2sd	%xmm0, %xmm0
	movl	$.L.str.9, %edi
	movl	%r14d, %esi
	movb	$1, %al
	callq	printf
.LBB7_4:                                #   in Loop: Header=BB7_1 Depth=1
	incq	%r14
	cmpq	$720, %r14                      # imm = 0x2D0
	jne	.LBB7_1
# %bb.5:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end7:
	.size	_Z12print_omegasPfi, .Lfunc_end7-_Z12print_omegasPfi
	.cfi_endproc
                                        # -- End function
	.globl	_Z20write_omegas_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPf # -- Begin function _Z20write_omegas_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPf
	.type	_Z20write_omegas_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPf,@function
_Z20write_omegas_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPf: # @_Z20write_omegas_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPf
.Lfunc_begin3:
	.cfi_startproc
	.cfi_personality 3, __gxx_personality_v0
	.cfi_lsda 3, .Lexception3
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	subq	$552, %rsp                      # imm = 0x228
	.cfi_def_cfa_offset 592
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%rsi, %rbx
	movq	%rdi, %r15
	leaq	40(%rsp), %rdi
	callq	_ZNSt14basic_ofstreamIcSt11char_traitsIcEEC1Ev
	movq	(%r15), %r14
	movq	8(%r15), %r15
	leaq	24(%rsp), %r12
	movq	%r12, -16(%r12)
	movq	$0, -8(%r12)
	movb	$0, (%r12)
	leaq	7(%r15), %rsi
.Ltmp62:
	leaq	8(%rsp), %rdi
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE7reserveEm
.Ltmp63:
# %bb.1:
.Ltmp64:
	leaq	8(%rsp), %rdi
	movl	$.L.str.13, %esi
	movl	$7, %edx
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE6appendEPKcm
.Ltmp65:
# %bb.2:
.Ltmp66:
	leaq	8(%rsp), %rdi
	movq	%r14, %rsi
	movq	%r15, %rdx
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE6appendEPKcm
.Ltmp67:
# %bb.3:                                # %_ZStplIcSt11char_traitsIcESaIcEENSt7__cxx1112basic_stringIT_T0_T1_EEPKS5_RKS8_.exit
.Ltmp69:
	leaq	40(%rsp), %rdi
	leaq	8(%rsp), %rsi
	movl	$16, %edx
	callq	_ZNSt14basic_ofstreamIcSt11char_traitsIcEE4openERKNSt7__cxx1112basic_stringIcS1_SaIcEEESt13_Ios_Openmode
.Ltmp70:
# %bb.4:
	movq	8(%rsp), %rdi
	cmpq	%r12, %rdi
	je	.LBB8_6
# %bb.5:                                # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i
	callq	_ZdlPv
.LBB8_6:                                # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit
	xorl	%r15d, %r15d
	leaq	40(%rsp), %r14
.LBB8_7:                                # =>This Inner Loop Header: Depth=1
	xorps	%xmm0, %xmm0
	cvtss2sd	(%rbx,%r15,4), %xmm0
.Ltmp72:
	movq	%r14, %rdi
	callq	_ZNSo9_M_insertIdEERSoT_
.Ltmp73:
# %bb.8:                                # %_ZNSolsEf.exit
                                        #   in Loop: Header=BB8_7 Depth=1
	cmpq	$719, %r15                      # imm = 0x2CF
	je	.LBB8_16
# %bb.9:                                #   in Loop: Header=BB8_7 Depth=1
.Ltmp74:
	movl	$.L.str.14, %esi
	movl	$1, %edx
	movq	%r14, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp75:
# %bb.10:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit
                                        #   in Loop: Header=BB8_7 Depth=1
	incq	%r15
	jmp	.LBB8_7
.LBB8_16:
.Ltmp77:
	leaq	40(%rsp), %rdi
	callq	_ZNSt14basic_ofstreamIcSt11char_traitsIcEE5closeEv
.Ltmp78:
# %bb.17:
	movq	_ZTTSt14basic_ofstreamIcSt11char_traitsIcEE(%rip), %rax
	leaq	48(%rsp), %rdi
	movq	%rax, -8(%rdi)
	movq	_ZTTSt14basic_ofstreamIcSt11char_traitsIcEE+24(%rip), %rcx
	movq	-24(%rax), %rax
	movq	%rcx, 40(%rsp,%rax)
	callq	_ZNSt13basic_filebufIcSt11char_traitsIcEED2Ev
	leaq	288(%rsp), %rdi
	callq	_ZNSt8ios_baseD2Ev
	addq	$552, %rsp                      # imm = 0x228
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.LBB8_18:
	.cfi_def_cfa_offset 592
.Ltmp79:
	jmp	.LBB8_19
.LBB8_12:
.Ltmp71:
	jmp	.LBB8_13
.LBB8_11:
.Ltmp68:
.LBB8_13:
	movq	%rax, %rbx
	movq	8(%rsp), %rdi
	cmpq	%r12, %rdi
	je	.LBB8_20
# %bb.14:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i12
	callq	_ZdlPv
	jmp	.LBB8_20
.LBB8_15:
.Ltmp76:
.LBB8_19:
	movq	%rax, %rbx
.LBB8_20:
	leaq	40(%rsp), %rdi
	callq	_ZNSt14basic_ofstreamIcSt11char_traitsIcEED1Ev
	movq	%rbx, %rdi
	callq	_Unwind_Resume@PLT
.Lfunc_end8:
	.size	_Z20write_omegas_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPf, .Lfunc_end8-_Z20write_omegas_to_fileNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPf
	.cfi_endproc
	.section	.gcc_except_table,"a",@progbits
	.p2align	2, 0x0
GCC_except_table8:
.Lexception3:
	.byte	255                             # @LPStart Encoding = omit
	.byte	255                             # @TType Encoding = omit
	.byte	1                               # Call site Encoding = uleb128
	.uleb128 .Lcst_end3-.Lcst_begin3
.Lcst_begin3:
	.uleb128 .Lfunc_begin3-.Lfunc_begin3    # >> Call Site 1 <<
	.uleb128 .Ltmp62-.Lfunc_begin3          #   Call between .Lfunc_begin3 and .Ltmp62
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp62-.Lfunc_begin3          # >> Call Site 2 <<
	.uleb128 .Ltmp67-.Ltmp62                #   Call between .Ltmp62 and .Ltmp67
	.uleb128 .Ltmp68-.Lfunc_begin3          #     jumps to .Ltmp68
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp69-.Lfunc_begin3          # >> Call Site 3 <<
	.uleb128 .Ltmp70-.Ltmp69                #   Call between .Ltmp69 and .Ltmp70
	.uleb128 .Ltmp71-.Lfunc_begin3          #     jumps to .Ltmp71
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp72-.Lfunc_begin3          # >> Call Site 4 <<
	.uleb128 .Ltmp75-.Ltmp72                #   Call between .Ltmp72 and .Ltmp75
	.uleb128 .Ltmp76-.Lfunc_begin3          #     jumps to .Ltmp76
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp77-.Lfunc_begin3          # >> Call Site 5 <<
	.uleb128 .Ltmp78-.Ltmp77                #   Call between .Ltmp77 and .Ltmp78
	.uleb128 .Ltmp79-.Lfunc_begin3          #     jumps to .Ltmp79
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp78-.Lfunc_begin3          # >> Call Site 6 <<
	.uleb128 .Lfunc_end8-.Ltmp78            #   Call between .Ltmp78 and .Lfunc_end8
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
.Lcst_end3:
	.p2align	2, 0x0
                                        # -- End function
	.section	.text._ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag,"axG",@progbits,_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag,comdat
	.weak	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag # -- Begin function _ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag
	.p2align	1, 0x90
	.type	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag,@function
_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag: # @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%rdx, %r14
	movq	%rsi, %r15
	movq	%rdi, %rbx
	subq	%rsi, %r14
	movq	%r14, (%rsp)
	cmpq	$15, %r14
	jbe	.LBB9_1
# %bb.2:
	movq	%rsp, %r12
	movq	%rbx, %rdi
	movq	%r12, %rsi
	xorl	%edx, %edx
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE9_M_createERmm
	movq	%rax, (%rbx)
	movq	(%r12), %rcx
	movq	%rcx, 16(%rbx)
	jmp	.LBB9_3
.LBB9_1:                                # %._crit_edge
	movq	(%rbx), %rax
.LBB9_3:
	testq	%r14, %r14
	je	.LBB9_7
# %bb.4:
	cmpq	$1, %r14
	jne	.LBB9_6
# %bb.5:
	movb	(%r15), %cl
	movb	%cl, (%rax)
	jmp	.LBB9_7
.LBB9_6:
	movq	%rax, %rdi
	movq	%r15, %rsi
	movq	%r14, %rdx
	callq	memcpy@PLT
.LBB9_7:                                # %_ZZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tagEN6_GuardD2Ev.exit
	movq	(%rsp), %rax
	movq	%rax, 8(%rbx)
	movq	(%rbx), %rcx
	movb	$0, (%rcx,%rax)
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end9:
	.size	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag, .Lfunc_end9-_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag
	.cfi_endproc
                                        # -- End function
	.text
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB10_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB10_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z22angle_between_galaxiesPfS_S_S_Pi, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end10:
	.size	__hip_module_ctor, .Lfunc_end10-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB11_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB11_2:
	retq
.Lfunc_end11:
	.size	__hip_module_dtor, .Lfunc_end11-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z22angle_between_galaxiesPfS_S_S_Pi,@object # @_Z22angle_between_galaxiesPfS_S_S_Pi
	.section	.rodata,"a",@progbits
	.globl	_Z22angle_between_galaxiesPfS_S_S_Pi
	.p2align	3, 0x0
_Z22angle_between_galaxiesPfS_S_S_Pi:
	.quad	_Z37__device_stub__angle_between_galaxiesPfS_S_S_Pi
	.size	_Z22angle_between_galaxiesPfS_S_S_Pi, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"test_data/flat_100k_arcmin.txt"
	.size	.L.str, 31

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"test_data/data_100k_arcmin.txt"
	.size	.L.str.1, 31

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

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

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

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"dd_histogram.txt"
	.size	.L.str.5, 17

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"dr_histogram.txt"
	.size	.L.str.6, 17

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"rr_histogram.txt"
	.size	.L.str.7, 17

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"omegas.txt"
	.size	.L.str.8, 11

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"omegas[%d]: %f\n"
	.size	.L.str.9, 16

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"[%f, %f]: %d\n"
	.size	.L.str.10, 14

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"Galaxy pairs counted in "
	.size	.L.str.11, 25

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

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

	.type	.L.str.14,@object               # @.str.14
.L.str.14:
	.asciz	"\n"
	.size	.L.str.14, 2

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z22angle_between_galaxiesPfS_S_S_Pi"
	.size	.L__unnamed_1, 37

	.type	__hip_fatbin_wrapper,@object    # @__hip_fatbin_wrapper
	.section	.hipFatBinSegment,"a",@progbits
	.p2align	3, 0x0
__hip_fatbin_wrapper:
	.long	1212764230                      # 0x48495046
	.long	1                               # 0x1
	.quad	__hip_fatbin
	.quad	0
	.size	__hip_fatbin_wrapper, 24

	.type	__hip_gpubin_handle,@object     # @__hip_gpubin_handle
	.local	__hip_gpubin_handle
	.comm	__hip_gpubin_handle,8,8
	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	__hip_module_ctor
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym _Z37__device_stub__angle_between_galaxiesPfS_S_S_Pi
	.addrsig_sym __gxx_personality_v0
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Unwind_Resume
	.addrsig_sym _Z22angle_between_galaxiesPfS_S_S_Pi
	.addrsig_sym _ZSt4cout
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z22angle_between_galaxiesPfS_S_S_Pi ; -- Begin function _Z22angle_between_galaxiesPfS_S_S_Pi
	.globl	_Z22angle_between_galaxiesPfS_S_S_Pi
	.p2align	8
	.type	_Z22angle_between_galaxiesPfS_S_S_Pi,@function
_Z22angle_between_galaxiesPfS_S_S_Pi:   ; @_Z22angle_between_galaxiesPfS_S_S_Pi
; %bb.0:
	s_load_b32 s2, s[0:1], 0x34
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_i32_e32 0x186a0, v1
	s_cbranch_execz .LBB0_29
; %bb.1:                                ; %.preheader
	s_load_b256 s[4:11], s[0:1], 0x0
	v_ashrrev_i32_e32 v2, 31, v1
	s_load_b64 s[2:3], s[0:1], 0x20
	s_mov_b64 s[12:13], 0
	s_mov_b32 s14, 0x800000
	v_mov_b32_e32 v6, 1
	v_lshlrev_b64 v[4:5], 2, v[1:2]
	v_mov_b32_e32 v1, 0
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v2, vcc_lo, s8, v4
	v_add_co_ci_u32_e32 v3, vcc_lo, s9, v5, vcc_lo
	v_add_co_u32 v4, vcc_lo, s10, v4
	v_add_co_ci_u32_e32 v5, vcc_lo, s11, v5, vcc_lo
	s_mov_b32 s8, 0xb94c1982
	s_mov_b32 s9, 0x37d75334
	s_mov_b32 s10, 0x3d1c21a7
.LBB0_2:                                ; =>This Inner Loop Header: Depth=1
	s_add_u32 s0, s4, s12
	s_addc_u32 s1, s5, s13
	global_load_b32 v11, v[2:3], off
	global_load_b32 v12, v1, s[0:1]
	v_mov_b32_e32 v0, 0
	s_mov_b32 s11, exec_lo
	s_waitcnt vmcnt(0)
	v_cmpx_neq_f32_e32 v12, v11
	s_cbranch_execz .LBB0_28
; %bb.3:                                ;   in Loop: Header=BB0_2 Depth=1
	s_add_u32 s0, s6, s12
	s_addc_u32 s1, s7, s13
	global_load_b32 v0, v1, s[0:1]
	global_load_b32 v7, v[4:5], off
	s_waitcnt vmcnt(1)
	v_readfirstlane_b32 s16, v0
	s_waitcnt vmcnt(0)
	v_cmp_neq_f32_e32 vcc_lo, v0, v7
	v_mov_b32_e32 v0, 0
	s_and_saveexec_b32 s15, vcc_lo
	s_cbranch_execz .LBB0_27
; %bb.4:                                ;   in Loop: Header=BB0_2 Depth=1
	v_cmp_ngt_f32_e64 s18, 0x48000000, |s16|
	s_and_b32 s17, s16, 0x7fffffff
	s_delay_alu instid0(VALU_DEP_1)
	s_and_b32 vcc_lo, exec_lo, s18
	s_cbranch_vccz .LBB0_6
; %bb.5:                                ;   in Loop: Header=BB0_2 Depth=1
	s_and_b32 s0, s17, 0x7fffff
	s_lshr_b32 s1, s17, 23
	s_bitset1_b32 s0, 23
	s_addk_i32 s1, 0xff88
	s_mul_hi_u32 s19, s0, 0xfe5163ab
	s_mul_i32 s20, s0, 0x3c439041
	s_mul_hi_u32 s21, s0, 0x3c439041
	s_add_u32 s19, s19, s20
	s_addc_u32 s20, 0, s21
	s_mul_i32 s21, s0, 0xdb629599
	s_mul_hi_u32 s22, s0, 0xdb629599
	s_add_u32 s20, s20, s21
	s_addc_u32 s21, 0, s22
	s_mul_i32 s22, s0, 0xf534ddc0
	s_mul_hi_u32 s23, s0, 0xf534ddc0
	s_add_u32 s21, s21, s22
	s_addc_u32 s22, 0, s23
	s_mul_i32 s23, s0, 0xfc2757d1
	s_mul_hi_u32 s24, s0, 0xfc2757d1
	s_add_u32 s22, s22, s23
	s_addc_u32 s23, 0, s24
	s_mul_i32 s24, s0, 0x4e441529
	s_mul_hi_u32 s25, s0, 0x4e441529
	s_add_u32 s23, s23, s24
	s_addc_u32 s24, 0, s25
	s_cmp_gt_u32 s1, 63
	s_mul_i32 s25, s0, 0xfe5163ab
	s_mul_hi_u32 s26, s0, 0xa2f9836e
	s_mul_i32 s0, s0, 0xa2f9836e
	s_cselect_b32 s27, s20, s22
	s_cselect_b32 s19, s19, s21
	s_cselect_b32 s20, s25, s20
	s_add_u32 s0, s24, s0
	s_addc_u32 s24, 0, s26
	s_cmp_gt_u32 s1, 63
	s_cselect_b32 s25, 0xffffffc0, 0
	s_cselect_b32 s21, s21, s23
	s_cselect_b32 s0, s22, s0
	s_cselect_b32 s22, s23, s24
	s_add_i32 s25, s25, s1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_gt_u32 s25, 31
	s_cselect_b32 s1, 0xffffffe0, 0
	s_cselect_b32 s23, s21, s0
	s_cselect_b32 s0, s0, s22
	s_cselect_b32 s21, s27, s21
	s_cselect_b32 s22, s19, s27
	s_cselect_b32 s19, s20, s19
	s_add_i32 s1, s1, s25
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_gt_u32 s1, 31
	s_cselect_b32 s20, 0xffffffe0, 0
	s_cselect_b32 s0, s23, s0
	s_cselect_b32 s23, s21, s23
	s_cselect_b32 s21, s22, s21
	s_cselect_b32 s19, s19, s22
	s_add_i32 s20, s20, s1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	s_sub_i32 s1, 32, s20
	s_cmp_eq_u32 s20, 0
	v_mov_b32_e32 v0, s1
	s_cselect_b32 s20, -1, 0
	v_alignbit_b32 v8, s0, s23, v0
	v_alignbit_b32 v9, s23, s21, v0
	v_alignbit_b32 v0, s21, s19, v0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_readfirstlane_b32 s1, v8
	v_cndmask_b32_e64 v8, v9, s23, s20
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v0, v0, s21, s20
	s_cselect_b32 s0, s0, s1
	s_delay_alu instid0(VALU_DEP_2) | instid1(SALU_CYCLE_1)
	v_alignbit_b32 v9, s0, v8, 30
	s_bfe_u32 s1, s0, 0x1001d
	s_delay_alu instid0(VALU_DEP_2)
	v_alignbit_b32 v8, v8, v0, 30
	s_sub_i32 s22, 0, s1
	v_alignbit_b32 v0, v0, s19, 30
	v_xor_b32_e32 v9, s22, v9
	s_lshr_b32 s19, s0, 29
	v_xor_b32_e32 v8, s22, v8
	s_lshl_b32 s19, s19, 31
	v_xor_b32_e32 v0, s22, v0
	v_clz_i32_u32_e32 v10, v9
	s_or_b32 s20, s19, 0.5
	s_lshr_b32 s0, s0, 30
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_add_i32 s1, s1, s0
	v_min_u32_e32 v10, 32, v10
	s_mov_b32 s0, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v13, 31, v10
	v_lshlrev_b32_e32 v14, 23, v10
	v_alignbit_b32 v9, v9, v8, v13
	v_alignbit_b32 v0, v8, v0, v13
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v14, s20, v14
	v_alignbit_b32 v8, v9, v0, 9
	v_lshrrev_b32_e32 v9, 9, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_clz_i32_u32_e32 v13, v8
	v_min_u32_e32 v13, 32, v13
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v15, 31, v13
	v_add_nc_u32_e32 v10, v13, v10
	v_alignbit_b32 v0, v8, v0, v15
	v_or_b32_e32 v8, v9, v14
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b32_e32 v9, 23, v10
	v_lshrrev_b32_e32 v0, 9, v0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v10, 0x3fc90fda, v8
	v_sub_nc_u32_e32 v0, v0, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v9, 0x3fc90fda, v8, -v10
	v_add_nc_u32_e32 v0, 0x33000000, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v9, 0x33a22168, v8
	v_or_b32_e32 v0, s19, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v9, 0x3fc90fda, v0
	v_add_f32_e32 v8, v10, v9
	s_branch .LBB0_7
.LBB0_6:                                ;   in Loop: Header=BB0_2 Depth=1
	s_mov_b32 s0, -1
                                        ; implicit-def: $sgpr1
                                        ; implicit-def: $vgpr8
.LBB0_7:                                ; %Flow93
                                        ;   in Loop: Header=BB0_2 Depth=1
	v_mul_f32_e64 v0, 0x3f22f983, |s16|
	v_mov_b32_e32 v10, s1
	s_and_not1_b32 vcc_lo, exec_lo, s0
	s_delay_alu instid0(VALU_DEP_2)
	v_rndne_f32_e32 v19, v0
	s_cbranch_vccnz .LBB0_9
; %bb.8:                                ;   in Loop: Header=BB0_2 Depth=1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fma_f32 v8, 0xbfc90fda, v19, |s16|
	v_cvt_i32_f32_e32 v10, v19
	v_fmac_f32_e32 v8, 0xb3a22168, v19
	s_delay_alu instid0(VALU_DEP_1)
	v_fmac_f32_e32 v8, 0xa7c234c4, v19
.LBB0_9:                                ; %_ZL3sinf.exit
                                        ;   in Loop: Header=BB0_2 Depth=1
	v_and_b32_e32 v9, 0x7fffffff, v7
	v_cmp_ngt_f32_e64 s19, 0x48000000, |v7|
                                        ; implicit-def: $vgpr14
                                        ; implicit-def: $vgpr13
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_lshrrev_b32_e32 v0, 23, v9
	v_and_or_b32 v16, 0x7fffff, v9, s14
	v_add_nc_u32_e32 v17, 0xffffff88, v0
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_saveexec_b32 s0, s19
	s_xor_b32 s20, exec_lo, s0
	s_cbranch_execz .LBB0_11
; %bb.10:                               ;   in Loop: Header=BB0_2 Depth=1
	v_mad_u64_u32 v[13:14], null, 0xfe5163ab, v16, 0
	v_cmp_lt_u32_e32 vcc_lo, 63, v17
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mov_b32_e32 v0, v14
	v_mad_u64_u32 v[14:15], null, 0x3c439041, v16, v[0:1]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mov_b32_e32 v0, v15
	v_cndmask_b32_e64 v15, 0, 0xffffffc0, vcc_lo
	v_add_nc_u32_e32 v15, v15, v17
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_lt_u32_e64 s0, 31, v15
	v_cndmask_b32_e64 v18, 0, 0xffffffe0, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v15, v18, v15
	v_mad_u64_u32 v[20:21], null, 0xdb629599, v16, v[0:1]
	v_cmp_lt_u32_e64 s1, 31, v15
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_mov_b32 v0, v21 :: v_dual_cndmask_b32 v13, v20, v13
	v_mad_u64_u32 v[21:22], null, 0xf534ddc0, v16, v[0:1]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mov_b32_e32 v0, v22
	v_cndmask_b32_e32 v14, v21, v14, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mad_u64_u32 v[22:23], null, 0xfc2757d1, v16, v[0:1]
	v_cndmask_b32_e64 v13, v14, v13, s0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mov_b32_e32 v0, v23
	v_mad_u64_u32 v[23:24], null, 0x4e441529, v16, v[0:1]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mov_b32_e32 v0, v24
	v_cndmask_b32_e32 v18, v23, v21, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[24:25], null, 0xa2f9836e, v16, v[0:1]
	v_cndmask_b32_e64 v0, 0, 0xffffffe0, s1
	v_add_nc_u32_e32 v0, v0, v15
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_dual_cndmask_b32 v24, v24, v22 :: v_dual_cndmask_b32 v23, v25, v23
	v_cndmask_b32_e32 v22, v22, v20, vcc_lo
	v_cmp_eq_u32_e32 vcc_lo, 0, v0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cndmask_b32_e64 v15, v24, v18, s0
	v_cndmask_b32_e64 v21, v23, v24, s0
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_cndmask_b32_e64 v18, v18, v22, s0
	v_sub_nc_u32_e32 v23, 32, v0
	v_cndmask_b32_e64 v22, v22, v14, s0
	v_cndmask_b32_e64 v21, v21, v15, s1
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v15, v15, v18, s1
	v_cndmask_b32_e64 v18, v18, v22, s1
	v_cndmask_b32_e64 v13, v22, v13, s1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_alignbit_b32 v24, v21, v15, v23
	v_alignbit_b32 v20, v15, v18, v23
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_alignbit_b32 v22, v18, v13, v23
	v_cndmask_b32_e32 v0, v24, v21, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cndmask_b32_e32 v14, v20, v15, vcc_lo
	v_cndmask_b32_e32 v18, v22, v18, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_bfe_u32 v15, v0, 29, 1
	v_alignbit_b32 v20, v0, v14, 30
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_alignbit_b32 v14, v14, v18, 30
	v_alignbit_b32 v13, v18, v13, 30
	v_sub_nc_u32_e32 v21, 0, v15
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_xor_b32_e32 v20, v20, v21
	v_xor_b32_e32 v14, v14, v21
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_xor_b32_e32 v13, v13, v21
	v_clz_i32_u32_e32 v22, v20
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_min_u32_e32 v22, 32, v22
	v_sub_nc_u32_e32 v18, 31, v22
	v_lshlrev_b32_e32 v24, 23, v22
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_3) | instid1(VALU_DEP_3)
	v_alignbit_b32 v20, v20, v14, v18
	v_alignbit_b32 v13, v14, v13, v18
	v_lshrrev_b32_e32 v18, 29, v0
	v_lshrrev_b32_e32 v0, 30, v0
	v_alignbit_b32 v14, v20, v13, 9
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_lshlrev_b32_e32 v18, 31, v18
	v_lshrrev_b32_e32 v20, 9, v20
	v_clz_i32_u32_e32 v21, v14
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_or_b32_e32 v23, 0.5, v18
	v_min_u32_e32 v21, 32, v21
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v23, v23, v24
	v_sub_nc_u32_e32 v25, 31, v21
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_alignbit_b32 v13, v14, v13, v25
	v_or_b32_e32 v14, v20, v23
	v_add_lshl_u32 v20, v21, v22, 23
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshrrev_b32_e32 v13, 9, v13
	v_mul_f32_e32 v21, 0x3fc90fda, v14
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v13, v13, v20
	v_fma_f32 v20, 0x3fc90fda, v14, -v21
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v13, 0x33000000, v13
	v_fmac_f32_e32 v20, 0x33a22168, v14
	v_add_nc_u32_e32 v14, v15, v0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_or_b32_e32 v13, v13, v18
	v_fmac_f32_e32 v20, 0x3fc90fda, v13
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f32_e32 v13, v21, v20
.LBB0_11:                               ; %Flow92
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_or_saveexec_b32 s0, s20
	v_mul_f32_e64 v0, 0x3f22f983, |v7|
	s_delay_alu instid0(VALU_DEP_1)
	v_rndne_f32_e32 v0, v0
	s_xor_b32 exec_lo, exec_lo, s0
; %bb.12:                               ;   in Loop: Header=BB0_2 Depth=1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fma_f32 v13, 0xbfc90fda, v0, |v7|
	v_cvt_i32_f32_e32 v14, v0
	v_fmac_f32_e32 v13, 0xb3a22168, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_fmac_f32_e32 v13, 0xa7c234c4, v0
; %bb.13:                               ; %_ZL3sinf.exit41
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 vcc_lo, exec_lo, s18
	s_cbranch_vccz .LBB0_15
; %bb.14:                               ;   in Loop: Header=BB0_2 Depth=1
	s_and_b32 s0, s17, 0x7fffff
	s_lshr_b32 s1, s17, 23
	s_bitset1_b32 s0, 23
	s_addk_i32 s1, 0xff88
	s_mul_hi_u32 s18, s0, 0xfe5163ab
	s_mul_i32 s20, s0, 0x3c439041
	s_mul_hi_u32 s21, s0, 0x3c439041
	s_add_u32 s18, s18, s20
	s_addc_u32 s20, 0, s21
	s_mul_i32 s21, s0, 0xdb629599
	s_mul_hi_u32 s22, s0, 0xdb629599
	s_add_u32 s20, s20, s21
	s_addc_u32 s21, 0, s22
	s_mul_i32 s22, s0, 0xf534ddc0
	s_mul_hi_u32 s23, s0, 0xf534ddc0
	s_add_u32 s21, s21, s22
	s_addc_u32 s22, 0, s23
	s_mul_i32 s23, s0, 0xfc2757d1
	s_mul_hi_u32 s24, s0, 0xfc2757d1
	s_add_u32 s22, s22, s23
	s_addc_u32 s23, 0, s24
	s_mul_i32 s24, s0, 0x4e441529
	s_mul_hi_u32 s25, s0, 0x4e441529
	s_add_u32 s23, s23, s24
	s_addc_u32 s24, 0, s25
	s_cmp_gt_u32 s1, 63
	s_mul_i32 s25, s0, 0xfe5163ab
	s_mul_hi_u32 s26, s0, 0xa2f9836e
	s_mul_i32 s0, s0, 0xa2f9836e
	s_cselect_b32 s27, s20, s22
	s_cselect_b32 s18, s18, s21
	s_cselect_b32 s20, s25, s20
	s_add_u32 s0, s24, s0
	s_addc_u32 s24, 0, s26
	s_cmp_gt_u32 s1, 63
	s_cselect_b32 s25, 0xffffffc0, 0
	s_cselect_b32 s21, s21, s23
	s_cselect_b32 s0, s22, s0
	s_cselect_b32 s22, s23, s24
	s_add_i32 s25, s25, s1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_gt_u32 s25, 31
	s_cselect_b32 s1, 0xffffffe0, 0
	s_cselect_b32 s23, s21, s0
	s_cselect_b32 s0, s0, s22
	s_cselect_b32 s21, s27, s21
	s_cselect_b32 s22, s18, s27
	s_cselect_b32 s18, s20, s18
	s_add_i32 s1, s1, s25
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_gt_u32 s1, 31
	s_cselect_b32 s20, 0xffffffe0, 0
	s_cselect_b32 s0, s23, s0
	s_cselect_b32 s23, s21, s23
	s_cselect_b32 s21, s22, s21
	s_cselect_b32 s18, s18, s22
	s_add_i32 s20, s20, s1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	s_sub_i32 s1, 32, s20
	s_cmp_eq_u32 s20, 0
	v_mov_b32_e32 v15, s1
	s_cselect_b32 s20, -1, 0
	v_alignbit_b32 v18, s0, s23, v15
	v_alignbit_b32 v20, s23, s21, v15
	v_alignbit_b32 v15, s21, s18, v15
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_readfirstlane_b32 s1, v18
	v_cndmask_b32_e64 v18, v20, s23, s20
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v15, v15, s21, s20
	s_cselect_b32 s0, s0, s1
	s_delay_alu instid0(VALU_DEP_2) | instid1(SALU_CYCLE_1)
	v_alignbit_b32 v20, s0, v18, 30
	s_bfe_u32 s1, s0, 0x1001d
	s_delay_alu instid0(VALU_DEP_2)
	v_alignbit_b32 v18, v18, v15, 30
	s_sub_i32 s22, 0, s1
	v_alignbit_b32 v15, v15, s18, 30
	v_xor_b32_e32 v20, s22, v20
	s_lshr_b32 s18, s0, 29
	v_xor_b32_e32 v18, s22, v18
	s_lshl_b32 s18, s18, 31
	v_xor_b32_e32 v15, s22, v15
	v_clz_i32_u32_e32 v21, v20
	s_or_b32 s20, s18, 0.5
	s_lshr_b32 s0, s0, 30
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_add_i32 s1, s1, s0
	v_min_u32_e32 v21, 32, v21
	s_mov_b32 s0, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v22, 31, v21
	v_lshlrev_b32_e32 v23, 23, v21
	v_alignbit_b32 v20, v20, v18, v22
	v_alignbit_b32 v15, v18, v15, v22
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v23, s20, v23
	v_alignbit_b32 v18, v20, v15, 9
	v_lshrrev_b32_e32 v20, 9, v20
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_clz_i32_u32_e32 v22, v18
	v_min_u32_e32 v22, 32, v22
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v24, 31, v22
	v_add_nc_u32_e32 v21, v22, v21
	v_alignbit_b32 v15, v18, v15, v24
	v_or_b32_e32 v18, v20, v23
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b32_e32 v20, 23, v21
	v_lshrrev_b32_e32 v15, 9, v15
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v21, 0x3fc90fda, v18
	v_sub_nc_u32_e32 v15, v15, v20
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v20, 0x3fc90fda, v18, -v21
	v_add_nc_u32_e32 v15, 0x33000000, v15
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v20, 0x33a22168, v18
	v_or_b32_e32 v15, s18, v15
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v20, 0x3fc90fda, v15
	v_add_f32_e32 v15, v21, v20
	s_branch .LBB0_16
.LBB0_15:                               ;   in Loop: Header=BB0_2 Depth=1
	s_mov_b32 s0, -1
                                        ; implicit-def: $sgpr1
                                        ; implicit-def: $vgpr15
.LBB0_16:                               ; %Flow91
                                        ;   in Loop: Header=BB0_2 Depth=1
	v_mov_b32_e32 v18, s1
	s_and_not1_b32 vcc_lo, exec_lo, s0
	s_cbranch_vccnz .LBB0_18
; %bb.17:                               ;   in Loop: Header=BB0_2 Depth=1
	v_fma_f32 v15, 0xbfc90fda, v19, |s16|
	v_cvt_i32_f32_e32 v18, v19
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v15, 0xb3a22168, v19
	v_fmac_f32_e32 v15, 0xa7c234c4, v19
.LBB0_18:                               ; %_ZL3cosf.exit
                                        ;   in Loop: Header=BB0_2 Depth=1
                                        ; implicit-def: $vgpr20
                                        ; implicit-def: $vgpr19
	s_and_saveexec_b32 s0, s19
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s18, exec_lo, s0
	s_cbranch_execz .LBB0_20
; %bb.19:                               ;   in Loop: Header=BB0_2 Depth=1
	v_mad_u64_u32 v[19:20], null, 0xfe5163ab, v16, 0
	v_cmp_lt_u32_e32 vcc_lo, 63, v17
	v_cndmask_b32_e64 v25, 0, 0xffffffc0, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_mov_b32 v0, v20 :: v_dual_add_nc_u32 v17, v25, v17
	v_mad_u64_u32 v[20:21], null, 0x3c439041, v16, v[0:1]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_lt_u32_e64 s0, 31, v17
	v_cndmask_b32_e64 v26, 0, 0xffffffe0, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_mov_b32 v0, v21 :: v_dual_add_nc_u32 v17, v26, v17
	v_mad_u64_u32 v[21:22], null, 0xdb629599, v16, v[0:1]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_lt_u32_e64 s1, 31, v17
	v_dual_mov_b32 v0, v22 :: v_dual_cndmask_b32 v19, v21, v19
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[22:23], null, 0xf534ddc0, v16, v[0:1]
	v_mov_b32_e32 v0, v23
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[23:24], null, 0xfc2757d1, v16, v[0:1]
	v_mov_b32_e32 v0, v24
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[24:25], null, 0x4e441529, v16, v[0:1]
	v_mov_b32_e32 v0, v25
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_mad_u64_u32 v[25:26], null, 0xa2f9836e, v16, v[0:1]
	v_cndmask_b32_e64 v0, 0, 0xffffffe0, s1
	v_cndmask_b32_e32 v16, v24, v22, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_dual_cndmask_b32 v25, v25, v23 :: v_dual_add_nc_u32 v0, v0, v17
	v_dual_cndmask_b32 v24, v26, v24 :: v_dual_cndmask_b32 v23, v23, v21
	v_cndmask_b32_e32 v17, v22, v20, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cmp_eq_u32_e32 vcc_lo, 0, v0
	v_cndmask_b32_e64 v20, v25, v16, s0
	s_delay_alu instid0(VALU_DEP_4)
	v_cndmask_b32_e64 v22, v24, v25, s0
	v_cndmask_b32_e64 v16, v16, v23, s0
	v_sub_nc_u32_e32 v24, 32, v0
	v_cndmask_b32_e64 v23, v23, v17, s0
	v_cndmask_b32_e64 v17, v17, v19, s0
	v_cndmask_b32_e64 v22, v22, v20, s1
	v_cndmask_b32_e64 v20, v20, v16, s1
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cndmask_b32_e64 v16, v16, v23, s1
	v_cndmask_b32_e64 v17, v23, v17, s1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_alignbit_b32 v25, v22, v20, v24
	v_alignbit_b32 v21, v20, v16, v24
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_alignbit_b32 v23, v16, v17, v24
	v_cndmask_b32_e32 v0, v25, v22, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cndmask_b32_e32 v19, v21, v20, vcc_lo
	v_cndmask_b32_e32 v16, v23, v16, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_bfe_u32 v20, v0, 29, 1
	v_alignbit_b32 v21, v0, v19, 30
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_alignbit_b32 v19, v19, v16, 30
	v_alignbit_b32 v16, v16, v17, 30
	v_sub_nc_u32_e32 v22, 0, v20
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_xor_b32_e32 v21, v21, v22
	v_xor_b32_e32 v17, v19, v22
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_xor_b32_e32 v16, v16, v22
	v_clz_i32_u32_e32 v23, v21
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_min_u32_e32 v23, 32, v23
	v_sub_nc_u32_e32 v19, 31, v23
	v_lshlrev_b32_e32 v25, 23, v23
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_alignbit_b32 v21, v21, v17, v19
	v_alignbit_b32 v16, v17, v16, v19
	v_lshrrev_b32_e32 v19, 29, v0
	v_alignbit_b32 v17, v21, v16, 9
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_lshlrev_b32_e32 v19, 31, v19
	v_lshrrev_b32_e32 v21, 9, v21
	v_clz_i32_u32_e32 v22, v17
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_or_b32_e32 v24, 0.5, v19
	v_min_u32_e32 v22, 32, v22
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v24, v24, v25
	v_sub_nc_u32_e32 v26, 31, v22
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_alignbit_b32 v16, v17, v16, v26
	v_or_b32_e32 v17, v21, v24
	v_add_lshl_u32 v21, v22, v23, 23
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshrrev_b32_e32 v16, 9, v16
	v_sub_nc_u32_e32 v16, v16, v21
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v16, 0x33000000, v16
	v_or_b32_e32 v16, v16, v19
	v_mul_f32_e32 v22, 0x3fc90fda, v17
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v21, 0x3fc90fda, v17, -v22
	v_fmac_f32_e32 v21, 0x33a22168, v17
	v_lshrrev_b32_e32 v0, 30, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v21, 0x3fc90fda, v16
	v_add_nc_u32_e32 v20, v20, v0
                                        ; implicit-def: $vgpr0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_f32_e32 v19, v22, v21
.LBB0_20:                               ; %Flow89
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_and_not1_saveexec_b32 s0, s18
; %bb.21:                               ;   in Loop: Header=BB0_2 Depth=1
	v_fma_f32 v19, 0xbfc90fda, v0, |v7|
	v_cvt_i32_f32_e32 v20, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v19, 0xb3a22168, v0
	v_fmac_f32_e32 v19, 0xa7c234c4, v0
; %bb.22:                               ; %_ZL3cosf.exit52
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s0
	v_sub_f32_e32 v11, v12, v11
                                        ; implicit-def: $vgpr12
                                        ; implicit-def: $vgpr0
	s_mov_b32 s1, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_ngt_f32_e64 0x48000000, |v11|
	s_xor_b32 s18, exec_lo, s1
	s_cbranch_execz .LBB0_24
; %bb.23:                               ;   in Loop: Header=BB0_2 Depth=1
	v_and_b32_e32 v12, 0x7fffffff, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_and_or_b32 v28, 0x7fffff, v12, s14
	v_lshrrev_b32_e32 v12, 23, v12
	v_mad_u64_u32 v[16:17], null, 0xfe5163ab, v28, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v12, 0xffffff88, v12
	v_cmp_lt_u32_e32 vcc_lo, 63, v12
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mov_b32_e32 v0, v17
	v_cndmask_b32_e64 v17, 0, 0xffffffc0, vcc_lo
	v_mad_u64_u32 v[21:22], null, 0x3c439041, v28, v[0:1]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v12, v17, v12
	v_mov_b32_e32 v0, v22
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_lt_u32_e64 s0, 31, v12
	v_mad_u64_u32 v[22:23], null, 0xdb629599, v28, v[0:1]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v17, 0, 0xffffffe0, s0
	v_add_nc_u32_e32 v12, v17, v12
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_mov_b32_e32 v0, v23
	v_cndmask_b32_e32 v16, v22, v16, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cmp_lt_u32_e64 s1, 31, v12
	v_mad_u64_u32 v[23:24], null, 0xf534ddc0, v28, v[0:1]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mov_b32_e32 v0, v24
	v_mad_u64_u32 v[24:25], null, 0xfc2757d1, v28, v[0:1]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mov_b32_e32 v0, v25
	v_mad_u64_u32 v[25:26], null, 0x4e441529, v28, v[0:1]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_mov_b32 v0, v26 :: v_dual_cndmask_b32 v17, v25, v23
	v_mad_u64_u32 v[26:27], null, 0xa2f9836e, v28, v[0:1]
	v_cndmask_b32_e64 v0, 0, 0xffffffe0, s1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v0, v0, v12
	v_dual_cndmask_b32 v26, v26, v24 :: v_dual_cndmask_b32 v25, v27, v25
	v_cndmask_b32_e32 v24, v24, v22, vcc_lo
	v_cndmask_b32_e32 v12, v23, v21, vcc_lo
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cmp_eq_u32_e32 vcc_lo, 0, v0
	v_cndmask_b32_e64 v21, v26, v17, s0
	v_cndmask_b32_e64 v23, v25, v26, s0
	v_cndmask_b32_e64 v17, v17, v24, s0
	v_sub_nc_u32_e32 v25, 32, v0
	v_cndmask_b32_e64 v24, v24, v12, s0
	v_cndmask_b32_e64 v12, v12, v16, s0
	v_cndmask_b32_e64 v23, v23, v21, s1
	v_cndmask_b32_e64 v21, v21, v17, s1
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cndmask_b32_e64 v17, v17, v24, s1
	v_cndmask_b32_e64 v12, v24, v12, s1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_alignbit_b32 v26, v23, v21, v25
	v_alignbit_b32 v22, v21, v17, v25
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_alignbit_b32 v24, v17, v12, v25
	v_cndmask_b32_e32 v0, v26, v23, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cndmask_b32_e32 v16, v22, v21, vcc_lo
	v_cndmask_b32_e32 v17, v24, v17, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_bfe_u32 v21, v0, 29, 1
	v_alignbit_b32 v22, v0, v16, 30
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_alignbit_b32 v16, v16, v17, 30
	v_alignbit_b32 v12, v17, v12, 30
	v_sub_nc_u32_e32 v23, 0, v21
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_xor_b32_e32 v22, v22, v23
	v_xor_b32_e32 v16, v16, v23
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_xor_b32_e32 v12, v12, v23
	v_clz_i32_u32_e32 v24, v22
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_min_u32_e32 v24, 32, v24
	v_sub_nc_u32_e32 v17, 31, v24
	v_lshlrev_b32_e32 v26, 23, v24
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_alignbit_b32 v22, v22, v16, v17
	v_alignbit_b32 v12, v16, v12, v17
	v_lshrrev_b32_e32 v17, 29, v0
	v_alignbit_b32 v16, v22, v12, 9
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_lshlrev_b32_e32 v17, 31, v17
	v_lshrrev_b32_e32 v22, 9, v22
	v_clz_i32_u32_e32 v23, v16
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_or_b32_e32 v25, 0.5, v17
	v_min_u32_e32 v23, 32, v23
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v25, v25, v26
	v_sub_nc_u32_e32 v27, 31, v23
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_alignbit_b32 v12, v16, v12, v27
	v_or_b32_e32 v16, v22, v25
	v_add_lshl_u32 v22, v23, v24, 23
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshrrev_b32_e32 v12, 9, v12
	v_mul_f32_e32 v23, 0x3fc90fda, v16
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v12, v12, v22
	v_fma_f32 v22, 0x3fc90fda, v16, -v23
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v12, 0x33000000, v12
	v_fmac_f32_e32 v22, 0x33a22168, v16
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_or_b32_e32 v12, v12, v17
	v_fmac_f32_e32 v22, 0x3fc90fda, v12
	v_lshrrev_b32_e32 v12, 30, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f32_e32 v0, v23, v22
	v_add_nc_u32_e32 v12, v21, v12
.LBB0_24:                               ; %Flow
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_and_not1_saveexec_b32 s0, s18
; %bb.25:                               ;   in Loop: Header=BB0_2 Depth=1
	v_mul_f32_e64 v0, 0x3f22f983, |v11|
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_rndne_f32_e32 v12, v0
	v_fma_f32 v0, 0xbfc90fda, v12, |v11|
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v0, 0xb3a22168, v12
	v_fmac_f32_e32 v0, 0xa7c234c4, v12
	v_cvt_i32_f32_e32 v12, v12
; %bb.26:                               ; %_ZL3cosf.exit58
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s0
	v_dual_mul_f32 v16, v15, v15 :: v_dual_and_b32 v17, 1, v18
	v_dual_mul_f32 v21, v19, v19 :: v_dual_lshlrev_b32 v18, 30, v18
	v_cmp_class_f32_e64 s0, v7, 0x1f8
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fmaak_f32 v22, s8, v16, 0x3c0881c4
	v_dual_fmaak_f32 v23, s9, v16, 0xbab64f3b :: v_dual_fmaak_f32 v26, s9, v21, 0xbab64f3b
	v_dual_fmaak_f32 v25, s8, v21, 0x3c0881c4 :: v_dual_and_b32 v24, 1, v20
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fmaak_f32 v22, v16, v22, 0xbe2aaa9d
	v_dual_fmaak_f32 v23, v16, v23, 0x3d2aabf7 :: v_dual_lshlrev_b32 v20, 30, v20
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_dual_fmaak_f32 v26, v21, v26, 0x3d2aabf7 :: v_dual_mul_f32 v27, v8, v8
	v_dual_fmaak_f32 v25, v21, v25, 0xbe2aaa9d :: v_dual_mul_f32 v22, v16, v22
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_fmaak_f32 v23, v16, v23, 0xbf000004
	v_cmp_eq_u32_e32 vcc_lo, 0, v17
	v_dual_mul_f32 v25, v21, v25 :: v_dual_and_b32 v18, 0x80000000, v18
	v_fmaak_f32 v26, v21, v26, 0xbf000004
	v_fmac_f32_e32 v15, v15, v22
	v_fma_f32 v16, v16, v23, 1.0
	v_xor_b32_e32 v7, v9, v7
	v_fmac_f32_e32 v19, v19, v25
	v_fma_f32 v21, v21, v26, 1.0
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v15, -v15, v16, vcc_lo
	v_fmaak_f32 v16, s8, v27, 0x3c0881c4
	v_cmp_eq_u32_e32 vcc_lo, 0, v24
	v_xor_b32_e32 v15, v18, v15
	v_cndmask_b32_e64 v17, -v19, v21, vcc_lo
	v_cmp_class_f32_e64 vcc_lo, s16, 0x1f8
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v15, 0x7fc00000, v15, vcc_lo
	v_and_b32_e32 v20, 0x80000000, v20
	v_xor_b32_e32 v17, v20, v17
	v_mul_f32_e32 v20, v0, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e64 v17, 0x7fc00000, v17, s0
	v_fmaak_f32 v21, s8, v20, 0x3c0881c4
	v_fmaak_f32 v16, v27, v16, 0xbe2aaa9d
	v_fmaak_f32 v22, s9, v20, 0xbab64f3b
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fmaak_f32 v21, v20, v21, 0xbe2aaa9d
	v_dual_fmaak_f32 v19, s9, v27, 0xbab64f3b :: v_dual_mul_f32 v16, v27, v16
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmaak_f32 v22, v20, v22, 0x3d2aabf7
	v_dual_mul_f32 v21, v20, v21 :: v_dual_fmaak_f32 v18, v27, v19, 0x3d2aabf7
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_4) | instid1(VALU_DEP_4)
	v_fmac_f32_e32 v8, v8, v16
	v_dual_mul_f32 v19, v13, v13 :: v_dual_and_b32 v16, 1, v10
	v_lshlrev_b32_e32 v10, 30, v10
	v_fmaak_f32 v22, v20, v22, 0xbf000004
	v_fmac_f32_e32 v0, v0, v21
	v_cmp_eq_u32_e64 s1, 0, v16
	v_fmaak_f32 v16, s8, v19, 0x3c0881c4
	v_dual_mul_f32 v9, v15, v17 :: v_dual_and_b32 v10, 0x80000000, v10
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmaak_f32 v16, v19, v16, 0xbe2aaa9d
	v_mul_f32_e32 v16, v19, v16
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_fmaak_f32 v18, v27, v18, 0xbf000004 :: v_dual_fmac_f32 v13, v13, v16
	v_fma_f32 v18, v27, v18, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	v_cndmask_b32_e64 v8, v18, v8, s1
	v_fmaak_f32 v18, s9, v19, 0xbab64f3b
	s_xor_b32 s1, s17, s16
	v_xor_b32_e32 v10, s1, v10
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmaak_f32 v18, v19, v18, 0x3d2aabf7
	v_xor_b32_e32 v8, v10, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmaak_f32 v18, v19, v18, 0xbf000004
	v_cndmask_b32_e32 v8, 0x7fc00000, v8, vcc_lo
	s_delay_alu instid0(VALU_DEP_2)
	v_fma_f32 v16, v19, v18, 1.0
	v_and_b32_e32 v18, 1, v12
	v_fma_f32 v19, v20, v22, 1.0
	v_lshlrev_b32_e32 v12, 30, v12
	v_and_b32_e32 v20, 1, v14
	v_lshlrev_b32_e32 v14, 30, v14
	v_cmp_eq_u32_e64 s1, 0, v18
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_and_b32_e32 v12, 0x80000000, v12
	v_and_b32_e32 v14, 0x80000000, v14
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e64 v0, -v0, v19, s1
	v_cmp_eq_u32_e64 s1, 0, v20
	v_xor_b32_e32 v0, v12, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e64 v13, v16, v13, s1
	v_cmp_class_f32_e64 s1, v11, 0x1f8
	v_xor3_b32 v7, v7, v14, v13
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e64 v0, 0x7fc00000, v0, s1
	v_cndmask_b32_e64 v7, 0x7fc00000, v7, s0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v0, v9, v0
	v_fmac_f32_e32 v0, v8, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_minmax_f32 v0, v0, 1.0, -1.0
	v_fma_f32 v7, |v0|, -0.5, 0.5
	v_mul_f32_e32 v8, v0, v0
	v_cmp_gt_f32_e64 vcc_lo, |v0|, 0.5
	v_cmp_gt_f32_e64 s0, 0, v0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v7, v8, v7, vcc_lo
	v_fmaak_f32 v8, s10, v7, 0x3c5fc5da
	v_sqrt_f32_e32 v9, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmaak_f32 v8, v7, v8, 0x3d034c3c
	v_fmaak_f32 v8, v7, v8, 0x3d3641b1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmaak_f32 v8, v7, v8, 0x3d999bc8
	v_fmaak_f32 v8, v7, v8, 0x3e2aaaac
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v7, v7, v8
	s_waitcnt_depctr 0xfff
	v_fmac_f32_e32 v9, v9, v7
	v_fmac_f32_e32 v0, v0, v7
	v_add_f32_e32 v8, v9, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v0, 0x3fc90fdb, v0
	v_sub_f32_e32 v9, 0x40490fdb, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v7, v8, v9, s0
	v_cndmask_b32_e32 v0, v0, v7, vcc_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_mul_f32_e32 v0, 0x42652ee1, v0
.LBB0_27:                               ; %Flow94
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s15
.LBB0_28:                               ;   in Loop: Header=BB0_2 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_or_b32 exec_lo, exec_lo, s11
	v_mul_f32_e32 v0, 4.0, v0
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	s_add_u32 s12, s12, 4
	v_floor_f32_e32 v0, v0
	s_addc_u32 s13, s13, 0
	s_cmp_lg_u32 s12, 0x61a80
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_i32_f32_e32 v0, v0
	v_mul_hi_i32 v7, 0xb60b60b7, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v7, v7, v0
	v_lshrrev_b32_e32 v8, 31, v7
	v_ashrrev_i32_e32 v7, 9, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v7, v7, v8
	v_mul_i32_i24_e32 v7, 0x2d0, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v7, v0, v7
	v_ashrrev_i32_e32 v8, 31, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[7:8], 2, v[7:8]
	v_add_co_u32 v7, vcc_lo, s2, v7
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v8, vcc_lo, s3, v8, vcc_lo
	global_atomic_add_u32 v[7:8], v6, off
	s_cbranch_scc1 .LBB0_2
.LBB0_29:                               ; %.loopexit
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z22angle_between_galaxiesPfS_S_S_Pi
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 296
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 29
		.amdhsa_next_free_sgpr 28
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z22angle_between_galaxiesPfS_S_S_Pi, .Lfunc_end0-_Z22angle_between_galaxiesPfS_S_S_Pi
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 4980
; NumSgprs: 30
; NumVgprs: 29
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 3
; VGPRBlocks: 3
; NumSGPRsForWavesPerEU: 30
; NumVGPRsForWavesPerEU: 29
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         44
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         52
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         54
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         56
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         58
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         60
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         62
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         104
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 296
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z22angle_between_galaxiesPfS_S_S_Pi
    .private_segment_fixed_size: 0
    .sgpr_count:     30
    .sgpr_spill_count: 0
    .symbol:         _Z22angle_between_galaxiesPfS_S_S_Pi.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     29
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
