//Based on the work of Andrew Krepps
#include <iostream>
#include <random>
#include <chrono>
#include <stdio.h>

__constant__ static const int VAL_A = 1;
__constant__ static const int VAL_B = 3;

// Device GPU add c[i] = a[i] + b[i]
__device__ void add(int * a, int * b, int * c)
{
    const unsigned int thread_idx = (blockIdx.x * blockDim.x) + threadIdx.x;
    
    c[thread_idx] = a[thread_idx] + b[thread_idx];
}

// Device GPU subtract c[i] = a[i] - b[i]
__device__ void subtract(int * a, int * b, int * c)
{
    const unsigned int thread_idx = (blockIdx.x * blockDim.x) + threadIdx.x;
    
    c[thread_idx] = a[thread_idx] - b[thread_idx];
}

// Device GPU multiply c[i] = a[i] * b[i]
__device__ void mult(int * a, int * b, int * c)
{
    const unsigned int thread_idx = (blockIdx.x * blockDim.x) + threadIdx.x;
    
    c[thread_idx] = a[thread_idx] * b[thread_idx];
}

// Device GPU div c[i] = a[i] / b[i]
__device__ void div(int *a, int * b, int * c)
{
    const unsigned int thread_idx = (blockIdx.x * blockDim.x) + threadIdx.x;
    
    c[thread_idx] = a[thread_idx] / b[thread_idx];
}

// Device GPU mod c[i] = a[i] % b[i]
__device__ void mod(int * a, int * b, int * c)
{
    const unsigned int thread_idx = (blockIdx.x * blockDim.x) + threadIdx.x;
    
    c[thread_idx] = a[thread_idx] % b[thread_idx];
}

// Executes all 5 shared math operations
__global__ void executeSharedMathOperations(int * a, int * b, int * addDest, int * subDest, int * multDest, int * divDest, int * modDest, const int size)
{
	const int tid = (blockIdx.x * blockDim.x) + threadIdx.x;
    extern __shared__ int sharedMem[];
    
    // Use offsets in the shared mem to create arrays.
    int * sharedA = &sharedMem[0];
    int * sharedB = &sharedMem[size];
    int * sharedRet = &sharedMem[2*size];
    
    sharedA[tid] = a[tid];
    sharedB[tid] = b[tid];
    
    // Add sharedA to sharedB and store in addDest
    add(sharedA, sharedB, sharedRet);
    addDest[tid] = sharedRet[tid];
    
    // Subtract sharedB from sharedA and store in subDest
    subtract(sharedA, sharedB, sharedRet);
    subDest[tid] = sharedRet[tid];
    // Multiply sharedA to sharedB and store in mutlDest
    mult(sharedA, sharedB, sharedRet);
    multDest[tid] = sharedRet[tid];
    
    // Divide sharedA by sharedB and store in divDest
    div(sharedA, sharedB, sharedRet);
    divDest[tid] = sharedRet[tid];
    
    // Mod sharedA by sharedB and store in modDest
    mod(sharedA, sharedB, sharedRet);
    modDest[tid] = sharedRet[tid];
}

// Executes all 5 global math operations
__global__ void executeGlobalMathOperations(int * a, int * b, int * addDest, int * subDest, int * multDest, int * divDest, int * modDest, const int size)
{
    // Add a to b and store in addDest
    add(a, b, addDest);
    
    // Subtract a from b and store in subDest
    subtract(a, b, subDest);
    
    // Multiply a to b and store in mutlDest
    mult(a, b, multDest);
    
    // Divide a by b and store in divDest
    div(a, b, divDest);
    
    // Mod a by b and store in modDest
    mod(a, b, modDest);
}

// Executes all 5 constant math operations
__global__ void executeConstantMathOperations(int * addDest, int * subDest, int * multDest, int * divDest, int * modDest, const int size)
{    
	const int tid = (blockIdx.x * blockDim.x) + threadIdx.x;
    // Add VAL_A to VAL_B and store in addDest
    addDest[tid] = VAL_A + VAL_B;
    
    // Subtract a from b and store in subDest
    subDest[tid] = VAL_A - VAL_B;
    
    // Multiply a to b and store in mutlDest
    multDest[tid] = VAL_A * VAL_B;
    
    // Divide a by b and store in divDest
    divDest[tid] = VAL_A / VAL_B; // B is chosen to not be 0.
    
    // Mod a by b and store in modDest
    modDest[tid] = VAL_A / VAL_B; // B is chosen to not be 0.
}


// Host (Cpu) add c[i] = a[i] + b[i]
void hostAdd(int * a, int * b, int *c, const int size)
{
    for (int i = 0; i < size; ++i)
    {
        c[i] = a[i] + b[i];
    }
}

// Host (Cpu) sub c[i] = a[i] - b[i]
void hostSub(int * a, int * b, int *c, const int size)
{
    for (int i = 0; i < size; ++i)
    {
        c[i] = a[i] - b[i];
    }
}

// Host (Cpu) multiply c[i] = a[i] * b[i]
void hostMult(int * a, int * b, int *c, const int size)
{
    for (int i = 0; i < size; ++i)
    {
        c[i] = a[i] * b[i];
    }
}

// Host (Cpu) divide c[i] = a[i] / b[i]
void hostDiv(int * a, int * b, int *c, const int size)
{
    for (int i = 0; i < size; ++i)
    {
        if (b[i] != 0)
        {
            c[i] = a[i] / b[i];
        }
        else
        {
            c[i] = 0;
        }
    }
}

// Host (Cpu) mod c[i] = a[i] % b[i]
void hostMod(int * a, int * b, int *c, const int size)
{
    for (int i = 0; i < size; ++i)
    {
        // Protect against divide by 0. 
        // cuda code catches this error and sets result to 0 by default.
        if (b[i] == 0)
        {
            c[i] = 0;
        }
        else
        {
            c[i] = a[i] % b[i];
        }
    }
}

// Executes each of the host (cpu) tests by creating local memory and executing all 5 math operations on the data.
// The data is filled with random numbers that uses the same seed as the GPU tests.
void executeHostTest(const int totalThreads, const int blockSize, const int numBlocks)
{
    int a[totalThreads], b[totalThreads], c[totalThreads];
    
    // Create a random generate that will generate random numbers from 0 to 4.
    // Use a set seed so output is deterministic
    unsigned seed = 12345;
    std::default_random_engine gen(seed);
    std::uniform_int_distribution<int> dist(0,4);
    
    for (size_t i = 0; i < totalThreads; ++i)
    {
        a[i] = i;
        b[i] = dist(gen);
    }
    
    // Add all of the numbers c[i] = a[i] + b[i];
    hostAdd(a,b,c, totalThreads);
    
    // Subtract all of the numbers c[i] = a[i] - b[i];
    hostSub(a,b,c, totalThreads);
    
    // Multiply all of the numbers c[i] = a[i] * b[i];
    hostMult(a,b,c, totalThreads);
    
    // Divides all of the numbers c[i] = a[i] / b[i]; if b[i] == 0, c[i] = 0
    hostDiv(a,b,c, totalThreads);
    
    // Mod all of the numbers c[i] = a[i] % b[i];
    hostMod(a,b,c, totalThreads);
}

// Executes each of the global memory gpu tests by creating local memory, copying it global memory, and then performing
//  all 5 math operations on the data. The data is filled with random numbers that uses the same seed as the CPU tests.
void executeGlobalTest(const int totalThreads, const int blockSize, const int numBlocks)
{
    int a[totalThreads], b[totalThreads], add_dest[totalThreads], sub_dest[totalThreads], mult_dest[totalThreads], div_dest[totalThreads], mod_dest[totalThreads];
    
    int *gpu_a, *gpu_b, *gpu_add_dest, *gpu_sub_dest, *gpu_mult_dest, *gpu_div_dest, *gpu_mod_dest;

    cudaMalloc((void**)&gpu_a,         totalThreads * sizeof(int));
    cudaMalloc((void**)&gpu_b,         totalThreads * sizeof(int));
    cudaMalloc((void**)&gpu_add_dest,  totalThreads * sizeof(int));
    cudaMalloc((void**)&gpu_sub_dest,  totalThreads * sizeof(int));
    cudaMalloc((void**)&gpu_mult_dest, totalThreads * sizeof(int));
    cudaMalloc((void**)&gpu_div_dest,  totalThreads * sizeof(int));
    cudaMalloc((void**)&gpu_mod_dest,  totalThreads * sizeof(int));
    
    // Create a random generate that will generate random numbers from 0 to 4.
    // Use a set seed so output is deterministic
    unsigned seed = 12345;
    std::default_random_engine gen(seed);
    std::uniform_int_distribution<int> dist(0,4);
    
    for (size_t i = 0; i < totalThreads; ++i)
    {
        a[i] = i;
        b[i] = dist(gen);
    }
    
    cudaMemcpy(gpu_a, a, totalThreads * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(gpu_b, b, totalThreads * sizeof(int), cudaMemcpyHostToDevice);
    
    executeGlobalMathOperations<<<numBlocks, blockSize>>>(gpu_a,gpu_b,gpu_add_dest,gpu_sub_dest,gpu_mult_dest,gpu_div_dest,gpu_mod_dest, numBlocks * blockSize);
    
    cudaMemcpy(add_dest,  gpu_add_dest,  totalThreads*sizeof(int), cudaMemcpyDeviceToHost);
    cudaMemcpy(sub_dest,  gpu_sub_dest,  totalThreads*sizeof(int), cudaMemcpyDeviceToHost);    
    cudaMemcpy(mult_dest, gpu_mult_dest, totalThreads*sizeof(int), cudaMemcpyDeviceToHost);
    cudaMemcpy(div_dest,  gpu_div_dest,  totalThreads*sizeof(int), cudaMemcpyDeviceToHost);
    cudaMemcpy(mod_dest,  gpu_mod_dest,  totalThreads*sizeof(int), cudaMemcpyDeviceToHost);
    
    cudaFree(gpu_a);
    cudaFree(gpu_b);
    cudaFree(gpu_add_dest);
    cudaFree(gpu_sub_dest);
    cudaFree(gpu_mult_dest);
    cudaFree(gpu_div_dest);
    cudaFree(gpu_mod_dest);
}

// Executes each of the shared memory gpu tests by creating local memory, copying it global memory, and then performing
// all 5 math operations on the data after creating shared memory. The data is filled with random numbers that uses the same seed as the CPU tests.
void executeSharedTest(const int totalThreads, const int blockSize, const int numBlocks)
{
    int a[totalThreads], b[totalThreads], add_dest[totalThreads], sub_dest[totalThreads], mult_dest[totalThreads], div_dest[totalThreads], mod_dest[totalThreads];
    
    int *gpu_a, *gpu_b, *gpu_add_dest, *gpu_sub_dest, *gpu_mult_dest, *gpu_div_dest, *gpu_mod_dest;

    cudaMalloc((void**)&gpu_a,         totalThreads * sizeof(int));
    cudaMalloc((void**)&gpu_b,         totalThreads * sizeof(int));
    cudaMalloc((void**)&gpu_add_dest,  totalThreads * sizeof(int));
    cudaMalloc((void**)&gpu_sub_dest,  totalThreads * sizeof(int));
    cudaMalloc((void**)&gpu_mult_dest, totalThreads * sizeof(int));
    cudaMalloc((void**)&gpu_div_dest,  totalThreads * sizeof(int));
    cudaMalloc((void**)&gpu_mod_dest,  totalThreads * sizeof(int));
    
    // Create a random generate that will generate random numbers from 0 to 4.
    // Use a set seed so output is deterministic
    unsigned seed = 12345;
    std::default_random_engine gen(seed);
    std::uniform_int_distribution<int> dist(0,4);
    
    for (size_t i = 0; i < totalThreads; ++i)
    {
        a[i] = i;
        b[i] = dist(gen);
    }
    
    cudaMemcpy(gpu_a, a, totalThreads * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(gpu_b, b, totalThreads * sizeof(int), cudaMemcpyHostToDevice);
    
    // The third parameter is the size of the shared memory
    // We multiply by 3 because we need to copy A and B and then also have room for the return in shared memory.
    executeSharedMathOperations<<<numBlocks, blockSize, 3 * totalThreads * sizeof(int)>>>(gpu_a,gpu_b,gpu_add_dest,gpu_sub_dest,gpu_mult_dest,gpu_div_dest,gpu_mod_dest, totalThreads);
    
    cudaMemcpy(add_dest,  gpu_add_dest,  totalThreads*sizeof(int), cudaMemcpyDeviceToHost);
    cudaMemcpy(sub_dest,  gpu_sub_dest,  totalThreads*sizeof(int), cudaMemcpyDeviceToHost);    
    cudaMemcpy(mult_dest, gpu_mult_dest, totalThreads*sizeof(int), cudaMemcpyDeviceToHost);
    cudaMemcpy(div_dest,  gpu_div_dest,  totalThreads*sizeof(int), cudaMemcpyDeviceToHost);
    cudaMemcpy(mod_dest,  gpu_mod_dest,  totalThreads*sizeof(int), cudaMemcpyDeviceToHost);
    
    cudaFree(gpu_a);
    cudaFree(gpu_b);
    cudaFree(gpu_add_dest);
    cudaFree(gpu_sub_dest);
    cudaFree(gpu_mult_dest);
    cudaFree(gpu_div_dest);
    cudaFree(gpu_mod_dest); 
}

// Executes each of the consnt memory gpu tests by creating local memory, copying it global memory, and then performing
// all 5 math operations on the data using constant values. The data is filled with random numbers that uses the same seed as the CPU tests.
void executeConstantTest(const int totalThreads, const int blockSize, const int numBlocks)
{
    int add_dest[totalThreads], sub_dest[totalThreads], mult_dest[totalThreads], div_dest[totalThreads], mod_dest[totalThreads];
    
    int *gpu_add_dest, *gpu_sub_dest, *gpu_mult_dest, *gpu_div_dest, *gpu_mod_dest;

    cudaMalloc((void**)&gpu_add_dest,  totalThreads * sizeof(int));
    cudaMalloc((void**)&gpu_sub_dest,  totalThreads * sizeof(int));
    cudaMalloc((void**)&gpu_mult_dest, totalThreads * sizeof(int));
    cudaMalloc((void**)&gpu_div_dest,  totalThreads * sizeof(int));
    cudaMalloc((void**)&gpu_mod_dest,  totalThreads * sizeof(int));
    
    // The third parameter is the size of the shared memory
    // We multiply by 3 because we need to copy A and B and then also have room for the return in shared memory.
    executeConstantMathOperations<<<numBlocks, blockSize>>>(gpu_add_dest,gpu_sub_dest,gpu_mult_dest,gpu_div_dest,gpu_mod_dest, totalThreads);
    
    cudaMemcpy(add_dest,  gpu_add_dest,  totalThreads*sizeof(int), cudaMemcpyDeviceToHost);
    cudaMemcpy(sub_dest,  gpu_sub_dest,  totalThreads*sizeof(int), cudaMemcpyDeviceToHost);    
    cudaMemcpy(mult_dest, gpu_mult_dest, totalThreads*sizeof(int), cudaMemcpyDeviceToHost);
    cudaMemcpy(div_dest,  gpu_div_dest,  totalThreads*sizeof(int), cudaMemcpyDeviceToHost);
    cudaMemcpy(mod_dest,  gpu_mod_dest,  totalThreads*sizeof(int), cudaMemcpyDeviceToHost);
    
    cudaFree(gpu_add_dest);
    cudaFree(gpu_sub_dest);
    cudaFree(gpu_mult_dest);
    cudaFree(gpu_div_dest);
    cudaFree(gpu_mod_dest); 
}

void printArray(const int * const arr, const int xSize, const int ySize)
{
    for (size_t i = 0; i < xSize; ++i)
    {
        for(size_t j = 0; j < ySize; ++j)
        {
            std::cout << arr[i * ySize + j] << " ";
        }
        std::cout << '\n';
    }
    
    std::cout << std::flush;
}

int main(int argc, char** argv)
{
    // read command line arguments
    int totalThreads = 256;
    int blockSize = 256;
    
    if (argc >= 2) {
        totalThreads = atoi(argv[1]);
    }
    if (argc >= 3) {
        blockSize = atoi(argv[2]);
    }
   
    int numBlocks = totalThreads/blockSize;

    // validate command line arguments
    if (totalThreads % blockSize != 0) {
        ++numBlocks;
        totalThreads = numBlocks*blockSize;
        
        printf("Warning: Total thread count is not evenly divisible by the block size\n");
        printf("The total number of threads will be rounded up to %d\n", totalThreads);
    }
    
    auto startTime = std::chrono::system_clock::now();
    executeHostTest(totalThreads, blockSize, numBlocks);
    auto endTime = std::chrono::system_clock::now();
    std::chrono::duration<double> totalTime = endTime-startTime;
    std::cout << "Host execution took: " << totalTime.count() << " seconds." << std::endl;
    
    startTime = std::chrono::system_clock::now();
    executeGlobalTest(totalThreads, blockSize, numBlocks);
    endTime = std::chrono::system_clock::now();
    totalTime = endTime-startTime;
    std::cout << "Global Memory execution took: " << totalTime.count() << " seconds." << std::endl;
    
    startTime = std::chrono::system_clock::now();
    executeSharedTest(totalThreads, blockSize, numBlocks);
    endTime = std::chrono::system_clock::now();
    totalTime = endTime-startTime;
    std::cout << "Shared Memory execution took: " << totalTime.count() << " seconds." << std::endl;
    
    startTime = std::chrono::system_clock::now();
    executeConstantTest(totalThreads, blockSize, numBlocks);
    endTime = std::chrono::system_clock::now();
    totalTime = endTime-startTime;
    std::cout << "Constant Memory execution took: " << totalTime.count() << " seconds." << std::endl;
    
    return 0;
}

	.file	"tmpxft_0020aa2d_00000000-6_assignment.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4840:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4840:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.align 2
	.type	_ZNSt24uniform_int_distributionIiEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEiRT_RKNS0_10param_typeE.isra.0, @function
_ZNSt24uniform_int_distributionIiEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEiRT_RKNS0_10param_typeE.isra.0:
.LFB5618:
	.cfi_startproc
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movslq	%esi, %rax
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	movq	%rdi, %r13
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	movq	%rax, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	movslq	%edx, %rbp
	subq	%rax, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	cmpq	$2147483644, %rbp
	ja	.L3
	incq	%rbp
	movl	$2147483645, %eax
	xorl	%edx, %edx
	movl	$2147483647, %esi
	divq	%rbp
	movq	(%rdi), %rdx
	imulq	%rax, %rbp
	movq	%rax, %rcx
.L4:
	imulq	$16807, %rdx, %rax
	xorl	%edx, %edx
	divq	%rsi
	leaq	-1(%rdx), %rax
	cmpq	%rbp, %rax
	jnb	.L4
	movq	%rdx, 0(%r13)
	xorl	%edx, %edx
	divq	%rcx
	jmp	.L5
.L3:
	cmpq	$2147483645, %rbp
	je	.L6
	movl	$2147483646, %ecx
	movq	%rbp, %rax
	xorl	%edx, %edx
	movl	$2147483647, %ebx
	divq	%rcx
	movq	%rax, %r14
.L11:
	movl	%r14d, %edx
	xorl	%esi, %esi
	movq	%r13, %rdi
	call	_ZNSt24uniform_int_distributionIiEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEiRT_RKNS0_10param_typeE.isra.0
	xorl	%edx, %edx
	movl	%eax, %ecx
	imulq	$16807, 0(%r13), %rax
	divq	%rbx
	movslq	%ecx, %rax
	imulq	$2147483646, %rax, %rax
	movq	%rdx, 0(%r13)
	decq	%rdx
	addq	%rdx, %rax
	setc	%dl
	movzbl	%dl, %edx
	cmpq	%rax, %rbp
	jb	.L11
	testq	%rdx, %rdx
	jne	.L11
	jmp	.L5
.L6:
	imulq	$16807, (%rdi), %rax
	movl	$2147483647, %ecx
	xorl	%edx, %edx
	divq	%rcx
	movq	%rdx, (%rdi)
	leaq	-1(%rdx), %rax
.L5:
	popq	%rbx
	.cfi_def_cfa_offset 40
	addl	%r12d, %eax
	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
.LFE5618:
	.size	_ZNSt24uniform_int_distributionIiEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEiRT_RKNS0_10param_typeE.isra.0, .-_ZNSt24uniform_int_distributionIiEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEiRT_RKNS0_10param_typeE.isra.0
	.align 2
	.type	_ZNSolsEPFRSoS_E.isra.0, @function
_ZNSolsEPFRSoS_E.isra.0:
.LFB5619:
	.cfi_startproc
	jmp	*%rsi
	.cfi_endproc
.LFE5619:
	.size	_ZNSolsEPFRSoS_E.isra.0, .-_ZNSolsEPFRSoS_E.isra.0
	.globl	_Z3addPiS_S_
	.type	_Z3addPiS_S_, @function
_Z3addPiS_S_:
.LFB4818:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4818:
	.size	_Z3addPiS_S_, .-_Z3addPiS_S_
	.globl	_Z8subtractPiS_S_
	.type	_Z8subtractPiS_S_, @function
_Z8subtractPiS_S_:
.LFB5611:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE5611:
	.size	_Z8subtractPiS_S_, .-_Z8subtractPiS_S_
	.globl	_Z4multPiS_S_
	.type	_Z4multPiS_S_, @function
_Z4multPiS_S_:
.LFB5613:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE5613:
	.size	_Z4multPiS_S_, .-_Z4multPiS_S_
	.globl	_Z3divPiS_S_
	.type	_Z3divPiS_S_, @function
_Z3divPiS_S_:
.LFB5615:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE5615:
	.size	_Z3divPiS_S_, .-_Z3divPiS_S_
	.globl	_Z3modPiS_S_
	.type	_Z3modPiS_S_, @function
_Z3modPiS_S_:
.LFB5617:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE5617:
	.size	_Z3modPiS_S_, .-_Z3modPiS_S_
	.globl	_Z7hostAddPiS_S_i
	.type	_Z7hostAddPiS_S_i, @function
_Z7hostAddPiS_S_i:
.LFB4823:
	.cfi_startproc
	endbr64
	xorl	%eax, %eax
.L26:
	cmpl	%eax, %ecx
	jle	.L28
	movl	(%rsi,%rax,4), %r8d
	addl	(%rdi,%rax,4), %r8d
	movl	%r8d, (%rdx,%rax,4)
	incq	%rax
	jmp	.L26
.L28:
	ret
	.cfi_endproc
.LFE4823:
	.size	_Z7hostAddPiS_S_i, .-_Z7hostAddPiS_S_i
	.globl	_Z7hostSubPiS_S_i
	.type	_Z7hostSubPiS_S_i, @function
_Z7hostSubPiS_S_i:
.LFB4824:
	.cfi_startproc
	endbr64
	xorl	%eax, %eax
.L30:
	cmpl	%eax, %ecx
	jle	.L32
	movl	(%rdi,%rax,4), %r8d
	subl	(%rsi,%rax,4), %r8d
	movl	%r8d, (%rdx,%rax,4)
	incq	%rax
	jmp	.L30
.L32:
	ret
	.cfi_endproc
.LFE4824:
	.size	_Z7hostSubPiS_S_i, .-_Z7hostSubPiS_S_i
	.globl	_Z8hostMultPiS_S_i
	.type	_Z8hostMultPiS_S_i, @function
_Z8hostMultPiS_S_i:
.LFB4825:
	.cfi_startproc
	endbr64
	xorl	%eax, %eax
.L34:
	cmpl	%eax, %ecx
	jle	.L36
	movl	(%rdi,%rax,4), %r8d
	imull	(%rsi,%rax,4), %r8d
	movl	%r8d, (%rdx,%rax,4)
	incq	%rax
	jmp	.L34
.L36:
	ret
	.cfi_endproc
.LFE4825:
	.size	_Z8hostMultPiS_S_i, .-_Z8hostMultPiS_S_i
	.globl	_Z7hostDivPiS_S_i
	.type	_Z7hostDivPiS_S_i, @function
_Z7hostDivPiS_S_i:
.LFB4826:
	.cfi_startproc
	endbr64
	movq	%rdx, %r10
	xorl	%r8d, %r8d
.L38:
	cmpl	%r8d, %ecx
	jle	.L44
	movl	(%rsi,%r8,4), %r9d
	testl	%r9d, %r9d
	je	.L39
	movl	(%rdi,%r8,4), %eax
	cltd
	idivl	%r9d
	movl	%eax, %r9d
.L39:
	movl	%r9d, (%r10,%r8,4)
	incq	%r8
	jmp	.L38
.L44:
	ret
	.cfi_endproc
.LFE4826:
	.size	_Z7hostDivPiS_S_i, .-_Z7hostDivPiS_S_i
	.globl	_Z7hostModPiS_S_i
	.type	_Z7hostModPiS_S_i, @function
_Z7hostModPiS_S_i:
.LFB4827:
	.cfi_startproc
	endbr64
	movq	%rdx, %r9
	xorl	%r8d, %r8d
.L46:
	cmpl	%r8d, %ecx
	jle	.L52
	movl	(%rsi,%r8,4), %edx
	testl	%edx, %edx
	je	.L47
	movl	(%rdi,%r8,4), %eax
	movl	%edx, %r10d
	cltd
	idivl	%r10d
.L47:
	movl	%edx, (%r9,%r8,4)
	incq	%r8
	jmp	.L46
.L52:
	ret
	.cfi_endproc
.LFE4827:
	.size	_Z7hostModPiS_S_i, .-_Z7hostModPiS_S_i
	.globl	_Z15executeHostTestiii
	.type	_Z15executeHostTestiii, @function
_Z15executeHostTestiii:
.LFB4828:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movslq	%edi, %r9
	movq	%r9, %r11
	movq	%rsp, %rbp
	.cfi_def_cfa_register 6
	pushq	%r14
	pushq	%r12
	pushq	%rbx
	subq	$24, %rsp
	.cfi_offset 14, -24
	.cfi_offset 12, -32
	.cfi_offset 3, -40
	movq	%fs:40, %rax
	movq	%rax, -40(%rbp)
	xorl	%eax, %eax
	leaq	15(,%r9,4), %rax
	movq	%rsp, %rcx
	movq	%rax, %rsi
	movq	%rax, %rdx
	andq	$-4096, %rsi
	andq	$-16, %rdx
	subq	%rsi, %rcx
.L54:
	cmpq	%rcx, %rsp
	je	.L55
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L54
.L55:
	andl	$4095, %edx
	subq	%rdx, %rsp
	testq	%rdx, %rdx
	je	.L56
	orq	$0, -8(%rsp,%rdx)
.L56:
	movq	%rax, %rsi
	movq	%rax, %rdx
	movq	%rsp, %rcx
	movq	%rsp, %rbx
	andq	$-4096, %rsi
	andq	$-16, %rdx
	subq	%rsi, %rcx
.L57:
	cmpq	%rcx, %rsp
	je	.L58
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L57
.L58:
	andl	$4095, %edx
	subq	%rdx, %rsp
	testq	%rdx, %rdx
	je	.L59
	orq	$0, -8(%rsp,%rdx)
.L59:
	movq	%rax, %rdx
	movq	%rsp, %rcx
	andq	$-4096, %rax
	movq	%rsp, %r12
	andq	$-16, %rdx
	subq	%rax, %rcx
.L60:
	cmpq	%rcx, %rsp
	je	.L61
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L60
.L61:
	movq	%rdx, %rax
	andl	$4095, %eax
	subq	%rax, %rsp
	testq	%rax, %rax
	je	.L62
	orq	$0, -8(%rsp,%rax)
.L62:
	movq	$12345, -48(%rbp)
	movq	%rsp, %r14
	xorl	%r8d, %r8d
	leaq	-48(%rbp), %r10
.L63:
	cmpq	%r9, %r8
	jnb	.L76
	movl	%r8d, (%rbx,%r8,4)
	movl	$4, %edx
	xorl	%esi, %esi
	movq	%r10, %rdi
	call	_ZNSt24uniform_int_distributionIiEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEiRT_RKNS0_10param_typeE.isra.0
	movl	%eax, (%r12,%r8,4)
	incq	%r8
	jmp	.L63
.L76:
	movq	%r14, %rdx
	movl	%r11d, %ecx
	movq	%r12, %rsi
	movq	%rbx, %rdi
	call	_Z7hostAddPiS_S_i
	call	_Z7hostSubPiS_S_i
	call	_Z8hostMultPiS_S_i
	call	_Z7hostDivPiS_S_i
	movq	%r14, %rdx
	call	_Z7hostModPiS_S_i
	movq	-40(%rbp), %rax
	subq	%fs:40, %rax
	je	.L65
	call	__stack_chk_fail@PLT
.L65:
	leaq	-24(%rbp), %rsp
	popq	%rbx
	popq	%r12
	popq	%r14
	popq	%rbp
	.cfi_def_cfa 7, 8
	ret
	.cfi_endproc
.LFE4828:
	.size	_Z15executeHostTestiii, .-_Z15executeHostTestiii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	" "
	.text
	.globl	_Z10printArrayPKiii
	.type	_Z10printArrayPKiii, @function
_Z10printArrayPKiii:
.LFB4832:
	.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
	movslq	%esi, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	xorl	%r13d, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	xorl	%r12d, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movslq	%edx, %rbx
	subq	$40, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 24(%rsp)
	xorl	%eax, %eax
.L78:
	cmpq	%r14, %r13
	leaq	_ZSt4cout(%rip), %r15
	jnb	.L82
	leaq	0(%rbp,%r12,4), %r8
	xorl	%edx, %edx
.L83:
	cmpq	%rbx, %rdx
	jnb	.L87
	movl	(%r8,%rdx,4), %esi
	movq	%r15, %rdi
	movq	%r8, 8(%rsp)
	movq	%rdx, (%rsp)
	call	_ZNSolsEi@PLT
	leaq	.LC0(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	(%rsp), %rdx
	movq	8(%rsp), %r8
	incq	%rdx
	jmp	.L83
.L87:
	movq	_ZSt4cout(%rip), %rax
	movb	$10, 23(%rsp)
	movq	-24(%rax), %rax
	cmpq	$0, 16(%r15,%rax)
	je	.L80
	leaq	23(%rsp), %rsi
	movl	$1, %edx
	movq	%r15, %rdi
	call	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l@PLT
	jmp	.L81
.L80:
	movl	$10, %esi
	movq	%r15, %rdi
	call	_ZNSo3putEc@PLT
.L81:
	incq	%r13
	addq	%rbx, %r12
	jmp	.L78
.L82:
	movq	24(%rsp), %rax
	subq	%fs:40, %rax
	je	.L84
	call	__stack_chk_fail@PLT
.L84:
	movq	_ZSt5flushIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	addq	$40, %rsp
	.cfi_def_cfa_offset 56
	movq	%r15, %rdi
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	jmp	_ZNSolsEPFRSoS_E.isra.0
	.cfi_endproc
.LFE4832:
	.size	_Z10printArrayPKiii, .-_Z10printArrayPKiii
	.globl	_Z60__device_stub__Z27executeSharedMathOperationsPiS_S_S_S_S_S_iPiS_S_S_S_S_S_i
	.type	_Z60__device_stub__Z27executeSharedMathOperationsPiS_S_S_S_S_S_iPiS_S_S_S_S_S_i, @function
_Z60__device_stub__Z27executeSharedMathOperationsPiS_S_S_S_S_S_iPiS_S_S_S_S_S_i:
.LFB4862:
	.cfi_startproc
	endbr64
	subq	$216, %rsp
	.cfi_def_cfa_offset 224
	movq	224(%rsp), %rax
	movq	%rdi, 56(%rsp)
	leaq	88(%rsp), %rdi
	movq	%rsi, 48(%rsp)
	leaq	100(%rsp), %rsi
	movq	%rdx, 40(%rsp)
	leaq	72(%rsp), %rdx
	movq	%rcx, 32(%rsp)
	leaq	80(%rsp), %rcx
	movq	%r8, 24(%rsp)
	movq	%r9, 16(%rsp)
	movq	%rax, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 200(%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	24(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	232(%rsp), %rax
	movq	%rax, 192(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 88(%rsp)
	movq	%rax, 100(%rsp)
	movl	$1, 108(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L88
	pushq	80(%rsp)
	.cfi_def_cfa_offset 232
	leaq	_Z27executeSharedMathOperationsPiS_S_S_S_S_S_i(%rip), %rdi
	pushq	80(%rsp)
	.cfi_def_cfa_offset 240
	movq	116(%rsp), %rcx
	movl	124(%rsp), %r8d
	movq	104(%rsp), %rsi
	movl	112(%rsp), %edx
	leaq	152(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 232
	popq	%rdx
	.cfi_def_cfa_offset 224
.L88:
	movq	200(%rsp), %rax
	subq	%fs:40, %rax
	je	.L90
	call	__stack_chk_fail@PLT
.L90:
	addq	$216, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4862:
	.size	_Z60__device_stub__Z27executeSharedMathOperationsPiS_S_S_S_S_S_iPiS_S_S_S_S_S_i, .-_Z60__device_stub__Z27executeSharedMathOperationsPiS_S_S_S_S_S_iPiS_S_S_S_S_S_i
	.globl	_Z27executeSharedMathOperationsPiS_S_S_S_S_S_i
	.type	_Z27executeSharedMathOperationsPiS_S_S_S_S_S_i, @function
_Z27executeSharedMathOperationsPiS_S_S_S_S_S_i:
.LFB4863:
	.cfi_startproc
	endbr64
	jmp	_Z60__device_stub__Z27executeSharedMathOperationsPiS_S_S_S_S_S_iPiS_S_S_S_S_S_i
	.cfi_endproc
.LFE4863:
	.size	_Z27executeSharedMathOperationsPiS_S_S_S_S_S_i, .-_Z27executeSharedMathOperationsPiS_S_S_S_S_S_i
	.globl	_Z17executeSharedTestiii
	.type	_Z17executeSharedTestiii, @function
_Z17executeSharedTestiii:
.LFB4830:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register 6
	pushq	%r15
	.cfi_offset 15, -24
	movslq	%edi, %r15
	pushq	%r14
	.cfi_offset 14, -32
	movl	%esi, %r14d
	pushq	%r13
	.cfi_offset 13, -40
	movl	%edx, %r13d
	pushq	%r12
	.cfi_offset 12, -48
	movq	%r15, %r12
	pushq	%rbx
	.cfi_offset 3, -56
	leaq	0(,%r15,4), %rbx
	subq	$168, %rsp
	movq	%fs:40, %rax
	movq	%rax, -56(%rbp)
	xorl	%eax, %eax
	leaq	15(%rbx), %rax
	movq	%rsp, %rcx
	movq	%rax, %rsi
	movq	%rax, %rdx
	andq	$-4096, %rsi
	andq	$-16, %rdx
	subq	%rsi, %rcx
.L94:
	cmpq	%rcx, %rsp
	je	.L95
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L94
.L95:
	andl	$4095, %edx
	subq	%rdx, %rsp
	testq	%rdx, %rdx
	je	.L96
	orq	$0, -8(%rsp,%rdx)
.L96:
	movq	%rax, %rsi
	movq	%rsp, -152(%rbp)
	movq	%rax, %rdx
	movq	%rsp, %rcx
	andq	$-4096, %rsi
	andq	$-16, %rdx
	subq	%rsi, %rcx
.L97:
	cmpq	%rcx, %rsp
	je	.L98
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L97
.L98:
	andl	$4095, %edx
	subq	%rdx, %rsp
	testq	%rdx, %rdx
	je	.L99
	orq	$0, -8(%rsp,%rdx)
.L99:
	movq	%rax, %rsi
	movq	%rsp, -160(%rbp)
	movq	%rax, %rdx
	movq	%rsp, %rcx
	andq	$-4096, %rsi
	andq	$-16, %rdx
	subq	%rsi, %rcx
.L100:
	cmpq	%rcx, %rsp
	je	.L101
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L100
.L101:
	andl	$4095, %edx
	subq	%rdx, %rsp
	testq	%rdx, %rdx
	je	.L102
	orq	$0, -8(%rsp,%rdx)
.L102:
	movq	%rax, %rsi
	movq	%rsp, -168(%rbp)
	movq	%rax, %rdx
	movq	%rsp, %rcx
	andq	$-4096, %rsi
	andq	$-16, %rdx
	subq	%rsi, %rcx
.L103:
	cmpq	%rcx, %rsp
	je	.L104
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L103
.L104:
	andl	$4095, %edx
	subq	%rdx, %rsp
	testq	%rdx, %rdx
	je	.L105
	orq	$0, -8(%rsp,%rdx)
.L105:
	movq	%rax, %rsi
	movq	%rsp, -176(%rbp)
	movq	%rax, %rdx
	movq	%rsp, %rcx
	andq	$-4096, %rsi
	andq	$-16, %rdx
	subq	%rsi, %rcx
.L106:
	cmpq	%rcx, %rsp
	je	.L107
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L106
.L107:
	andl	$4095, %edx
	subq	%rdx, %rsp
	testq	%rdx, %rdx
	je	.L108
	orq	$0, -8(%rsp,%rdx)
.L108:
	movq	%rax, %rsi
	movq	%rsp, -184(%rbp)
	movq	%rax, %rdx
	movq	%rsp, %rcx
	andq	$-4096, %rsi
	andq	$-16, %rdx
	subq	%rsi, %rcx
.L109:
	cmpq	%rcx, %rsp
	je	.L110
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L109
.L110:
	andl	$4095, %edx
	subq	%rdx, %rsp
	testq	%rdx, %rdx
	je	.L111
	orq	$0, -8(%rsp,%rdx)
.L111:
	movq	%rsp, -192(%rbp)
	movq	%rax, %rdx
	movq	%rsp, %rcx
	andq	$-4096, %rax
	andq	$-16, %rdx
	subq	%rax, %rcx
.L112:
	cmpq	%rcx, %rsp
	je	.L113
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L112
.L113:
	movq	%rdx, %rax
	andl	$4095, %eax
	subq	%rax, %rsp
	testq	%rax, %rax
	je	.L114
	orq	$0, -8(%rsp,%rax)
.L114:
	leaq	-144(%rbp), %rdi
	movq	%rbx, %rsi
	movq	%rsp, -200(%rbp)
	call	cudaMalloc@PLT
	leaq	-136(%rbp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	leaq	-128(%rbp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	leaq	-120(%rbp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	leaq	-112(%rbp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	leaq	-104(%rbp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	leaq	-96(%rbp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	movq	$12345, -88(%rbp)
	xorl	%r8d, %r8d
	leaq	-88(%rbp), %r9
.L115:
	cmpq	%r15, %r8
	jnb	.L141
	movq	-152(%rbp), %rax
	movq	%r9, %rdi
	movl	$4, %edx
	xorl	%esi, %esi
	movl	%r8d, (%rax,%r8,4)
	call	_ZNSt24uniform_int_distributionIiEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEiRT_RKNS0_10param_typeE.isra.0
	movq	-160(%rbp), %rdi
	movl	%eax, (%rdi,%r8,4)
	incq	%r8
	jmp	.L115
.L141:
	movq	-152(%rbp), %rsi
	movq	-144(%rbp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rdx
	call	cudaMemcpy@PLT
	movq	-160(%rbp), %rsi
	movl	$1, %ecx
	movq	%rbx, %rdx
	movq	-136(%rbp), %rdi
	call	cudaMemcpy@PLT
	movl	%r14d, -68(%rbp)
	xorl	%r9d, %r9d
	movabsq	$4294967297, %rax
	movq	%rax, -64(%rbp)
	movl	-60(%rbp), %ecx
	movq	%rax, -76(%rbp)
	leal	(%r12,%r12,2), %eax
	movq	-68(%rbp), %rdx
	movl	%r13d, -80(%rbp)
	cltq
	movq	-80(%rbp), %rdi
	movl	-72(%rbp), %esi
	leaq	0(,%rax,4), %r8
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L117
	pushq	%r12
	pushq	-96(%rbp)
	movq	-128(%rbp), %rdx
	movq	-104(%rbp), %r9
	movq	-112(%rbp), %r8
	movq	-120(%rbp), %rcx
	movq	-136(%rbp), %rsi
	movq	-144(%rbp), %rdi
	call	_Z60__device_stub__Z27executeSharedMathOperationsPiS_S_S_S_S_S_iPiS_S_S_S_S_S_i
	popq	%rax
	popq	%rdx
.L117:
	movq	-128(%rbp), %rsi
	movq	-168(%rbp), %rdi
	movl	$2, %ecx
	movq	%rbx, %rdx
	call	cudaMemcpy@PLT
	movq	-120(%rbp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	-176(%rbp), %rdi
	call	cudaMemcpy@PLT
	movq	-112(%rbp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	-184(%rbp), %rdi
	call	cudaMemcpy@PLT
	movq	-104(%rbp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	-192(%rbp), %rdi
	call	cudaMemcpy@PLT
	movq	-96(%rbp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	-200(%rbp), %rdi
	call	cudaMemcpy@PLT
	movq	-144(%rbp), %rdi
	call	cudaFree@PLT
	movq	-136(%rbp), %rdi
	call	cudaFree@PLT
	movq	-128(%rbp), %rdi
	call	cudaFree@PLT
	movq	-120(%rbp), %rdi
	call	cudaFree@PLT
	movq	-112(%rbp), %rdi
	call	cudaFree@PLT
	movq	-104(%rbp), %rdi
	call	cudaFree@PLT
	movq	-96(%rbp), %rdi
	call	cudaFree@PLT
	movq	-56(%rbp), %rax
	subq	%fs:40, %rax
	je	.L118
	call	__stack_chk_fail@PLT
.L118:
	leaq	-40(%rbp), %rsp
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa 7, 8
	ret
	.cfi_endproc
.LFE4830:
	.size	_Z17executeSharedTestiii, .-_Z17executeSharedTestiii
	.globl	_Z60__device_stub__Z27executeGlobalMathOperationsPiS_S_S_S_S_S_iPiS_S_S_S_S_S_i
	.type	_Z60__device_stub__Z27executeGlobalMathOperationsPiS_S_S_S_S_S_iPiS_S_S_S_S_S_i, @function
_Z60__device_stub__Z27executeGlobalMathOperationsPiS_S_S_S_S_S_iPiS_S_S_S_S_S_i:
.LFB4864:
	.cfi_startproc
	endbr64
	subq	$216, %rsp
	.cfi_def_cfa_offset 224
	movq	224(%rsp), %rax
	movq	%rdi, 56(%rsp)
	leaq	88(%rsp), %rdi
	movq	%rsi, 48(%rsp)
	leaq	100(%rsp), %rsi
	movq	%rdx, 40(%rsp)
	leaq	72(%rsp), %rdx
	movq	%rcx, 32(%rsp)
	leaq	80(%rsp), %rcx
	movq	%r8, 24(%rsp)
	movq	%r9, 16(%rsp)
	movq	%rax, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 200(%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	24(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	232(%rsp), %rax
	movq	%rax, 192(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 88(%rsp)
	movq	%rax, 100(%rsp)
	movl	$1, 108(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L142
	pushq	80(%rsp)
	.cfi_def_cfa_offset 232
	leaq	_Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i(%rip), %rdi
	pushq	80(%rsp)
	.cfi_def_cfa_offset 240
	movq	116(%rsp), %rcx
	movl	124(%rsp), %r8d
	movq	104(%rsp), %rsi
	movl	112(%rsp), %edx
	leaq	152(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 232
	popq	%rdx
	.cfi_def_cfa_offset 224
.L142:
	movq	200(%rsp), %rax
	subq	%fs:40, %rax
	je	.L144
	call	__stack_chk_fail@PLT
.L144:
	addq	$216, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4864:
	.size	_Z60__device_stub__Z27executeGlobalMathOperationsPiS_S_S_S_S_S_iPiS_S_S_S_S_S_i, .-_Z60__device_stub__Z27executeGlobalMathOperationsPiS_S_S_S_S_S_iPiS_S_S_S_S_S_i
	.globl	_Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i
	.type	_Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i, @function
_Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i:
.LFB4865:
	.cfi_startproc
	endbr64
	jmp	_Z60__device_stub__Z27executeGlobalMathOperationsPiS_S_S_S_S_S_iPiS_S_S_S_S_S_i
	.cfi_endproc
.LFE4865:
	.size	_Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i, .-_Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i
	.globl	_Z17executeGlobalTestiii
	.type	_Z17executeGlobalTestiii, @function
_Z17executeGlobalTestiii:
.LFB4829:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register 6
	pushq	%r15
	pushq	%r14
	.cfi_offset 15, -24
	.cfi_offset 14, -32
	movslq	%edi, %r14
	pushq	%r13
	.cfi_offset 13, -40
	movl	%edx, %r13d
	pushq	%r12
	.cfi_offset 12, -48
	movl	%esi, %r12d
	pushq	%rbx
	.cfi_offset 3, -56
	leaq	0(,%r14,4), %rbx
	subq	$152, %rsp
	movq	%fs:40, %rax
	movq	%rax, -56(%rbp)
	xorl	%eax, %eax
	leaq	15(%rbx), %rax
	movq	%rsp, %rcx
	movq	%rax, %rsi
	movq	%rax, %rdx
	andq	$-4096, %rsi
	andq	$-16, %rdx
	subq	%rsi, %rcx
.L148:
	cmpq	%rcx, %rsp
	je	.L149
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L148
.L149:
	andl	$4095, %edx
	subq	%rdx, %rsp
	testq	%rdx, %rdx
	je	.L150
	orq	$0, -8(%rsp,%rdx)
.L150:
	movq	%rax, %rsi
	movq	%rax, %rdx
	movq	%rsp, %rcx
	movq	%rsp, %r15
	andq	$-4096, %rsi
	andq	$-16, %rdx
	subq	%rsi, %rcx
.L151:
	cmpq	%rcx, %rsp
	je	.L152
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L151
.L152:
	andl	$4095, %edx
	subq	%rdx, %rsp
	testq	%rdx, %rdx
	je	.L153
	orq	$0, -8(%rsp,%rdx)
.L153:
	movq	%rax, %rsi
	movq	%rsp, -152(%rbp)
	movq	%rax, %rdx
	movq	%rsp, %rcx
	andq	$-4096, %rsi
	andq	$-16, %rdx
	subq	%rsi, %rcx
.L154:
	cmpq	%rcx, %rsp
	je	.L155
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L154
.L155:
	andl	$4095, %edx
	subq	%rdx, %rsp
	testq	%rdx, %rdx
	je	.L156
	orq	$0, -8(%rsp,%rdx)
.L156:
	movq	%rax, %rsi
	movq	%rsp, -160(%rbp)
	movq	%rax, %rdx
	movq	%rsp, %rcx
	andq	$-4096, %rsi
	andq	$-16, %rdx
	subq	%rsi, %rcx
.L157:
	cmpq	%rcx, %rsp
	je	.L158
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L157
.L158:
	andl	$4095, %edx
	subq	%rdx, %rsp
	testq	%rdx, %rdx
	je	.L159
	orq	$0, -8(%rsp,%rdx)
.L159:
	movq	%rax, %rsi
	movq	%rsp, -168(%rbp)
	movq	%rax, %rdx
	movq	%rsp, %rcx
	andq	$-4096, %rsi
	andq	$-16, %rdx
	subq	%rsi, %rcx
.L160:
	cmpq	%rcx, %rsp
	je	.L161
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L160
.L161:
	andl	$4095, %edx
	subq	%rdx, %rsp
	testq	%rdx, %rdx
	je	.L162
	orq	$0, -8(%rsp,%rdx)
.L162:
	movq	%rax, %rsi
	movq	%rsp, -176(%rbp)
	movq	%rax, %rdx
	movq	%rsp, %rcx
	andq	$-4096, %rsi
	andq	$-16, %rdx
	subq	%rsi, %rcx
.L163:
	cmpq	%rcx, %rsp
	je	.L164
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L163
.L164:
	andl	$4095, %edx
	subq	%rdx, %rsp
	testq	%rdx, %rdx
	je	.L165
	orq	$0, -8(%rsp,%rdx)
.L165:
	movq	%rsp, -184(%rbp)
	movq	%rax, %rdx
	movq	%rsp, %rcx
	andq	$-4096, %rax
	andq	$-16, %rdx
	subq	%rax, %rcx
.L166:
	cmpq	%rcx, %rsp
	je	.L167
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L166
.L167:
	movq	%rdx, %rax
	andl	$4095, %eax
	subq	%rax, %rsp
	testq	%rax, %rax
	je	.L168
	orq	$0, -8(%rsp,%rax)
.L168:
	leaq	-144(%rbp), %rdi
	movq	%rbx, %rsi
	movq	%rsp, -192(%rbp)
	call	cudaMalloc@PLT
	leaq	-136(%rbp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	leaq	-128(%rbp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	leaq	-120(%rbp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	leaq	-112(%rbp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	leaq	-104(%rbp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	leaq	-96(%rbp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	movq	$12345, -88(%rbp)
	xorl	%r8d, %r8d
	leaq	-88(%rbp), %r9
.L169:
	cmpq	%r14, %r8
	jnb	.L195
	movl	%r8d, (%r15,%r8,4)
	movq	%r9, %rdi
	movl	$4, %edx
	xorl	%esi, %esi
	call	_ZNSt24uniform_int_distributionIiEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEiRT_RKNS0_10param_typeE.isra.0
	movq	-152(%rbp), %rdi
	movl	%eax, (%rdi,%r8,4)
	incq	%r8
	jmp	.L169
.L195:
	movq	-144(%rbp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rdx
	movq	%r15, %rsi
	call	cudaMemcpy@PLT
	movq	-152(%rbp), %rsi
	movl	$1, %ecx
	movq	%rbx, %rdx
	movq	-136(%rbp), %rdi
	call	cudaMemcpy@PLT
	movl	%r12d, -68(%rbp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movabsq	$4294967297, %rax
	movl	%r13d, -80(%rbp)
	movq	%rax, -64(%rbp)
	movl	-60(%rbp), %ecx
	movq	%rax, -76(%rbp)
	movq	-68(%rbp), %rdx
	movq	-80(%rbp), %rdi
	movl	-72(%rbp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L171
	imull	%r13d, %r12d
	pushq	%r12
	pushq	-96(%rbp)
	movq	-128(%rbp), %rdx
	movq	-104(%rbp), %r9
	movq	-112(%rbp), %r8
	movq	-120(%rbp), %rcx
	movq	-136(%rbp), %rsi
	movq	-144(%rbp), %rdi
	call	_Z60__device_stub__Z27executeGlobalMathOperationsPiS_S_S_S_S_S_iPiS_S_S_S_S_S_i
	popq	%rax
	popq	%rdx
.L171:
	movq	-128(%rbp), %rsi
	movq	-160(%rbp), %rdi
	movl	$2, %ecx
	movq	%rbx, %rdx
	call	cudaMemcpy@PLT
	movq	-120(%rbp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	-168(%rbp), %rdi
	call	cudaMemcpy@PLT
	movq	-112(%rbp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	-176(%rbp), %rdi
	call	cudaMemcpy@PLT
	movq	-104(%rbp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	-184(%rbp), %rdi
	call	cudaMemcpy@PLT
	movq	-96(%rbp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	-192(%rbp), %rdi
	call	cudaMemcpy@PLT
	movq	-144(%rbp), %rdi
	call	cudaFree@PLT
	movq	-136(%rbp), %rdi
	call	cudaFree@PLT
	movq	-128(%rbp), %rdi
	call	cudaFree@PLT
	movq	-120(%rbp), %rdi
	call	cudaFree@PLT
	movq	-112(%rbp), %rdi
	call	cudaFree@PLT
	movq	-104(%rbp), %rdi
	call	cudaFree@PLT
	movq	-96(%rbp), %rdi
	call	cudaFree@PLT
	movq	-56(%rbp), %rax
	subq	%fs:40, %rax
	je	.L172
	call	__stack_chk_fail@PLT
.L172:
	leaq	-40(%rbp), %rsp
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa 7, 8
	ret
	.cfi_endproc
.LFE4829:
	.size	_Z17executeGlobalTestiii, .-_Z17executeGlobalTestiii
	.globl	_Z58__device_stub__Z29executeConstantMathOperationsPiS_S_S_S_iPiS_S_S_S_i
	.type	_Z58__device_stub__Z29executeConstantMathOperationsPiS_S_S_S_iPiS_S_S_S_i, @function
_Z58__device_stub__Z29executeConstantMathOperationsPiS_S_S_S_iPiS_S_S_S_i:
.LFB4866:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movq	%rdi, 40(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	56(%rsp), %rdx
	movq	%rcx, 16(%rsp)
	leaq	64(%rsp), %rcx
	movq	%r8, 8(%rsp)
	movl	%r9d, 4(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%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	16(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L196
	pushq	64(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z29executeConstantMathOperationsPiS_S_S_S_i(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 208
	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 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L196:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L198
	call	__stack_chk_fail@PLT
.L198:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4866:
	.size	_Z58__device_stub__Z29executeConstantMathOperationsPiS_S_S_S_iPiS_S_S_S_i, .-_Z58__device_stub__Z29executeConstantMathOperationsPiS_S_S_S_iPiS_S_S_S_i
	.globl	_Z29executeConstantMathOperationsPiS_S_S_S_i
	.type	_Z29executeConstantMathOperationsPiS_S_S_S_i, @function
_Z29executeConstantMathOperationsPiS_S_S_S_i:
.LFB4867:
	.cfi_startproc
	endbr64
	jmp	_Z58__device_stub__Z29executeConstantMathOperationsPiS_S_S_S_iPiS_S_S_S_i
	.cfi_endproc
.LFE4867:
	.size	_Z29executeConstantMathOperationsPiS_S_S_S_i, .-_Z29executeConstantMathOperationsPiS_S_S_S_i
	.globl	_Z19executeConstantTestiii
	.type	_Z19executeConstantTestiii, @function
_Z19executeConstantTestiii:
.LFB4831:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register 6
	pushq	%r15
	pushq	%r14
	.cfi_offset 15, -24
	.cfi_offset 14, -32
	movl	%edx, %r14d
	pushq	%r13
	.cfi_offset 13, -40
	movl	%esi, %r13d
	pushq	%r12
	pushq	%rbx
	.cfi_offset 12, -48
	.cfi_offset 3, -56
	movslq	%edi, %rbx
	movq	%rbx, %r12
	salq	$2, %rbx
	subq	$120, %rsp
	movq	%fs:40, %rax
	movq	%rax, -56(%rbp)
	xorl	%eax, %eax
	leaq	15(%rbx), %rax
	movq	%rsp, %rcx
	movq	%rax, %rsi
	movq	%rax, %rdx
	andq	$-4096, %rsi
	andq	$-16, %rdx
	subq	%rsi, %rcx
.L202:
	cmpq	%rcx, %rsp
	je	.L203
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L202
.L203:
	andl	$4095, %edx
	subq	%rdx, %rsp
	testq	%rdx, %rdx
	je	.L204
	orq	$0, -8(%rsp,%rdx)
.L204:
	movq	%rax, %rsi
	movq	%rax, %rdx
	movq	%rsp, %rcx
	movq	%rsp, %r15
	andq	$-4096, %rsi
	andq	$-16, %rdx
	subq	%rsi, %rcx
.L205:
	cmpq	%rcx, %rsp
	je	.L206
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L205
.L206:
	andl	$4095, %edx
	subq	%rdx, %rsp
	testq	%rdx, %rdx
	je	.L207
	orq	$0, -8(%rsp,%rdx)
.L207:
	movq	%rax, %rsi
	movq	%rsp, -136(%rbp)
	movq	%rax, %rdx
	movq	%rsp, %rcx
	andq	$-4096, %rsi
	andq	$-16, %rdx
	subq	%rsi, %rcx
.L208:
	cmpq	%rcx, %rsp
	je	.L209
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L208
.L209:
	andl	$4095, %edx
	subq	%rdx, %rsp
	testq	%rdx, %rdx
	je	.L210
	orq	$0, -8(%rsp,%rdx)
.L210:
	movq	%rax, %rsi
	movq	%rsp, -144(%rbp)
	movq	%rax, %rdx
	movq	%rsp, %rcx
	andq	$-4096, %rsi
	andq	$-16, %rdx
	subq	%rsi, %rcx
.L211:
	cmpq	%rcx, %rsp
	je	.L212
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L211
.L212:
	andl	$4095, %edx
	subq	%rdx, %rsp
	testq	%rdx, %rdx
	je	.L213
	orq	$0, -8(%rsp,%rdx)
.L213:
	movq	%rsp, -152(%rbp)
	movq	%rax, %rdx
	movq	%rsp, %rcx
	andq	$-4096, %rax
	andq	$-16, %rdx
	subq	%rax, %rcx
.L214:
	cmpq	%rcx, %rsp
	je	.L215
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L214
.L215:
	movq	%rdx, %rax
	andl	$4095, %eax
	subq	%rax, %rsp
	testq	%rax, %rax
	je	.L216
	orq	$0, -8(%rsp,%rax)
.L216:
	leaq	-120(%rbp), %rdi
	movq	%rbx, %rsi
	movq	%rsp, -160(%rbp)
	call	cudaMalloc@PLT
	leaq	-112(%rbp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	leaq	-104(%rbp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	leaq	-96(%rbp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	leaq	-88(%rbp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	movl	%r13d, -68(%rbp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movabsq	$4294967297, %rax
	movl	%r14d, -80(%rbp)
	movq	%rax, -64(%rbp)
	movl	-60(%rbp), %ecx
	movq	%rax, -76(%rbp)
	movq	-68(%rbp), %rdx
	movq	-80(%rbp), %rdi
	movl	-72(%rbp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L217
	movq	-88(%rbp), %r8
	movq	-96(%rbp), %rcx
	movl	%r12d, %r9d
	movq	-104(%rbp), %rdx
	movq	-112(%rbp), %rsi
	movq	-120(%rbp), %rdi
	call	_Z58__device_stub__Z29executeConstantMathOperationsPiS_S_S_S_iPiS_S_S_S_i
.L217:
	movq	-120(%rbp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	%r15, %rdi
	call	cudaMemcpy@PLT
	movq	-112(%rbp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	-136(%rbp), %rdi
	call	cudaMemcpy@PLT
	movq	-104(%rbp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	-144(%rbp), %rdi
	call	cudaMemcpy@PLT
	movq	-96(%rbp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	-152(%rbp), %rdi
	call	cudaMemcpy@PLT
	movq	-88(%rbp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	-160(%rbp), %rdi
	call	cudaMemcpy@PLT
	movq	-120(%rbp), %rdi
	call	cudaFree@PLT
	movq	-112(%rbp), %rdi
	call	cudaFree@PLT
	movq	-104(%rbp), %rdi
	call	cudaFree@PLT
	movq	-96(%rbp), %rdi
	call	cudaFree@PLT
	movq	-88(%rbp), %rdi
	call	cudaFree@PLT
	movq	-56(%rbp), %rax
	subq	%fs:40, %rax
	je	.L218
	call	__stack_chk_fail@PLT
.L218:
	leaq	-40(%rbp), %rsp
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa 7, 8
	ret
	.cfi_endproc
.LFE4831:
	.size	_Z19executeConstantTestiii, .-_Z19executeConstantTestiii
	.section	.rodata.str1.1
.LC1:
	.string	"_Z29executeConstantMathOperationsPiS_S_S_S_i"
.LC2:
	.string	"_Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i"
.LC3:
	.string	"_Z27executeSharedMathOperationsPiS_S_S_S_S_S_i"
.LC4:
	.string	"VAL_A"
.LC5:
	.string	"VAL_B"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4869:
	.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	.LC1(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z29executeConstantMathOperationsPiS_S_S_S_i(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC2(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC3(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z27executeSharedMathOperationsPiS_S_S_S_S_S_i(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC4(%rip), %rdx
	movl	$4, %r9d
	leaq	_ZL5VAL_A(%rip), %rsi
	pushq	$1
	.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	$1
	.cfi_def_cfa_offset 32
	leaq	.LC5(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$4, %r9d
	leaq	_ZL5VAL_B(%rip), %rsi
	call	__cudaRegisterVar@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.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
.LFE4869:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.text._ZNSt6chrono8durationIdSt5ratioILl1ELl1EEEC2IlS1_ILl1ELl1000000000EEvEERKNS0_IT_T0_EE,"axG",@progbits,_ZNSt6chrono8durationIdSt5ratioILl1ELl1EEEC5IlS1_ILl1ELl1000000000EEvEERKNS0_IT_T0_EE,comdat
	.align 2
	.weak	_ZNSt6chrono8durationIdSt5ratioILl1ELl1EEEC2IlS1_ILl1ELl1000000000EEvEERKNS0_IT_T0_EE
	.type	_ZNSt6chrono8durationIdSt5ratioILl1ELl1EEEC2IlS1_ILl1ELl1000000000EEvEERKNS0_IT_T0_EE, @function
_ZNSt6chrono8durationIdSt5ratioILl1ELl1EEEC2IlS1_ILl1ELl1000000000EEvEERKNS0_IT_T0_EE:
.LFB5224:
	.cfi_startproc
	endbr64
	cvtsi2sdq	(%rsi), %xmm0
	divsd	.LC6(%rip), %xmm0
	movsd	%xmm0, (%rdi)
	ret
	.cfi_endproc
.LFE5224:
	.size	_ZNSt6chrono8durationIdSt5ratioILl1ELl1EEEC2IlS1_ILl1ELl1000000000EEvEERKNS0_IT_T0_EE, .-_ZNSt6chrono8durationIdSt5ratioILl1ELl1EEEC2IlS1_ILl1ELl1000000000EEvEERKNS0_IT_T0_EE
	.weak	_ZNSt6chrono8durationIdSt5ratioILl1ELl1EEEC1IlS1_ILl1ELl1000000000EEvEERKNS0_IT_T0_EE
	.set	_ZNSt6chrono8durationIdSt5ratioILl1ELl1EEEC1IlS1_ILl1ELl1000000000EEvEERKNS0_IT_T0_EE,_ZNSt6chrono8durationIdSt5ratioILl1ELl1EEEC2IlS1_ILl1ELl1000000000EEvEERKNS0_IT_T0_EE
	.section	.rodata.str1.1
.LC7:
	.string	"Warning: Total thread count is not evenly divisible by the block size\n"
.LC8:
	.string	"The total number of threads will be rounded up to %d\n"
.LC9:
	.string	"Host execution took: "
.LC10:
	.string	" seconds."
.LC11:
	.string	"Global Memory execution took: "
.LC12:
	.string	"Shared Memory execution took: "
.LC13:
	.string	"Constant Memory execution took: "
	.section	.text.startup
	.globl	main
	.type	main, @function
main:
.LFB4833:
	.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	$40, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 24(%rsp)
	xorl	%eax, %eax
	cmpl	$1, %edi
	jle	.L242
	movl	%edi, %r13d
	movq	8(%rsi), %rdi
	movq	%rsi, %r12
	movl	$256, %ebp
	call	atoi@PLT
	movl	%eax, %ebx
	cmpl	$2, %r13d
	je	.L239
	movq	16(%r12), %rdi
	call	atoi@PLT
	movl	%eax, %ebp
	jmp	.L239
.L242:
	movl	$256, %ebx
	movl	$256, %ebp
.L239:
	movl	%ebx, %eax
	cltd
	idivl	%ebp
	movl	%eax, %r12d
	testl	%edx, %edx
	je	.L240
	incl	%r12d
	movl	%ebp, %ebx
	leaq	.LC7(%rip), %rsi
	xorl	%eax, %eax
	imull	%r12d, %ebx
	movl	$2, %edi
	call	__printf_chk@PLT
	leaq	.LC8(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	%ebx, %edx
	call	__printf_chk@PLT
.L240:
	call	_ZNSt6chrono3_V212system_clock3nowEv@PLT
	movl	%r12d, %edx
	movl	%ebp, %esi
	movl	%ebx, %edi
	movq	%rax, %r13
	leaq	16(%rsp), %r14
	leaq	8(%rsp), %r15
	call	_Z15executeHostTestiii
	call	_ZNSt6chrono3_V212system_clock3nowEv@PLT
	movq	%rsp, %rdi
	movq	%r14, %rsi
	subq	%r13, %rax
	movq	%rax, 16(%rsp)
	call	_ZNSt6chrono8durationIdSt5ratioILl1ELl1EEEC1IlS1_ILl1ELl1000000000EEvEERKNS0_IT_T0_EE
	leaq	.LC9(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	leaq	.LC10(%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
	call	_ZNSt6chrono3_V212system_clock3nowEv@PLT
	movl	%r12d, %edx
	movl	%ebp, %esi
	movl	%ebx, %edi
	movq	%rax, %r13
	call	_Z17executeGlobalTestiii
	call	_ZNSt6chrono3_V212system_clock3nowEv@PLT
	movq	%r15, %rsi
	movq	%r14, %rdi
	subq	%r13, %rax
	movq	%rax, 8(%rsp)
	call	_ZNSt6chrono8durationIdSt5ratioILl1ELl1EEEC1IlS1_ILl1ELl1000000000EEvEERKNS0_IT_T0_EE
	movsd	16(%rsp), %xmm0
	leaq	.LC11(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	movsd	%xmm0, (%rsp)
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	leaq	.LC10(%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
	call	_ZNSt6chrono3_V212system_clock3nowEv@PLT
	movl	%r12d, %edx
	movl	%ebp, %esi
	movl	%ebx, %edi
	movq	%rax, %r13
	call	_Z17executeSharedTestiii
	call	_ZNSt6chrono3_V212system_clock3nowEv@PLT
	movq	%r15, %rsi
	movq	%r14, %rdi
	subq	%r13, %rax
	movq	%rax, 8(%rsp)
	call	_ZNSt6chrono8durationIdSt5ratioILl1ELl1EEEC1IlS1_ILl1ELl1000000000EEvEERKNS0_IT_T0_EE
	movsd	16(%rsp), %xmm0
	leaq	.LC12(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	movsd	%xmm0, (%rsp)
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	leaq	.LC10(%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
	call	_ZNSt6chrono3_V212system_clock3nowEv@PLT
	movl	%r12d, %edx
	movl	%ebp, %esi
	movl	%ebx, %edi
	movq	%rax, %r13
	call	_Z19executeConstantTestiii
	call	_ZNSt6chrono3_V212system_clock3nowEv@PLT
	movq	%r15, %rsi
	movq	%r14, %rdi
	subq	%r13, %rax
	movq	%rax, 8(%rsp)
	call	_ZNSt6chrono8durationIdSt5ratioILl1ELl1EEEC1IlS1_ILl1ELl1000000000EEvEERKNS0_IT_T0_EE
	movsd	16(%rsp), %xmm0
	leaq	.LC13(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	movsd	%xmm0, (%rsp)
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	leaq	.LC10(%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
	movq	24(%rsp), %rax
	subq	%fs:40, %rax
	je	.L241
	call	__stack_chk_fail@PLT
.L241:
	addq	$40, %rsp
	.cfi_def_cfa_offset 56
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4833:
	.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
	.align 4
	.type	_ZL5VAL_B, @object
	.size	_ZL5VAL_B, 4
_ZL5VAL_B:
	.long	3
	.align 4
	.type	_ZL5VAL_A, @object
	.size	_ZL5VAL_A, 4
_ZL5VAL_A:
	.long	1
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC6:
	.long	0
	.long	1104006501
	.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:

//Based on the work of Andrew Krepps

#include <hip/hip_runtime.h>
#include <iostream>
#include <random>
#include <chrono>
#include <stdio.h>

__constant__ static const int VAL_A = 1;
__constant__ static const int VAL_B = 3;

// Device GPU add c[i] = a[i] + b[i]
__device__ void add(int * a, int * b, int * c)
{
    const unsigned int thread_idx = (blockIdx.x * blockDim.x) + threadIdx.x;
    
    c[thread_idx] = a[thread_idx] + b[thread_idx];
}

// Device GPU subtract c[i] = a[i] - b[i]
__device__ void subtract(int * a, int * b, int * c)
{
    const unsigned int thread_idx = (blockIdx.x * blockDim.x) + threadIdx.x;
    
    c[thread_idx] = a[thread_idx] - b[thread_idx];
}

// Device GPU multiply c[i] = a[i] * b[i]
__device__ void mult(int * a, int * b, int * c)
{
    const unsigned int thread_idx = (blockIdx.x * blockDim.x) + threadIdx.x;
    
    c[thread_idx] = a[thread_idx] * b[thread_idx];
}

// Device GPU div c[i] = a[i] / b[i]
__device__ void div(int *a, int * b, int * c)
{
    const unsigned int thread_idx = (blockIdx.x * blockDim.x) + threadIdx.x;
    
    c[thread_idx] = a[thread_idx] / b[thread_idx];
}

// Device GPU mod c[i] = a[i] % b[i]
__device__ void mod(int * a, int * b, int * c)
{
    const unsigned int thread_idx = (blockIdx.x * blockDim.x) + threadIdx.x;
    
    c[thread_idx] = a[thread_idx] % b[thread_idx];
}

// Executes all 5 shared math operations
__global__ void executeSharedMathOperations(int * a, int * b, int * addDest, int * subDest, int * multDest, int * divDest, int * modDest, const int size)
{
	const int tid = (blockIdx.x * blockDim.x) + threadIdx.x;
    extern __shared__ int sharedMem[];
    
    // Use offsets in the shared mem to create arrays.
    int * sharedA = &sharedMem[0];
    int * sharedB = &sharedMem[size];
    int * sharedRet = &sharedMem[2*size];
    
    sharedA[tid] = a[tid];
    sharedB[tid] = b[tid];
    
    // Add sharedA to sharedB and store in addDest
    add(sharedA, sharedB, sharedRet);
    addDest[tid] = sharedRet[tid];
    
    // Subtract sharedB from sharedA and store in subDest
    subtract(sharedA, sharedB, sharedRet);
    subDest[tid] = sharedRet[tid];
    // Multiply sharedA to sharedB and store in mutlDest
    mult(sharedA, sharedB, sharedRet);
    multDest[tid] = sharedRet[tid];
    
    // Divide sharedA by sharedB and store in divDest
    div(sharedA, sharedB, sharedRet);
    divDest[tid] = sharedRet[tid];
    
    // Mod sharedA by sharedB and store in modDest
    mod(sharedA, sharedB, sharedRet);
    modDest[tid] = sharedRet[tid];
}

// Executes all 5 global math operations
__global__ void executeGlobalMathOperations(int * a, int * b, int * addDest, int * subDest, int * multDest, int * divDest, int * modDest, const int size)
{
    // Add a to b and store in addDest
    add(a, b, addDest);
    
    // Subtract a from b and store in subDest
    subtract(a, b, subDest);
    
    // Multiply a to b and store in mutlDest
    mult(a, b, multDest);
    
    // Divide a by b and store in divDest
    div(a, b, divDest);
    
    // Mod a by b and store in modDest
    mod(a, b, modDest);
}

// Executes all 5 constant math operations
__global__ void executeConstantMathOperations(int * addDest, int * subDest, int * multDest, int * divDest, int * modDest, const int size)
{    
	const int tid = (blockIdx.x * blockDim.x) + threadIdx.x;
    // Add VAL_A to VAL_B and store in addDest
    addDest[tid] = VAL_A + VAL_B;
    
    // Subtract a from b and store in subDest
    subDest[tid] = VAL_A - VAL_B;
    
    // Multiply a to b and store in mutlDest
    multDest[tid] = VAL_A * VAL_B;
    
    // Divide a by b and store in divDest
    divDest[tid] = VAL_A / VAL_B; // B is chosen to not be 0.
    
    // Mod a by b and store in modDest
    modDest[tid] = VAL_A / VAL_B; // B is chosen to not be 0.
}


// Host (Cpu) add c[i] = a[i] + b[i]
void hostAdd(int * a, int * b, int *c, const int size)
{
    for (int i = 0; i < size; ++i)
    {
        c[i] = a[i] + b[i];
    }
}

// Host (Cpu) sub c[i] = a[i] - b[i]
void hostSub(int * a, int * b, int *c, const int size)
{
    for (int i = 0; i < size; ++i)
    {
        c[i] = a[i] - b[i];
    }
}

// Host (Cpu) multiply c[i] = a[i] * b[i]
void hostMult(int * a, int * b, int *c, const int size)
{
    for (int i = 0; i < size; ++i)
    {
        c[i] = a[i] * b[i];
    }
}

// Host (Cpu) divide c[i] = a[i] / b[i]
void hostDiv(int * a, int * b, int *c, const int size)
{
    for (int i = 0; i < size; ++i)
    {
        if (b[i] != 0)
        {
            c[i] = a[i] / b[i];
        }
        else
        {
            c[i] = 0;
        }
    }
}

// Host (Cpu) mod c[i] = a[i] % b[i]
void hostMod(int * a, int * b, int *c, const int size)
{
    for (int i = 0; i < size; ++i)
    {
        // Protect against divide by 0. 
        // cuda code catches this error and sets result to 0 by default.
        if (b[i] == 0)
        {
            c[i] = 0;
        }
        else
        {
            c[i] = a[i] % b[i];
        }
    }
}

// Executes each of the host (cpu) tests by creating local memory and executing all 5 math operations on the data.
// The data is filled with random numbers that uses the same seed as the GPU tests.
void executeHostTest(const int totalThreads, const int blockSize, const int numBlocks)
{
    int a[totalThreads], b[totalThreads], c[totalThreads];
    
    // Create a random generate that will generate random numbers from 0 to 4.
    // Use a set seed so output is deterministic
    unsigned seed = 12345;
    std::default_random_engine gen(seed);
    std::uniform_int_distribution<int> dist(0,4);
    
    for (size_t i = 0; i < totalThreads; ++i)
    {
        a[i] = i;
        b[i] = dist(gen);
    }
    
    // Add all of the numbers c[i] = a[i] + b[i];
    hostAdd(a,b,c, totalThreads);
    
    // Subtract all of the numbers c[i] = a[i] - b[i];
    hostSub(a,b,c, totalThreads);
    
    // Multiply all of the numbers c[i] = a[i] * b[i];
    hostMult(a,b,c, totalThreads);
    
    // Divides all of the numbers c[i] = a[i] / b[i]; if b[i] == 0, c[i] = 0
    hostDiv(a,b,c, totalThreads);
    
    // Mod all of the numbers c[i] = a[i] % b[i];
    hostMod(a,b,c, totalThreads);
}

// Executes each of the global memory gpu tests by creating local memory, copying it global memory, and then performing
//  all 5 math operations on the data. The data is filled with random numbers that uses the same seed as the CPU tests.
void executeGlobalTest(const int totalThreads, const int blockSize, const int numBlocks)
{
    int a[totalThreads], b[totalThreads], add_dest[totalThreads], sub_dest[totalThreads], mult_dest[totalThreads], div_dest[totalThreads], mod_dest[totalThreads];
    
    int *gpu_a, *gpu_b, *gpu_add_dest, *gpu_sub_dest, *gpu_mult_dest, *gpu_div_dest, *gpu_mod_dest;

    hipMalloc((void**)&gpu_a,         totalThreads * sizeof(int));
    hipMalloc((void**)&gpu_b,         totalThreads * sizeof(int));
    hipMalloc((void**)&gpu_add_dest,  totalThreads * sizeof(int));
    hipMalloc((void**)&gpu_sub_dest,  totalThreads * sizeof(int));
    hipMalloc((void**)&gpu_mult_dest, totalThreads * sizeof(int));
    hipMalloc((void**)&gpu_div_dest,  totalThreads * sizeof(int));
    hipMalloc((void**)&gpu_mod_dest,  totalThreads * sizeof(int));
    
    // Create a random generate that will generate random numbers from 0 to 4.
    // Use a set seed so output is deterministic
    unsigned seed = 12345;
    std::default_random_engine gen(seed);
    std::uniform_int_distribution<int> dist(0,4);
    
    for (size_t i = 0; i < totalThreads; ++i)
    {
        a[i] = i;
        b[i] = dist(gen);
    }
    
    hipMemcpy(gpu_a, a, totalThreads * sizeof(int), hipMemcpyHostToDevice);
    hipMemcpy(gpu_b, b, totalThreads * sizeof(int), hipMemcpyHostToDevice);
    
    executeGlobalMathOperations<<<numBlocks, blockSize>>>(gpu_a,gpu_b,gpu_add_dest,gpu_sub_dest,gpu_mult_dest,gpu_div_dest,gpu_mod_dest, numBlocks * blockSize);
    
    hipMemcpy(add_dest,  gpu_add_dest,  totalThreads*sizeof(int), hipMemcpyDeviceToHost);
    hipMemcpy(sub_dest,  gpu_sub_dest,  totalThreads*sizeof(int), hipMemcpyDeviceToHost);    
    hipMemcpy(mult_dest, gpu_mult_dest, totalThreads*sizeof(int), hipMemcpyDeviceToHost);
    hipMemcpy(div_dest,  gpu_div_dest,  totalThreads*sizeof(int), hipMemcpyDeviceToHost);
    hipMemcpy(mod_dest,  gpu_mod_dest,  totalThreads*sizeof(int), hipMemcpyDeviceToHost);
    
    hipFree(gpu_a);
    hipFree(gpu_b);
    hipFree(gpu_add_dest);
    hipFree(gpu_sub_dest);
    hipFree(gpu_mult_dest);
    hipFree(gpu_div_dest);
    hipFree(gpu_mod_dest);
}

// Executes each of the shared memory gpu tests by creating local memory, copying it global memory, and then performing
// all 5 math operations on the data after creating shared memory. The data is filled with random numbers that uses the same seed as the CPU tests.
void executeSharedTest(const int totalThreads, const int blockSize, const int numBlocks)
{
    int a[totalThreads], b[totalThreads], add_dest[totalThreads], sub_dest[totalThreads], mult_dest[totalThreads], div_dest[totalThreads], mod_dest[totalThreads];
    
    int *gpu_a, *gpu_b, *gpu_add_dest, *gpu_sub_dest, *gpu_mult_dest, *gpu_div_dest, *gpu_mod_dest;

    hipMalloc((void**)&gpu_a,         totalThreads * sizeof(int));
    hipMalloc((void**)&gpu_b,         totalThreads * sizeof(int));
    hipMalloc((void**)&gpu_add_dest,  totalThreads * sizeof(int));
    hipMalloc((void**)&gpu_sub_dest,  totalThreads * sizeof(int));
    hipMalloc((void**)&gpu_mult_dest, totalThreads * sizeof(int));
    hipMalloc((void**)&gpu_div_dest,  totalThreads * sizeof(int));
    hipMalloc((void**)&gpu_mod_dest,  totalThreads * sizeof(int));
    
    // Create a random generate that will generate random numbers from 0 to 4.
    // Use a set seed so output is deterministic
    unsigned seed = 12345;
    std::default_random_engine gen(seed);
    std::uniform_int_distribution<int> dist(0,4);
    
    for (size_t i = 0; i < totalThreads; ++i)
    {
        a[i] = i;
        b[i] = dist(gen);
    }
    
    hipMemcpy(gpu_a, a, totalThreads * sizeof(int), hipMemcpyHostToDevice);
    hipMemcpy(gpu_b, b, totalThreads * sizeof(int), hipMemcpyHostToDevice);
    
    // The third parameter is the size of the shared memory
    // We multiply by 3 because we need to copy A and B and then also have room for the return in shared memory.
    executeSharedMathOperations<<<numBlocks, blockSize, 3 * totalThreads * sizeof(int)>>>(gpu_a,gpu_b,gpu_add_dest,gpu_sub_dest,gpu_mult_dest,gpu_div_dest,gpu_mod_dest, totalThreads);
    
    hipMemcpy(add_dest,  gpu_add_dest,  totalThreads*sizeof(int), hipMemcpyDeviceToHost);
    hipMemcpy(sub_dest,  gpu_sub_dest,  totalThreads*sizeof(int), hipMemcpyDeviceToHost);    
    hipMemcpy(mult_dest, gpu_mult_dest, totalThreads*sizeof(int), hipMemcpyDeviceToHost);
    hipMemcpy(div_dest,  gpu_div_dest,  totalThreads*sizeof(int), hipMemcpyDeviceToHost);
    hipMemcpy(mod_dest,  gpu_mod_dest,  totalThreads*sizeof(int), hipMemcpyDeviceToHost);
    
    hipFree(gpu_a);
    hipFree(gpu_b);
    hipFree(gpu_add_dest);
    hipFree(gpu_sub_dest);
    hipFree(gpu_mult_dest);
    hipFree(gpu_div_dest);
    hipFree(gpu_mod_dest); 
}

// Executes each of the consnt memory gpu tests by creating local memory, copying it global memory, and then performing
// all 5 math operations on the data using constant values. The data is filled with random numbers that uses the same seed as the CPU tests.
void executeConstantTest(const int totalThreads, const int blockSize, const int numBlocks)
{
    int add_dest[totalThreads], sub_dest[totalThreads], mult_dest[totalThreads], div_dest[totalThreads], mod_dest[totalThreads];
    
    int *gpu_add_dest, *gpu_sub_dest, *gpu_mult_dest, *gpu_div_dest, *gpu_mod_dest;

    hipMalloc((void**)&gpu_add_dest,  totalThreads * sizeof(int));
    hipMalloc((void**)&gpu_sub_dest,  totalThreads * sizeof(int));
    hipMalloc((void**)&gpu_mult_dest, totalThreads * sizeof(int));
    hipMalloc((void**)&gpu_div_dest,  totalThreads * sizeof(int));
    hipMalloc((void**)&gpu_mod_dest,  totalThreads * sizeof(int));
    
    // The third parameter is the size of the shared memory
    // We multiply by 3 because we need to copy A and B and then also have room for the return in shared memory.
    executeConstantMathOperations<<<numBlocks, blockSize>>>(gpu_add_dest,gpu_sub_dest,gpu_mult_dest,gpu_div_dest,gpu_mod_dest, totalThreads);
    
    hipMemcpy(add_dest,  gpu_add_dest,  totalThreads*sizeof(int), hipMemcpyDeviceToHost);
    hipMemcpy(sub_dest,  gpu_sub_dest,  totalThreads*sizeof(int), hipMemcpyDeviceToHost);    
    hipMemcpy(mult_dest, gpu_mult_dest, totalThreads*sizeof(int), hipMemcpyDeviceToHost);
    hipMemcpy(div_dest,  gpu_div_dest,  totalThreads*sizeof(int), hipMemcpyDeviceToHost);
    hipMemcpy(mod_dest,  gpu_mod_dest,  totalThreads*sizeof(int), hipMemcpyDeviceToHost);
    
    hipFree(gpu_add_dest);
    hipFree(gpu_sub_dest);
    hipFree(gpu_mult_dest);
    hipFree(gpu_div_dest);
    hipFree(gpu_mod_dest); 
}

void printArray(const int * const arr, const int xSize, const int ySize)
{
    for (size_t i = 0; i < xSize; ++i)
    {
        for(size_t j = 0; j < ySize; ++j)
        {
            std::cout << arr[i * ySize + j] << " ";
        }
        std::cout << '\n';
    }
    
    std::cout << std::flush;
}

int main(int argc, char** argv)
{
    // read command line arguments
    int totalThreads = 256;
    int blockSize = 256;
    
    if (argc >= 2) {
        totalThreads = atoi(argv[1]);
    }
    if (argc >= 3) {
        blockSize = atoi(argv[2]);
    }
   
    int numBlocks = totalThreads/blockSize;

    // validate command line arguments
    if (totalThreads % blockSize != 0) {
        ++numBlocks;
        totalThreads = numBlocks*blockSize;
        
        printf("Warning: Total thread count is not evenly divisible by the block size\n");
        printf("The total number of threads will be rounded up to %d\n", totalThreads);
    }
    
    auto startTime = std::chrono::system_clock::now();
    executeHostTest(totalThreads, blockSize, numBlocks);
    auto endTime = std::chrono::system_clock::now();
    std::chrono::duration<double> totalTime = endTime-startTime;
    std::cout << "Host execution took: " << totalTime.count() << " seconds." << std::endl;
    
    startTime = std::chrono::system_clock::now();
    executeGlobalTest(totalThreads, blockSize, numBlocks);
    endTime = std::chrono::system_clock::now();
    totalTime = endTime-startTime;
    std::cout << "Global Memory execution took: " << totalTime.count() << " seconds." << std::endl;
    
    startTime = std::chrono::system_clock::now();
    executeSharedTest(totalThreads, blockSize, numBlocks);
    endTime = std::chrono::system_clock::now();
    totalTime = endTime-startTime;
    std::cout << "Shared Memory execution took: " << totalTime.count() << " seconds." << std::endl;
    
    startTime = std::chrono::system_clock::now();
    executeConstantTest(totalThreads, blockSize, numBlocks);
    endTime = std::chrono::system_clock::now();
    totalTime = endTime-startTime;
    std::cout << "Constant Memory execution took: " << totalTime.count() << " seconds." << std::endl;
    
    return 0;
}

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

                                        # End of file scope inline assembly
	.globl	_Z42__device_stub__executeSharedMathOperationsPiS_S_S_S_S_S_i # -- Begin function _Z42__device_stub__executeSharedMathOperationsPiS_S_S_S_S_S_i
	.type	_Z42__device_stub__executeSharedMathOperationsPiS_S_S_S_S_S_i,@function
_Z42__device_stub__executeSharedMathOperationsPiS_S_S_S_S_S_i: # @_Z42__device_stub__executeSharedMathOperationsPiS_S_S_S_S_S_i
	.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	$160, %rsp
	.cfi_def_cfa_offset 208
	.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	16(%rsp), %r8
	movq	%r9, (%r8)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	208(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	216(%rsp), %rax
	movq	%rax, 56(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%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	$_Z27executeSharedMathOperationsPiS_S_S_S_S_S_i, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$176, %rsp
	.cfi_adjust_cfa_offset -176
	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	_Z42__device_stub__executeSharedMathOperationsPiS_S_S_S_S_S_i, .Lfunc_end0-_Z42__device_stub__executeSharedMathOperationsPiS_S_S_S_S_S_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z42__device_stub__executeGlobalMathOperationsPiS_S_S_S_S_S_i # -- Begin function _Z42__device_stub__executeGlobalMathOperationsPiS_S_S_S_S_S_i
	.type	_Z42__device_stub__executeGlobalMathOperationsPiS_S_S_S_S_S_i,@function
_Z42__device_stub__executeGlobalMathOperationsPiS_S_S_S_S_S_i: # @_Z42__device_stub__executeGlobalMathOperationsPiS_S_S_S_S_S_i
	.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	$160, %rsp
	.cfi_def_cfa_offset 208
	.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	16(%rsp), %r8
	movq	%r9, (%r8)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	208(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	216(%rsp), %rax
	movq	%rax, 56(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%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	$_Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$176, %rsp
	.cfi_adjust_cfa_offset -176
	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_end1:
	.size	_Z42__device_stub__executeGlobalMathOperationsPiS_S_S_S_S_S_i, .Lfunc_end1-_Z42__device_stub__executeGlobalMathOperationsPiS_S_S_S_S_S_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z44__device_stub__executeConstantMathOperationsPiS_S_S_S_i # -- Begin function _Z44__device_stub__executeConstantMathOperationsPiS_S_S_S_i
	.type	_Z44__device_stub__executeConstantMathOperationsPiS_S_S_S_i,@function
_Z44__device_stub__executeConstantMathOperationsPiS_S_S_S_i: # @_Z44__device_stub__executeConstantMathOperationsPiS_S_S_S_i
	.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	4(%rsp), %r8
	movl	%r9d, (%r8)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%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	$_Z29executeConstantMathOperationsPiS_S_S_S_i, %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_end2:
	.size	_Z44__device_stub__executeConstantMathOperationsPiS_S_S_S_i, .Lfunc_end2-_Z44__device_stub__executeConstantMathOperationsPiS_S_S_S_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z7hostAddPiS_S_i               # -- Begin function _Z7hostAddPiS_S_i
	.type	_Z7hostAddPiS_S_i,@function
_Z7hostAddPiS_S_i:                      # @_Z7hostAddPiS_S_i
	.cfi_startproc
# %bb.0:
	testl	%ecx, %ecx
	jle	.LBB3_3
# %bb.1:                                # %.lr.ph.preheader
	movl	%ecx, %eax
	xorl	%ecx, %ecx
.LBB3_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	(%rsi,%rcx,4), %r8d
	addl	(%rdi,%rcx,4), %r8d
	movl	%r8d, (%rdx,%rcx,4)
	incq	%rcx
	cmpq	%rcx, %rax
	jne	.LBB3_2
.LBB3_3:                                # %._crit_edge
	retq
.Lfunc_end3:
	.size	_Z7hostAddPiS_S_i, .Lfunc_end3-_Z7hostAddPiS_S_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z7hostSubPiS_S_i               # -- Begin function _Z7hostSubPiS_S_i
	.type	_Z7hostSubPiS_S_i,@function
_Z7hostSubPiS_S_i:                      # @_Z7hostSubPiS_S_i
	.cfi_startproc
# %bb.0:
	testl	%ecx, %ecx
	jle	.LBB4_3
# %bb.1:                                # %.lr.ph.preheader
	movl	%ecx, %eax
	xorl	%ecx, %ecx
.LBB4_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	(%rdi,%rcx,4), %r8d
	subl	(%rsi,%rcx,4), %r8d
	movl	%r8d, (%rdx,%rcx,4)
	incq	%rcx
	cmpq	%rcx, %rax
	jne	.LBB4_2
.LBB4_3:                                # %._crit_edge
	retq
.Lfunc_end4:
	.size	_Z7hostSubPiS_S_i, .Lfunc_end4-_Z7hostSubPiS_S_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z8hostMultPiS_S_i              # -- Begin function _Z8hostMultPiS_S_i
	.type	_Z8hostMultPiS_S_i,@function
_Z8hostMultPiS_S_i:                     # @_Z8hostMultPiS_S_i
	.cfi_startproc
# %bb.0:
	testl	%ecx, %ecx
	jle	.LBB5_3
# %bb.1:                                # %.lr.ph.preheader
	movl	%ecx, %eax
	xorl	%ecx, %ecx
.LBB5_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	(%rsi,%rcx,4), %r8d
	imull	(%rdi,%rcx,4), %r8d
	movl	%r8d, (%rdx,%rcx,4)
	incq	%rcx
	cmpq	%rcx, %rax
	jne	.LBB5_2
.LBB5_3:                                # %._crit_edge
	retq
.Lfunc_end5:
	.size	_Z8hostMultPiS_S_i, .Lfunc_end5-_Z8hostMultPiS_S_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z7hostDivPiS_S_i               # -- Begin function _Z7hostDivPiS_S_i
	.type	_Z7hostDivPiS_S_i,@function
_Z7hostDivPiS_S_i:                      # @_Z7hostDivPiS_S_i
	.cfi_startproc
# %bb.0:
	testl	%ecx, %ecx
	jle	.LBB6_5
# %bb.1:                                # %.lr.ph.preheader
	movq	%rdx, %r8
	movl	%ecx, %ecx
	xorl	%r9d, %r9d
.LBB6_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	(%rsi,%r9,4), %r10d
	movl	$0, %eax
	testl	%r10d, %r10d
	je	.LBB6_4
# %bb.3:                                #   in Loop: Header=BB6_2 Depth=1
	movl	(%rdi,%r9,4), %eax
	cltd
	idivl	%r10d
.LBB6_4:                                #   in Loop: Header=BB6_2 Depth=1
	movl	%eax, (%r8,%r9,4)
	incq	%r9
	cmpq	%r9, %rcx
	jne	.LBB6_2
.LBB6_5:                                # %._crit_edge
	retq
.Lfunc_end6:
	.size	_Z7hostDivPiS_S_i, .Lfunc_end6-_Z7hostDivPiS_S_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z7hostModPiS_S_i               # -- Begin function _Z7hostModPiS_S_i
	.type	_Z7hostModPiS_S_i,@function
_Z7hostModPiS_S_i:                      # @_Z7hostModPiS_S_i
	.cfi_startproc
# %bb.0:
	testl	%ecx, %ecx
	jle	.LBB7_5
# %bb.1:                                # %.lr.ph.preheader
	movq	%rdx, %r8
	movl	%ecx, %ecx
	xorl	%r9d, %r9d
.LBB7_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	(%rsi,%r9,4), %r10d
	movl	$0, %edx
	testl	%r10d, %r10d
	je	.LBB7_4
# %bb.3:                                #   in Loop: Header=BB7_2 Depth=1
	movl	(%rdi,%r9,4), %eax
	cltd
	idivl	%r10d
.LBB7_4:                                #   in Loop: Header=BB7_2 Depth=1
	movl	%edx, (%r8,%r9,4)
	incq	%r9
	cmpq	%r9, %rcx
	jne	.LBB7_2
.LBB7_5:                                # %._crit_edge
	retq
.Lfunc_end7:
	.size	_Z7hostModPiS_S_i, .Lfunc_end7-_Z7hostModPiS_S_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z15executeHostTestiii          # -- Begin function _Z15executeHostTestiii
	.type	_Z15executeHostTestiii,@function
_Z15executeHostTestiii:                 # @_Z15executeHostTestiii
	.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
	subq	$16, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	$12345, (%rsp)                  # imm = 0x3039
	movabsq	$17179869184, %rax              # imm = 0x400000000
	movq	%rax, 8(%rsp)
	testl	%edi, %edi
	je	.LBB8_3
# %bb.1:
	movslq	%edi, %r15
	movq	%rsp, %rbx
	leaq	8(%rsp), %r14
.LBB8_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movq	%r14, %rdi
	movq	%rbx, %rsi
	movq	%r14, %rdx
	callq	_ZNSt24uniform_int_distributionIiEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEiRT_RKNS0_10param_typeE
	decq	%r15
	jne	.LBB8_2
.LBB8_3:                                # %_Z7hostModPiS_S_i.exit
	addq	$16, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end8:
	.size	_Z15executeHostTestiii, .Lfunc_end8-_Z15executeHostTestiii
	.cfi_endproc
                                        # -- End function
	.globl	_Z17executeGlobalTestiii        # -- Begin function _Z17executeGlobalTestiii
	.type	_Z17executeGlobalTestiii,@function
_Z17executeGlobalTestiii:               # @_Z17executeGlobalTestiii
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset %rbp, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register %rbp
	pushq	%r15
	pushq	%r14
	pushq	%r13
	pushq	%r12
	pushq	%rbx
	subq	$152, %rsp
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
                                        # kill: def $edx killed $edx def $rdx
	movq	%rdx, -168(%rbp)                # 8-byte Spill
	movl	%esi, -100(%rbp)                # 4-byte Spill
	movl	%edi, %eax
	movq	%rsp, %rcx
	leaq	15(,%rax,4), %rax
	andq	$-16, %rax
	subq	%rax, %rcx
	movq	%rcx, -112(%rbp)                # 8-byte Spill
	movq	%rcx, %rsp
	movq	%rsp, %r12
	subq	%rax, %r12
	movq	%r12, %rsp
	movq	%rsp, %rcx
	subq	%rax, %rcx
	movq	%rcx, -120(%rbp)                # 8-byte Spill
	movq	%rcx, %rsp
	movq	%rsp, %rcx
	subq	%rax, %rcx
	movq	%rcx, -128(%rbp)                # 8-byte Spill
	movq	%rcx, %rsp
	movq	%rsp, %rcx
	subq	%rax, %rcx
	movq	%rcx, -136(%rbp)                # 8-byte Spill
	movq	%rcx, %rsp
	movq	%rsp, %rcx
	subq	%rax, %rcx
	movq	%rcx, -152(%rbp)                # 8-byte Spill
	movq	%rcx, %rsp
	movq	%rsp, %rcx
	subq	%rax, %rcx
	movq	%rcx, -160(%rbp)                # 8-byte Spill
	movq	%rcx, %rsp
	movslq	%edi, %r13
	leaq	(,%r13,4), %rbx
	leaq	-96(%rbp), %rdi
	movq	%rbx, %rsi
	callq	hipMalloc
	leaq	-88(%rbp), %rdi
	movq	%rbx, %rsi
	callq	hipMalloc
	leaq	-80(%rbp), %rdi
	movq	%rbx, %rsi
	callq	hipMalloc
	leaq	-72(%rbp), %rdi
	movq	%rbx, %rsi
	callq	hipMalloc
	leaq	-64(%rbp), %rdi
	movq	%rbx, %rsi
	callq	hipMalloc
	leaq	-56(%rbp), %rdi
	movq	%rbx, %rsi
	callq	hipMalloc
	leaq	-48(%rbp), %rdi
	movq	%rbx, -144(%rbp)                # 8-byte Spill
	movq	%rbx, %rsi
	callq	hipMalloc
	movq	$12345, -176(%rbp)              # imm = 0x3039
	movabsq	$17179869184, %rax              # imm = 0x400000000
	movq	%rax, -184(%rbp)
	testl	%r13d, %r13d
	je	.LBB9_3
# %bb.1:                                # %.lr.ph.preheader
	xorl	%r15d, %r15d
	leaq	-176(%rbp), %rbx
	leaq	-184(%rbp), %r14
.LBB9_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movq	-112(%rbp), %rax                # 8-byte Reload
	movl	%r15d, (%rax,%r15,4)
	movq	%r14, %rdi
	movq	%rbx, %rsi
	movq	%r14, %rdx
	callq	_ZNSt24uniform_int_distributionIiEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEiRT_RKNS0_10param_typeE
	movl	%eax, (%r12,%r15,4)
	incq	%r15
	cmpq	%r15, %r13
	jne	.LBB9_2
.LBB9_3:                                # %._crit_edge
	movq	-96(%rbp), %rdi
	movq	-112(%rbp), %rsi                # 8-byte Reload
	movq	-144(%rbp), %r15                # 8-byte Reload
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	-88(%rbp), %rdi
	movq	%r12, %rsi
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	-168(%rbp), %rbx                # 8-byte Reload
	movl	%ebx, %edi
	btsq	$32, %rdi
	movl	-100(%rbp), %r14d               # 4-byte Reload
	movl	%r14d, %edx
	btsq	$32, %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB9_5
# %bb.4:
	movq	-96(%rbp), %rdi
	movq	-88(%rbp), %rsi
	movq	-80(%rbp), %rdx
	movq	-72(%rbp), %rcx
	movq	-64(%rbp), %r8
	movq	-56(%rbp), %r9
	imull	%r14d, %ebx
	pushq	%rbx
	pushq	-48(%rbp)
	callq	_Z42__device_stub__executeGlobalMathOperationsPiS_S_S_S_S_S_i
	addq	$16, %rsp
.LBB9_5:
	movq	-80(%rbp), %rsi
	movq	-120(%rbp), %rdi                # 8-byte Reload
	movq	%r15, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	-72(%rbp), %rsi
	movq	-128(%rbp), %rdi                # 8-byte Reload
	movq	%r15, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	-64(%rbp), %rsi
	movq	-136(%rbp), %rdi                # 8-byte Reload
	movq	%r15, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	-56(%rbp), %rsi
	movq	-152(%rbp), %rdi                # 8-byte Reload
	movq	%r15, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	-48(%rbp), %rsi
	movq	-160(%rbp), %rdi                # 8-byte Reload
	movq	%r15, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	-96(%rbp), %rdi
	callq	hipFree
	movq	-88(%rbp), %rdi
	callq	hipFree
	movq	-80(%rbp), %rdi
	callq	hipFree
	movq	-72(%rbp), %rdi
	callq	hipFree
	movq	-64(%rbp), %rdi
	callq	hipFree
	movq	-56(%rbp), %rdi
	callq	hipFree
	movq	-48(%rbp), %rdi
	callq	hipFree
	leaq	-40(%rbp), %rsp
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa %rsp, 8
	retq
.Lfunc_end9:
	.size	_Z17executeGlobalTestiii, .Lfunc_end9-_Z17executeGlobalTestiii
	.cfi_endproc
                                        # -- End function
	.globl	_Z17executeSharedTestiii        # -- Begin function _Z17executeSharedTestiii
	.type	_Z17executeSharedTestiii,@function
_Z17executeSharedTestiii:               # @_Z17executeSharedTestiii
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset %rbp, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register %rbp
	pushq	%r15
	pushq	%r14
	pushq	%r13
	pushq	%r12
	pushq	%rbx
	subq	$152, %rsp
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	movl	%edx, -100(%rbp)                # 4-byte Spill
	movl	%esi, -104(%rbp)                # 4-byte Spill
                                        # kill: def $edi killed $edi def $rdi
	movl	%edi, %eax
	movq	%rsp, %rcx
	leaq	15(,%rax,4), %rax
	andq	$-16, %rax
	subq	%rax, %rcx
	movq	%rcx, -112(%rbp)                # 8-byte Spill
	movq	%rcx, %rsp
	movq	%rsp, %r13
	subq	%rax, %r13
	movq	%r13, %rsp
	movq	%rsp, %rcx
	subq	%rax, %rcx
	movq	%rcx, -120(%rbp)                # 8-byte Spill
	movq	%rcx, %rsp
	movq	%rsp, %rcx
	subq	%rax, %rcx
	movq	%rcx, -128(%rbp)                # 8-byte Spill
	movq	%rcx, %rsp
	movq	%rsp, %rcx
	subq	%rax, %rcx
	movq	%rcx, -136(%rbp)                # 8-byte Spill
	movq	%rcx, %rsp
	movq	%rsp, %rcx
	subq	%rax, %rcx
	movq	%rcx, -152(%rbp)                # 8-byte Spill
	movq	%rcx, %rsp
	movq	%rsp, %rcx
	subq	%rax, %rcx
	movq	%rcx, -160(%rbp)                # 8-byte Spill
	movq	%rcx, %rsp
	movq	%rdi, -168(%rbp)                # 8-byte Spill
	movslq	%edi, %r12
	leaq	(,%r12,4), %rbx
	leaq	-96(%rbp), %rdi
	movq	%rbx, %rsi
	callq	hipMalloc
	leaq	-88(%rbp), %rdi
	movq	%rbx, %rsi
	callq	hipMalloc
	leaq	-80(%rbp), %rdi
	movq	%rbx, %rsi
	callq	hipMalloc
	leaq	-72(%rbp), %rdi
	movq	%rbx, %rsi
	callq	hipMalloc
	leaq	-64(%rbp), %rdi
	movq	%rbx, %rsi
	callq	hipMalloc
	leaq	-56(%rbp), %rdi
	movq	%rbx, %rsi
	callq	hipMalloc
	leaq	-48(%rbp), %rdi
	movq	%rbx, -144(%rbp)                # 8-byte Spill
	movq	%rbx, %rsi
	callq	hipMalloc
	movq	$12345, -176(%rbp)              # imm = 0x3039
	movabsq	$17179869184, %rax              # imm = 0x400000000
	movq	%rax, -184(%rbp)
	testl	%r12d, %r12d
	je	.LBB10_3
# %bb.1:                                # %.lr.ph.preheader
	xorl	%r15d, %r15d
	leaq	-176(%rbp), %rbx
	leaq	-184(%rbp), %r14
.LBB10_2:                               # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movq	-112(%rbp), %rax                # 8-byte Reload
	movl	%r15d, (%rax,%r15,4)
	movq	%r14, %rdi
	movq	%rbx, %rsi
	movq	%r14, %rdx
	callq	_ZNSt24uniform_int_distributionIiEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEiRT_RKNS0_10param_typeE
	movl	%eax, (%r13,%r15,4)
	incq	%r15
	cmpq	%r15, %r12
	jne	.LBB10_2
.LBB10_3:                               # %._crit_edge
	movq	-96(%rbp), %rdi
	movq	-112(%rbp), %rsi                # 8-byte Reload
	movq	-144(%rbp), %r14                # 8-byte Reload
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	-88(%rbp), %rdi
	movq	%r13, %rsi
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	-168(%rbp), %rbx                # 8-byte Reload
	leal	(%rbx,%rbx,2), %eax
	movslq	%eax, %r8
	shlq	$2, %r8
	movl	-100(%rbp), %edi                # 4-byte Reload
	btsq	$32, %rdi
	movl	-104(%rbp), %edx                # 4-byte Reload
	btsq	$32, %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB10_5
# %bb.4:
	movq	-96(%rbp), %rdi
	movq	-88(%rbp), %rsi
	movq	-80(%rbp), %rdx
	movq	-72(%rbp), %rcx
	movq	-64(%rbp), %r8
	movq	-56(%rbp), %r9
	pushq	%rbx
	pushq	-48(%rbp)
	callq	_Z42__device_stub__executeSharedMathOperationsPiS_S_S_S_S_S_i
	addq	$16, %rsp
.LBB10_5:
	movq	-80(%rbp), %rsi
	movq	-120(%rbp), %rdi                # 8-byte Reload
	movq	%r14, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	-72(%rbp), %rsi
	movq	-128(%rbp), %rdi                # 8-byte Reload
	movq	%r14, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	-64(%rbp), %rsi
	movq	-136(%rbp), %rdi                # 8-byte Reload
	movq	%r14, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	-56(%rbp), %rsi
	movq	-152(%rbp), %rdi                # 8-byte Reload
	movq	%r14, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	-48(%rbp), %rsi
	movq	-160(%rbp), %rdi                # 8-byte Reload
	movq	%r14, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	-96(%rbp), %rdi
	callq	hipFree
	movq	-88(%rbp), %rdi
	callq	hipFree
	movq	-80(%rbp), %rdi
	callq	hipFree
	movq	-72(%rbp), %rdi
	callq	hipFree
	movq	-64(%rbp), %rdi
	callq	hipFree
	movq	-56(%rbp), %rdi
	callq	hipFree
	movq	-48(%rbp), %rdi
	callq	hipFree
	leaq	-40(%rbp), %rsp
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa %rsp, 8
	retq
.Lfunc_end10:
	.size	_Z17executeSharedTestiii, .Lfunc_end10-_Z17executeSharedTestiii
	.cfi_endproc
                                        # -- End function
	.globl	_Z19executeConstantTestiii      # -- Begin function _Z19executeConstantTestiii
	.type	_Z19executeConstantTestiii,@function
_Z19executeConstantTestiii:             # @_Z19executeConstantTestiii
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset %rbp, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register %rbp
	pushq	%r15
	pushq	%r14
	pushq	%r13
	pushq	%r12
	pushq	%rbx
	subq	$72, %rsp
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	movl	%edx, -88(%rbp)                 # 4-byte Spill
	movl	%esi, -92(%rbp)                 # 4-byte Spill
	movl	%edi, %eax
	movq	%rsp, %r13
	leaq	15(,%rax,4), %rax
	andq	$-16, %rax
	subq	%rax, %r13
	movq	%r13, %rsp
	movq	%rsp, %rbx
	subq	%rax, %rbx
	movq	%rbx, %rsp
	movq	%rsp, %r12
	subq	%rax, %r12
	movq	%r12, %rsp
	movq	%rsp, %r14
	subq	%rax, %r14
	movq	%r14, %rsp
	movq	%rsp, %rcx
	subq	%rax, %rcx
	movq	%rcx, -104(%rbp)                # 8-byte Spill
	movq	%rcx, %rsp
	movl	%edi, -84(%rbp)                 # 4-byte Spill
	movslq	%edi, %r15
	shlq	$2, %r15
	leaq	-80(%rbp), %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	leaq	-72(%rbp), %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	leaq	-64(%rbp), %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	leaq	-56(%rbp), %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	leaq	-48(%rbp), %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	movl	-88(%rbp), %edi                 # 4-byte Reload
	btsq	$32, %rdi
	movl	-92(%rbp), %edx                 # 4-byte Reload
	btsq	$32, %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB11_2
# %bb.1:
	movq	-80(%rbp), %rdi
	movq	-72(%rbp), %rsi
	movq	-64(%rbp), %rdx
	movq	-56(%rbp), %rcx
	movq	-48(%rbp), %r8
	movl	-84(%rbp), %r9d                 # 4-byte Reload
	callq	_Z44__device_stub__executeConstantMathOperationsPiS_S_S_S_i
.LBB11_2:
	movq	-80(%rbp), %rsi
	movq	%r13, %rdi
	movq	%r15, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	-72(%rbp), %rsi
	movq	%rbx, %rdi
	movq	%r15, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	-64(%rbp), %rsi
	movq	%r12, %rdi
	movq	%r15, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	-56(%rbp), %rsi
	movq	%r14, %rdi
	movq	%r15, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	-48(%rbp), %rsi
	movq	-104(%rbp), %rdi                # 8-byte Reload
	movq	%r15, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	-80(%rbp), %rdi
	callq	hipFree
	movq	-72(%rbp), %rdi
	callq	hipFree
	movq	-64(%rbp), %rdi
	callq	hipFree
	movq	-56(%rbp), %rdi
	callq	hipFree
	movq	-48(%rbp), %rdi
	callq	hipFree
	leaq	-40(%rbp), %rsp
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa %rsp, 8
	retq
.Lfunc_end11:
	.size	_Z19executeConstantTestiii, .Lfunc_end11-_Z19executeConstantTestiii
	.cfi_endproc
                                        # -- End function
	.globl	_Z10printArrayPKiii             # -- Begin function _Z10printArrayPKiii
	.type	_Z10printArrayPKiii,@function
_Z10printArrayPKiii:                    # @_Z10printArrayPKiii
	.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
	movl	%edx, 4(%rsp)                   # 4-byte Spill
	testl	%esi, %esi
	je	.LBB12_9
# %bb.1:                                # %.preheader.lr.ph
	movq	%rdi, %r14
	movslq	%esi, %r12
	movslq	4(%rsp), %r13                   # 4-byte Folded Reload
	leaq	(,%r13,4), %rbp
	xorl	%r15d, %r15d
.LBB12_2:                               # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB12_4 Depth 2
	cmpl	$0, 4(%rsp)                     # 4-byte Folded Reload
	je	.LBB12_5
# %bb.3:                                # %.lr.ph
                                        #   in Loop: Header=BB12_2 Depth=1
	xorl	%ebx, %ebx
.LBB12_4:                               #   Parent Loop BB12_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	(%r14,%rbx,4), %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSolsEi
	movl	$.L.str, %esi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	incq	%rbx
	cmpq	%rbx, %r13
	jne	.LBB12_4
.LBB12_5:                               # %._crit_edge
                                        #   in Loop: Header=BB12_2 Depth=1
	movb	$10, 3(%rsp)
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rax
	cmpq	$0, _ZSt4cout+16(%rax)
	je	.LBB12_7
# %bb.6:                                #   in Loop: Header=BB12_2 Depth=1
	movl	$_ZSt4cout, %edi
	movl	$1, %edx
	leaq	3(%rsp), %rsi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	jmp	.LBB12_8
.LBB12_7:                               #   in Loop: Header=BB12_2 Depth=1
	movl	$_ZSt4cout, %edi
	movl	$10, %esi
	callq	_ZNSo3putEc
.LBB12_8:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.exit
                                        #   in Loop: Header=BB12_2 Depth=1
	incq	%r15
	addq	%rbp, %r14
	cmpq	%r12, %r15
	jne	.LBB12_2
.LBB12_9:                               # %._crit_edge13
	movl	$_ZSt4cout, %edi
	callq	_ZNSo5flushEv
	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
	retq
.Lfunc_end12:
	.size	_Z10printArrayPKiii, .Lfunc_end12-_Z10printArrayPKiii
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI13_0:
	.quad	0x41cdcd6500000000              # double 1.0E+9
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.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	%rbx
	.cfi_def_cfa_offset 40
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	$256, %ebx                      # imm = 0x100
	cmpl	$2, %edi
	jl	.LBB13_1
# %bb.2:
	movq	%rsi, %r14
	movl	%edi, %r15d
	movq	8(%rsi), %rdi
	callq	atoi
	movl	%eax, %ebp
	cmpl	$2, %r15d
	je	.LBB13_4
# %bb.3:
	movq	16(%r14), %rdi
	callq	atoi
	movl	%eax, %ebx
	jmp	.LBB13_4
.LBB13_1:
	movl	$256, %ebp                      # imm = 0x100
.LBB13_4:                               # %.thread
	movl	%ebp, %eax
	cltd
	idivl	%ebx
	movl	%eax, %r14d
	testl	%edx, %edx
	je	.LBB13_6
# %bb.5:
	incl	%r14d
	movl	%r14d, %ebp
	imull	%ebx, %ebp
	movl	$.Lstr, %edi
	callq	puts@PLT
	movl	$.L.str.2, %edi
	movl	%ebp, %esi
	xorl	%eax, %eax
	callq	printf
.LBB13_6:
	callq	_ZNSt6chrono3_V212system_clock3nowEv
	movq	%rax, %r15
	movl	%ebp, %edi
	callq	_Z15executeHostTestiii
	callq	_ZNSt6chrono3_V212system_clock3nowEv
	subq	%r15, %rax
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI13_0(%rip), %xmm0
	movsd	%xmm0, (%rsp)                   # 8-byte Spill
	movl	$_ZSt4cout, %edi
	movl	$.L.str.3, %esi
	movl	$21, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movsd	(%rsp), %xmm0                   # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r15
	movl	$.L.str.4, %esi
	movl	$9, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	(%r15), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	callq	_ZNSt6chrono3_V212system_clock3nowEv
	movq	%rax, %r15
	movl	%ebp, %edi
	movl	%ebx, %esi
	movl	%r14d, %edx
	callq	_Z17executeGlobalTestiii
	callq	_ZNSt6chrono3_V212system_clock3nowEv
	subq	%r15, %rax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI13_0(%rip), %xmm0
	movsd	%xmm0, (%rsp)                   # 8-byte Spill
	movl	$_ZSt4cout, %edi
	movl	$.L.str.5, %esi
	movl	$30, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movsd	(%rsp), %xmm0                   # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r15
	movl	$.L.str.4, %esi
	movl	$9, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	(%r15), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	callq	_ZNSt6chrono3_V212system_clock3nowEv
	movq	%rax, %r15
	movl	%ebp, %edi
	movl	%ebx, %esi
	movl	%r14d, %edx
	callq	_Z17executeSharedTestiii
	callq	_ZNSt6chrono3_V212system_clock3nowEv
	subq	%r15, %rax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI13_0(%rip), %xmm0
	movsd	%xmm0, (%rsp)                   # 8-byte Spill
	movl	$_ZSt4cout, %edi
	movl	$.L.str.6, %esi
	movl	$30, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movsd	(%rsp), %xmm0                   # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r15
	movl	$.L.str.4, %esi
	movl	$9, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	(%r15), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	callq	_ZNSt6chrono3_V212system_clock3nowEv
	movq	%rax, %r15
	movl	%ebp, %edi
	movl	%ebx, %esi
	movl	%r14d, %edx
	callq	_Z19executeConstantTestiii
	callq	_ZNSt6chrono3_V212system_clock3nowEv
	subq	%r15, %rax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI13_0(%rip), %xmm0
	movsd	%xmm0, (%rsp)                   # 8-byte Spill
	movl	$_ZSt4cout, %edi
	movl	$.L.str.7, %esi
	movl	$32, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movsd	(%rsp), %xmm0                   # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %rbx
	movl	$.L.str.4, %esi
	movl	$9, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	(%rbx), %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
	callq	_ZNSo5flushEv
	xorl	%eax, %eax
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.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_end13:
	.size	main, .Lfunc_end13-main
	.cfi_endproc
                                        # -- End function
	.section	.text._ZNSt24uniform_int_distributionIiEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEiRT_RKNS0_10param_typeE,"axG",@progbits,_ZNSt24uniform_int_distributionIiEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEiRT_RKNS0_10param_typeE,comdat
	.weak	_ZNSt24uniform_int_distributionIiEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEiRT_RKNS0_10param_typeE # -- Begin function _ZNSt24uniform_int_distributionIiEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEiRT_RKNS0_10param_typeE
	.p2align	1, 0x90
	.type	_ZNSt24uniform_int_distributionIiEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEiRT_RKNS0_10param_typeE,@function
_ZNSt24uniform_int_distributionIiEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEiRT_RKNS0_10param_typeE: # @_ZNSt24uniform_int_distributionIiEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEiRT_RKNS0_10param_typeE
	.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	$24, %rsp
	.cfi_def_cfa_offset 80
	.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, %rbx
	movq	%rdi, %r15
	movslq	4(%rdx), %rbp
	movslq	(%rdx), %rdi
	subq	%rdi, %rbp
	cmpq	$2147483644, %rbp               # imm = 0x7FFFFFFC
	ja	.LBB14_4
# %bb.1:
	incq	%rbp
	movl	$2147483645, %eax               # imm = 0x7FFFFFFD
	xorl	%edx, %edx
	divl	%ebp
	movl	%eax, %ecx
	imulq	%rcx, %rbp
	movq	(%rbx), %rdx
	movabsq	$8589934597, %r8                # imm = 0x200000005
.LBB14_2:                               # =>This Inner Loop Header: Depth=1
	imulq	$16807, %rdx, %rsi              # imm = 0x41A7
	movq	%rsi, %rax
	mulq	%r8
	movq	%rsi, %rax
	subq	%rdx, %rax
	shrq	%rax
	addq	%rdx, %rax
	shrq	$30, %rax
	movq	%rax, %rdx
	shlq	$31, %rdx
	subq	%rdx, %rax
	leaq	(%rsi,%rax), %rdx
	addq	%rsi, %rax
	decq	%rax
	cmpq	%rbp, %rax
	jae	.LBB14_2
# %bb.3:
	movq	%rdx, (%rbx)
	xorl	%edx, %edx
	divq	%rcx
	jmp	.LBB14_9
.LBB14_4:
	cmpq	$2147483645, %rbp               # imm = 0x7FFFFFFD
	jne	.LBB14_5
# %bb.10:
	imulq	$16807, (%rbx), %rcx            # imm = 0x41A7
	movabsq	$8589934597, %rdx               # imm = 0x200000005
	movq	%rcx, %rax
	mulq	%rdx
	movq	%rcx, %rax
	subq	%rdx, %rax
	shrq	%rax
	addq	%rdx, %rax
	shrq	$30, %rax
	movq	%rax, %rdx
	shlq	$31, %rdx
	subq	%rdx, %rax
	leaq	(%rcx,%rax), %rdx
	movq	%rdx, (%rbx)
	addq	%rcx, %rax
	decq	%rax
	jmp	.LBB14_9
.LBB14_5:                               # %.preheader
	movq	%rdx, 16(%rsp)                  # 8-byte Spill
	movq	%rbp, %rax
	shrq	%rax
	movabsq	$-9223372028264841207, %rcx     # imm = 0x8000000200000009
	mulq	%rcx
	movq	%rdx, %r12
	shrq	$29, %r12
	leaq	8(%rsp), %r13
	movabsq	$8589934597, %r14               # imm = 0x200000005
.LBB14_6:                               # =>This Inner Loop Header: Depth=1
	movl	$0, 8(%rsp)
	movl	%r12d, 12(%rsp)
	movq	%r15, %rdi
	movq	%rbx, %rsi
	movq	%r13, %rdx
	callq	_ZNSt24uniform_int_distributionIiEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEiRT_RKNS0_10param_typeE
	movslq	%eax, %rsi
	leaq	(%rsi,%rsi), %rax
	shlq	$31, %rsi
	subq	%rax, %rsi
	imulq	$16807, (%rbx), %rcx            # imm = 0x41A7
	movq	%rcx, %rax
	mulq	%r14
	movq	%rcx, %rax
	subq	%rdx, %rax
	shrq	%rax
	addq	%rdx, %rax
	shrq	$30, %rax
	movq	%rax, %rdx
	shlq	$31, %rdx
	subq	%rdx, %rax
	addq	%rcx, %rax
	movq	%rax, (%rbx)
	addq	%rsi, %rax
	decq	%rax
	cmpq	%rbp, %rax
	ja	.LBB14_6
# %bb.7:                                #   in Loop: Header=BB14_6 Depth=1
	cmpq	%rsi, %rax
	jb	.LBB14_6
# %bb.8:                                # %.loopexit.loopexit
	movq	16(%rsp), %rcx                  # 8-byte Reload
	movl	(%rcx), %edi
.LBB14_9:                               # %.loopexit
	addl	%edi, %eax
                                        # kill: def $eax killed $eax killed $rax
	addq	$24, %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_end14:
	.size	_ZNSt24uniform_int_distributionIiEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEiRT_RKNS0_10param_typeE, .Lfunc_end14-_ZNSt24uniform_int_distributionIiEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEiRT_RKNS0_10param_typeE
	.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:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB15_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB15_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z27executeSharedMathOperationsPiS_S_S_S_S_S_i, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z29executeConstantMathOperationsPiS_S_S_S_i, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end15:
	.size	__hip_module_ctor, .Lfunc_end15-__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	.LBB16_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
.LBB16_2:
	retq
.Lfunc_end16:
	.size	__hip_module_dtor, .Lfunc_end16-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z27executeSharedMathOperationsPiS_S_S_S_S_S_i,@object # @_Z27executeSharedMathOperationsPiS_S_S_S_S_S_i
	.section	.rodata,"a",@progbits
	.globl	_Z27executeSharedMathOperationsPiS_S_S_S_S_S_i
	.p2align	3, 0x0
_Z27executeSharedMathOperationsPiS_S_S_S_S_S_i:
	.quad	_Z42__device_stub__executeSharedMathOperationsPiS_S_S_S_S_S_i
	.size	_Z27executeSharedMathOperationsPiS_S_S_S_S_S_i, 8

	.type	_Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i,@object # @_Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i
	.globl	_Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i
	.p2align	3, 0x0
_Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i:
	.quad	_Z42__device_stub__executeGlobalMathOperationsPiS_S_S_S_S_S_i
	.size	_Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i, 8

	.type	_Z29executeConstantMathOperationsPiS_S_S_S_i,@object # @_Z29executeConstantMathOperationsPiS_S_S_S_i
	.globl	_Z29executeConstantMathOperationsPiS_S_S_S_i
	.p2align	3, 0x0
_Z29executeConstantMathOperationsPiS_S_S_S_i:
	.quad	_Z44__device_stub__executeConstantMathOperationsPiS_S_S_S_i
	.size	_Z29executeConstantMathOperationsPiS_S_S_S_i, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	" "
	.size	.L.str, 2

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"The total number of threads will be rounded up to %d\n"
	.size	.L.str.2, 54

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"Host execution took: "
	.size	.L.str.3, 22

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	" seconds."
	.size	.L.str.4, 10

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"Global Memory execution took: "
	.size	.L.str.5, 31

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"Shared Memory execution took: "
	.size	.L.str.6, 31

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"Constant Memory execution took: "
	.size	.L.str.7, 33

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z27executeSharedMathOperationsPiS_S_S_S_S_S_i"
	.size	.L__unnamed_1, 47

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i"
	.size	.L__unnamed_2, 47

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"_Z29executeConstantMathOperationsPiS_S_S_S_i"
	.size	.L__unnamed_3, 45

	.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

	.type	.Lstr,@object                   # @str
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr:
	.asciz	"Warning: Total thread count is not evenly divisible by the block size"
	.size	.Lstr, 70

	.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 _Z42__device_stub__executeSharedMathOperationsPiS_S_S_S_S_S_i
	.addrsig_sym _Z42__device_stub__executeGlobalMathOperationsPiS_S_S_S_S_S_i
	.addrsig_sym _Z44__device_stub__executeConstantMathOperationsPiS_S_S_S_i
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z27executeSharedMathOperationsPiS_S_S_S_S_S_i
	.addrsig_sym _Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i
	.addrsig_sym _Z29executeConstantMathOperationsPiS_S_S_S_i
	.addrsig_sym _ZSt4cout
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z27executeSharedMathOperationsPiS_S_S_S_S_S_i ; -- Begin function _Z27executeSharedMathOperationsPiS_S_S_S_S_S_i
	.globl	_Z27executeSharedMathOperationsPiS_S_S_S_S_S_i
	.p2align	8
	.type	_Z27executeSharedMathOperationsPiS_S_S_S_S_S_i,@function
_Z27executeSharedMathOperationsPiS_S_S_S_S_S_i: ; @_Z27executeSharedMathOperationsPiS_S_S_S_S_S_i
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x4c
	s_load_b256 s[4:11], s[0:1], 0x0
	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_load_b32 s2, s[0:1], 0x38
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[2:3], 2, v[1:2]
	v_lshlrev_b32_e32 v1, 2, v1
	v_add_co_u32 v4, vcc_lo, s4, v2
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v5, vcc_lo, s5, v3, vcc_lo
	v_add_co_u32 v6, vcc_lo, s6, v2
	v_add_co_ci_u32_e32 v7, vcc_lo, s7, v3, vcc_lo
	s_waitcnt lgkmcnt(0)
	s_lshl_b32 s2, s2, 2
	global_load_b32 v0, v[4:5], off
	global_load_b32 v4, v[6:7], off
	s_add_i32 s3, s2, 0
	v_add_nc_u32_e32 v5, 0, v1
	v_add_nc_u32_e32 v6, s3, v1
	v_add3_u32 v10, s3, s2, v1
	s_load_b256 s[0:7], s[0:1], 0x20
	s_waitcnt vmcnt(1)
	ds_store_b32 v5, v0
	s_waitcnt vmcnt(0)
	ds_store_b32 v6, v4
	ds_load_b32 v0, v5
	s_waitcnt lgkmcnt(0)
	v_add_nc_u32_e32 v11, v0, v4
	ds_store_b32 v10, v11
	ds_load_b32 v0, v5
	ds_load_b32 v1, v6
	s_waitcnt lgkmcnt(0)
	v_sub_nc_u32_e32 v12, v0, v1
	ds_store_b32 v10, v12
	ds_load_b32 v0, v5
	ds_load_b32 v1, v6
	s_waitcnt lgkmcnt(0)
	v_mul_lo_u32 v13, v1, v0
	ds_store_b32 v10, v13
	ds_load_b32 v0, v6
	ds_load_b32 v7, v5
	s_waitcnt lgkmcnt(1)
	v_ashrrev_i32_e32 v1, 31, v0
	s_waitcnt lgkmcnt(0)
	v_ashrrev_i32_e32 v9, 31, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v0, v0, v1
	v_add_nc_u32_e32 v7, v7, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v0, v0, v1
	v_xor_b32_e32 v7, v7, v9
	v_xor_b32_e32 v1, v9, v1
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cvt_f32_u32_e32 v4, v0
	v_sub_nc_u32_e32 v8, 0, v0
	v_rcp_iflag_f32_e32 v4, v4
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v4, 0x4f7ffffe, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v4, v4
	v_mul_lo_u32 v8, v8, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v8, v4, v8
	v_add_nc_u32_e32 v4, v4, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v4, v7, v4
	v_mul_lo_u32 v8, v4, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v7, v7, v8
	v_add_nc_u32_e32 v8, 1, v4
	v_sub_nc_u32_e32 v14, v7, v0
	v_cmp_ge_u32_e32 vcc_lo, v7, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v4, v4, v8 :: v_dual_cndmask_b32 v7, v7, v14
	v_add_nc_u32_e32 v8, 1, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_ge_u32_e32 vcc_lo, v7, v0
	v_cndmask_b32_e32 v0, v4, v8, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v0, v0, v1
	v_sub_nc_u32_e32 v14, v0, v1
	ds_store_b32 v10, v14
	ds_load_b32 v0, v6
	s_waitcnt lgkmcnt(0)
	v_ashrrev_i32_e32 v1, 31, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v0, v0, v1
	v_xor_b32_e32 v8, v0, v1
	ds_load_b32 v1, v5
	v_cvt_f32_u32_e32 v0, v8
	v_sub_nc_u32_e32 v4, 0, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v0, v0
	s_waitcnt lgkmcnt(0)
	v_ashrrev_i32_e32 v15, 31, v1
	s_waitcnt_depctr 0xfff
	v_dual_mul_f32 v0, 0x4f7ffffe, v0 :: v_dual_add_nc_u32 v1, v1, v15
	v_cvt_u32_f32_e32 v0, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v1, v1, v15
	v_mul_lo_u32 v4, v4, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v4, v0, v4
	v_add_nc_u32_e32 v0, v0, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v0, v1, v0
	v_mul_lo_u32 v0, v0, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_sub_nc_u32_e32 v4, v1, v0
	v_add_co_u32 v0, vcc_lo, s8, v2
	v_add_co_ci_u32_e32 v1, vcc_lo, s9, v3, vcc_lo
	v_sub_nc_u32_e32 v5, v4, v8
	v_cmp_ge_u32_e32 vcc_lo, v4, v8
	s_delay_alu instid0(VALU_DEP_2)
	v_cndmask_b32_e32 v9, v4, v5, vcc_lo
	v_add_co_u32 v4, vcc_lo, s10, v2
	v_add_co_ci_u32_e32 v5, vcc_lo, s11, v3, vcc_lo
	v_add_co_u32 v6, vcc_lo, s0, v2
	v_add_co_ci_u32_e32 v7, vcc_lo, s1, v3, vcc_lo
	v_sub_nc_u32_e32 v16, v9, v8
	v_cmp_ge_u32_e32 vcc_lo, v9, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_cndmask_b32_e32 v16, v9, v16, vcc_lo
	v_add_co_u32 v8, vcc_lo, s2, v2
	v_add_co_ci_u32_e32 v9, vcc_lo, s3, v3, vcc_lo
	v_xor_b32_e32 v16, v16, v15
	v_add_co_u32 v2, vcc_lo, s4, v2
	v_add_co_ci_u32_e32 v3, vcc_lo, s5, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_3)
	v_sub_nc_u32_e32 v15, v16, v15
	global_store_b32 v[0:1], v11, off
	global_store_b32 v[4:5], v12, off
	global_store_b32 v[6:7], v13, off
	global_store_b32 v[8:9], v14, off
	ds_store_b32 v10, v15
	global_store_b32 v[2:3], v15, off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z27executeSharedMathOperationsPiS_S_S_S_S_S_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 320
		.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 17
		.amdhsa_next_free_sgpr 16
		.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	_Z27executeSharedMathOperationsPiS_S_S_S_S_S_i, .Lfunc_end0-_Z27executeSharedMathOperationsPiS_S_S_S_S_S_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 780
; NumSgprs: 18
; NumVgprs: 17
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 2
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 17
; 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
	.protected	_Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i ; -- Begin function _Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i
	.globl	_Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i
	.p2align	8
	.type	_Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i,@function
_Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i: ; @_Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x4c
	s_load_b256 s[4:11], s[0:1], 0x0
	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]
	v_mov_b32_e32 v2, 0
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v2, vcc_lo, s4, v0
	v_add_co_ci_u32_e32 v3, vcc_lo, s5, v1, vcc_lo
	v_add_co_u32 v4, vcc_lo, s6, v0
	v_add_co_ci_u32_e32 v5, vcc_lo, s7, v1, vcc_lo
	v_add_co_u32 v6, vcc_lo, s8, v0
	global_load_b32 v8, v[2:3], off
	global_load_b32 v9, v[4:5], off
	v_add_co_ci_u32_e32 v7, vcc_lo, s9, v1, vcc_lo
	s_load_b256 s[0:7], s[0:1], 0x20
	s_waitcnt vmcnt(0)
	v_add_nc_u32_e32 v8, v9, v8
	global_store_b32 v[6:7], v8, off
	global_load_b32 v8, v[2:3], off
	global_load_b32 v9, v[4:5], off
	v_add_co_u32 v6, vcc_lo, s10, v0
	v_add_co_ci_u32_e32 v7, vcc_lo, s11, v1, vcc_lo
	s_waitcnt vmcnt(0)
	v_sub_nc_u32_e32 v8, v8, v9
	global_store_b32 v[6:7], v8, off
	global_load_b32 v6, v[2:3], off
	global_load_b32 v7, v[4:5], off
	s_waitcnt vmcnt(0)
	v_mul_lo_u32 v8, v7, v6
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v6, vcc_lo, s0, v0
	v_add_co_ci_u32_e32 v7, vcc_lo, s1, v1, vcc_lo
	global_store_b32 v[6:7], v8, off
	global_load_b32 v6, v[4:5], off
	global_load_b32 v7, v[2:3], off
	s_waitcnt vmcnt(1)
	v_ashrrev_i32_e32 v8, 31, v6
	s_waitcnt vmcnt(0)
	v_ashrrev_i32_e32 v11, 31, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v6, v6, v8
	v_add_nc_u32_e32 v7, v7, v11
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v6, v6, v8
	v_xor_b32_e32 v7, v7, v11
	v_xor_b32_e32 v8, v11, v8
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cvt_f32_u32_e32 v9, v6
	v_sub_nc_u32_e32 v10, 0, v6
	v_rcp_iflag_f32_e32 v9, v9
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v9, 0x4f7ffffe, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v9, v9
	v_mul_lo_u32 v10, v10, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v10, v9, v10
	v_add_nc_u32_e32 v9, v9, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v9, v7, v9
	v_mul_lo_u32 v10, v9, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v7, v7, v10
	v_sub_nc_u32_e32 v12, v7, v6
	v_cmp_ge_u32_e32 vcc_lo, v7, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v7, v7, v12 :: v_dual_add_nc_u32 v10, 1, v9
	v_cndmask_b32_e32 v9, v9, v10, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_ge_u32_e32 vcc_lo, v7, v6
	v_add_nc_u32_e32 v10, 1, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v6, v9, v10, vcc_lo
	v_xor_b32_e32 v9, v6, v8
	v_add_co_u32 v6, vcc_lo, s2, v0
	v_add_co_ci_u32_e32 v7, vcc_lo, s3, v1, vcc_lo
	s_delay_alu instid0(VALU_DEP_3)
	v_sub_nc_u32_e32 v8, v9, v8
	global_store_b32 v[6:7], v8, off
	global_load_b32 v4, v[4:5], off
	global_load_b32 v2, v[2:3], off
	s_waitcnt vmcnt(1)
	v_ashrrev_i32_e32 v3, 31, v4
	s_waitcnt vmcnt(0)
	v_ashrrev_i32_e32 v6, 31, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v4, v4, v3
	v_add_nc_u32_e32 v2, v2, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v3, v4, v3
	v_xor_b32_e32 v2, v2, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cvt_f32_u32_e32 v4, v3
	v_sub_nc_u32_e32 v5, 0, v3
	v_rcp_iflag_f32_e32 v4, v4
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v4, 0x4f7ffffe, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v4, v4
	v_mul_lo_u32 v5, v5, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v5, v4, v5
	v_add_nc_u32_e32 v4, v4, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v4, v2, v4
	v_mul_lo_u32 v4, v4, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v2, v2, v4
	v_sub_nc_u32_e32 v4, v2, v3
	v_cmp_ge_u32_e32 vcc_lo, v2, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v2, v2, v4, vcc_lo
	v_sub_nc_u32_e32 v4, v2, v3
	v_cmp_ge_u32_e32 vcc_lo, v2, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_cndmask_b32_e32 v2, v2, v4, vcc_lo
	v_add_co_u32 v0, vcc_lo, s4, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s5, v1, vcc_lo
	v_xor_b32_e32 v2, v2, v6
	s_delay_alu instid0(VALU_DEP_1)
	v_sub_nc_u32_e32 v2, v2, v6
	global_store_b32 v[0:1], v2, off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 320
		.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 13
		.amdhsa_next_free_sgpr 16
		.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_end1:
	.size	_Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i, .Lfunc_end1-_Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 676
; NumSgprs: 18
; NumVgprs: 13
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 13
; 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
	.protected	_Z29executeConstantMathOperationsPiS_S_S_S_i ; -- Begin function _Z29executeConstantMathOperationsPiS_S_S_S_i
	.globl	_Z29executeConstantMathOperationsPiS_S_S_S_i
	.p2align	8
	.type	_Z29executeConstantMathOperationsPiS_S_S_S_i,@function
_Z29executeConstantMathOperationsPiS_S_S_S_i: ; @_Z29executeConstantMathOperationsPiS_S_S_S_i
; %bb.0:
	s_clause 0x2
	s_load_b32 s2, s[0:1], 0x3c
	s_load_b256 s[4:11], s[0:1], 0x0
	s_load_b64 s[0:1], s[0:1], 0x20
	v_dual_mov_b32 v10, 4 :: v_dual_mov_b32 v11, -2
	v_dual_mov_b32 v12, 3 :: v_dual_mov_b32 v13, 0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	v_add_co_u32 v2, vcc_lo, s4, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v3, vcc_lo, s5, v1, vcc_lo
	v_add_co_u32 v4, vcc_lo, s6, v0
	v_add_co_ci_u32_e32 v5, vcc_lo, s7, v1, vcc_lo
	v_add_co_u32 v6, vcc_lo, s8, v0
	v_add_co_ci_u32_e32 v7, vcc_lo, s9, v1, vcc_lo
	v_add_co_u32 v8, vcc_lo, s10, v0
	v_add_co_ci_u32_e32 v9, vcc_lo, s11, v1, vcc_lo
	v_add_co_u32 v0, vcc_lo, s0, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_store_b32 v[2:3], v10, off
	global_store_b32 v[4:5], v11, off
	global_store_b32 v[6:7], v12, off
	global_store_b32 v[8:9], v13, off
	global_store_b32 v[0:1], v13, off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z29executeConstantMathOperationsPiS_S_S_S_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 304
		.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 14
		.amdhsa_next_free_sgpr 16
		.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_end2:
	.size	_Z29executeConstantMathOperationsPiS_S_S_S_i, .Lfunc_end2-_Z29executeConstantMathOperationsPiS_S_S_S_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 200
; NumSgprs: 18
; NumVgprs: 14
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 14
; 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
      - .address_space:  global
        .offset:         40
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         48
        .size:           8
        .value_kind:     global_buffer
      - .offset:         56
        .size:           4
        .value_kind:     by_value
      - .offset:         64
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         68
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         72
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         76
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         78
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         80
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         82
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         84
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         86
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         104
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         112
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         120
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         128
        .size:           2
        .value_kind:     hidden_grid_dims
      - .offset:         184
        .size:           4
        .value_kind:     hidden_dynamic_lds_size
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 320
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z27executeSharedMathOperationsPiS_S_S_S_S_S_i
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z27executeSharedMathOperationsPiS_S_S_S_S_S_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     17
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .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
      - .address_space:  global
        .offset:         40
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         48
        .size:           8
        .value_kind:     global_buffer
      - .offset:         56
        .size:           4
        .value_kind:     by_value
      - .offset:         64
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         68
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         72
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         76
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         78
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         80
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         82
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         84
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         86
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         104
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         112
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         120
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         128
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 320
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z27executeGlobalMathOperationsPiS_S_S_S_S_S_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     13
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .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:     by_value
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         52
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         56
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         60
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         62
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         64
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         66
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         68
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         70
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         104
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         112
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 304
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z29executeConstantMathOperationsPiS_S_S_S_i
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z29executeConstantMathOperationsPiS_S_S_S_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     14
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
