#include <stdio.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <cuda_runtime_api.h>
#include <curand.h>
#include <curand_kernel.h>
#include "kernels.cuh"

__device__ __forceinline__ 
int get_polarity(int id) {

    // If id is an even number, 1 will be returned.
    // If id is an odd number, -1 will be returned.
    return 1 - (2 * (id & 1));
}

__device__ __forceinline__
bool automata_action(unsigned int automata_state, unsigned int max_state) {

    // Returns true if the automata is in an include state, while false if the automata is in an exclude state
    return (automata_state > (static_cast<unsigned int>(max_state / 2))); 
}

__global__
void validate_clauses(unsigned int* model, bool* clauses_output, unsigned int* x_data, unsigned int sample_id, unsigned int clauses_amount, unsigned int features_amount, unsigned int automatas_amount, unsigned int class_id, unsigned int max_state, bool prediction) 
{
    // Declare some shared variables
    // shared[0] = The output of the clause
    // shared[1] = Boolean flag if all is in exclude mode
    __shared__ bool shared[2];

    // Calculate the clause id to work on
    const int thread_id = threadIdx.x;

    // Initialize some "private variables"
    unsigned int sample_value;
    unsigned int automata_value;
    bool action;
    int automata_polarity;

    for (unsigned int clause_id = blockIdx.x; clause_id < clauses_amount; clause_id += gridDim.x) { 
    
        // Set the clause output to be true
        if(thread_id == 0) {
            shared[0] = true;
            shared[1] = true;
        }

        // Wait until all threads are ready
        __syncthreads();

        // Loop over each of the automata and "stride" through
        for(unsigned int automata_id = thread_id; automata_id < automatas_amount; automata_id += blockDim.x) {
        
            // Check if any of the other threads have evaluated the clause to false. This way we could skip checking.
            if(shared[0] == false) {
                break;
            }

            // Get the automatas value
            automata_value = model[(class_id * clauses_amount * automatas_amount) + (clause_id * automatas_amount) + automata_id];

            // Get the action of the automata
            action = automata_action(automata_value, max_state);

            // Check if the automata is in an include state, if so, investigate further...
            if(action == true) {

                // Calculate the polarity of the automata
                automata_polarity = get_polarity(automata_id);

                // Get the sample's value
                sample_value = x_data[(sample_id * features_amount) + (automata_id / 2)]; 

                // Flip the flag that says that all automatas are in exclude mode
                shared[1] = false;

                // Since the automata is in an include state, lets check if the DOES NOT match the desired value
                if(((automata_polarity == 1) && (sample_value != 1)) || ((automata_polarity == -1) && (sample_value != 0))){
                
                    // A condition has been met that would falsify the entire clause. Therefore, evaluate the entire clause to false
                    shared[0] = false;
                    break;
                }
            }
        }

        // Wait until all threads to evaluate until finished
        __syncthreads();

        // Check if we are thread id 0
        if(thread_id == 0)
        {
            // Check if the clause was, when finished evaluating, evaluated to false
            if(shared[0] == false || (prediction == true && shared[1] == true)) {
                clauses_output[clause_id] = false;
            }
            // Assuming it was not false, then it is true
            else {
                clauses_output[clause_id] = true;
            }
        }
    }
}

__global__
void reduce_votes(int* scores, unsigned int scores_index, bool* clauses_output, unsigned int clauses_amount, unsigned int threshold) {

    // Tempromary shared results
    extern __shared__ int results[];

    // Declare some private variables
    int thread_result = 0;

    for(unsigned int clause_id = threadIdx.x; clause_id < clauses_amount; clause_id += blockDim.x) {
        
        // Add the score to this threads tempromary score
        thread_result += (get_polarity(clause_id) * clauses_output[clause_id]);
    }

    // Move the threads result into shared memory
    results[threadIdx.x] = thread_result;

    // Wait until all the threads have completed the summation of all clause outputs
    __syncthreads();

    // Start to reduce the threads and score
    for(unsigned int offset = blockDim.x / 2; offset > 0; offset /= 2) {
    
        // Check if this thread is doing some reduction
        if(threadIdx.x < offset) {
            results[threadIdx.x] += results[threadIdx.x + offset];
        }

        __syncthreads();
    }

    // Thread 0 will store the result in the scores list
    if(threadIdx.x == 0)
    {
        if(threshold != 0) {

            if(results[threadIdx.x] > threshold) {
                results[threadIdx.x] = static_cast<int>(threshold);
            }
            else if(results[threadIdx.x] < -threshold) {
                results[threadIdx.x] = -static_cast<int>(threshold);
            }
        }

        scores[scores_index] = results[threadIdx.x];
    }
}

__global__
void calculate_feedback(unsigned int* clauses_feedback, int* scores, unsigned int threshold, float s, unsigned int class_id, bool correct_class, unsigned int clauses_amount, curandState* random_states) {
    
    // Calculate the position of the thread
    unsigned int global_thread_id = (blockIdx.x * blockDim.x) + threadIdx.x;

    // Declare some private variables
    curandState rnd_state = random_states[global_thread_id];
    float clause_polarity;
    int class_score = scores[0];

    // Loop all clauses 
    for (unsigned int clause_id = global_thread_id; clause_id < clauses_amount; clause_id += gridDim.x) {
   
        // Determine the polarity of the clause
        clause_polarity = static_cast<float>(get_polarity(clause_id));

        // Check if we are on the correct class
        if (correct_class == true) {
            
            // Check if we are to skip feedback for this clause
            if(curand_uniform(&rnd_state) > (((1.0f * threshold) - class_score) / (2.0f * threshold))) {

                // No feedback will be given to this clause
                clauses_feedback[clause_id] = 0;
            }
            else {
                // A small performant operation that calculates that will return the following
                // Clauses for = Type 1 feedback
                // Clauses against = Type 2 feedback
                clauses_feedback[clause_id] = 1 + static_cast<int>(signbit(clause_polarity));
            }

        }
        else {

            // Check if we are to skip feedback for this clause
            if(curand_uniform(&rnd_state) > (((1.0f * threshold) + class_score) / (2.0f * threshold))) {

                // No feedback will be given to this clause
                clauses_feedback[clause_id] = 0;
            }
            else {
                // A small performant operation that calculates that will return the following
                // Clauses for = Type 2 feedback
                // Clauses against = Type 1 feedback
                clauses_feedback[clause_id] = 2 - static_cast<int>(signbit(clause_polarity));
            }
            
        }
    }

    // Copy the random state back to global memory
    random_states[global_thread_id] = rnd_state;
}


__global__ 
void give_feedback_to_clauses(unsigned int* model, unsigned int* clauses_feedback, unsigned int* x_data, bool* clauses_output, unsigned int class_id, unsigned int sample_id, const bool correct_class, unsigned int clauses_amount, unsigned int features_amount, unsigned int automatas_amount, unsigned int max_state, unsigned int threshold, float s, curandState* random_states) {
    
    // Calculate and declare some "private variables"
    // Get the clause id, based on the block id in the grid
    unsigned int global_thread_id = (blockIdx.x * blockDim.x) + threadIdx.x;
   
    // Used to calculate the absolute index of an automata
    unsigned int automata_model_index;
    unsigned int automata_temp;

    // Used to tempromary store whether an automata is in include or exclude state 
    bool action;

    // Used to tempromary store the polarity of an automata
    int automata_polarity;

    // Used to tempromary store the feature id of which feature an automata is associated with
    unsigned int sample_value;

    // Get the random state from the random values matrix (used to generate "random" numbers)
    curandState rnd_state = random_states[global_thread_id];

    // In case there are more clauses than blocks, we need to loop them 
    for(unsigned int clause_id = blockIdx.x; clause_id < clauses_amount; clause_id += gridDim.x) {
        
        // Check if we are to do type 1 feedback
        if(clauses_feedback[clause_id] == 1){
        
            // If the clause output was evaluated to false
            if(clauses_output[clause_id] == 0) {
            
                // Loop and potentially punish all automatas
                for(unsigned int automata_index = threadIdx.x; automata_index < automatas_amount; automata_index += blockDim.x) {
                
                    // Calculate the position of the current automata
                    automata_model_index = (class_id * clauses_amount * automatas_amount) + (clause_id * automatas_amount) + automata_index;

                    // Get the value for the automata
                    automata_temp = model[automata_model_index];

                    if((automata_temp > 1) && (curand_uniform(&rnd_state) <= (1.0 / s))) {
                        model[automata_model_index] = automata_temp - 1;
                    }
                }
            }
            else {
            
                // Loop over each of the automatas
                for(unsigned int automata_index = threadIdx.x; automata_index < automatas_amount; automata_index += blockDim.x){
                
                    // Calculate the position of the current automata
                    automata_model_index = (class_id * clauses_amount * automatas_amount) + (clause_id * automatas_amount) + automata_index;
            
                    // Get the value of the sample for the current automata
                    sample_value = x_data[(sample_id * features_amount) + static_cast<unsigned int>(automata_index / 2)];

                    // Calculate the polarity of the automata
                    automata_polarity = get_polarity(automata_index);

                    // Get the value for the automata
                    automata_temp = model[automata_model_index];

                    // Check if the sample was False
                    if(sample_value == 0) {
                
                        // Check if the automata is an against automata
                        if(automata_polarity == -1){
                        
                            // Increment state
                            if((curand_uniform(&rnd_state) <= ((s - 1.0) / s)) && (automata_temp < max_state)) {
                                model[automata_model_index] = automata_temp + 1;
                            }
                        }
                        // Assumes that the automata is a for automata (since it is not an against automata)
                        else {

                            // Decrement state
                            if((curand_uniform(&rnd_state) <= (1.0 / s)) && automata_temp > 1) {
                                model[automata_model_index] = automata_temp - 1;
                            }
                        }

                    }
                    // Assumes that the sample is 1 (since it was not 0)
                    else {
                    
                        // Check if the automata is a for automata
                        if(automata_polarity == 1) {
                    
                            // Decrement the state 
                            if((curand_uniform(&rnd_state) <= ((s - 1.0) / s)) && (automata_temp < max_state)) {
                                model[automata_model_index] = automata_temp + 1;
                            }
                        }
                        // Assumes that the automata is an against automata (since it is not an for automata)
                        else {
                        
                            // Decrement state
                            if((curand_uniform(&rnd_state) <= (1.0 / s)) && automata_temp > 1) {
                                model[automata_model_index] = automata_temp - 1;
                            }
                        }
                    }
                }
            }
        }
        // Check if we are to do type 2 feedback
        else if(clauses_feedback[clause_id] == 2) {
        
            // Check if the clause was evaluated to true in the evaluation phase. 
            if(clauses_output[clause_id] == 1) {
            
                // Loop over all the automatas
                for(unsigned int automata_id = threadIdx.x; automata_id < automatas_amount; automata_id += blockDim.x) {
            
                    // Calculate the automata model index
                    automata_model_index = (class_id * clauses_amount * automatas_amount) + (clause_id * automatas_amount) + automata_id;

                    // Get the automata value
                    automata_temp = model[automata_model_index];

                    // Get the sample's value
                    sample_value = x_data[(sample_id * features_amount) + (automata_id / 2)];

                    // Calculate the polarity of the automata
                    automata_polarity = get_polarity(automata_id);

                    // Get the include/exclude action for the automata
                    action = automata_action(automata_temp, max_state);

                    // Check if the automata is an for automata and that the feature is 0
                    if((automata_polarity == 1) && (sample_value == 0) && (action == false) && (automata_temp < max_state)){
                        
                        model[automata_model_index] = automata_temp + 1;
                    }
                    else if((automata_polarity == -1) && (sample_value == 1) && (action == false) && (automata_temp < max_state)){

                        model[automata_model_index] = automata_temp + 1;
                    }
                }
            }
        }
    }

    // Some cleanup and persistence before exiting
    // Copy back the random state
    random_states[global_thread_id] = rnd_state;
}


__global__ 
void improved_feedback(unsigned int* model, unsigned int* clauses_feedback, unsigned int* x_data, bool* clauses_output, unsigned int class_id, unsigned int sample_id, const bool correct_class, unsigned int clauses_amount, unsigned int features_amount, unsigned int automatas_amount, unsigned int max_state, unsigned int threshold, float s, curandState* random_states) {
    
    // Data indexing
    unsigned int global_thread_id = (blockIdx.x * blockDim.x) + threadIdx.x;
    unsigned int clause_id;
    unsigned int automata_id;
    unsigned int feature_id;
   
    // Model indexing
    unsigned int model_index;

    // Used to tempromary store the polarity of an automata
    int automata_polarity;
    unsigned int clause_feedback_type;
    bool clause_output;
    unsigned int automata_value;
    unsigned int sample_value;

    // Get the random state from the random values matrix (used to generate "random" numbers)
    curandState rnd_state = random_states[global_thread_id];

    // In case there are more clauses than blocks, we need to loop them 
    for(unsigned int model_relative_index = global_thread_id; model_relative_index < (clauses_amount * automatas_amount); model_relative_index += gridDim.x) {

        // Calculate the position in the model
        model_index = (class_id * clauses_amount * automatas_amount) + model_relative_index;
        clause_id = model_relative_index / automatas_amount;
        automata_id = model_relative_index % automatas_amount;
        feature_id = automata_id / 2;

        // Get the feedback type
        clause_feedback_type = clauses_feedback[clause_id];
        clause_output = clauses_output[clause_id];

        // Get the clause value and sample value
        automata_polarity = get_polarity(automata_id);
        automata_value = model[model_index];
        sample_value = x_data[(sample_id * features_amount) + feature_id];

        // Check if we are to do type 1 feedback
        if(clause_feedback_type == 1){
        
            // If the clause output was evaluated to false
            if(clause_output == false) {

                // Punish the automata if the following conditions are met
                model[model_index] = automata_value - static_cast<int>((curand_uniform(&rnd_state) <= (1.0f / s)) && (automata_value > 1));
            }
            else {                

                // Check if the sample was False
                if(sample_value == 0) {
            
                    // Check if the automata is an against automata
                    if(automata_polarity == -1){
                    
                        // Increment state if the conditions are met
                        model[model_index] = automata_value + static_cast<int>((curand_uniform(&rnd_state) <= ((s - 1.0f) / s)) && (automata_value < max_state));
                    }
                    // Assumes that the automata is a for automata (since it is not an against automata)
                    else {

                        // Decrement the state if the conditions are met
                        model[model_index] = automata_value - static_cast<int>((curand_uniform(&rnd_state) <= (1.0f / s)) && automata_value > 1);
                    }

                }
                // Assumes that the sample is 1 (since it was not 0)
                else {
                
                    // Check if the automata is a for automata
                    if(automata_polarity == 1) {
                
                        // Decrement the state if the conditions are met
                        model[model_index] = automata_value + static_cast<int>((curand_uniform(&rnd_state) <= ((s - 1.0f) / s)) && (automata_value < max_state));
                    }
                    // Assumes that the automata is an against automata (since it is not an for automata)
                    else {
                    
                        // Increment state if the conditions are met
                        model[model_index] = automata_value - static_cast<int>((curand_uniform(&rnd_state) <= (1.0f / s)) && automata_value > 1);
                    }
                }
            }
            
        }
        // Check if we are to do type 2 feedback
        else if(clause_feedback_type == 2) {
        
            // Check if the clause was evaluated to true in the evaluation phase. 
            if(clause_output == true) {

                // Increment the state if the conditions are met
                model[model_index] = automata_value + static_cast<int>(((automata_polarity == 1) && (sample_value == 0)) || ((automata_polarity == -1) && (sample_value == 1)) && (automata_action(automata_value, max_state) == false) && (automata_value < max_state));
            }
        }
    }

    // Some cleanup and persistence before exiting
    // Copy back the random state
    random_states[global_thread_id] = rnd_state;
}

__global__ 
void initialize_random_states(curandState* states, int seed, unsigned int amount_of_states) {

    // Calculate the global thread id
    unsigned int global_thread_id = ((blockIdx.x * blockDim.x) + threadIdx.x);

    // Calculate the offset (to make it a "bit more random")
    int offset = seed+global_thread_id;

    for(unsigned int index = global_thread_id; index < amount_of_states; index += gridDim.x) {
    
        // Initialize the random state
        curand_init(seed, index, offset, &states[index]);
    }
}

	.file	"tmpxft_0023eb89_00000000-6_kernels.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2245:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2245:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z47__device_stub__Z16validate_clausesPjPbS_jjjjjjbPjPbS_jjjjjjb
	.type	_Z47__device_stub__Z16validate_clausesPjPbS_jjjjjjbPjPbS_jjjjjjb, @function
_Z47__device_stub__Z16validate_clausesPjPbS_jjjjjjbPjPbS_jjjjjjb:
.LFB2267:
	.cfi_startproc
	endbr64
	subq	$216, %rsp
	.cfi_def_cfa_offset 224
	movl	248(%rsp), %eax
	movq	%rdi, 40(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	56(%rsp), %rdx
	movl	%ecx, 20(%rsp)
	leaq	64(%rsp), %rcx
	movl	%r8d, 16(%rsp)
	movl	%r9d, 12(%rsp)
	movb	%al, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 200(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 120(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	224(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	232(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	240(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 192(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	64(%rsp)
	.cfi_def_cfa_offset 232
	leaq	_Z16validate_clausesPjPbS_jjjjjjb(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 240
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	136(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 232
	popq	%rdx
	.cfi_def_cfa_offset 224
.L2:
	movq	200(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$216, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2267:
	.size	_Z47__device_stub__Z16validate_clausesPjPbS_jjjjjjbPjPbS_jjjjjjb, .-_Z47__device_stub__Z16validate_clausesPjPbS_jjjjjjbPjPbS_jjjjjjb
	.globl	_Z16validate_clausesPjPbS_jjjjjjb
	.type	_Z16validate_clausesPjPbS_jjjjjjb, @function
_Z16validate_clausesPjPbS_jjjjjjb:
.LFB2268:
	.cfi_startproc
	endbr64
	movzbl	32(%rsp), %eax
	movl	%eax, 32(%rsp)
	jmp	_Z47__device_stub__Z16validate_clausesPjPbS_jjjjjjbPjPbS_jjjjjjb
	.cfi_endproc
.LFE2268:
	.size	_Z16validate_clausesPjPbS_jjjjjjb, .-_Z16validate_clausesPjPbS_jjjjjjb
	.globl	_Z37__device_stub__Z12reduce_votesPijPbjjPijPbjj
	.type	_Z37__device_stub__Z12reduce_votesPijPbjjPijPbjj, @function
_Z37__device_stub__Z12reduce_votesPijPbjjPijPbjj:
.LFB2269:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rdi
	movl	%esi, 20(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movl	%ecx, 16(%rsp)
	leaq	40(%rsp), %rcx
	movl	%r8d, 4(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L8
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z12reduce_votesPijPbjj(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 176
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L8:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2269:
	.size	_Z37__device_stub__Z12reduce_votesPijPbjjPijPbjj, .-_Z37__device_stub__Z12reduce_votesPijPbjjPijPbjj
	.globl	_Z12reduce_votesPijPbjj
	.type	_Z12reduce_votesPijPbjj, @function
_Z12reduce_votesPijPbjj:
.LFB2270:
	.cfi_startproc
	endbr64
	jmp	_Z37__device_stub__Z12reduce_votesPijPbjjPijPbjj
	.cfi_endproc
.LFE2270:
	.size	_Z12reduce_votesPijPbjj, .-_Z12reduce_votesPijPbjj
	.globl	_Z65__device_stub__Z18calculate_feedbackPjPijfjbjP17curandStateXORWOWPjPijfjbjP17curandStateXORWOW
	.type	_Z65__device_stub__Z18calculate_feedbackPjPijfjbjP17curandStateXORWOWPjPijfjbjP17curandStateXORWOW, @function
_Z65__device_stub__Z18calculate_feedbackPjPijfjbjP17curandStateXORWOWPjPijfjbjP17curandStateXORWOW:
.LFB2271:
	.cfi_startproc
	endbr64
	subq	$200, %rsp
	.cfi_def_cfa_offset 208
	movq	208(%rsp), %rax
	movq	%rdi, 40(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movl	%edx, 28(%rsp)
	leaq	56(%rsp), %rdx
	movl	%ecx, 20(%rsp)
	leaq	64(%rsp), %rcx
	movl	%r9d, 12(%rsp)
	movb	%r8b, 16(%rsp)
	movq	%rax, (%rsp)
	movss	%xmm0, 24(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 184(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 120(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 168(%rsp)
	movq	%rsp, %rax
	movq	%rax, 176(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L13
	pushq	64(%rsp)
	.cfi_def_cfa_offset 216
	leaq	_Z18calculate_feedbackPjPijfjbjP17curandStateXORWOW(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 224
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	136(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 216
	popq	%rdx
	.cfi_def_cfa_offset 208
.L13:
	movq	184(%rsp), %rax
	subq	%fs:40, %rax
	je	.L15
	call	__stack_chk_fail@PLT
.L15:
	addq	$200, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2271:
	.size	_Z65__device_stub__Z18calculate_feedbackPjPijfjbjP17curandStateXORWOWPjPijfjbjP17curandStateXORWOW, .-_Z65__device_stub__Z18calculate_feedbackPjPijfjbjP17curandStateXORWOWPjPijfjbjP17curandStateXORWOW
	.globl	_Z18calculate_feedbackPjPijfjbjP17curandStateXORWOW
	.type	_Z18calculate_feedbackPjPijfjbjP17curandStateXORWOW, @function
_Z18calculate_feedbackPjPijfjbjP17curandStateXORWOW:
.LFB2272:
	.cfi_startproc
	endbr64
	movzbl	%r8b, %r8d
	jmp	_Z65__device_stub__Z18calculate_feedbackPjPijfjbjP17curandStateXORWOWPjPijfjbjP17curandStateXORWOW
	.cfi_endproc
.LFE2272:
	.size	_Z18calculate_feedbackPjPijfjbjP17curandStateXORWOW, .-_Z18calculate_feedbackPjPijfjbjP17curandStateXORWOW
	.globl	_Z79__device_stub__Z24give_feedback_to_clausesPjS_S_PbjjbjjjjjfP17curandStateXORWOWPjS_S_PbjjbjjjjjfP17curandStateXORWOW
	.type	_Z79__device_stub__Z24give_feedback_to_clausesPjS_S_PbjjbjjjjjfP17curandStateXORWOWPjS_S_PbjjbjjjjjfP17curandStateXORWOW, @function
_Z79__device_stub__Z24give_feedback_to_clausesPjS_S_PbjjbjjjjjfP17curandStateXORWOWPjS_S_PbjjbjjjjjfP17curandStateXORWOW:
.LFB2273:
	.cfi_startproc
	endbr64
	subq	$264, %rsp
	.cfi_def_cfa_offset 272
	movl	272(%rsp), %eax
	movq	%rdi, 56(%rsp)
	leaq	88(%rsp), %rdi
	movq	%rsi, 48(%rsp)
	leaq	100(%rsp), %rsi
	movb	%al, 20(%rsp)
	movq	320(%rsp), %rax
	movq	%rdx, 40(%rsp)
	leaq	72(%rsp), %rdx
	movq	%rcx, 32(%rsp)
	leaq	80(%rsp), %rcx
	movl	%r8d, 28(%rsp)
	movl	%r9d, 24(%rsp)
	movq	%rax, 8(%rsp)
	movss	%xmm0, 16(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 248(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rax
	movl	$1, 96(%rsp)
	movq	%rax, 136(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	280(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	288(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	296(%rsp), %rax
	movq	%rax, 208(%rsp)
	leaq	304(%rsp), %rax
	movq	%rax, 216(%rsp)
	leaq	312(%rsp), %rax
	movq	%rax, 224(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 232(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 240(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 88(%rsp)
	movq	%rax, 100(%rsp)
	movl	$1, 108(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L18
	pushq	80(%rsp)
	.cfi_def_cfa_offset 280
	leaq	_Z24give_feedback_to_clausesPjS_S_PbjjbjjjjjfP17curandStateXORWOW(%rip), %rdi
	pushq	80(%rsp)
	.cfi_def_cfa_offset 288
	movq	116(%rsp), %rcx
	movl	124(%rsp), %r8d
	movq	104(%rsp), %rsi
	movl	112(%rsp), %edx
	leaq	152(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 280
	popq	%rdx
	.cfi_def_cfa_offset 272
.L18:
	movq	248(%rsp), %rax
	subq	%fs:40, %rax
	je	.L20
	call	__stack_chk_fail@PLT
.L20:
	addq	$264, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2273:
	.size	_Z79__device_stub__Z24give_feedback_to_clausesPjS_S_PbjjbjjjjjfP17curandStateXORWOWPjS_S_PbjjbjjjjjfP17curandStateXORWOW, .-_Z79__device_stub__Z24give_feedback_to_clausesPjS_S_PbjjbjjjjjfP17curandStateXORWOWPjS_S_PbjjbjjjjjfP17curandStateXORWOW
	.globl	_Z24give_feedback_to_clausesPjS_S_PbjjbjjjjjfP17curandStateXORWOW
	.type	_Z24give_feedback_to_clausesPjS_S_PbjjbjjjjjfP17curandStateXORWOW, @function
_Z24give_feedback_to_clausesPjS_S_PbjjbjjjjjfP17curandStateXORWOW:
.LFB2274:
	.cfi_startproc
	endbr64
	movzbl	8(%rsp), %eax
	movl	%eax, 8(%rsp)
	jmp	_Z79__device_stub__Z24give_feedback_to_clausesPjS_S_PbjjbjjjjjfP17curandStateXORWOWPjS_S_PbjjbjjjjjfP17curandStateXORWOW
	.cfi_endproc
.LFE2274:
	.size	_Z24give_feedback_to_clausesPjS_S_PbjjbjjjjjfP17curandStateXORWOW, .-_Z24give_feedback_to_clausesPjS_S_PbjjbjjjjjfP17curandStateXORWOW
	.globl	_Z72__device_stub__Z17improved_feedbackPjS_S_PbjjbjjjjjfP17curandStateXORWOWPjS_S_PbjjbjjjjjfP17curandStateXORWOW
	.type	_Z72__device_stub__Z17improved_feedbackPjS_S_PbjjbjjjjjfP17curandStateXORWOWPjS_S_PbjjbjjjjjfP17curandStateXORWOW, @function
_Z72__device_stub__Z17improved_feedbackPjS_S_PbjjbjjjjjfP17curandStateXORWOWPjS_S_PbjjbjjjjjfP17curandStateXORWOW:
.LFB2275:
	.cfi_startproc
	endbr64
	subq	$264, %rsp
	.cfi_def_cfa_offset 272
	movl	272(%rsp), %eax
	movq	%rdi, 56(%rsp)
	leaq	88(%rsp), %rdi
	movq	%rsi, 48(%rsp)
	leaq	100(%rsp), %rsi
	movb	%al, 20(%rsp)
	movq	320(%rsp), %rax
	movq	%rdx, 40(%rsp)
	leaq	72(%rsp), %rdx
	movq	%rcx, 32(%rsp)
	leaq	80(%rsp), %rcx
	movl	%r8d, 28(%rsp)
	movl	%r9d, 24(%rsp)
	movq	%rax, 8(%rsp)
	movss	%xmm0, 16(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 248(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rax
	movl	$1, 96(%rsp)
	movq	%rax, 136(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	280(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	288(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	296(%rsp), %rax
	movq	%rax, 208(%rsp)
	leaq	304(%rsp), %rax
	movq	%rax, 216(%rsp)
	leaq	312(%rsp), %rax
	movq	%rax, 224(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 232(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 240(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 88(%rsp)
	movq	%rax, 100(%rsp)
	movl	$1, 108(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L23
	pushq	80(%rsp)
	.cfi_def_cfa_offset 280
	leaq	_Z17improved_feedbackPjS_S_PbjjbjjjjjfP17curandStateXORWOW(%rip), %rdi
	pushq	80(%rsp)
	.cfi_def_cfa_offset 288
	movq	116(%rsp), %rcx
	movl	124(%rsp), %r8d
	movq	104(%rsp), %rsi
	movl	112(%rsp), %edx
	leaq	152(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 280
	popq	%rdx
	.cfi_def_cfa_offset 272
.L23:
	movq	248(%rsp), %rax
	subq	%fs:40, %rax
	je	.L25
	call	__stack_chk_fail@PLT
.L25:
	addq	$264, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2275:
	.size	_Z72__device_stub__Z17improved_feedbackPjS_S_PbjjbjjjjjfP17curandStateXORWOWPjS_S_PbjjbjjjjjfP17curandStateXORWOW, .-_Z72__device_stub__Z17improved_feedbackPjS_S_PbjjbjjjjjfP17curandStateXORWOWPjS_S_PbjjbjjjjjfP17curandStateXORWOW
	.globl	_Z17improved_feedbackPjS_S_PbjjbjjjjjfP17curandStateXORWOW
	.type	_Z17improved_feedbackPjS_S_PbjjbjjjjjfP17curandStateXORWOW, @function
_Z17improved_feedbackPjS_S_PbjjbjjjjjfP17curandStateXORWOW:
.LFB2276:
	.cfi_startproc
	endbr64
	movzbl	8(%rsp), %eax
	movl	%eax, 8(%rsp)
	jmp	_Z72__device_stub__Z17improved_feedbackPjS_S_PbjjbjjjjjfP17curandStateXORWOWPjS_S_PbjjbjjjjjfP17curandStateXORWOW
	.cfi_endproc
.LFE2276:
	.size	_Z17improved_feedbackPjS_S_PbjjbjjjjjfP17curandStateXORWOW, .-_Z17improved_feedbackPjS_S_PbjjbjjjjjfP17curandStateXORWOW
	.globl	_Z64__device_stub__Z24initialize_random_statesP17curandStateXORWOWijP17curandStateXORWOWij
	.type	_Z64__device_stub__Z24initialize_random_statesP17curandStateXORWOWijP17curandStateXORWOWij, @function
_Z64__device_stub__Z24initialize_random_statesP17curandStateXORWOWijP17curandStateXORWOWij:
.LFB2277:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	24(%rsp), %rcx
	leaq	32(%rsp), %rdi
	movl	%esi, 4(%rsp)
	leaq	44(%rsp), %rsi
	movl	%edx, (%rsp)
	leaq	16(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 40(%rsp)
	movq	%rax, 80(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L28
	pushq	24(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z24initialize_random_statesP17curandStateXORWOWij(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 144
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	96(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L28:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L30
	call	__stack_chk_fail@PLT
.L30:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2277:
	.size	_Z64__device_stub__Z24initialize_random_statesP17curandStateXORWOWijP17curandStateXORWOWij, .-_Z64__device_stub__Z24initialize_random_statesP17curandStateXORWOWijP17curandStateXORWOWij
	.globl	_Z24initialize_random_statesP17curandStateXORWOWij
	.type	_Z24initialize_random_statesP17curandStateXORWOWij, @function
_Z24initialize_random_statesP17curandStateXORWOWij:
.LFB2278:
	.cfi_startproc
	endbr64
	jmp	_Z64__device_stub__Z24initialize_random_statesP17curandStateXORWOWijP17curandStateXORWOWij
	.cfi_endproc
.LFE2278:
	.size	_Z24initialize_random_statesP17curandStateXORWOWij, .-_Z24initialize_random_statesP17curandStateXORWOWij
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z24initialize_random_statesP17curandStateXORWOWij"
.LC1:
	.string	"_Z17improved_feedbackPjS_S_PbjjbjjjjjfP17curandStateXORWOW"
.LC2:
	.string	"_Z24give_feedback_to_clausesPjS_S_PbjjbjjjjjfP17curandStateXORWOW"
.LC3:
	.string	"_Z18calculate_feedbackPjPijfjbjP17curandStateXORWOW"
.LC4:
	.string	"_Z12reduce_votesPijPbjj"
.LC5:
	.string	"_Z16validate_clausesPjPbS_jjjjjjb"
.LC6:
	.string	"precalc_xorwow_matrix"
.LC7:
	.string	"precalc_xorwow_offset_matrix"
.LC8:
	.string	"mrg32k3aM1"
.LC9:
	.string	"mrg32k3aM2"
.LC10:
	.string	"mrg32k3aM1SubSeq"
.LC11:
	.string	"mrg32k3aM2SubSeq"
.LC12:
	.string	"mrg32k3aM1Seq"
.LC13:
	.string	"mrg32k3aM2Seq"
.LC14:
	.string	"__cr_lgamma_table"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2280:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC0(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z24initialize_random_statesP17curandStateXORWOWij(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC1(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z17improved_feedbackPjS_S_PbjjbjjjjjfP17curandStateXORWOW(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC2(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z24give_feedback_to_clausesPjS_S_PbjjbjjjjjfP17curandStateXORWOW(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC3(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z18calculate_feedbackPjPijfjbjP17curandStateXORWOW(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC4(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z12reduce_votesPijPbjj(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC5(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z16validate_clausesPjPbS_jjjjjjb(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC6(%rip), %rdx
	movl	$102400, %r9d
	leaq	_ZL21precalc_xorwow_matrix(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC7(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$102400, %r9d
	leaq	_ZL28precalc_xorwow_offset_matrix(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC8(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$2304, %r9d
	leaq	_ZL10mrg32k3aM1(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rdi
	.cfi_def_cfa_offset 24
	popq	%r8
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC9(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$2304, %r9d
	leaq	_ZL10mrg32k3aM2(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%r9
	.cfi_def_cfa_offset 24
	popq	%r10
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC10(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$2016, %r9d
	leaq	_ZL16mrg32k3aM1SubSeq(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%r11
	.cfi_def_cfa_offset 24
	popq	%rax
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC11(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$2016, %r9d
	leaq	_ZL16mrg32k3aM2SubSeq(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC12(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$2304, %r9d
	leaq	_ZL13mrg32k3aM1Seq(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC13(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$2304, %r9d
	leaq	_ZL13mrg32k3aM2Seq(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rdi
	.cfi_def_cfa_offset 24
	popq	%r8
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC14(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movl	$72, %r9d
	movq	%rdx, %rcx
	leaq	_ZL17__cr_lgamma_table(%rip), %rsi
	call	__cudaRegisterVar@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%r9
	.cfi_def_cfa_offset 24
	popq	%r10
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2280:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.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
	.local	_ZL17__cr_lgamma_table
	.comm	_ZL17__cr_lgamma_table,72,32
	.local	_ZL13mrg32k3aM2Seq
	.comm	_ZL13mrg32k3aM2Seq,2304,32
	.local	_ZL13mrg32k3aM1Seq
	.comm	_ZL13mrg32k3aM1Seq,2304,32
	.local	_ZL16mrg32k3aM2SubSeq
	.comm	_ZL16mrg32k3aM2SubSeq,2016,32
	.local	_ZL16mrg32k3aM1SubSeq
	.comm	_ZL16mrg32k3aM1SubSeq,2016,32
	.local	_ZL10mrg32k3aM2
	.comm	_ZL10mrg32k3aM2,2304,32
	.local	_ZL10mrg32k3aM1
	.comm	_ZL10mrg32k3aM1,2304,32
	.local	_ZL28precalc_xorwow_offset_matrix
	.comm	_ZL28precalc_xorwow_offset_matrix,102400,32
	.local	_ZL21precalc_xorwow_matrix
	.comm	_ZL21precalc_xorwow_matrix,102400,32
	.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:

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

#include <hip/hip_runtime_api.h>
#include <hiprand/hiprand.h>
#include <hiprand/hiprand_kernel.h>

__device__ __forceinline__
int get_polarity(int id);

__device__ __forceinline__
bool automata_action(unsigned int automata_state, unsigned int max_state);

__global__
void reduce_votes(int* scores, unsigned int scores_index, bool* clauses_output, unsigned int clauses_amount, unsigned int threshold);

__global__
void validate_clauses(unsigned int* model, bool* clauses_output, unsigned int* x_data, unsigned int sample_id, unsigned int clauses_amount, unsigned int features_amount, unsigned int automatas_amount, unsigned int class_id, unsigned int max_state, bool prediction);

__global__ 
void calculate_feedback(unsigned int* clauses_feedback, int* scores, unsigned int threshold, float s, unsigned int class_id, bool correct_class, unsigned int clauses_amount, hiprandState* random_states); 

__global__
void give_feedback_to_clauses(unsigned int* model, unsigned int* clauses_feedback, unsigned int* x_data, bool* clauses_output, unsigned int class_id, unsigned int sample_id, bool correct_class, unsigned int clauses_amount, unsigned int features_amount, unsigned int automatas_amount, unsigned int max_state, unsigned int threshold, float s, hiprandState* random_states);

__global__
void improved_feedback(unsigned int* model, unsigned int* clauses_feedback, unsigned int* x_data, bool* clauses_output, unsigned int class_id, unsigned int sample_id, bool correct_class, unsigned int clauses_amount, unsigned int features_amount, unsigned int automatas_amount, unsigned int max_state, unsigned int threshold, float s, hiprandState* random_states);

__global__
void initialize_random_states(hiprandState* states, int seed, unsigned int amount_of_states);


	.text
	.file	"kernels.hip"
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_

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

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

	.end_amdgpu_metadata
