//
// Created by alex on 7/15/20.
//

#include "Processor.cuh"

inline cudaError_t checkCuda(cudaError_t result)
{
    if (result != cudaSuccess) {
        fprintf(stderr, "CUDA Runtime Error: %s\n", cudaGetErrorString(result));
        assert(result == cudaSuccess);
    }
    return result;
}

__global__ void gpu_count_zeros(Message** flow, int* sum, int flowLength)
{
    int indx = blockDim.x * blockIdx.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;

    for(int i = indx; i < flowLength; i += stride)
    {
        for(int j = 0; j < flow[i]->bufferSize; j++)
        {
            if(flow[i]->buffer[j] == 0)
            {
                sum[i] += 1;
                //cout << "found a zero at msg[" << i << "] byte[" << j << "]" << endl;
            }
        }
    }
}


void cpu_count_zeros(Message** flow, int& sum, int flowLength)
{
    for(int i = 0; i < flowLength; i++)
    {
        for(int j = 0; j < flow[i]->bufferSize; j++)
        {
            if(flow[i]->buffer[j] == 0)
            {
                sum += 1;
                //cout << "found a zero at msg[" << i << "] byte[" << j << "]" << endl;
            }
        }
    }
}


Processor::Processor(ITransport* t) {
    transport = t;
}

void Processor::procCountZerosGPU(int minMessageToProcess) {
    timer t;

    int deviceId;
    int numberOfSMs;

    cudaGetDevice(&deviceId);
    cudaDeviceGetAttribute(&numberOfSMs, cudaDevAttrMultiProcessorCount, deviceId);

    size_t threadsPerBlock;
    size_t numberOfBlocks;

    threadsPerBlock = 256;
    numberOfBlocks = 32 * numberOfSMs;

    int msgCountReturned = 0;
    int processedMessages = 0;
    int sum =0;

    Message* m[MSG_BLOCK_SIZE];//Create array that is max message block size
    uint8_t * d; //The data we will store.
    size_t msgBlockSize = MSG_BLOCK_SIZE * sizeof(Message);
    size_t msgDataSize = MSG_MAX_SIZE * MSG_BLOCK_SIZE;
    checkCuda( cudaMallocManaged(m, msgBlockSize));
    checkCuda( cudaMallocManaged(&d, msgDataSize));

    int* blockSum;   //Array with sum of zeros for this message
    size_t sumArraySize = MSG_BLOCK_SIZE * sizeof(int);
    checkCuda( cudaMallocManaged(&blockSum, sumArraySize));
   // cout << "Processing on GPU using " <<  numberOfBlocks << " blocks with " << threadsPerBlock << " threads per block" << endl;

    while (processedMessages < minMessageToProcess) {

        if (0 != transport->pop(m, MSG_BLOCK_SIZE, msgCountReturned, eTransportDest::DEVICE)) {
            exit(EXIT_FAILURE);
        }

        cudaMemPrefetchAsync(m, msgBlockSize, deviceId);

        if(msgCountReturned > 0) //If there are new messages process them
        {
            std::cerr << "\rProcessed " << processedMessages << " messages";
            gpu_count_zeros <<< threadsPerBlock, numberOfBlocks >>>(m, blockSum, msgCountReturned);

            checkCuda( cudaGetLastError() );
            checkCuda( cudaDeviceSynchronize() ); //Wait for GPU threads to complete

            cudaMemPrefetchAsync(blockSum, sumArraySize, cudaCpuDeviceId);

            for(int k = 0; k < msgCountReturned; k++)
            {
                sum += blockSum[k]; //Add all the counts to the accumulator
                blockSum[k] = 0;
            }

            processedMessages += msgCountReturned;
        }
        //m.clear();
        msgCountReturned=0;

    }

    checkCuda( cudaFree(m));
    checkCuda( cudaFree(blockSum));

    std::cout << "\n Processing Completed: " << std::endl;
    std::cout << "\t processed " << processedMessages << " in " << t.seconds_elapsed() << " sec" << std::endl;
    std::cout << "\t total zero's in messages = " << sum << std::endl;

    exit(EXIT_SUCCESS);
}

int Processor::procCountZerosCPU(int minMessageToProcess) {
    timer t;

    Message* m[MSG_BLOCK_SIZE];
    int msgCountReturned = 0;
    int sum = 0;
    int processedMessages = 0;

    while (processedMessages < minMessageToProcess) {

        if (0 != transport->pop(m, MSG_BLOCK_SIZE, msgCountReturned, eTransportDest::HOST)) {
            exit(EXIT_FAILURE);
        }

        if(msgCountReturned > 0) //If there are new messages process them
        {
            std::cerr << "\rProcessed " << processedMessages << " messages";
            cpu_count_zeros(m, sum, msgCountReturned);
            processedMessages += msgCountReturned;
        }
        msgCountReturned=0;

    }

    std::cout << "\nProcessing Completed: " << std::endl;
    std::cout << "\t processed " << processedMessages << " in " << t.seconds_elapsed() << " sec" << std::endl;
    std::cout << "\t total zero's in messages = " << sum << std::endl;
    exit(EXIT_SUCCESS);
}

void Processor::procDropMsg(int minMessageToProcess) {
    timer t;

    Message* m[MSG_BLOCK_SIZE];
    int msgCountReturned = 0;
    int processedMessages = 0;

    while (processedMessages < minMessageToProcess) {

        if (0 != transport->pop(m, MSG_BLOCK_SIZE, msgCountReturned, eTransportDest::HOST)) {
            exit(EXIT_FAILURE);
        }

        if(msgCountReturned > 0) //If there are new messages process them
        {
            std::cerr << "\rProcessed " << processedMessages << " messages";
            processedMessages += msgCountReturned;
        }
        msgCountReturned=0;

    }

    std::cout << "\nProcessing Completed: " << std::endl;
    std::cout << "\t processed " << processedMessages << " in " << t.seconds_elapsed() << " sec" << std::endl;
    exit(EXIT_SUCCESS);
}

int Processor::procPrintMessages(int minMessageToProcess) {
    Message* m[MSG_BLOCK_SIZE];
    int processedCount = 0;
    int r = 0;

    do {

        if (0 != transport->pop(m, MSG_BLOCK_SIZE, r, eTransportDest::HOST)) {
            exit(EXIT_FAILURE);
        }

        processedCount += r;

        std::cout << "Printing first bytes of " << min(r,minMessageToProcess) << " messages" << std::endl;
        for(int i = 0; i<min(r,minMessageToProcess); i++)
        {
            transport->printMessage(m[i], 32);
            std::cout << std::endl;
        }
    } while (processedCount < minMessageToProcess);

    //Simple process (i.e. print)
    std::cout << "Processing Completed: found " << processedCount << " messages" << std::endl;
    exit(EXIT_SUCCESS);
}

	.file	"tmpxft_00358b81_00000000-6_Processor.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3769:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3769:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.align 2
	.type	_ZNSolsEPFRSoS_E.isra.0, @function
_ZNSolsEPFRSoS_E.isra.0:
.LFB4478:
	.cfi_startproc
	jmp	*%rsi
	.cfi_endproc
.LFE4478:
	.size	_ZNSolsEPFRSoS_E.isra.0, .-_ZNSolsEPFRSoS_E.isra.0
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"CUDA Runtime Error: %s\n"
	.text
	.type	_Z9checkCuda9cudaError.isra.0, @function
_Z9checkCuda9cudaError.isra.0:
.LFB4479:
	.cfi_startproc
	testl	%edi, %edi
	je	.L3
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	cudaGetErrorString@PLT
	movq	stderr(%rip), %rdi
	movl	$2, %esi
	popq	%r8
	.cfi_def_cfa_offset 8
	movq	%rax, %rcx
	leaq	.LC0(%rip), %rdx
	xorl	%eax, %eax
	jmp	__fprintf_chk@PLT
.L3:
	ret
	.cfi_endproc
.LFE4479:
	.size	_Z9checkCuda9cudaError.isra.0, .-_Z9checkCuda9cudaError.isra.0
	.section	.text._ZNK5timer15seconds_elapsedEv,"axG",@progbits,_ZNK5timer15seconds_elapsedEv,comdat
	.align 2
	.weak	_ZNK5timer15seconds_elapsedEv
	.type	_ZNK5timer15seconds_elapsedEv, @function
_ZNK5timer15seconds_elapsedEv:
.LFB3748:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rdi, %rbx
	call	_ZNSt6chrono3_V212steady_clock3nowEv@PLT
	subq	(%rbx), %rax
	movl	$1000000000, %ecx
	popq	%rbx
	.cfi_def_cfa_offset 8
	cqto
	idivq	%rcx
	ret
	.cfi_endproc
.LFE3748:
	.size	_ZNK5timer15seconds_elapsedEv, .-_ZNK5timer15seconds_elapsedEv
	.section	.rodata._ZN10ITransport12printMessageEP7Messagei.str1.1,"aMS",@progbits,1
.LC1:
	.string	"[Message Seq #: "
.LC2:
	.string	"\tsize: "
.LC3:
	.string	"\tintreval: "
.LC4:
	.string	"]"
.LC5:
	.string	"%02hhX "
	.section	.text._ZN10ITransport12printMessageEP7Messagei,"axG",@progbits,_ZN10ITransport12printMessageEP7Messagei,comdat
	.weak	_ZN10ITransport12printMessageEP7Messagei
	.type	_ZN10ITransport12printMessageEP7Messagei, @function
_ZN10ITransport12printMessageEP7Messagei:
.LFB3753:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	leaq	_ZSt4cout(%rip), %r13
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	movl	%esi, %ebp
	leaq	.LC1(%rip), %rsi
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	movq	%rdi, %rbx
	movq	%r13, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	8(%rbx), %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC2(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	4(%rbx), %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC3(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	(%rbx), %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC4(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	4(%rbx), %r12d
	testl	%ebp, %ebp
	jle	.L12
	cmpl	%ebp, %r12d
	cmovg	%ebp, %r12d
.L12:
	xorl	%ebp, %ebp
	leaq	.LC5(%rip), %r14
.L13:
	cmpl	%ebp, %r12d
	jle	.L17
	testl	$15, %ebp
	jne	.L14
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%r13, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
.L14:
	movzbl	12(%rbx,%rbp), %edx
	movq	%r14, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	incq	%rbp
	call	__printf_chk@PLT
	jmp	.L13
.L17:
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	popq	%rbx
	.cfi_def_cfa_offset 40
	movq	%r13, %rdi
	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
	jmp	_ZNSolsEPFRSoS_E.isra.0
	.cfi_endproc
.LFE3753:
	.size	_ZN10ITransport12printMessageEP7Messagei, .-_ZN10ITransport12printMessageEP7Messagei
	.text
	.globl	_Z15cpu_count_zerosPP7MessageRii
	.type	_Z15cpu_count_zerosPP7MessageRii, @function
_Z15cpu_count_zerosPP7MessageRii:
.LFB3755:
	.cfi_startproc
	endbr64
	xorl	%ecx, %ecx
.L19:
	cmpl	%ecx, %edx
	jle	.L18
	movq	(%rdi,%rcx,8), %r8
	xorl	%eax, %eax
.L23:
	cmpl	%eax, 4(%r8)
	jle	.L24
	cmpb	$0, 12(%r8,%rax)
	jne	.L20
	incl	(%rsi)
.L20:
	incq	%rax
	jmp	.L23
.L24:
	incq	%rcx
	jmp	.L19
.L18:
	ret
	.cfi_endproc
.LFE3755:
	.size	_Z15cpu_count_zerosPP7MessageRii, .-_Z15cpu_count_zerosPP7MessageRii
	.align 2
	.globl	_ZN9ProcessorC2EP10ITransport
	.type	_ZN9ProcessorC2EP10ITransport, @function
_ZN9ProcessorC2EP10ITransport:
.LFB3757:
	.cfi_startproc
	endbr64
	movq	%rsi, (%rdi)
	ret
	.cfi_endproc
.LFE3757:
	.size	_ZN9ProcessorC2EP10ITransport, .-_ZN9ProcessorC2EP10ITransport
	.globl	_ZN9ProcessorC1EP10ITransport
	.set	_ZN9ProcessorC1EP10ITransport,_ZN9ProcessorC2EP10ITransport
	.section	.rodata.str1.1
.LC6:
	.string	"\rProcessed "
.LC7:
	.string	" messages"
.LC8:
	.string	"\nProcessing Completed: "
.LC9:
	.string	"\t processed "
.LC10:
	.string	" in "
.LC11:
	.string	" sec"
.LC12:
	.string	"\t total zero's in messages = "
	.text
	.align 2
	.globl	_ZN9Processor17procCountZerosCPUEi
	.type	_ZN9Processor17procCountZerosCPUEi, @function
_ZN9Processor17procCountZerosCPUEi:
.LFB3764:
	.cfi_startproc
	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
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	subq	$4096, %rsp
	.cfi_def_cfa_offset 4144
	orq	$0, (%rsp)
	subq	$4096, %rsp
	.cfi_def_cfa_offset 8240
	orq	$0, (%rsp)
	subq	$32, %rsp
	.cfi_def_cfa_offset 8272
	xorl	%ebx, %ebx
	leaq	.LC6(%rip), %r14
	movq	%fs:40, %rax
	movq	%rax, 8216(%rsp)
	xorl	%eax, %eax
	movq	%rdi, %r12
	movl	%esi, %ebp
	call	_ZNSt6chrono3_V212steady_clock3nowEv@PLT
	xorl	%edx, %edx
	movq	%rax, 16(%rsp)
	movl	%edx, 8(%rsp)
	movl	%edx, 12(%rsp)
.L27:
	cmpl	%ebp, %ebx
	jge	.L34
	movq	(%r12), %rdi
	leaq	24(%rsp), %r13
	xorl	%r8d, %r8d
	leaq	8(%rsp), %rcx
	movl	$1024, %edx
	movq	%r13, %rsi
	movq	(%rdi), %rax
	call	*8(%rax)
	testl	%eax, %eax
	je	.L28
	movl	$1, %edi
	jmp	.L33
.L28:
	cmpl	$0, 8(%rsp)
	jle	.L29
	movq	%r14, %rsi
	leaq	_ZSt4cerr(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC7(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	8(%rsp), %edx
	leaq	12(%rsp), %rsi
	movq	%r13, %rdi
	call	_Z15cpu_count_zerosPP7MessageRii
	addl	%edx, %ebx
.L29:
	xorl	%eax, %eax
	movl	%eax, 8(%rsp)
	jmp	.L27
.L34:
	leaq	_ZSt4cout(%rip), %rbp
	leaq	.LC8(%rip), %rsi
	movq	%rbp, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC9(%rip), %rsi
	movq	%rbp, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC10(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	leaq	16(%rsp), %rdi
	movq	%rax, %rbx
	call	_ZNK5timer15seconds_elapsedEv
	movq	%rbx, %rdi
	movq	%rax, %rsi
	call	_ZNSo9_M_insertIyEERSoT_@PLT
	leaq	.LC11(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC12(%rip), %rsi
	movq	%rbp, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	12(%rsp), %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	xorl	%edi, %edi
.L33:
	call	exit@PLT
	.cfi_endproc
.LFE3764:
	.size	_ZN9Processor17procCountZerosCPUEi, .-_ZN9Processor17procCountZerosCPUEi
	.align 2
	.globl	_ZN9Processor11procDropMsgEi
	.type	_ZN9Processor11procDropMsgEi, @function
_ZN9Processor11procDropMsgEi:
.LFB3765:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	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
	subq	$4096, %rsp
	.cfi_def_cfa_offset 4136
	orq	$0, (%rsp)
	subq	$4096, %rsp
	.cfi_def_cfa_offset 8232
	orq	$0, (%rsp)
	subq	$40, %rsp
	.cfi_def_cfa_offset 8272
	xorl	%ebx, %ebx
	leaq	.LC6(%rip), %r13
	movq	%fs:40, %rax
	movq	%rax, 8216(%rsp)
	xorl	%eax, %eax
	movq	%rdi, %r12
	movl	%esi, %ebp
	call	_ZNSt6chrono3_V212steady_clock3nowEv@PLT
	movq	%rax, 16(%rsp)
.L36:
	xorl	%eax, %eax
	movl	%eax, 12(%rsp)
	cmpl	%ebp, %ebx
	jge	.L43
	movq	(%r12), %rdi
	xorl	%r8d, %r8d
	leaq	12(%rsp), %rcx
	leaq	24(%rsp), %rsi
	movl	$1024, %edx
	movq	(%rdi), %rax
	call	*8(%rax)
	testl	%eax, %eax
	je	.L37
	movl	$1, %edi
	jmp	.L42
.L37:
	cmpl	$0, 12(%rsp)
	jle	.L36
	movq	%r13, %rsi
	leaq	_ZSt4cerr(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC7(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	addl	12(%rsp), %ebx
	jmp	.L36
.L43:
	leaq	_ZSt4cout(%rip), %rbp
	leaq	.LC8(%rip), %rsi
	movq	%rbp, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC9(%rip), %rsi
	movq	%rbp, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC10(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	leaq	16(%rsp), %rdi
	movq	%rax, %rbx
	call	_ZNK5timer15seconds_elapsedEv
	movq	%rbx, %rdi
	movq	%rax, %rsi
	call	_ZNSo9_M_insertIyEERSoT_@PLT
	leaq	.LC11(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	xorl	%edi, %edi
.L42:
	call	exit@PLT
	.cfi_endproc
.LFE3765:
	.size	_ZN9Processor11procDropMsgEi, .-_ZN9Processor11procDropMsgEi
	.section	.rodata.str1.1
.LC13:
	.string	"Printing first bytes of "
.LC14:
	.string	"Processing Completed: found "
	.text
	.align 2
	.globl	_ZN9Processor17procPrintMessagesEi
	.type	_ZN9Processor17procPrintMessagesEi, @function
_ZN9Processor17procPrintMessagesEi:
.LFB3766:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$4096, %rsp
	.cfi_def_cfa_offset 4152
	orq	$0, (%rsp)
	subq	$4096, %rsp
	.cfi_def_cfa_offset 8248
	orq	$0, (%rsp)
	subq	$24, %rsp
	.cfi_def_cfa_offset 8272
	xorl	%r13d, %r13d
	movq	%fs:40, %rax
	movq	%rax, 8200(%rsp)
	xorl	%eax, %eax
	xorl	%eax, %eax
	movq	%rdi, %r12
	movl	%esi, %ebx
	movl	%eax, 4(%rsp)
.L48:
	movq	(%r12), %rdi
	leaq	8(%rsp), %r14
	xorl	%r8d, %r8d
	leaq	4(%rsp), %rcx
	movl	$1024, %edx
	movq	%r14, %rsi
	movq	(%rdi), %rax
	call	*8(%rax)
	testl	%eax, %eax
	je	.L45
	movl	$1, %edi
	jmp	.L52
.L45:
	leaq	.LC13(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	movl	4(%rsp), %ebp
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	4(%rsp), %esi
	addl	%r13d, %ebp
	movq	%rax, %rdi
	cmpl	%esi, %ebx
	movl	%ebp, %r13d
	cmovle	%ebx, %esi
	xorl	%r15d, %r15d
	call	_ZNSolsEi@PLT
	leaq	.LC7(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
.L46:
	movl	4(%rsp), %edx
	cmpl	%edx, %ebx
	cmovle	%ebx, %edx
	cmpl	%r15d, %edx
	jle	.L53
	movq	(%r14,%r15,8), %rdi
	movl	$32, %esi
	incq	%r15
	call	_ZN10ITransport12printMessageEP7Messagei
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	jmp	.L46
.L53:
	cmpl	%ebx, %ebp
	jl	.L48
	leaq	.LC14(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebp, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC7(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	xorl	%edi, %edi
.L52:
	call	exit@PLT
	.cfi_endproc
.LFE3766:
	.size	_ZN9Processor17procPrintMessagesEi, .-_ZN9Processor17procPrintMessagesEi
	.globl	_Z46__device_stub__Z15gpu_count_zerosPP7MessagePiiPP7MessagePii
	.type	_Z46__device_stub__Z15gpu_count_zerosPP7MessagePiiPP7MessagePii, @function
_Z46__device_stub__Z15gpu_count_zerosPP7MessagePiiPP7MessagePii:
.LFB3791:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L54
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z15gpu_count_zerosPP7MessagePii(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	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 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L54:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L56
	call	__stack_chk_fail@PLT
.L56:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3791:
	.size	_Z46__device_stub__Z15gpu_count_zerosPP7MessagePiiPP7MessagePii, .-_Z46__device_stub__Z15gpu_count_zerosPP7MessagePiiPP7MessagePii
	.globl	_Z15gpu_count_zerosPP7MessagePii
	.type	_Z15gpu_count_zerosPP7MessagePii, @function
_Z15gpu_count_zerosPP7MessagePii:
.LFB3792:
	.cfi_startproc
	endbr64
	jmp	_Z46__device_stub__Z15gpu_count_zerosPP7MessagePiiPP7MessagePii
	.cfi_endproc
.LFE3792:
	.size	_Z15gpu_count_zerosPP7MessagePii, .-_Z15gpu_count_zerosPP7MessagePii
	.section	.rodata.str1.1
.LC15:
	.string	"\n Processing Completed: "
	.text
	.align 2
	.globl	_ZN9Processor17procCountZerosGPUEi
	.type	_ZN9Processor17procCountZerosGPUEi, @function
_ZN9Processor17procCountZerosGPUEi:
.LFB3759:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$4096, %rsp
	.cfi_def_cfa_offset 4152
	orq	$0, (%rsp)
	subq	$4096, %rsp
	.cfi_def_cfa_offset 8248
	orq	$0, (%rsp)
	subq	$104, %rsp
	.cfi_def_cfa_offset 8352
	xorl	%r12d, %r12d
	xorl	%ebx, %ebx
	leaq	.LC6(%rip), %r15
	movq	%fs:40, %rax
	movq	%rax, 8280(%rsp)
	xorl	%eax, %eax
	movl	%esi, %r13d
	movq	%rdi, %r14
	call	_ZNSt6chrono3_V212steady_clock3nowEv@PLT
	leaq	28(%rsp), %rdi
	movq	%rax, 40(%rsp)
	call	cudaGetDevice@PLT
	movl	28(%rsp), %edx
	leaq	32(%rsp), %rdi
	movl	$16, %esi
	call	cudaDeviceGetAttribute@PLT
	leaq	88(%rsp), %rax
	xorl	%ecx, %ecx
	movl	32(%rsp), %ebp
	movl	$1, %edx
	movl	$1548288, %esi
	movq	%rax, %rdi
	movl	%ecx, 36(%rsp)
	movq	%rax, 8(%rsp)
	sall	$5, %ebp
	call	cudaMallocManaged@PLT
	movslq	%ebp, %rbp
	movl	%eax, %edi
	call	_Z9checkCuda9cudaError.isra.0
	movl	$1, %edx
	movl	$1536000, %esi
	leaq	48(%rsp), %rdi
	call	cudaMallocManaged@PLT
	movl	%eax, %edi
	call	_Z9checkCuda9cudaError.isra.0
	leaq	56(%rsp), %rdi
	movl	$1, %edx
	movl	$4096, %esi
	call	cudaMallocManaged@PLT
	movl	%eax, %edi
	call	_Z9checkCuda9cudaError.isra.0
.L60:
	cmpl	%r13d, %ebx
	jge	.L70
	movq	(%r14), %rdi
	movq	8(%rsp), %rsi
	leaq	36(%rsp), %rcx
	movl	$1, %r8d
	movl	$1024, %edx
	movq	(%rdi), %rax
	call	*8(%rax)
	testl	%eax, %eax
	je	.L61
	movl	$1, %edi
	jmp	.L69
.L61:
	movl	28(%rsp), %edx
	movq	8(%rsp), %rdi
	xorl	%ecx, %ecx
	movl	$1548288, %esi
	call	cudaMemPrefetchAsync@PLT
	cmpl	$0, 36(%rsp)
	jle	.L62
	movq	%r15, %rsi
	leaq	_ZSt4cerr(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC7(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebp, 76(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movabsq	$4294967297, %rax
	movl	$16777217, %edi
	movl	$1, %esi
	movl	$1, 72(%rsp)
	movq	%rax, 80(%rsp)
	movl	84(%rsp), %ecx
	salq	$8, %rdi
	movq	76(%rsp), %rdx
	movq	%rdi, 64(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L63
	movl	36(%rsp), %edx
	movq	56(%rsp), %rsi
	movq	8(%rsp), %rdi
	call	_Z46__device_stub__Z15gpu_count_zerosPP7MessagePiiPP7MessagePii
.L63:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	call	_Z9checkCuda9cudaError.isra.0
	call	cudaDeviceSynchronize@PLT
	movl	%eax, %edi
	call	_Z9checkCuda9cudaError.isra.0
	movq	56(%rsp), %rdi
	xorl	%ecx, %ecx
	orl	$-1, %edx
	movl	$4096, %esi
	call	cudaMemPrefetchAsync@PLT
	movq	56(%rsp), %rcx
	xorl	%eax, %eax
.L64:
	movl	36(%rsp), %edx
	cmpl	%eax, %edx
	jle	.L71
	xorl	%edx, %edx
	addl	(%rcx,%rax,4), %r12d
	movl	%edx, (%rcx,%rax,4)
	incq	%rax
	jmp	.L64
.L71:
	addl	%edx, %ebx
.L62:
	xorl	%eax, %eax
	movl	%eax, 36(%rsp)
	jmp	.L60
.L70:
	movq	8(%rsp), %rdi
	leaq	_ZSt4cout(%rip), %rbp
	call	cudaFree@PLT
	movl	%eax, %edi
	call	_Z9checkCuda9cudaError.isra.0
	movq	56(%rsp), %rdi
	call	cudaFree@PLT
	movl	%eax, %edi
	call	_Z9checkCuda9cudaError.isra.0
	leaq	.LC15(%rip), %rsi
	movq	%rbp, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC9(%rip), %rsi
	movq	%rbp, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC10(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	leaq	40(%rsp), %rdi
	movq	%rax, %rbx
	call	_ZNK5timer15seconds_elapsedEv
	movq	%rbx, %rdi
	movq	%rax, %rsi
	call	_ZNSo9_M_insertIyEERSoT_@PLT
	leaq	.LC11(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC12(%rip), %rsi
	movq	%rbp, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%r12d, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	xorl	%edi, %edi
.L69:
	call	exit@PLT
	.cfi_endproc
.LFE3759:
	.size	_ZN9Processor17procCountZerosGPUEi, .-_ZN9Processor17procCountZerosGPUEi
	.section	.rodata.str1.1
.LC16:
	.string	"_Z15gpu_count_zerosPP7MessagePii"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3794:
	.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	.LC16(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z15gpu_count_zerosPP7MessagePii(%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
.LFE3794:
	.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
	.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:

//
// Created by alex on 7/15/20.
//

#ifndef SENSORSIM_PROCESSOR_CUH
#define SENSORSIM_PROCESSOR_CUH

#include <unistd.h>
#include <assert.h>

#include "../common.cuh"
#include "../transport/itransport.cuh"

__global__ void gpu_count_zeros(Message* flow, int* sum, int flowLength);
void cpu_count_zeros(Message* flow, int* sum, int flowLength);

class Processor {
public:

    // Constructor declaration
    explicit Processor(ITransport *t);

    //pop a message and process if you get one.
    int procPrintMessages(int minMsg);
    int procCountZerosCPU(int minMsg);
    void procCountZerosGPU(int i);
    void procDropMsg(int i);

private:
    ITransport *transport;


};


#endif //SENSORSIM_PROCESSOR_CUH

	.text
	.file	"Processor.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.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
