/*
    CS 4402 Distributed and Parallel Systems
    Assignment 2 Question 1: N thread blocks and N threads per thread block
    Sarah Whelan 250778849

    TO RUN: nvcc q1_swhela2.cu -o q1_swhela2
            ./q1_swhela2
*/

#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <math.h>

int modBy = 103; // common prime num used for modding coefficient values during generation, multiplication, and addition

void genPolynomials(int *polyA, int *polyB, int size);
void multPolynomialsSerial(int *polyA, int *polyB, int polySize, int *product, int productSize);
__global__ void multPolynomialsParallel(int *polyA, int *polyB, int *product, int polySize, int modBy);
__global__ void sumProductsParallel(int prodSize, int threadsPerBlock, int *summedProduct, int *products, int numBlocks, int modBy);
void checkCUDAError(const char* msg);

int main() {
    srand(time(NULL));
    int numTerms;

    // get user desired input on length of polynomials
    printf("Specify the number of terms in the polynomial by specifying the exponent on base 2, UP TO 10, e.g. enter '3' if you want 2^3 terms (AKA 8 terms) per polynomial: ");
    scanf("%d", &numTerms);

    printf("\nYou entered '%d'.\n", numTerms);
    if (numTerms > 10) {
        printf("Invalid entry. The maximum number of terms is 2^10. Please enter a term less than or equal to 10 next time.");
        return 1;
    }
    
    // then bitshift by input value to determine actual value of numTerms
    numTerms = 1 << numTerms;
    printf("Number of terms per polynomial = %d, hence each polynomial will have degree = %d.\n\n", numTerms, numTerms-1);

    // use numTerms as the number of blocks per thread and the number of blocks
    int threadsPerBlock = numTerms;
    int blocks = numTerms;

    // instantiate and allocate host memory blocks to store each polynomial of size numTerms
    int *host_polyA, *host_polyB;
    host_polyA = (int *) malloc(numTerms * sizeof(int));
    host_polyB = (int *) malloc(numTerms * sizeof(int));

    // generate random polynomials of size numTerms
    printf("Generating polynomials...\n\n");
    genPolynomials(host_polyA, host_polyB, numTerms);

    printf("polyA:\n");
    for (int i = 0; i < numTerms; i++) {
        printf("%dx^%d ", host_polyA[i], i);
        if (i != numTerms-1) {
            printf("+ ");
        }
    }

    printf("\n\npolyB:\n");
    for (int i = 0; i < numTerms; i++) {
        printf("%dx^%d ", host_polyB[i], i);
        if (i != numTerms-1) {
            printf("+ ");
        }
    }

    printf("\n\n");

    // determine degree of product
    int degreeOfProduct = (numTerms - 1) * 2; // e.g. degree(polyA, polyB) = 3 then x^3 * x^3 = x^6 and degree = numTerms - 1

    // allocate blocks of memory on the host for storing the product with size degreeOfProduct + 1 (serial)
    // and numTerms*numTerms for the intermediary parallel product, as well asthe final parallel product
    // two different allocations in order to verify results at the end!
    int *host_product_serial, *host_product_parallel, *host_final_product;
    host_product_serial = (int *) malloc((degreeOfProduct+1) * sizeof(int)); // sum of products is intrinsic
    host_product_parallel = (int *) malloc(numTerms * numTerms * sizeof(int)); // because of n threads in each n thread blocks
    host_final_product = (int *) malloc((degreeOfProduct+1) * sizeof(int)); // final product from parallel version once summed

    // ensure all vals in host_product_parallel are 0 (this is done within the serial version so we don't need to worry about that one)
    for (int i = 0; i < numTerms*numTerms; i++) {
        host_product_parallel[i] = 0;
    }
    // ensure all vals in host_final_product are 0
    for (int i = 0; i < degreeOfProduct+1; i++) {
        host_final_product[i] = 0;
    }

    // initialize and allocate memory on the devices for storing dev_polyA, dev_polyB, and dev_product
    int *dev_polyA, *dev_polyB, *dev_product;
    cudaMalloc( (void **) &dev_polyA, numTerms * sizeof(int));
    cudaMalloc( (void **) &dev_polyB, numTerms * sizeof(int));
    cudaMalloc( (void **) &dev_product, numTerms * numTerms * sizeof(int));

    // copy polynomials: host -> device (dest, src, size, direction)
    cudaMemcpy(dev_polyA, host_polyA, numTerms * sizeof(int), cudaMemcpyHostToDevice); 
    cudaMemcpy(dev_polyB, host_polyB, numTerms * sizeof(int), cudaMemcpyHostToDevice); 
    cudaMemcpy(dev_product, host_product_parallel, numTerms * numTerms * sizeof(int), cudaMemcpyHostToDevice);

    // setup kernel params & launch
    dim3 dimGrid(blocks);
    dim3 dimBlock(threadsPerBlock);
    multPolynomialsParallel<<<dimGrid, dimBlock>>>(dev_polyA, dev_polyB, dev_product, numTerms, modBy);

    cudaThreadSynchronize(); // wait for ALL threads from all blocks to complete
    checkCUDAError("kernel invocation");

    // copy dev_product back into host_product_parallel (dest, src, size, direction)
    cudaMemcpy(host_product_parallel, dev_product, numTerms * numTerms * sizeof(int), cudaMemcpyDeviceToHost);
    
    /* ~~~ now we need to deal with the summation of intermediary products ~~~ */

    // allocate device mem for final product
    int *dev_final;
    cudaMalloc( (void **) &dev_final, (degreeOfProduct+1) * sizeof(int));

    // copy zero'd host_final_product to dev_final and host_product_parallel to dev_product
    // (dest, src, size, direction)
    cudaMemcpy(dev_final, host_final_product, (degreeOfProduct+1) * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_product, host_product_parallel, numTerms * numTerms * sizeof(int), cudaMemcpyHostToDevice);

    // parameters are (int prodSize, int threadsPerBlock, int *summedProduct, int *products, int numBlocks, int modBy)
    sumProductsParallel<<<dimGrid, dimBlock>>>(degreeOfProduct+1, threadsPerBlock, dev_final, dev_product, blocks, modBy);

    cudaThreadSynchronize(); // wait for ALL threads from all blocks to complete
    checkCUDAError("kernel invocation");

    // copy summation of products back to host (dest, src, size, direction)
    cudaMemcpy(host_final_product, dev_final, (degreeOfProduct+1) * sizeof(int), cudaMemcpyDeviceToHost);

    // multiply polynomials in serial and write to host_product_serial for verification
    multPolynomialsSerial(host_polyA, host_polyB, numTerms, host_product_serial, degreeOfProduct+1);

    printf("Serial result:\n");
    for (int i = 0; i < degreeOfProduct+1; i++) {
        printf("%dx^%d ", host_product_serial[i], i);
        if (i != degreeOfProduct) {
            printf("+ ");
        }
    }
    printf("\n\nParallel result:\n");
    for (int i = 0; i < degreeOfProduct+1; i++) {
        printf("%dx^%d ", host_final_product[i], i);
        if (i != degreeOfProduct) {
            printf("+ ");
        }
    }
    printf("\n\n");
    bool allRight = 1;
    for (int i = 0; i < degreeOfProduct+1; i++) {
        if (host_product_serial[i] == host_final_product[i]) {
            continue;
        } else {
            printf("Coefficients at degree %d are not equivalent: serial!=parallel (%d!=%d)\n", i, host_product_serial[i], host_final_product[i]);
            allRight = 0;
        }
    }
    if (allRight) {
        printf("Verification successful. The serial and parallel polynomial multiplications produced the same result!\n\n");
    } else {
        printf("Looks like there were some discrepancies. Verification failed.\n\n");
    }

    // free host and device memory
    free(host_polyA);
    free(host_polyB);
    free(host_product_serial);
    free(host_product_parallel);
    free(host_final_product);

    cudaFree(dev_polyA);
    cudaFree(dev_polyB);
    cudaFree(dev_product);
    cudaFree(dev_final);

    return 0;
}

// genPolynomials takes two polynomials and their size (number of terms per polynomial),
// and generates random coefficients for each term mod p
void genPolynomials(int *polyA, int *polyB, int size) {

    // coefficient generation using rand mod p where p = 103
    for (int i = 0; i < size; i++) {
        polyA[i] = rand() % modBy;
        if (polyA[i] == 0) { // we don't want any zeros!!!
            polyA[i] = 1;
        }

        polyB[i] = rand() % modBy;
        if (polyB[i] == 0) {
            polyB[i] = 1;
        }
    }
}

// multPolynomialsSerial takes two polynomials and their size, in addition to a memory block to place 
// the sum of products into, as well as the size of the product polynomial
void multPolynomialsSerial(int *polyA, int *polyB, int polySize, int *product, int productSize) {
    int degreeOfTerms;

    // ensure all coefficients of product are 0
    for (int i = 0; i < productSize; i++) {
        product[i] = 0;
    }

    // calculate sum of products
    for (int a = 0; a < polySize; a++) { // iterate through terms in A
        for (int b = 0; b < polySize; b++) { // for each term in A, iterate through all terms in B
            // add degrees (indices) to determine which index this product belongs to in the product array block
            degreeOfTerms = a + b;

            // add product of terms to previous sum and mod by 103
            product[degreeOfTerms] = (product[degreeOfTerms] + polyA[a] * polyB[b]) % modBy;
        }
    }
}

// multPolynomialsParallel determines the intermediary products of the polynomial multiplication problem
__global__ void multPolynomialsParallel(int *polyA, int *polyB, int *product, int polySize, int modBy) {
    int a = blockIdx.x; // all threads in the same block will access the same polyA element
    int b = threadIdx.x; // but all threads will access individual polyB elements
    int myIndex = blockDim.x * blockIdx.x + threadIdx.x; // where to write this thread's product
    product[myIndex] = (polyA[a] * polyB[b]) % modBy;
}

// sumProductsParallel uses prodSize threads, each thread corresponding to a degree, to sum common terms and determine the actual product of polyA and polyB
__global__ void sumProductsParallel(int prodSize, int threadsPerBlock, int *summedProduct, int *products, int numBlocks, int modBy) {
    int responsibleFor = blockIdx.x * blockDim.x + threadIdx.x; // used to check which threads are going to be active during this step

    if (responsibleFor < prodSize) { // e.g. 1 < 7 so thread 1 is going to be in charge of summing the x^1 terms, threads >= prodSize will be inactive for remainder
        for (int blockNum = 0; blockNum < numBlocks; blockNum++) { // loop through blocks
            for (int indexInBlock = 0; indexInBlock < threadsPerBlock; indexInBlock++) { // loop through each index per block
                
                int degreeOfElement = blockNum + indexInBlock; // the degree related to the coefficient stored at each products[] index is equal to the block number + the relative index in the block
                
                if (indexInBlock == 0 && degreeOfElement > responsibleFor) {
                    return; // this thread is done summing its common terms
                } 
                else if (degreeOfElement == responsibleFor) { // if this thread is responsible for the degree we just calculated
                    int spotInProducts = blockNum * blockDim.x + indexInBlock; // get its actual index in products[]
                    summedProduct[responsibleFor] = (summedProduct[responsibleFor] + products[spotInProducts]) % modBy; // and write that value into the final summedProduct[our degree]
                }
            }
        }
    }
}

void checkCUDAError(const char *msg) {
    cudaError_t err = cudaGetLastError();
    if(cudaSuccess != err)
    {
        fprintf(stderr, "CUDA error: %s: %s.\n", msg, cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }
}


	.file	"tmpxft_002e41c2_00000000-6_q1_swhela2.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2033:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2033:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z14genPolynomialsPiS_i
	.type	_Z14genPolynomialsPiS_i, @function
_Z14genPolynomialsPiS_i:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movq	%rsi, %r14
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	movl	$1, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	movq	%rdi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	movl	%edx, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	xorl	%ebx, %ebx
.L3:
	cmpl	%ebx, %ebp
	jle	.L9
	call	rand@PLT
	cltd
	idivl	modBy(%rip)
	testl	%edx, %edx
	cmove	%r13d, %edx
	movl	%edx, (%r12,%rbx,4)
	call	rand@PLT
	cltd
	idivl	modBy(%rip)
	testl	%edx, %edx
	cmove	%r13d, %edx
	movl	%edx, (%r14,%rbx,4)
	incq	%rbx
	jmp	.L3
.L9:
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2028:
	.size	_Z14genPolynomialsPiS_i, .-_Z14genPolynomialsPiS_i
	.globl	_Z21multPolynomialsSerialPiS_iS_i
	.type	_Z21multPolynomialsSerialPiS_iS_i, @function
_Z21multPolynomialsSerialPiS_iS_i:
.LFB2029:
	.cfi_startproc
	endbr64
	movq	%rdi, %r10
	movq	%rsi, %r11
	movq	%rcx, %rdi
	movl	%edx, %esi
	xorl	%eax, %eax
.L11:
	cmpl	%eax, %r8d
	jle	.L17
	xorl	%edx, %edx
	movl	%edx, (%rdi,%rax,4)
	incq	%rax
	jmp	.L11
.L17:
	xorl	%ecx, %ecx
.L13:
	cmpl	%ecx, %esi
	jle	.L18
	leaq	(%rdi,%rcx,4), %r9
	xorl	%r8d, %r8d
.L14:
	movl	(%r10,%rcx,4), %eax
	imull	(%r11,%r8,4), %eax
	incq	%r8
	addq	$4, %r9
	addl	-4(%r9), %eax
	cltd
	idivl	modBy(%rip)
	movl	%edx, -4(%r9)
	cmpl	%r8d, %esi
	jg	.L14
	incq	%rcx
	jmp	.L13
.L18:
	ret
	.cfi_endproc
.LFE2029:
	.size	_Z21multPolynomialsSerialPiS_iS_i, .-_Z21multPolynomialsSerialPiS_iS_i
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"CUDA error: %s: %s.\n"
	.text
	.globl	_Z14checkCUDAErrorPKc
	.type	_Z14checkCUDAErrorPKc, @function
_Z14checkCUDAErrorPKc:
.LFB2030:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rdi, %rbx
	call	cudaGetLastError@PLT
	testl	%eax, %eax
	je	.L19
	movl	%eax, %edi
	call	cudaGetErrorString@PLT
	movq	stderr(%rip), %rdi
	movq	%rbx, %rcx
	leaq	.LC0(%rip), %rdx
	movq	%rax, %r8
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	movl	$1, %edi
	call	exit@PLT
.L19:
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2030:
	.size	_Z14checkCUDAErrorPKc, .-_Z14checkCUDAErrorPKc
	.globl	_Z49__device_stub__Z23multPolynomialsParallelPiS_S_iiPiS_S_ii
	.type	_Z49__device_stub__Z23multPolynomialsParallelPiS_S_iiPiS_S_ii, @function
_Z49__device_stub__Z23multPolynomialsParallelPiS_S_iiPiS_S_ii:
.LFB2055:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movl	%ecx, 4(%rsp)
	leaq	40(%rsp), %rcx
	movl	%r8d, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	%rsp, %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L22
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z23multPolynomialsParallelPiS_S_ii(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 176
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L22:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L24
	call	__stack_chk_fail@PLT
.L24:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2055:
	.size	_Z49__device_stub__Z23multPolynomialsParallelPiS_S_iiPiS_S_ii, .-_Z49__device_stub__Z23multPolynomialsParallelPiS_S_iiPiS_S_ii
	.globl	_Z23multPolynomialsParallelPiS_S_ii
	.type	_Z23multPolynomialsParallelPiS_S_ii, @function
_Z23multPolynomialsParallelPiS_S_ii:
.LFB2056:
	.cfi_startproc
	endbr64
	jmp	_Z49__device_stub__Z23multPolynomialsParallelPiS_S_iiPiS_S_ii
	.cfi_endproc
.LFE2056:
	.size	_Z23multPolynomialsParallelPiS_S_ii, .-_Z23multPolynomialsParallelPiS_S_ii
	.globl	_Z45__device_stub__Z19sumProductsParalleliiPiS_iiiiPiS_ii
	.type	_Z45__device_stub__Z19sumProductsParalleliiPiS_iiiiPiS_ii, @function
_Z45__device_stub__Z19sumProductsParalleliiPiS_iiiiPiS_ii:
.LFB2057:
	.cfi_startproc
	endbr64
	subq	$168, %rsp
	.cfi_def_cfa_offset 176
	movl	%edi, 28(%rsp)
	leaq	56(%rsp), %rdi
	movl	%esi, 24(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 16(%rsp)
	leaq	40(%rsp), %rdx
	movq	%rcx, 8(%rsp)
	leaq	48(%rsp), %rcx
	movl	%r8d, 4(%rsp)
	movl	%r9d, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 136(%rsp)
	movq	%rsp, %rax
	movq	%rax, 144(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L27
	pushq	48(%rsp)
	.cfi_def_cfa_offset 184
	leaq	_Z19sumProductsParalleliiPiS_ii(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 192
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 184
	popq	%rdx
	.cfi_def_cfa_offset 176
.L27:
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L29
	call	__stack_chk_fail@PLT
.L29:
	addq	$168, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2057:
	.size	_Z45__device_stub__Z19sumProductsParalleliiPiS_iiiiPiS_ii, .-_Z45__device_stub__Z19sumProductsParalleliiPiS_iiiiPiS_ii
	.globl	_Z19sumProductsParalleliiPiS_ii
	.type	_Z19sumProductsParalleliiPiS_ii, @function
_Z19sumProductsParalleliiPiS_ii:
.LFB2058:
	.cfi_startproc
	endbr64
	jmp	_Z45__device_stub__Z19sumProductsParalleliiPiS_iiiiPiS_ii
	.cfi_endproc
.LFE2058:
	.size	_Z19sumProductsParalleliiPiS_ii, .-_Z19sumProductsParalleliiPiS_ii
	.section	.rodata.str1.1
.LC1:
	.string	"Specify the number of terms in the polynomial by specifying the exponent on base 2, UP TO 10, e.g. enter '3' if you want 2^3 terms (AKA 8 terms) per polynomial: "
.LC2:
	.string	"%d"
.LC3:
	.string	"\nYou entered '%d'.\n"
.LC4:
	.string	"Invalid entry. The maximum number of terms is 2^10. Please enter a term less than or equal to 10 next time."
.LC5:
	.string	"Number of terms per polynomial = %d, hence each polynomial will have degree = %d.\n\n"
.LC6:
	.string	"Generating polynomials...\n\n"
.LC7:
	.string	"polyA:\n"
.LC8:
	.string	"%dx^%d "
.LC9:
	.string	"+ "
.LC10:
	.string	"\n\npolyB:\n"
.LC11:
	.string	"\n\n"
.LC12:
	.string	"kernel invocation"
.LC13:
	.string	"Serial result:\n"
.LC14:
	.string	"\n\nParallel result:\n"
.LC15:
	.string	"Coefficients at degree %d are not equivalent: serial!=parallel (%d!=%d)\n"
.LC16:
	.string	"Verification successful. The serial and parallel polynomial multiplications produced the same result!\n\n"
.LC17:
	.string	"Looks like there were some discrepancies. Verification failed.\n\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	xorl	%edi, %edi
	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	$120, %rsp
	.cfi_def_cfa_offset 176
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	call	time@PLT
	movl	%eax, %edi
	call	srand@PLT
	leaq	.LC1(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	44(%rsp), %rsi
	leaq	.LC2(%rip), %rdi
	xorl	%eax, %eax
	call	__isoc23_scanf@PLT
	movl	44(%rsp), %edx
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC3(%rip), %rsi
	call	__printf_chk@PLT
	movl	44(%rsp), %ecx
	cmpl	$10, %ecx
	jle	.L33
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$1, %eax
	jmp	.L32
.L33:
	movl	$1, %edx
	leaq	.LC5(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	sall	%cl, %edx
	leaq	.LC8(%rip), %r12
	leal	-1(%rdx), %ecx
	movl	%edx, 44(%rsp)
	call	__printf_chk@PLT
	movslq	44(%rsp), %rbx
	movq	%rbx, %r14
	salq	$2, %rbx
	movq	%rbx, %rdi
	call	malloc@PLT
	movq	%rbx, %rdi
	xorl	%ebx, %ebx
	movq	%rax, %rbp
	call	malloc@PLT
	leaq	.LC6(%rip), %rsi
	movl	$2, %edi
	movq	%rax, 8(%rsp)
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	44(%rsp), %edx
	movq	8(%rsp), %rsi
	movq	%rbp, %rdi
	call	_Z14genPolynomialsPiS_i
	leaq	.LC7(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L35:
	movl	%ebx, %ecx
	cmpl	%ebx, 44(%rsp)
	jle	.L61
	movl	0(%rbp,%rbx,4), %edx
	movq	%r12, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	44(%rsp), %eax
	decl	%eax
	cmpl	%ebx, %eax
	je	.L36
	leaq	.LC9(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L36:
	incq	%rbx
	jmp	.L35
.L61:
	leaq	.LC10(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	xorl	%ebx, %ebx
	call	__printf_chk@PLT
	leaq	.LC8(%rip), %r12
.L38:
	movl	%ebx, %ecx
	cmpl	%ebx, 44(%rsp)
	jle	.L62
	movq	8(%rsp), %rax
	movq	%r12, %rsi
	movl	$2, %edi
	movl	(%rax,%rbx,4), %edx
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	44(%rsp), %eax
	decl	%eax
	cmpl	%ebx, %eax
	je	.L39
	leaq	.LC9(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L39:
	incq	%rbx
	jmp	.L38
.L62:
	leaq	.LC11(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	44(%rsp), %r15d
	leal	-2(%r15,%r15), %r12d
	leal	1(%r12), %eax
	movl	%eax, 24(%rsp)
	cltq
	salq	$2, %rax
	movq	%rax, %rdi
	movq	%rax, (%rsp)
	call	malloc@PLT
	movl	%r15d, %edx
	imull	%r15d, %edx
	movq	%rax, %r13
	movslq	%edx, %rdi
	movl	%edx, 28(%rsp)
	salq	$2, %rdi
	call	malloc@PLT
	movq	(%rsp), %rdi
	movq	%rax, 16(%rsp)
	call	malloc@PLT
	movl	28(%rsp), %edx
	movq	%rax, %rbx
	xorl	%eax, %eax
.L41:
	cmpl	%eax, %edx
	jle	.L63
	movq	16(%rsp), %rcx
	xorl	%esi, %esi
	movl	%esi, (%rcx,%rax,4)
	incq	%rax
	jmp	.L41
.L63:
	xorl	%eax, %eax
.L43:
	cmpl	%eax, %r12d
	jl	.L64
	xorl	%edx, %edx
	movl	%edx, (%rbx,%rax,4)
	incq	%rax
	jmp	.L43
.L64:
	movslq	%r15d, %rsi
	leaq	48(%rsp), %rdi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	movslq	44(%rsp), %rsi
	leaq	56(%rsp), %rdi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	movl	44(%rsp), %esi
	leaq	64(%rsp), %rdi
	imull	%esi, %esi
	movslq	%esi, %rsi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	movslq	44(%rsp), %rdx
	movq	48(%rsp), %rdi
	movq	%rbp, %rsi
	movl	$1, %ecx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movslq	44(%rsp), %rdx
	movq	8(%rsp), %rsi
	movl	$1, %ecx
	movq	56(%rsp), %rdi
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movl	44(%rsp), %edx
	movq	16(%rsp), %rsi
	movl	$1, %ecx
	movq	64(%rsp), %rdi
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movl	%r14d, 80(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movabsq	$4294967297, %rax
	movl	%r14d, 92(%rsp)
	movq	%rax, 84(%rsp)
	movq	80(%rsp), %rdi
	movq	%rax, 96(%rsp)
	movl	88(%rsp), %esi
	movl	100(%rsp), %ecx
	movq	92(%rsp), %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L45
	movl	modBy(%rip), %r8d
	movl	44(%rsp), %ecx
	movq	64(%rsp), %rdx
	movq	56(%rsp), %rsi
	movq	48(%rsp), %rdi
	call	_Z49__device_stub__Z23multPolynomialsParallelPiS_S_iiPiS_S_ii
.L45:
	call	cudaThreadSynchronize@PLT
	leaq	.LC12(%rip), %rdi
	call	_Z14checkCUDAErrorPKc
	movl	44(%rsp), %edx
	movq	64(%rsp), %rsi
	movl	$2, %ecx
	movq	16(%rsp), %rdi
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movq	(%rsp), %rsi
	leaq	72(%rsp), %rdi
	call	cudaMalloc@PLT
	movq	(%rsp), %rdx
	movq	72(%rsp), %rdi
	movq	%rbx, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movl	44(%rsp), %edx
	movq	16(%rsp), %rsi
	movl	$1, %ecx
	movq	64(%rsp), %rdi
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movl	100(%rsp), %ecx
	movl	88(%rsp), %esi
	xorl	%r9d, %r9d
	movq	92(%rsp), %rdx
	movq	80(%rsp), %rdi
	xorl	%r8d, %r8d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L46
	movq	64(%rsp), %rcx
	movq	72(%rsp), %rdx
	movl	%r14d, %r8d
	movl	%r14d, %esi
	movl	modBy(%rip), %r9d
	movl	24(%rsp), %edi
	call	_Z45__device_stub__Z19sumProductsParalleliiPiS_iiiiPiS_ii
.L46:
	call	cudaThreadSynchronize@PLT
	leaq	.LC12(%rip), %rdi
	xorl	%r15d, %r15d
	leaq	.LC8(%rip), %r14
	call	_Z14checkCUDAErrorPKc
	movq	(%rsp), %rdx
	movq	72(%rsp), %rsi
	movq	%rbx, %rdi
	movl	$2, %ecx
	call	cudaMemcpy@PLT
	movl	24(%rsp), %r8d
	movq	%r13, %rcx
	movq	%rbp, %rdi
	movl	44(%rsp), %edx
	movq	8(%rsp), %rsi
	call	_Z21multPolynomialsSerialPiS_iS_i
	leaq	.LC13(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L47:
	movl	%r15d, %ecx
	cmpl	%r15d, %r12d
	jl	.L65
	movl	0(%r13,%r15,4), %edx
	xorl	%eax, %eax
	movq	%r14, %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	cmpl	%r15d, %r12d
	je	.L48
	leaq	.LC9(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L48:
	incq	%r15
	jmp	.L47
.L65:
	leaq	.LC14(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	xorl	%r15d, %r15d
	call	__printf_chk@PLT
	leaq	.LC8(%rip), %r14
.L50:
	movl	%r15d, %ecx
	cmpl	%r15d, %r12d
	jl	.L66
	movl	(%rbx,%r15,4), %edx
	xorl	%eax, %eax
	movq	%r14, %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	cmpl	%r15d, %r12d
	je	.L51
	leaq	.LC9(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L51:
	incq	%r15
	jmp	.L50
.L66:
	leaq	.LC11(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	xorl	%r15d, %r15d
	call	__printf_chk@PLT
	leaq	.LC15(%rip), %r14
	movb	$1, %al
.L53:
	cmpl	%r15d, %r12d
	jl	.L67
	movl	0(%r13,%r15,4), %ecx
	movl	(%rbx,%r15,4), %r8d
	cmpl	%r8d, %ecx
	je	.L54
	movl	%r15d, %edx
	movq	%r14, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	xorl	%eax, %eax
.L54:
	incq	%r15
	jmp	.L53
.L67:
	leaq	.LC16(%rip), %rsi
	testb	%al, %al
	jne	.L60
	leaq	.LC17(%rip), %rsi
.L60:
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	%rbp, %rdi
	call	free@PLT
	movq	8(%rsp), %rdi
	call	free@PLT
	movq	%r13, %rdi
	call	free@PLT
	movq	16(%rsp), %rdi
	call	free@PLT
	movq	%rbx, %rdi
	call	free@PLT
	movq	48(%rsp), %rdi
	call	cudaFree@PLT
	movq	56(%rsp), %rdi
	call	cudaFree@PLT
	movq	64(%rsp), %rdi
	call	cudaFree@PLT
	movq	72(%rsp), %rdi
	call	cudaFree@PLT
	xorl	%eax, %eax
.L32:
	movq	104(%rsp), %rdx
	subq	%fs:40, %rdx
	je	.L58
	call	__stack_chk_fail@PLT
.L58:
	addq	$120, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	main, .-main
	.section	.rodata.str1.1
.LC18:
	.string	"_Z19sumProductsParalleliiPiS_ii"
.LC19:
	.string	"_Z23multPolynomialsParallelPiS_S_ii"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2060:
	.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	.LC18(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z19sumProductsParalleliiPiS_ii(%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	.LC19(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z23multPolynomialsParallelPiS_S_ii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2060:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.globl	modBy
	.data
	.align 4
	.type	modBy, @object
	.size	modBy, 4
modBy:
	.long	103
	.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:

/*
    CS 4402 Distributed and Parallel Systems
    Assignment 2 Question 1: N thread blocks and N threads per thread block
    Sarah Whelan 250778849

    TO RUN: nvcc q1_swhela2.cu -o q1_swhela2
            ./q1_swhela2
*/


#include <hip/hip_runtime.h>
#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <math.h>

int modBy = 103; // common prime num used for modding coefficient values during generation, multiplication, and addition

void genPolynomials(int *polyA, int *polyB, int size);
void multPolynomialsSerial(int *polyA, int *polyB, int polySize, int *product, int productSize);
__global__ void multPolynomialsParallel(int *polyA, int *polyB, int *product, int polySize, int modBy);
__global__ void sumProductsParallel(int prodSize, int threadsPerBlock, int *summedProduct, int *products, int numBlocks, int modBy);
void checkCUDAError(const char* msg);

int main() {
    srand(time(NULL));
    int numTerms;

    // get user desired input on length of polynomials
    printf("Specify the number of terms in the polynomial by specifying the exponent on base 2, UP TO 10, e.g. enter '3' if you want 2^3 terms (AKA 8 terms) per polynomial: ");
    scanf("%d", &numTerms);

    printf("\nYou entered '%d'.\n", numTerms);
    if (numTerms > 10) {
        printf("Invalid entry. The maximum number of terms is 2^10. Please enter a term less than or equal to 10 next time.");
        return 1;
    }
    
    // then bitshift by input value to determine actual value of numTerms
    numTerms = 1 << numTerms;
    printf("Number of terms per polynomial = %d, hence each polynomial will have degree = %d.\n\n", numTerms, numTerms-1);

    // use numTerms as the number of blocks per thread and the number of blocks
    int threadsPerBlock = numTerms;
    int blocks = numTerms;

    // instantiate and allocate host memory blocks to store each polynomial of size numTerms
    int *host_polyA, *host_polyB;
    host_polyA = (int *) malloc(numTerms * sizeof(int));
    host_polyB = (int *) malloc(numTerms * sizeof(int));

    // generate random polynomials of size numTerms
    printf("Generating polynomials...\n\n");
    genPolynomials(host_polyA, host_polyB, numTerms);

    printf("polyA:\n");
    for (int i = 0; i < numTerms; i++) {
        printf("%dx^%d ", host_polyA[i], i);
        if (i != numTerms-1) {
            printf("+ ");
        }
    }

    printf("\n\npolyB:\n");
    for (int i = 0; i < numTerms; i++) {
        printf("%dx^%d ", host_polyB[i], i);
        if (i != numTerms-1) {
            printf("+ ");
        }
    }

    printf("\n\n");

    // determine degree of product
    int degreeOfProduct = (numTerms - 1) * 2; // e.g. degree(polyA, polyB) = 3 then x^3 * x^3 = x^6 and degree = numTerms - 1

    // allocate blocks of memory on the host for storing the product with size degreeOfProduct + 1 (serial)
    // and numTerms*numTerms for the intermediary parallel product, as well asthe final parallel product
    // two different allocations in order to verify results at the end!
    int *host_product_serial, *host_product_parallel, *host_final_product;
    host_product_serial = (int *) malloc((degreeOfProduct+1) * sizeof(int)); // sum of products is intrinsic
    host_product_parallel = (int *) malloc(numTerms * numTerms * sizeof(int)); // because of n threads in each n thread blocks
    host_final_product = (int *) malloc((degreeOfProduct+1) * sizeof(int)); // final product from parallel version once summed

    // ensure all vals in host_product_parallel are 0 (this is done within the serial version so we don't need to worry about that one)
    for (int i = 0; i < numTerms*numTerms; i++) {
        host_product_parallel[i] = 0;
    }
    // ensure all vals in host_final_product are 0
    for (int i = 0; i < degreeOfProduct+1; i++) {
        host_final_product[i] = 0;
    }

    // initialize and allocate memory on the devices for storing dev_polyA, dev_polyB, and dev_product
    int *dev_polyA, *dev_polyB, *dev_product;
    hipMalloc( (void **) &dev_polyA, numTerms * sizeof(int));
    hipMalloc( (void **) &dev_polyB, numTerms * sizeof(int));
    hipMalloc( (void **) &dev_product, numTerms * numTerms * sizeof(int));

    // copy polynomials: host -> device (dest, src, size, direction)
    hipMemcpy(dev_polyA, host_polyA, numTerms * sizeof(int), hipMemcpyHostToDevice); 
    hipMemcpy(dev_polyB, host_polyB, numTerms * sizeof(int), hipMemcpyHostToDevice); 
    hipMemcpy(dev_product, host_product_parallel, numTerms * numTerms * sizeof(int), hipMemcpyHostToDevice);

    // setup kernel params & launch
    dim3 dimGrid(blocks);
    dim3 dimBlock(threadsPerBlock);
    multPolynomialsParallel<<<dimGrid, dimBlock>>>(dev_polyA, dev_polyB, dev_product, numTerms, modBy);

    hipDeviceSynchronize(); // wait for ALL threads from all blocks to complete
    checkCUDAError("kernel invocation");

    // copy dev_product back into host_product_parallel (dest, src, size, direction)
    hipMemcpy(host_product_parallel, dev_product, numTerms * numTerms * sizeof(int), hipMemcpyDeviceToHost);
    
    /* ~~~ now we need to deal with the summation of intermediary products ~~~ */

    // allocate device mem for final product
    int *dev_final;
    hipMalloc( (void **) &dev_final, (degreeOfProduct+1) * sizeof(int));

    // copy zero'd host_final_product to dev_final and host_product_parallel to dev_product
    // (dest, src, size, direction)
    hipMemcpy(dev_final, host_final_product, (degreeOfProduct+1) * sizeof(int), hipMemcpyHostToDevice);
    hipMemcpy(dev_product, host_product_parallel, numTerms * numTerms * sizeof(int), hipMemcpyHostToDevice);

    // parameters are (int prodSize, int threadsPerBlock, int *summedProduct, int *products, int numBlocks, int modBy)
    sumProductsParallel<<<dimGrid, dimBlock>>>(degreeOfProduct+1, threadsPerBlock, dev_final, dev_product, blocks, modBy);

    hipDeviceSynchronize(); // wait for ALL threads from all blocks to complete
    checkCUDAError("kernel invocation");

    // copy summation of products back to host (dest, src, size, direction)
    hipMemcpy(host_final_product, dev_final, (degreeOfProduct+1) * sizeof(int), hipMemcpyDeviceToHost);

    // multiply polynomials in serial and write to host_product_serial for verification
    multPolynomialsSerial(host_polyA, host_polyB, numTerms, host_product_serial, degreeOfProduct+1);

    printf("Serial result:\n");
    for (int i = 0; i < degreeOfProduct+1; i++) {
        printf("%dx^%d ", host_product_serial[i], i);
        if (i != degreeOfProduct) {
            printf("+ ");
        }
    }
    printf("\n\nParallel result:\n");
    for (int i = 0; i < degreeOfProduct+1; i++) {
        printf("%dx^%d ", host_final_product[i], i);
        if (i != degreeOfProduct) {
            printf("+ ");
        }
    }
    printf("\n\n");
    bool allRight = 1;
    for (int i = 0; i < degreeOfProduct+1; i++) {
        if (host_product_serial[i] == host_final_product[i]) {
            continue;
        } else {
            printf("Coefficients at degree %d are not equivalent: serial!=parallel (%d!=%d)\n", i, host_product_serial[i], host_final_product[i]);
            allRight = 0;
        }
    }
    if (allRight) {
        printf("Verification successful. The serial and parallel polynomial multiplications produced the same result!\n\n");
    } else {
        printf("Looks like there were some discrepancies. Verification failed.\n\n");
    }

    // free host and device memory
    free(host_polyA);
    free(host_polyB);
    free(host_product_serial);
    free(host_product_parallel);
    free(host_final_product);

    hipFree(dev_polyA);
    hipFree(dev_polyB);
    hipFree(dev_product);
    hipFree(dev_final);

    return 0;
}

// genPolynomials takes two polynomials and their size (number of terms per polynomial),
// and generates random coefficients for each term mod p
void genPolynomials(int *polyA, int *polyB, int size) {

    // coefficient generation using rand mod p where p = 103
    for (int i = 0; i < size; i++) {
        polyA[i] = rand() % modBy;
        if (polyA[i] == 0) { // we don't want any zeros!!!
            polyA[i] = 1;
        }

        polyB[i] = rand() % modBy;
        if (polyB[i] == 0) {
            polyB[i] = 1;
        }
    }
}

// multPolynomialsSerial takes two polynomials and their size, in addition to a memory block to place 
// the sum of products into, as well as the size of the product polynomial
void multPolynomialsSerial(int *polyA, int *polyB, int polySize, int *product, int productSize) {
    int degreeOfTerms;

    // ensure all coefficients of product are 0
    for (int i = 0; i < productSize; i++) {
        product[i] = 0;
    }

    // calculate sum of products
    for (int a = 0; a < polySize; a++) { // iterate through terms in A
        for (int b = 0; b < polySize; b++) { // for each term in A, iterate through all terms in B
            // add degrees (indices) to determine which index this product belongs to in the product array block
            degreeOfTerms = a + b;

            // add product of terms to previous sum and mod by 103
            product[degreeOfTerms] = (product[degreeOfTerms] + polyA[a] * polyB[b]) % modBy;
        }
    }
}

// multPolynomialsParallel determines the intermediary products of the polynomial multiplication problem
__global__ void multPolynomialsParallel(int *polyA, int *polyB, int *product, int polySize, int modBy) {
    int a = blockIdx.x; // all threads in the same block will access the same polyA element
    int b = threadIdx.x; // but all threads will access individual polyB elements
    int myIndex = blockDim.x * blockIdx.x + threadIdx.x; // where to write this thread's product
    product[myIndex] = (polyA[a] * polyB[b]) % modBy;
}

// sumProductsParallel uses prodSize threads, each thread corresponding to a degree, to sum common terms and determine the actual product of polyA and polyB
__global__ void sumProductsParallel(int prodSize, int threadsPerBlock, int *summedProduct, int *products, int numBlocks, int modBy) {
    int responsibleFor = blockIdx.x * blockDim.x + threadIdx.x; // used to check which threads are going to be active during this step

    if (responsibleFor < prodSize) { // e.g. 1 < 7 so thread 1 is going to be in charge of summing the x^1 terms, threads >= prodSize will be inactive for remainder
        for (int blockNum = 0; blockNum < numBlocks; blockNum++) { // loop through blocks
            for (int indexInBlock = 0; indexInBlock < threadsPerBlock; indexInBlock++) { // loop through each index per block
                
                int degreeOfElement = blockNum + indexInBlock; // the degree related to the coefficient stored at each products[] index is equal to the block number + the relative index in the block
                
                if (indexInBlock == 0 && degreeOfElement > responsibleFor) {
                    return; // this thread is done summing its common terms
                } 
                else if (degreeOfElement == responsibleFor) { // if this thread is responsible for the degree we just calculated
                    int spotInProducts = blockNum * blockDim.x + indexInBlock; // get its actual index in products[]
                    summedProduct[responsibleFor] = (summedProduct[responsibleFor] + products[spotInProducts]) % modBy; // and write that value into the final summedProduct[our degree]
                }
            }
        }
    }
}

void checkCUDAError(const char *msg) {
    hipError_t err = hipGetLastError();
    if(hipSuccess != err)
    {
        fprintf(stderr, "CUDA error: %s: %s.\n", msg, hipGetErrorString(err));
        exit(EXIT_FAILURE);
    }
}


	.text
	.file	"q1_swhela2.hip"
	.globl	main                            # -- Begin function 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	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	xorl	%edi, %edi
	callq	time
	movl	%eax, %edi
	callq	srand
	movl	$.L.str, %edi
	xorl	%eax, %eax
	callq	printf
	movq	%rsp, %rbx
	movl	$.L.str.1, %edi
	movq	%rbx, %rsi
	xorl	%eax, %eax
	callq	__isoc23_scanf
	movl	(%rbx), %esi
	movl	$.L.str.2, %edi
	xorl	%eax, %eax
	callq	printf
	movl	(%rbx), %ecx
	cmpl	$11, %ecx
	jl	.LBB0_2
# %bb.1:
	movl	$.L.str.3, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$1, %eax
	jmp	.LBB0_47
.LBB0_2:
	movl	$1, %esi
                                        # kill: def $cl killed $cl killed $ecx
	shll	%cl, %esi
	movl	%esi, (%rsp)
	leal	-1(%rsi), %edx
	movl	$.L.str.4, %edi
                                        # kill: def $esi killed $esi killed $rsi
	xorl	%eax, %eax
	callq	printf
	movslq	(%rsp), %r12
	leaq	(,%r12,4), %r14
	movq	%r14, %rdi
	callq	malloc
	movq	%rax, %rbx
	movq	%r14, %rdi
	callq	malloc
	movq	%rax, %r14
	movl	$.Lstr, %edi
	callq	puts@PLT
	movl	(%rsp), %edx
	movq	%rbx, %rdi
	movq	%r14, %rsi
	callq	_Z14genPolynomialsPiS_i
	movl	$.Lstr.1, %edi
	callq	puts@PLT
	cmpl	$0, (%rsp)
	jle	.LBB0_7
# %bb.3:                                # %.lr.ph.preheader
	xorl	%r15d, %r15d
.LBB0_4:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	(%rbx,%r15,4), %esi
	movl	$.L.str.7, %edi
	movl	%r15d, %edx
	xorl	%eax, %eax
	callq	printf
	movl	(%rsp), %eax
	leal	-1(%rax), %ecx
	cmpq	%rcx, %r15
	je	.LBB0_6
# %bb.5:                                #   in Loop: Header=BB0_4 Depth=1
	movl	$.L.str.8, %edi
	xorl	%eax, %eax
	callq	printf
	movl	(%rsp), %eax
.LBB0_6:                                #   in Loop: Header=BB0_4 Depth=1
	incq	%r15
	cltq
	cmpq	%rax, %r15
	jl	.LBB0_4
.LBB0_7:                                # %._crit_edge
	movl	$.Lstr.2, %edi
	callq	puts@PLT
	cmpl	$0, (%rsp)
	jle	.LBB0_12
# %bb.8:                                # %.lr.ph127.preheader
	xorl	%r15d, %r15d
.LBB0_9:                                # %.lr.ph127
                                        # =>This Inner Loop Header: Depth=1
	movl	(%r14,%r15,4), %esi
	movl	$.L.str.7, %edi
	movl	%r15d, %edx
	xorl	%eax, %eax
	callq	printf
	movl	(%rsp), %eax
	leal	-1(%rax), %ecx
	cmpq	%rcx, %r15
	je	.LBB0_11
# %bb.10:                               #   in Loop: Header=BB0_9 Depth=1
	movl	$.L.str.8, %edi
	xorl	%eax, %eax
	callq	printf
	movl	(%rsp), %eax
.LBB0_11:                               #   in Loop: Header=BB0_9 Depth=1
	incq	%r15
	cltq
	cmpq	%rax, %r15
	jl	.LBB0_9
.LBB0_12:                               # %._crit_edge128
	movl	%r12d, %eax
	movq	%rax, 48(%rsp)                  # 8-byte Spill
	movl	$.Lstr.6, %edi
	callq	puts@PLT
	movslq	(%rsp), %r12
	leal	-1(,%r12,2), %eax
	movl	%eax, 4(%rsp)                   # 4-byte Spill
	movslq	%eax, %r13
	shlq	$2, %r13
	movq	%r13, %rdi
	callq	malloc
	movq	%rax, %r15
	movl	%r12d, %ebp
	imull	%ebp, %ebp
	leaq	(,%rbp,4), %rdi
	callq	malloc
	movq	%r13, 56(%rsp)                  # 8-byte Spill
	movq	%r13, %rdi
	movq	%rax, %r13
	callq	malloc
	movq	%rax, 8(%rsp)                   # 8-byte Spill
	leal	-2(,%r12,2), %eax
	testq	%r12, %r12
	je	.LBB0_14
# %bb.13:                               # %.lr.ph131.preheader
	cmpq	$1, %rbp
	adcq	$0, %rbp
	shlq	$2, %rbp
	movq	%r13, %rdi
	xorl	%esi, %esi
	movq	%rbp, %rdx
	movl	%eax, %ebp
	callq	memset@PLT
	movl	%ebp, %eax
.LBB0_14:                               # %.preheader
	movq	%r13, 32(%rsp)                  # 8-byte Spill
	movl	%eax, %r13d
	cmpl	$0, 4(%rsp)                     # 4-byte Folded Reload
	jle	.LBB0_16
# %bb.15:                               # %.lr.ph133.preheader
	leaq	4(,%r13,4), %rdx
	movq	8(%rsp), %rdi                   # 8-byte Reload
	xorl	%esi, %esi
	callq	memset@PLT
.LBB0_16:                               # %._crit_edge134
	shlq	$2, %r12
	leaq	64(%rsp), %rbp
	movq	%rbp, %rdi
	movq	%r12, %rsi
	callq	hipMalloc
	movslq	(%rsp), %rsi
	shlq	$2, %rsi
	leaq	40(%rsp), %rdi
	callq	hipMalloc
	movl	(%rsp), %esi
	imull	%esi, %esi
	shlq	$2, %rsi
	leaq	16(%rsp), %r12
	movq	%r12, %rdi
	callq	hipMalloc
	movq	(%rbp), %rdi
	movslq	(%rsp), %rdx
	shlq	$2, %rdx
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	40(%rsp), %rax
	movq	(%rax), %rdi
	movslq	(%rsp), %rdx
	shlq	$2, %rdx
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r12), %rdi
	movl	(%rsp), %edx
	imull	%edx, %edx
	shlq	$2, %rdx
	movq	32(%rsp), %r12                  # 8-byte Reload
	movq	%r12, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	48(%rsp), %rbp                  # 8-byte Reload
	btsq	$32, %rbp
	movq	%rbp, %rdi
	movl	$1, %esi
	movq	%rbp, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_18
# %bb.17:
	movq	64(%rsp), %rdi
	movq	40(%rsp), %rsi
	movq	16(%rsp), %rdx
	movl	(%rsp), %ecx
	movl	modBy(%rip), %r8d
	callq	_Z38__device_stub__multPolynomialsParallelPiS_S_ii
.LBB0_18:
	callq	hipDeviceSynchronize
	movl	$.L.str.11, %edi
	callq	_Z14checkCUDAErrorPKc
	movq	16(%rsp), %rsi
	movl	(%rsp), %edx
	imull	%edx, %edx
	shlq	$2, %rdx
	movq	%r12, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	leaq	24(%rsp), %rdi
	movq	56(%rsp), %r12                  # 8-byte Reload
	movq	%r12, %rsi
	callq	hipMalloc
	leaq	24(%rsp), %rax
	movq	(%rax), %rdi
	movq	8(%rsp), %rsi                   # 8-byte Reload
	movq	%r12, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	16(%rsp), %rdi
	movl	(%rsp), %edx
	imull	%edx, %edx
	shlq	$2, %rdx
	movq	32(%rsp), %rsi                  # 8-byte Reload
	movl	$1, %ecx
	callq	hipMemcpy
	movq	%rbp, %rdi
	movl	$1, %esi
	movq	%rbp, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	movl	4(%rsp), %ebp                   # 4-byte Reload
	jne	.LBB0_20
# %bb.19:
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rcx
	movl	modBy(%rip), %r9d
	movl	%ebp, %edi
	movq	48(%rsp), %rsi                  # 8-byte Reload
	movl	%esi, %r8d
	callq	_Z34__device_stub__sumProductsParalleliiPiS_ii
.LBB0_20:
	callq	hipDeviceSynchronize
	movl	$.L.str.11, %edi
	callq	_Z14checkCUDAErrorPKc
	movq	24(%rsp), %rsi
	movq	8(%rsp), %rdi                   # 8-byte Reload
	movq	56(%rsp), %rdx                  # 8-byte Reload
	movl	$2, %ecx
	callq	hipMemcpy
	movl	(%rsp), %r12d
	testl	%ebp, %ebp
	jle	.LBB0_22
# %bb.21:                               # %.lr.ph.preheader.i
	movl	%ebp, %edx
	shlq	$2, %rdx
	movq	%r15, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
.LBB0_22:                               # %.preheader23.i
	testl	%r12d, %r12d
	jle	.LBB0_27
# %bb.23:                               # %.preheader.lr.ph.i
	movl	modBy(%rip), %ecx
	xorl	%esi, %esi
	movq	%r15, %rdi
.LBB0_24:                               # %.lr.ph26.i
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB0_25 Depth 2
	movl	(%rbx,%rsi,4), %r8d
	xorl	%r9d, %r9d
.LBB0_25:                               #   Parent Loop BB0_24 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	(%r14,%r9,4), %eax
	imull	%r8d, %eax
	addl	(%rdi,%r9,4), %eax
	cltd
	idivl	%ecx
	movl	%edx, (%rdi,%r9,4)
	incq	%r9
	cmpq	%r9, %r12
	jne	.LBB0_25
# %bb.26:                               # %._crit_edge.i
                                        #   in Loop: Header=BB0_24 Depth=1
	incq	%rsi
	addq	$4, %rdi
	cmpq	%r12, %rsi
	jne	.LBB0_24
.LBB0_27:                               # %_Z21multPolynomialsSerialPiS_iS_i.exit
	movl	$.Lstr.4, %edi
	callq	puts@PLT
	testl	%ebp, %ebp
	jle	.LBB0_28
# %bb.37:                               # %.lr.ph136.preheader
	leaq	1(%r13), %rbp
	xorl	%r12d, %r12d
.LBB0_38:                               # %.lr.ph136
                                        # =>This Inner Loop Header: Depth=1
	movl	(%r15,%r12,4), %esi
	movl	$.L.str.7, %edi
	movl	%r12d, %edx
	xorl	%eax, %eax
	callq	printf
	cmpq	%r12, %r13
	je	.LBB0_40
# %bb.39:                               #   in Loop: Header=BB0_38 Depth=1
	movl	$.L.str.8, %edi
	xorl	%eax, %eax
	callq	printf
.LBB0_40:                               #   in Loop: Header=BB0_38 Depth=1
	incq	%r12
	cmpq	%r12, %rbp
	jne	.LBB0_38
# %bb.30:                               # %._crit_edge137
	movl	$.Lstr.5, %edi
	callq	puts@PLT
	cmpl	$0, 4(%rsp)                     # 4-byte Folded Reload
	jle	.LBB0_29
# %bb.31:                               # %.lr.ph140.preheader
	leaq	1(%r13), %rbp
	xorl	%r12d, %r12d
	movq	8(%rsp), %rax                   # 8-byte Reload
.LBB0_32:                               # %.lr.ph140
                                        # =>This Inner Loop Header: Depth=1
	movl	(%rax,%r12,4), %esi
	movl	$.L.str.7, %edi
	movl	%r12d, %edx
	xorl	%eax, %eax
	callq	printf
	cmpq	%r12, %r13
	je	.LBB0_34
# %bb.33:                               #   in Loop: Header=BB0_32 Depth=1
	movl	$.L.str.8, %edi
	xorl	%eax, %eax
	callq	printf
.LBB0_34:                               #   in Loop: Header=BB0_32 Depth=1
	incq	%r12
	cmpq	%r12, %rbp
	movq	8(%rsp), %rax                   # 8-byte Reload
	jne	.LBB0_32
# %bb.35:                               # %._crit_edge141
	movl	$.Lstr.6, %edi
	callq	puts@PLT
	movl	4(%rsp), %eax                   # 4-byte Reload
	testl	%eax, %eax
	jle	.LBB0_36
# %bb.41:                               # %.lr.ph145.preheader
	movl	%eax, %r13d
	movb	$1, %bpl
	xorl	%r12d, %r12d
	movq	8(%rsp), %rax                   # 8-byte Reload
.LBB0_42:                               # %.lr.ph145
                                        # =>This Inner Loop Header: Depth=1
	movl	(%r15,%r12,4), %edx
	movl	(%rax,%r12,4), %ecx
	cmpl	%ecx, %edx
	je	.LBB0_44
# %bb.43:                               #   in Loop: Header=BB0_42 Depth=1
	xorl	%ebp, %ebp
	movl	$.L.str.14, %edi
	movl	%r12d, %esi
	xorl	%eax, %eax
	callq	printf
	movq	8(%rsp), %rax                   # 8-byte Reload
.LBB0_44:                               #   in Loop: Header=BB0_42 Depth=1
	incq	%r12
	cmpq	%r12, %r13
	jne	.LBB0_42
# %bb.45:                               # %._crit_edge146
	testb	$1, %bpl
	movl	$.Lstr.7, %eax
	movl	$.Lstr.8, %edi
	cmoveq	%rax, %rdi
	jmp	.LBB0_46
.LBB0_28:                               # %._crit_edge137.thread
	movl	$.Lstr.5, %edi
	callq	puts@PLT
.LBB0_29:                               # %._crit_edge141.thread
	movl	$.Lstr.6, %edi
	callq	puts@PLT
.LBB0_36:
	movl	$.Lstr.8, %edi
.LBB0_46:                               # %._crit_edge146.thread
	callq	puts@PLT
	movq	%rbx, %rdi
	callq	free
	movq	%r14, %rdi
	callq	free
	movq	%r15, %rdi
	callq	free
	movq	32(%rsp), %rdi                  # 8-byte Reload
	callq	free
	movq	8(%rsp), %rdi                   # 8-byte Reload
	callq	free
	movq	64(%rsp), %rdi
	callq	hipFree
	movq	40(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	xorl	%eax, %eax
.LBB0_47:
	addq	$72, %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_end0:
	.size	main, .Lfunc_end0-main
	.cfi_endproc
                                        # -- End function
	.globl	_Z14genPolynomialsPiS_i         # -- Begin function _Z14genPolynomialsPiS_i
	.type	_Z14genPolynomialsPiS_i,@function
_Z14genPolynomialsPiS_i:                # @_Z14genPolynomialsPiS_i
	.cfi_startproc
# %bb.0:
	testl	%edx, %edx
	jle	.LBB1_4
# %bb.1:                                # %.lr.ph.preheader
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%rsi, %rbx
	movq	%rdi, %r14
	movl	%edx, %r15d
	xorl	%r12d, %r12d
.LBB1_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	callq	rand
	cltd
	idivl	modBy(%rip)
	cmpl	$1, %edx
	adcl	$0, %edx
	movl	%edx, (%r14,%r12,4)
	callq	rand
	cltd
	idivl	modBy(%rip)
	cmpl	$1, %edx
	adcl	$0, %edx
	movl	%edx, (%rbx,%r12,4)
	incq	%r12
	cmpq	%r12, %r15
	jne	.LBB1_2
# %bb.3:
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	.cfi_restore %rbx
	.cfi_restore %r12
	.cfi_restore %r14
	.cfi_restore %r15
.LBB1_4:                                # %._crit_edge
	retq
.Lfunc_end1:
	.size	_Z14genPolynomialsPiS_i, .Lfunc_end1-_Z14genPolynomialsPiS_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z38__device_stub__multPolynomialsParallelPiS_S_ii # -- Begin function _Z38__device_stub__multPolynomialsParallelPiS_S_ii
	.type	_Z38__device_stub__multPolynomialsParallelPiS_S_ii,@function
_Z38__device_stub__multPolynomialsParallelPiS_S_ii: # @_Z38__device_stub__multPolynomialsParallelPiS_S_ii
	.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	$128, %rsp
	.cfi_def_cfa_offset 176
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	4(%rsp), %rdx
	movl	%ecx, (%rdx)
	movq	%rsp, %rcx
	movl	%r8d, (%rcx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%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	$_Z23multPolynomialsParallelPiS_S_ii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$144, %rsp
	.cfi_adjust_cfa_offset -144
	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	_Z38__device_stub__multPolynomialsParallelPiS_S_ii, .Lfunc_end2-_Z38__device_stub__multPolynomialsParallelPiS_S_ii
	.cfi_endproc
                                        # -- End function
	.globl	_Z14checkCUDAErrorPKc           # -- Begin function _Z14checkCUDAErrorPKc
	.type	_Z14checkCUDAErrorPKc,@function
_Z14checkCUDAErrorPKc:                  # @_Z14checkCUDAErrorPKc
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	pushq	%rax
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	movq	%rdi, %rbx
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB3_2
# %bb.1:
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.LBB3_2:
	.cfi_def_cfa_offset 32
	movq	stderr(%rip), %r14
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.17, %esi
	movq	%r14, %rdi
	movq	%rbx, %rdx
	movq	%rax, %rcx
	xorl	%eax, %eax
	callq	fprintf
	movl	$1, %edi
	callq	exit
.Lfunc_end3:
	.size	_Z14checkCUDAErrorPKc, .Lfunc_end3-_Z14checkCUDAErrorPKc
	.cfi_endproc
                                        # -- End function
	.globl	_Z34__device_stub__sumProductsParalleliiPiS_ii # -- Begin function _Z34__device_stub__sumProductsParalleliiPiS_ii
	.type	_Z34__device_stub__sumProductsParalleliiPiS_ii,@function
_Z34__device_stub__sumProductsParalleliiPiS_ii: # @_Z34__device_stub__sumProductsParalleliiPiS_ii
	.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	$128, %rsp
	.cfi_def_cfa_offset 176
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	12(%rsp), %rax
	movl	%edi, (%rax)
	leaq	8(%rsp), %rdi
	movl	%esi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	4(%rsp), %rcx
	movl	%r8d, (%rcx)
	movq	%rsp, %r8
	movl	%r9d, (%r8)
	leaq	80(%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	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%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	$_Z19sumProductsParalleliiPiS_ii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$144, %rsp
	.cfi_adjust_cfa_offset -144
	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_end4:
	.size	_Z34__device_stub__sumProductsParalleliiPiS_ii, .Lfunc_end4-_Z34__device_stub__sumProductsParalleliiPiS_ii
	.cfi_endproc
                                        # -- End function
	.globl	_Z21multPolynomialsSerialPiS_iS_i # -- Begin function _Z21multPolynomialsSerialPiS_iS_i
	.type	_Z21multPolynomialsSerialPiS_iS_i,@function
_Z21multPolynomialsSerialPiS_iS_i:      # @_Z21multPolynomialsSerialPiS_iS_i
	.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
	movq	%rcx, %rbx
	movl	%edx, %ebp
	movq	%rsi, %r14
	movq	%rdi, %r15
	testl	%r8d, %r8d
	jle	.LBB5_2
# %bb.1:                                # %.lr.ph.preheader
	movl	%r8d, %edx
	shlq	$2, %rdx
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
.LBB5_2:                                # %.preheader23
	testl	%ebp, %ebp
	jle	.LBB5_7
# %bb.3:                                # %.preheader.lr.ph
	movl	%ebp, %ecx
	xorl	%esi, %esi
.LBB5_4:                                # %.lr.ph26
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB5_5 Depth 2
	xorl	%edi, %edi
.LBB5_5:                                #   Parent Loop BB5_4 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	(%r14,%rdi,4), %eax
	imull	(%r15,%rsi,4), %eax
	addl	(%rbx,%rdi,4), %eax
	cltd
	idivl	modBy(%rip)
	movl	%edx, (%rbx,%rdi,4)
	incq	%rdi
	cmpq	%rdi, %rcx
	jne	.LBB5_5
# %bb.6:                                # %._crit_edge
                                        #   in Loop: Header=BB5_4 Depth=1
	incq	%rsi
	addq	$4, %rbx
	cmpq	%rcx, %rsi
	jne	.LBB5_4
.LBB5_7:                                # %._crit_edge28
	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_end5:
	.size	_Z21multPolynomialsSerialPiS_iS_i, .Lfunc_end5-_Z21multPolynomialsSerialPiS_iS_i
	.cfi_endproc
                                        # -- End function
	.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	.LBB6_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB6_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z23multPolynomialsParallelPiS_S_ii, %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	$_Z19sumProductsParalleliiPiS_ii, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %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_end6:
	.size	__hip_module_ctor, .Lfunc_end6-__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	.LBB7_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
.LBB7_2:
	retq
.Lfunc_end7:
	.size	__hip_module_dtor, .Lfunc_end7-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	modBy,@object                   # @modBy
	.data
	.globl	modBy
	.p2align	2, 0x0
modBy:
	.long	103                             # 0x67
	.size	modBy, 4

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Specify the number of terms in the polynomial by specifying the exponent on base 2, UP TO 10, e.g. enter '3' if you want 2^3 terms (AKA 8 terms) per polynomial: "
	.size	.L.str, 162

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"%d"
	.size	.L.str.1, 3

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"\nYou entered '%d'.\n"
	.size	.L.str.2, 20

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"Invalid entry. The maximum number of terms is 2^10. Please enter a term less than or equal to 10 next time."
	.size	.L.str.3, 108

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"Number of terms per polynomial = %d, hence each polynomial will have degree = %d.\n\n"
	.size	.L.str.4, 84

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"%dx^%d "
	.size	.L.str.7, 8

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

	.type	_Z23multPolynomialsParallelPiS_S_ii,@object # @_Z23multPolynomialsParallelPiS_S_ii
	.section	.rodata,"a",@progbits
	.globl	_Z23multPolynomialsParallelPiS_S_ii
	.p2align	3, 0x0
_Z23multPolynomialsParallelPiS_S_ii:
	.quad	_Z38__device_stub__multPolynomialsParallelPiS_S_ii
	.size	_Z23multPolynomialsParallelPiS_S_ii, 8

	.type	.L.str.11,@object               # @.str.11
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.11:
	.asciz	"kernel invocation"
	.size	.L.str.11, 18

	.type	_Z19sumProductsParalleliiPiS_ii,@object # @_Z19sumProductsParalleliiPiS_ii
	.section	.rodata,"a",@progbits
	.globl	_Z19sumProductsParalleliiPiS_ii
	.p2align	3, 0x0
_Z19sumProductsParalleliiPiS_ii:
	.quad	_Z34__device_stub__sumProductsParalleliiPiS_ii
	.size	_Z19sumProductsParalleliiPiS_ii, 8

	.type	.L.str.14,@object               # @.str.14
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.14:
	.asciz	"Coefficients at degree %d are not equivalent: serial!=parallel (%d!=%d)\n"
	.size	.L.str.14, 73

	.type	.L.str.17,@object               # @.str.17
.L.str.17:
	.asciz	"CUDA error: %s: %s.\n"
	.size	.L.str.17, 21

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z23multPolynomialsParallelPiS_S_ii"
	.size	.L__unnamed_1, 36

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z19sumProductsParalleliiPiS_ii"
	.size	.L__unnamed_2, 32

	.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	"Generating polynomials...\n"
	.size	.Lstr, 27

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"polyA:"
	.size	.Lstr.1, 7

	.type	.Lstr.2,@object                 # @str.2
.Lstr.2:
	.asciz	"\n\npolyB:"
	.size	.Lstr.2, 9

	.type	.Lstr.4,@object                 # @str.4
.Lstr.4:
	.asciz	"Serial result:"
	.size	.Lstr.4, 15

	.type	.Lstr.5,@object                 # @str.5
.Lstr.5:
	.asciz	"\n\nParallel result:"
	.size	.Lstr.5, 19

	.type	.Lstr.6,@object                 # @str.6
.Lstr.6:
	.asciz	"\n"
	.size	.Lstr.6, 2

	.type	.Lstr.7,@object                 # @str.7
.Lstr.7:
	.asciz	"Looks like there were some discrepancies. Verification failed.\n"
	.size	.Lstr.7, 64

	.type	.Lstr.8,@object                 # @str.8
.Lstr.8:
	.asciz	"Verification successful. The serial and parallel polynomial multiplications produced the same result!\n"
	.size	.Lstr.8, 103

	.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 _Z38__device_stub__multPolynomialsParallelPiS_S_ii
	.addrsig_sym _Z34__device_stub__sumProductsParalleliiPiS_ii
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z23multPolynomialsParallelPiS_S_ii
	.addrsig_sym _Z19sumProductsParalleliiPiS_ii
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z23multPolynomialsParallelPiS_S_ii ; -- Begin function _Z23multPolynomialsParallelPiS_S_ii
	.globl	_Z23multPolynomialsParallelPiS_S_ii
	.p2align	8
	.type	_Z23multPolynomialsParallelPiS_S_ii,@function
_Z23multPolynomialsParallelPiS_S_ii:    ; @_Z23multPolynomialsParallelPiS_S_ii
; %bb.0:
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[8:9], s[0:1], 0x10
	v_lshlrev_b32_e32 v1, 2, v0
	s_mov_b32 s2, s15
	s_ashr_i32 s3, s15, 31
	s_waitcnt lgkmcnt(0)
	global_load_b32 v1, v1, s[6:7]
	s_clause 0x1
	s_load_b32 s6, s[0:1], 0x2c
	s_load_b32 s7, s[0:1], 0x1c
	s_lshl_b64 s[0:1], s[2:3], 2
	s_waitcnt lgkmcnt(0)
	s_and_b32 s3, s6, 0xffff
	s_add_u32 s0, s4, s0
	s_addc_u32 s1, s5, s1
	s_ashr_i32 s4, s7, 31
	s_load_b32 s0, s[0:1], 0x0
	s_add_i32 s7, s7, s4
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_xor_b32 s4, s7, s4
	v_cvt_f32_u32_e32 v2, s4
	s_sub_i32 s1, 0, s4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v2, v2
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v2, 0x4f7ffffe, v2
	v_cvt_u32_f32_e32 v2, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v3, s1, v2
	v_mul_hi_u32 v3, v2, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v2, v2, v3
	s_waitcnt vmcnt(0) lgkmcnt(0)
	v_mul_lo_u32 v1, v1, s0
	v_ashrrev_i32_e32 v4, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v1, v1, v4
	v_xor_b32_e32 v1, v1, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v2, v1, v2
	v_mul_lo_u32 v2, v2, s4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v1, v1, v2
	v_subrev_nc_u32_e32 v2, s4, v1
	v_cmp_le_u32_e32 vcc_lo, s4, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v3, v1, v2, vcc_lo
	v_mad_u64_u32 v[1:2], null, s2, s3, v[0:1]
	v_subrev_nc_u32_e32 v0, s4, v3
	v_cmp_le_u32_e32 vcc_lo, s4, v3
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_ashrrev_i32_e32 v2, 31, v1
	v_cndmask_b32_e32 v0, v3, v0, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_xor_b32_e32 v3, v0, v4
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v2, v3, v4
	v_add_co_u32 v0, vcc_lo, s8, v0
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v1, vcc_lo, s9, v1, vcc_lo
	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 _Z23multPolynomialsParallelPiS_S_ii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.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 5
		.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	_Z23multPolynomialsParallelPiS_S_ii, .Lfunc_end0-_Z23multPolynomialsParallelPiS_S_ii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 332
; NumSgprs: 18
; NumVgprs: 5
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 5
; 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	_Z19sumProductsParalleliiPiS_ii ; -- Begin function _Z19sumProductsParalleliiPiS_ii
	.globl	_Z19sumProductsParalleliiPiS_ii
	.p2align	8
	.type	_Z19sumProductsParalleliiPiS_ii,@function
_Z19sumProductsParalleliiPiS_ii:        ; @_Z19sumProductsParalleliiPiS_ii
; %bb.0:
	s_clause 0x2
	s_load_b32 s2, s[0:1], 0x2c
	s_load_b64 s[4:5], s[0:1], 0x0
	s_load_b64 s[6:7], s[0:1], 0x18
	s_waitcnt lgkmcnt(0)
	s_and_b32 s8, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s8, v[0:1]
	s_cmp_gt_i32 s6, 0
	s_cselect_b32 s2, -1, 0
	v_cmp_gt_i32_e32 vcc_lo, s4, v1
	s_mov_b32 s4, 0
	s_and_b32 s2, vcc_lo, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s3, s2
	s_cbranch_execz .LBB1_13
; %bb.1:                                ; %.preheader.lr.ph
	s_cmp_gt_i32 s5, 0
	v_ashrrev_i32_e32 v2, 31, v1
	s_cselect_b32 s9, -1, 0
	s_ashr_i32 s2, s7, 31
	v_sub_nc_u32_e32 v4, 0, v1
	s_add_i32 s3, s7, s2
	v_lshlrev_b64 v[2:3], 2, v[1:2]
	s_xor_b32 s7, s3, s2
	s_load_b128 s[0:3], s[0:1], 0x8
	v_cvt_f32_u32_e32 v0, s7
	s_sub_i32 s10, 0, s7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_4) | instid1(VALU_DEP_2)
	v_rcp_iflag_f32_e32 v0, v0
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v0, 0x4f7ffffe, v0
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s0, v2
	v_cvt_u32_f32_e32 v5, v0
	v_add_co_ci_u32_e32 v3, vcc_lo, s1, v3, vcc_lo
	s_mov_b32 s0, 0
	s_mov_b32 s1, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v0, s10, v5
	v_mul_hi_u32 v6, v5, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_dual_mov_b32 v0, 0 :: v_dual_add_nc_u32 v5, v5, v6
.LBB1_2:                                ; %.preheader
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB1_4 Depth 2
	s_and_not1_b32 vcc_lo, exec_lo, s9
	s_mov_b32 s10, -1
	s_cbranch_vccnz .LBB1_10
; %bb.3:                                ; %.lr.ph
                                        ;   in Loop: Header=BB1_2 Depth=1
	v_add_nc_u32_e32 v6, s1, v4
	s_mov_b32 s11, 0
	s_mov_b32 s12, 0
                                        ; implicit-def: $sgpr10
                                        ; implicit-def: $sgpr14
                                        ; implicit-def: $sgpr13
.LBB1_4:                                ;   Parent Loop BB1_2 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_i32 s15, s1, s12
	s_cmp_lg_u32 s12, 0
	v_cmp_le_i32_e32 vcc_lo, s15, v1
	s_cselect_b32 s15, -1, 0
	s_or_b32 s13, s13, exec_lo
	s_or_b32 s14, s14, exec_lo
	s_or_b32 s16, s15, vcc_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s15, s16
	s_cbranch_execz .LBB1_8
; %bb.5:                                ;   in Loop: Header=BB1_4 Depth=2
	v_add_nc_u32_e32 v7, s12, v6
	s_mov_b32 s16, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v7
	s_cbranch_execz .LBB1_7
; %bb.6:                                ;   in Loop: Header=BB1_4 Depth=2
	s_add_i32 s18, s0, s12
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_ashr_i32 s19, s18, 31
	s_lshl_b64 s[18:19], s[18:19], 2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s18, s2, s18
	s_addc_u32 s19, s3, s19
	global_load_b32 v7, v[2:3], off
	global_load_b32 v8, v0, s[18:19]
	s_waitcnt vmcnt(0)
	v_add_nc_u32_e32 v7, v8, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v8, 31, v7
	v_add_nc_u32_e32 v7, v7, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v7, v7, v8
	v_mul_hi_u32 v9, v7, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v9, v9, s7
	v_sub_nc_u32_e32 v7, v7, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v9, s7, v7
	v_cmp_le_u32_e32 vcc_lo, s7, v7
	v_cndmask_b32_e32 v7, v7, v9, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v9, s7, v7
	v_cmp_le_u32_e32 vcc_lo, s7, v7
	v_cndmask_b32_e32 v7, v7, v9, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v7, v7, v8
	v_sub_nc_u32_e32 v7, v7, v8
	global_store_b32 v[2:3], v7, off
.LBB1_7:                                ;   in Loop: Header=BB1_4 Depth=2
	s_or_b32 exec_lo, exec_lo, s16
	s_add_i32 s12, s12, 1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_eq_u32 s5, s12
	s_cselect_b32 s16, -1, 0
	s_and_not1_b32 s14, s14, exec_lo
	s_and_b32 s16, s16, exec_lo
	s_and_not1_b32 s13, s13, exec_lo
	s_or_b32 s14, s14, s16
.LBB1_8:                                ; %Flow
                                        ;   in Loop: Header=BB1_4 Depth=2
	s_or_b32 exec_lo, exec_lo, s15
	s_xor_b32 s15, s13, -1
	s_and_b32 s16, exec_lo, s14
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	s_or_b32 s11, s16, s11
	s_and_not1_b32 s10, s10, exec_lo
	s_and_b32 s15, s15, exec_lo
	s_or_b32 s10, s10, s15
	s_and_not1_b32 exec_lo, exec_lo, s11
	s_cbranch_execnz .LBB1_4
; %bb.9:                                ; %loop.exit.guard
                                        ;   in Loop: Header=BB1_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s11
.LBB1_10:                               ; %Flow66
                                        ;   in Loop: Header=BB1_2 Depth=1
	s_mov_b32 s11, -1
	s_and_saveexec_b32 s12, s10
; %bb.11:                               ; %.critedge34
                                        ;   in Loop: Header=BB1_2 Depth=1
	s_add_i32 s1, s1, 1
	s_add_i32 s0, s0, s8
	s_cmp_eq_u32 s1, s6
	s_cselect_b32 s10, -1, 0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_not1_b32 s11, s10, exec_lo
; %bb.12:                               ; %Flow67
                                        ;   in Loop: Header=BB1_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s12
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s10, exec_lo, s11
	s_or_b32 s4, s10, s4
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s4
	s_cbranch_execnz .LBB1_2
.LBB1_13:                               ; %.critedge36
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z19sumProductsParalleliiPiS_ii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.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 10
		.amdhsa_next_free_sgpr 20
		.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	_Z19sumProductsParalleliiPiS_ii, .Lfunc_end1-_Z19sumProductsParalleliiPiS_ii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 604
; NumSgprs: 22
; NumVgprs: 10
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 22
; NumVGPRsForWavesPerEU: 10
; 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
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         28
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z23multPolynomialsParallelPiS_S_ii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z23multPolynomialsParallelPiS_S_ii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     5
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .offset:         4
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         28
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z19sumProductsParalleliiPiS_ii
    .private_segment_fixed_size: 0
    .sgpr_count:     22
    .sgpr_spill_count: 0
    .symbol:         _Z19sumProductsParalleliiPiS_ii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     10
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
