#pragma once

// #include "fixnum/word_fixnum.cu"

template < typename fixnum> 
class cubic_ext_element {
public:
    typedef fixnum modnum;
    modnum a0;
    modnum a1;
    modnum a2;

    __device__ cubic_ext_element() { }
};

template < typename fixnum, typename monty >
class cubic_ext {
public:
    typedef fixnum modnum;
    monty mod;
    modnum alpha; 

    typedef cubic_ext_element<fixnum> ext_element;

    __device__ cubic_ext(fixnum modulus, fixnum _alpha) : mod(modulus), alpha(_alpha) {
        modnum t;
        mod.to_modnum(t, alpha);
        alpha = t;
    }

    __device__ void to_modnum(ext_element &z) {
        modnum t0, t1, t2;
        mod.to_modnum(t0, z.a0);
        mod.to_modnum(t1, z.a1);
        mod.to_modnum(t2, z.a2);

        z.a0 = t0; z.a1 = t1; z.a2 = t2;
    }

    __device__ void from_modnum(ext_element &z) {
        fixnum t0, t1, t2;
        mod.from_modnum(t0, z.a0);
        mod.from_modnum(t1, z.a1);
        mod.from_modnum(t2, z.a2);

        z.a0 = t0; z.a1 = t1; z.a2 = t2;
    }

    __device__ void add(ext_element &z, ext_element &x, ext_element &y) {
        modnum t0, t1, t2;
        mod.add(t0, x.a0, y.a0);
        mod.add(t1, x.a1, y.a1);
        mod.add(t2, x.a2, y.a2);

        z.a0 = t0; z.a1 = t1; z.a2 = t2;
    }

    __device__ void mul(ext_element &z, ext_element x, ext_element y) {
        modnum t0, t1, t2;
        modnum a0b0, a0b1, a0b2;
        modnum a1b0, a1b1, a1b2;
        modnum a2b0, a2b1, a2b2;

        mod.mul(a0b0, x.a0, y.a0);
        mod.mul(a0b1, x.a0, y.a1);
        mod.mul(a0b2, x.a0, y.a2);
        mod.mul(a1b0, x.a1, y.a0);
        mod.mul(a1b1, x.a1, y.a1);
        mod.mul(a1b2, x.a1, y.a2);
        mod.mul(a2b0, x.a2, y.a0);
        mod.mul(a2b1, x.a2, y.a1);
        mod.mul(a2b2, x.a2, y.a2);

        // c0 = a0*b0 + 11*(a1*b2 + a2*b1)
        mod.add(t0, a1b2, a2b1);
        mod.mul(t1, alpha, t0);
        mod.add(z.a0, a0b0, t1);

        // c1 = a0*b1 + a1*b0 + alpha*a2*b2
        mod.mul(t0, alpha, a2b2);
        mod.add(t1, t0, a0b1);
        mod.add(z.a1, t1, a1b0);

        // c2 = a0*b2 + a1*b1 + a2*b0
        mod.add(t0, a0b2, a1b1);
        mod.add(z.a2, t0, a2b0);
    }
};

	.file	"tmpxft_002ed131_00000000-6_cubic_ext.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2016:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2016:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2039:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	movq	%rax, %rdi
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2039:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

#pragma once
#include <hip/hip_runtime.h>


// #include "fixnum/word_fixnum.cu"

template < typename fixnum> 
class cubic_ext_element {
public:
    typedef fixnum modnum;
    modnum a0;
    modnum a1;
    modnum a2;

    __device__ cubic_ext_element() { }
};

template < typename fixnum, typename monty >
class cubic_ext {
public:
    typedef fixnum modnum;
    monty mod;
    modnum alpha; 

    typedef cubic_ext_element<fixnum> ext_element;

    __device__ cubic_ext(fixnum modulus, fixnum _alpha) : mod(modulus), alpha(_alpha) {
        modnum t;
        mod.to_modnum(t, alpha);
        alpha = t;
    }

    __device__ void to_modnum(ext_element &z) {
        modnum t0, t1, t2;
        mod.to_modnum(t0, z.a0);
        mod.to_modnum(t1, z.a1);
        mod.to_modnum(t2, z.a2);

        z.a0 = t0; z.a1 = t1; z.a2 = t2;
    }

    __device__ void from_modnum(ext_element &z) {
        fixnum t0, t1, t2;
        mod.from_modnum(t0, z.a0);
        mod.from_modnum(t1, z.a1);
        mod.from_modnum(t2, z.a2);

        z.a0 = t0; z.a1 = t1; z.a2 = t2;
    }

    __device__ void add(ext_element &z, ext_element &x, ext_element &y) {
        modnum t0, t1, t2;
        mod.add(t0, x.a0, y.a0);
        mod.add(t1, x.a1, y.a1);
        mod.add(t2, x.a2, y.a2);

        z.a0 = t0; z.a1 = t1; z.a2 = t2;
    }

    __device__ void mul(ext_element &z, ext_element x, ext_element y) {
        modnum t0, t1, t2;
        modnum a0b0, a0b1, a0b2;
        modnum a1b0, a1b1, a1b2;
        modnum a2b0, a2b1, a2b2;

        mod.mul(a0b0, x.a0, y.a0);
        mod.mul(a0b1, x.a0, y.a1);
        mod.mul(a0b2, x.a0, y.a2);
        mod.mul(a1b0, x.a1, y.a0);
        mod.mul(a1b1, x.a1, y.a1);
        mod.mul(a1b2, x.a1, y.a2);
        mod.mul(a2b0, x.a2, y.a0);
        mod.mul(a2b1, x.a2, y.a1);
        mod.mul(a2b2, x.a2, y.a2);

        // c0 = a0*b0 + 11*(a1*b2 + a2*b1)
        mod.add(t0, a1b2, a2b1);
        mod.mul(t1, alpha, t0);
        mod.add(z.a0, a0b0, t1);

        // c1 = a0*b1 + a1*b0 + alpha*a2*b2
        mod.mul(t0, alpha, a2b2);
        mod.add(t1, t0, a0b1);
        mod.add(z.a1, t1, a1b0);

        // c2 = a0*b2 + a1*b1 + a2*b0
        mod.add(t0, a0b2, a1b1);
        mod.add(z.a2, t0, a2b0);
    }
};

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

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

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

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

	.end_amdgpu_metadata
