#include "gpuVector4D.cu"

#include <iosfwd>

class gpuMatrix4x4 {

  public:


    // The default constructor.
    __device__ __host__ gpuMatrix4x4(void) { }

    // Constructor for row major form data.
    // Transposes to the internal column major form.
    // REQUIRES: data should be of size 16.
    __device__ __host__ gpuMatrix4x4(float * data)
    {
      for( int i = 0; i < 4; i++ )
        for( int j = 0; j < 4; j++ )
        {
          // Transpostion happens within the () query.
          (*this)(i,j) = data[i*4 + j];
        }

    }


    /**
     * Sets all elements to val.
     */
    __device__ void zero(float val = 0.0);

    /**
     * Returns the determinant of A.
     */
    __device__ float det( void ) const;

    /**
     * Returns the Frobenius norm of A.
     */
    __device__ float norm( void ) const;

    /**
     * Returns a fresh 4x4 identity matrix.
     */
    __device__ static gpuMatrix4x4 identity( void );

    // No Cross products for 4 by 4 matrix.

    /**
     * Returns the ith column.
     */
    __device__     gpuVector4D& column( int i );
    __device__ const gpuVector4D& column( int i ) const;

    /**
     * Returns the transpose of A.
     */
    __device__ gpuMatrix4x4 T( void ) const;

    /**
     * Returns the inverse of A.
     */
    __device__ gpuMatrix4x4 inv( void ) const;

    // accesses element (i,j) of A using 0-based indexing
    // where (i, j) is (row, column).
    __device__ __host__      float& operator()( int i, int j );
    __device__ __host__ const float& operator()( int i, int j ) const;

    // accesses the ith column of A
    __device__ __host__    gpuVector4D& operator[]( int i );
    __device__ __host__ const gpuVector4D& operator[]( int i ) const;

    // increments by B
    __device__ void operator+=( const gpuMatrix4x4& B );

    // returns -A
    __device__ gpuMatrix4x4 operator-( void ) const;

    // returns A-B
    __device__ gpuMatrix4x4 operator-( const gpuMatrix4x4& B ) const;

    // returns c*A
    __device__ gpuMatrix4x4 operator*( float c ) const;

    // returns A*B
    __device__ gpuMatrix4x4 operator*( const gpuMatrix4x4& B ) const;

    __device__ // returns A*x
      gpuVector4D operator*( const gpuVector4D& x ) const;

    // divides each element by x
    __device__ void operator/=( float x );

  protected:

    // 4 by 4 matrices are represented by an array of 4 column vectors.
    gpuVector4D entries[4];

}; // class Matrix3x3

// returns the outer product of u and v.
__device__ gpuMatrix4x4 outer( const gpuVector4D& u, const gpuVector4D& v );

// returns c*A
__device__ gpuMatrix4x4 operator*( float c, const gpuMatrix4x4& A );

__device__
void gpuMatrix4x4::zero( float val ) {
  // sets all elements to val
  entries[0] =
    entries[1] =
    entries[2] =
    entries[3] = gpuVector4D( val, val, val, val );
}

__device__
float gpuMatrix4x4::det( void ) const {
  const gpuMatrix4x4& A( *this );

  return
    A(0,3)*A(1,2)*A(2,1)*A(3,0) - A(0,2)*A(1,3)*A(2,1)*A(3,0) -
    A(0,3)*A(1,1)*A(2,2)*A(3,0) + A(0,1)*A(1,3)*A(2,2)*A(3,0) +
    A(0,2)*A(1,1)*A(2,3)*A(3,0) - A(0,1)*A(1,2)*A(2,3)*A(3,0) -
    A(0,3)*A(1,2)*A(2,0)*A(3,1) + A(0,2)*A(1,3)*A(2,0)*A(3,1) +
    A(0,3)*A(1,0)*A(2,2)*A(3,1) - A(0,0)*A(1,3)*A(2,2)*A(3,1) -
    A(0,2)*A(1,0)*A(2,3)*A(3,1) + A(0,0)*A(1,2)*A(2,3)*A(3,1) +
    A(0,3)*A(1,1)*A(2,0)*A(3,2) - A(0,1)*A(1,3)*A(2,0)*A(3,2) -
    A(0,3)*A(1,0)*A(2,1)*A(3,2) + A(0,0)*A(1,3)*A(2,1)*A(3,2) +
    A(0,1)*A(1,0)*A(2,3)*A(3,2) - A(0,0)*A(1,1)*A(2,3)*A(3,2) -
    A(0,2)*A(1,1)*A(2,0)*A(3,3) + A(0,1)*A(1,2)*A(2,0)*A(3,3) +
    A(0,2)*A(1,0)*A(2,1)*A(3,3) - A(0,0)*A(1,2)*A(2,1)*A(3,3) -
    A(0,1)*A(1,0)*A(2,2)*A(3,3) + A(0,0)*A(1,1)*A(2,2)*A(3,3);

}

__device__
float gpuMatrix4x4::norm( void ) const {
  return sqrt( entries[0].norm2() +
      entries[1].norm2() +
      entries[2].norm2() +
      entries[3].norm2());
}

__device__
gpuMatrix4x4 gpuMatrix4x4::operator-( void ) const {

  // returns -A (Negation).
  const gpuMatrix4x4& A( *this );
  gpuMatrix4x4 B;

  B(0,0) = -A(0,0); B(0,1) = -A(0,1); B(0,2) = -A(0,2); B(0,3) = -A(0,3);
  B(1,0) = -A(1,0); B(1,1) = -A(1,1); B(1,2) = -A(1,2); B(1,3) = -A(1,3);
  B(2,0) = -A(2,0); B(2,1) = -A(2,1); B(2,2) = -A(2,2); B(2,3) = -A(2,3);
  B(3,0) = -A(3,0); B(3,1) = -A(3,1); B(3,2) = -A(3,2); B(3,3) = -A(3,3);

  return B;
}

__device__
void gpuMatrix4x4::operator+=( const gpuMatrix4x4& B ) {

  gpuMatrix4x4& A( *this );
  float* Aij = (float*) &A;
  const float* Bij = (const float*) &B;

  // Add the 16 contigous vector packed float values.
  *Aij++ += *Bij++;//0
  *Aij++ += *Bij++;
  *Aij++ += *Bij++;
  *Aij++ += *Bij++;
  *Aij++ += *Bij++;//4
  *Aij++ += *Bij++;
  *Aij++ += *Bij++;
  *Aij++ += *Bij++;
  *Aij++ += *Bij++;//8
  *Aij++ += *Bij++;
  *Aij++ += *Bij++;
  *Aij++ += *Bij++;
  *Aij++ += *Bij++;//12
  *Aij++ += *Bij++;
  *Aij++ += *Bij++;
  *Aij++ += *Bij++;//15
  //16.

}

__device__
gpuMatrix4x4 gpuMatrix4x4::operator-( const gpuMatrix4x4& B ) const {
  const gpuMatrix4x4& A( *this );
  gpuMatrix4x4 C;

  for( int i = 0; i < 4; i++ )
    for( int j = 0; j < 4; j++ )
    {
      C(i,j) = A(i,j) - B(i,j);
    }

  return C;
}

__device__
gpuMatrix4x4 gpuMatrix4x4::operator*( float c ) const {
  const gpuMatrix4x4& A( *this );
  gpuMatrix4x4 B;

  for( int i = 0; i < 4; i++ )
    for( int j = 0; j < 4; j++ )
    {
      B(i,j) = c*A(i,j);
    }

  return B;
}

// Returns c*A.
__device__
gpuMatrix4x4 operator*( float c, const gpuMatrix4x4& A ) {

  gpuMatrix4x4 cA;
  const float* Aij = (const float*) &A;
  float* cAij = (float*) &cA;

  *cAij++ = c * (*Aij++);//0
  *cAij++ = c * (*Aij++);
  *cAij++ = c * (*Aij++);
  *cAij++ = c * (*Aij++);
  *cAij++ = c * (*Aij++);//4
  *cAij++ = c * (*Aij++);
  *cAij++ = c * (*Aij++);
  *cAij++ = c * (*Aij++);
  *cAij++ = c * (*Aij++);//8
  *cAij++ = c * (*Aij++);
  *cAij++ = c * (*Aij++);
  *cAij++ = c * (*Aij++);
  *cAij++ = c * (*Aij++);//12
  *cAij++ = c * (*Aij++);
  *cAij++ = c * (*Aij++);
  *cAij++ = c * (*Aij++);//15
  //16
  return cA;
}

// Tradiational Grade School Multiplication. N^3
__device__
gpuMatrix4x4 gpuMatrix4x4::operator*( const gpuMatrix4x4& B ) const {
  const gpuMatrix4x4& A( *this );
  gpuMatrix4x4 C;

  for( int i = 0; i < 4; i++ )
    for( int j = 0; j < 4; j++ )
    {
      C(i,j) = 0.;

      for( int k = 0; k < 4; k++ )
      {
        C(i,j) += A(i,k)*B(k,j);
      }
    }

  return C;
}


__device__
gpuVector4D gpuMatrix4x4::operator*( const gpuVector4D& x ) const {
  return x[0]*entries[0] + // Add up products for each matrix column.
    x[1]*entries[1] +
    x[2]*entries[2] +
    x[3]*entries[3];
}

// Naive Transposition.
__device__
gpuMatrix4x4 gpuMatrix4x4::T( void ) const {
  const gpuMatrix4x4& A( *this );
  gpuMatrix4x4 B;

  for( int i = 0; i < 4; i++ )
    for( int j = 0; j < 4; j++ )
    {
      B(i,j) = A(j,i);
    }

  return B;
}

__device__
gpuMatrix4x4 gpuMatrix4x4::inv( void ) const {
  const gpuMatrix4x4& A( *this );
  gpuMatrix4x4 B;

  // Hardcoded in Fully Symbolic computation.

  B(0,0) = A(1,2)*A(2,3)*A(3,1) - A(1,3)*A(2,2)*A(3,1) + A(1,3)*A(2,1)*A(3,2) - A(1,1)*A(2,3)*A(3,2) - A(1,2)*A(2,1)*A(3,3) + A(1,1)*A(2,2)*A(3,3);
  B(0,1) = A(0,3)*A(2,2)*A(3,1) - A(0,2)*A(2,3)*A(3,1) - A(0,3)*A(2,1)*A(3,2) + A(0,1)*A(2,3)*A(3,2) + A(0,2)*A(2,1)*A(3,3) - A(0,1)*A(2,2)*A(3,3);
  B(0,2) = A(0,2)*A(1,3)*A(3,1) - A(0,3)*A(1,2)*A(3,1) + A(0,3)*A(1,1)*A(3,2) - A(0,1)*A(1,3)*A(3,2) - A(0,2)*A(1,1)*A(3,3) + A(0,1)*A(1,2)*A(3,3);
  B(0,3) = A(0,3)*A(1,2)*A(2,1) - A(0,2)*A(1,3)*A(2,1) - A(0,3)*A(1,1)*A(2,2) + A(0,1)*A(1,3)*A(2,2) + A(0,2)*A(1,1)*A(2,3) - A(0,1)*A(1,2)*A(2,3);
  B(1,0) = A(1,3)*A(2,2)*A(3,0) - A(1,2)*A(2,3)*A(3,0) - A(1,3)*A(2,0)*A(3,2) + A(1,0)*A(2,3)*A(3,2) + A(1,2)*A(2,0)*A(3,3) - A(1,0)*A(2,2)*A(3,3);
  B(1,1) = A(0,2)*A(2,3)*A(3,0) - A(0,3)*A(2,2)*A(3,0) + A(0,3)*A(2,0)*A(3,2) - A(0,0)*A(2,3)*A(3,2) - A(0,2)*A(2,0)*A(3,3) + A(0,0)*A(2,2)*A(3,3);
  B(1,2) = A(0,3)*A(1,2)*A(3,0) - A(0,2)*A(1,3)*A(3,0) - A(0,3)*A(1,0)*A(3,2) + A(0,0)*A(1,3)*A(3,2) + A(0,2)*A(1,0)*A(3,3) - A(0,0)*A(1,2)*A(3,3);
  B(1,3) = A(0,2)*A(1,3)*A(2,0) - A(0,3)*A(1,2)*A(2,0) + A(0,3)*A(1,0)*A(2,2) - A(0,0)*A(1,3)*A(2,2) - A(0,2)*A(1,0)*A(2,3) + A(0,0)*A(1,2)*A(2,3);
  B(2,0) = A(1,1)*A(2,3)*A(3,0) - A(1,3)*A(2,1)*A(3,0) + A(1,3)*A(2,0)*A(3,1) - A(1,0)*A(2,3)*A(3,1) - A(1,1)*A(2,0)*A(3,3) + A(1,0)*A(2,1)*A(3,3);
  B(2,1) = A(0,3)*A(2,1)*A(3,0) - A(0,1)*A(2,3)*A(3,0) - A(0,3)*A(2,0)*A(3,1) + A(0,0)*A(2,3)*A(3,1) + A(0,1)*A(2,0)*A(3,3) - A(0,0)*A(2,1)*A(3,3);
  B(2,2) = A(0,1)*A(1,3)*A(3,0) - A(0,3)*A(1,1)*A(3,0) + A(0,3)*A(1,0)*A(3,1) - A(0,0)*A(1,3)*A(3,1) - A(0,1)*A(1,0)*A(3,3) + A(0,0)*A(1,1)*A(3,3);
  B(2,3) = A(0,3)*A(1,1)*A(2,0) - A(0,1)*A(1,3)*A(2,0) - A(0,3)*A(1,0)*A(2,1) + A(0,0)*A(1,3)*A(2,1) + A(0,1)*A(1,0)*A(2,3) - A(0,0)*A(1,1)*A(2,3);
  B(3,0) = A(1,2)*A(2,1)*A(3,0) - A(1,1)*A(2,2)*A(3,0) - A(1,2)*A(2,0)*A(3,1) + A(1,0)*A(2,2)*A(3,1) + A(1,1)*A(2,0)*A(3,2) - A(1,0)*A(2,1)*A(3,2);
  B(3,1) = A(0,1)*A(2,2)*A(3,0) - A(0,2)*A(2,1)*A(3,0) + A(0,2)*A(2,0)*A(3,1) - A(0,0)*A(2,2)*A(3,1) - A(0,1)*A(2,0)*A(3,2) + A(0,0)*A(2,1)*A(3,2);
  B(3,2) = A(0,2)*A(1,1)*A(3,0) - A(0,1)*A(1,2)*A(3,0) - A(0,2)*A(1,0)*A(3,1) + A(0,0)*A(1,2)*A(3,1) + A(0,1)*A(1,0)*A(3,2) - A(0,0)*A(1,1)*A(3,2);
  B(3,3) = A(0,1)*A(1,2)*A(2,0) - A(0,2)*A(1,1)*A(2,0) + A(0,2)*A(1,0)*A(2,1) - A(0,0)*A(1,2)*A(2,1) - A(0,1)*A(1,0)*A(2,2) + A(0,0)*A(1,1)*A(2,2);

  // Invertable iff the determinant is not equal to zero.
  B /= det();

  return B;
}

__device__
void gpuMatrix4x4::operator/=( float x ) {
  gpuMatrix4x4& A( *this );
  float rx = 1./x;

  for( int i = 0; i < 4; i++ )
    for( int j = 0; j < 4; j++ )
    {
      A( i, j ) *= rx;
    }
}

__device__
gpuMatrix4x4 gpuMatrix4x4::identity( void ) {
  gpuMatrix4x4 B;

  B(0,0) = 1.; B(0,1) = 0.; B(0,2) = 0.; B(0,3) = 0.;
  B(1,0) = 0.; B(1,1) = 1.; B(1,2) = 0.; B(1,3) = 0.;
  B(2,0) = 0.; B(2,1) = 0.; B(2,2) = 1.; B(2,3) = 0.;
  B(3,0) = 0.; B(3,1) = 0.; B(3,2) = 0.; B(3,3) = 1.;

  return B;
}

__device__
gpuMatrix4x4 outer( const gpuVector4D& u, const gpuVector4D& v ) {
  gpuMatrix4x4 B;

  // Opposite of an inner product.
  for( int i = 0; i < 4; i++ )
    for( int j = 0; j < 4; j++ )
    {
      B( i, j ) = u[i]*v[j];
    }

  return B;
}

	.file	"tmpxft_002d9a89_00000000-6_gpuMatrix4x4.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2136:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2136:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.align 2
	.globl	_ZN12gpuMatrix4x44zeroEf
	.type	_ZN12gpuMatrix4x44zeroEf, @function
_ZN12gpuMatrix4x44zeroEf:
.LFB2119:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2119:
	.size	_ZN12gpuMatrix4x44zeroEf, .-_ZN12gpuMatrix4x44zeroEf
	.align 2
	.globl	_ZNK12gpuMatrix4x43detEv
	.type	_ZNK12gpuMatrix4x43detEv, @function
_ZNK12gpuMatrix4x43detEv:
.LFB2120:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2120:
	.size	_ZNK12gpuMatrix4x43detEv, .-_ZNK12gpuMatrix4x43detEv
	.align 2
	.globl	_ZNK12gpuMatrix4x44normEv
	.type	_ZNK12gpuMatrix4x44normEv, @function
_ZNK12gpuMatrix4x44normEv:
.LFB2279:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_ZNK12gpuMatrix4x43detEv
	.cfi_endproc
.LFE2279:
	.size	_ZNK12gpuMatrix4x44normEv, .-_ZNK12gpuMatrix4x44normEv
	.align 2
	.globl	_ZNK12gpuMatrix4x4ngEv
	.type	_ZNK12gpuMatrix4x4ngEv, @function
_ZNK12gpuMatrix4x4ngEv:
.LFB2283:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2283:
	.size	_ZNK12gpuMatrix4x4ngEv, .-_ZNK12gpuMatrix4x4ngEv
	.align 2
	.globl	_ZN12gpuMatrix4x4pLERKS_
	.type	_ZN12gpuMatrix4x4pLERKS_, @function
_ZN12gpuMatrix4x4pLERKS_:
.LFB2123:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2123:
	.size	_ZN12gpuMatrix4x4pLERKS_, .-_ZN12gpuMatrix4x4pLERKS_
	.align 2
	.globl	_ZNK12gpuMatrix4x4miERKS_
	.type	_ZNK12gpuMatrix4x4miERKS_, @function
_ZNK12gpuMatrix4x4miERKS_:
.LFB2124:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2124:
	.size	_ZNK12gpuMatrix4x4miERKS_, .-_ZNK12gpuMatrix4x4miERKS_
	.align 2
	.globl	_ZNK12gpuMatrix4x4mlEf
	.type	_ZNK12gpuMatrix4x4mlEf, @function
_ZNK12gpuMatrix4x4mlEf:
.LFB2125:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2125:
	.size	_ZNK12gpuMatrix4x4mlEf, .-_ZNK12gpuMatrix4x4mlEf
	.globl	_ZmlfRK12gpuMatrix4x4
	.type	_ZmlfRK12gpuMatrix4x4, @function
_ZmlfRK12gpuMatrix4x4:
.LFB2126:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2126:
	.size	_ZmlfRK12gpuMatrix4x4, .-_ZmlfRK12gpuMatrix4x4
	.align 2
	.globl	_ZNK12gpuMatrix4x4mlERKS_
	.type	_ZNK12gpuMatrix4x4mlERKS_, @function
_ZNK12gpuMatrix4x4mlERKS_:
.LFB2277:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2277:
	.size	_ZNK12gpuMatrix4x4mlERKS_, .-_ZNK12gpuMatrix4x4mlERKS_
	.align 2
	.globl	_ZNK12gpuMatrix4x4mlERK11gpuVector4D
	.type	_ZNK12gpuMatrix4x4mlERK11gpuVector4D, @function
_ZNK12gpuMatrix4x4mlERK11gpuVector4D:
.LFB2128:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2128:
	.size	_ZNK12gpuMatrix4x4mlERK11gpuVector4D, .-_ZNK12gpuMatrix4x4mlERK11gpuVector4D
	.align 2
	.globl	_ZNK12gpuMatrix4x41TEv
	.type	_ZNK12gpuMatrix4x41TEv, @function
_ZNK12gpuMatrix4x41TEv:
.LFB2129:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2129:
	.size	_ZNK12gpuMatrix4x41TEv, .-_ZNK12gpuMatrix4x41TEv
	.align 2
	.globl	_ZNK12gpuMatrix4x43invEv
	.type	_ZNK12gpuMatrix4x43invEv, @function
_ZNK12gpuMatrix4x43invEv:
.LFB2281:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2281:
	.size	_ZNK12gpuMatrix4x43invEv, .-_ZNK12gpuMatrix4x43invEv
	.align 2
	.globl	_ZN12gpuMatrix4x4dVEf
	.type	_ZN12gpuMatrix4x4dVEf, @function
_ZN12gpuMatrix4x4dVEf:
.LFB2275:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_ZN12gpuMatrix4x44zeroEf
	.cfi_endproc
.LFE2275:
	.size	_ZN12gpuMatrix4x4dVEf, .-_ZN12gpuMatrix4x4dVEf
	.align 2
	.globl	_ZN12gpuMatrix4x48identityEv
	.type	_ZN12gpuMatrix4x48identityEv, @function
_ZN12gpuMatrix4x48identityEv:
.LFB2132:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2132:
	.size	_ZN12gpuMatrix4x48identityEv, .-_ZN12gpuMatrix4x48identityEv
	.globl	_Z5outerRK11gpuVector4DS1_
	.type	_Z5outerRK11gpuVector4DS1_, @function
_Z5outerRK11gpuVector4DS1_:
.LFB2133:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2133:
	.size	_Z5outerRK11gpuVector4DS1_, .-_Z5outerRK11gpuVector4DS1_
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2159:
	.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
.LFE2159:
	.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:


#include <hip/hip_runtime.h>
#include "gpuVector4D.cu"

#include <iosfwd>

class gpuMatrix4x4 {

  public:


    // The default constructor.
    __device__ __host__ gpuMatrix4x4(void) { }

    // Constructor for row major form data.
    // Transposes to the internal column major form.
    // REQUIRES: data should be of size 16.
    __device__ __host__ gpuMatrix4x4(float * data)
    {
      for( int i = 0; i < 4; i++ )
        for( int j = 0; j < 4; j++ )
        {
          // Transpostion happens within the () query.
          (*this)(i,j) = data[i*4 + j];
        }

    }


    /**
     * Sets all elements to val.
     */
    __device__ void zero(float val = 0.0);

    /**
     * Returns the determinant of A.
     */
    __device__ float det( void ) const;

    /**
     * Returns the Frobenius norm of A.
     */
    __device__ float norm( void ) const;

    /**
     * Returns a fresh 4x4 identity matrix.
     */
    __device__ static gpuMatrix4x4 identity( void );

    // No Cross products for 4 by 4 matrix.

    /**
     * Returns the ith column.
     */
    __device__     gpuVector4D& column( int i );
    __device__ const gpuVector4D& column( int i ) const;

    /**
     * Returns the transpose of A.
     */
    __device__ gpuMatrix4x4 T( void ) const;

    /**
     * Returns the inverse of A.
     */
    __device__ gpuMatrix4x4 inv( void ) const;

    // accesses element (i,j) of A using 0-based indexing
    // where (i, j) is (row, column).
    __device__ __host__      float& operator()( int i, int j );
    __device__ __host__ const float& operator()( int i, int j ) const;

    // accesses the ith column of A
    __device__ __host__    gpuVector4D& operator[]( int i );
    __device__ __host__ const gpuVector4D& operator[]( int i ) const;

    // increments by B
    __device__ void operator+=( const gpuMatrix4x4& B );

    // returns -A
    __device__ gpuMatrix4x4 operator-( void ) const;

    // returns A-B
    __device__ gpuMatrix4x4 operator-( const gpuMatrix4x4& B ) const;

    // returns c*A
    __device__ gpuMatrix4x4 operator*( float c ) const;

    // returns A*B
    __device__ gpuMatrix4x4 operator*( const gpuMatrix4x4& B ) const;

    __device__ // returns A*x
      gpuVector4D operator*( const gpuVector4D& x ) const;

    // divides each element by x
    __device__ void operator/=( float x );

  protected:

    // 4 by 4 matrices are represented by an array of 4 column vectors.
    gpuVector4D entries[4];

}; // class Matrix3x3

// returns the outer product of u and v.
__device__ gpuMatrix4x4 outer( const gpuVector4D& u, const gpuVector4D& v );

// returns c*A
__device__ gpuMatrix4x4 operator*( float c, const gpuMatrix4x4& A );

__device__
void gpuMatrix4x4::zero( float val ) {
  // sets all elements to val
  entries[0] =
    entries[1] =
    entries[2] =
    entries[3] = gpuVector4D( val, val, val, val );
}

__device__
float gpuMatrix4x4::det( void ) const {
  const gpuMatrix4x4& A( *this );

  return
    A(0,3)*A(1,2)*A(2,1)*A(3,0) - A(0,2)*A(1,3)*A(2,1)*A(3,0) -
    A(0,3)*A(1,1)*A(2,2)*A(3,0) + A(0,1)*A(1,3)*A(2,2)*A(3,0) +
    A(0,2)*A(1,1)*A(2,3)*A(3,0) - A(0,1)*A(1,2)*A(2,3)*A(3,0) -
    A(0,3)*A(1,2)*A(2,0)*A(3,1) + A(0,2)*A(1,3)*A(2,0)*A(3,1) +
    A(0,3)*A(1,0)*A(2,2)*A(3,1) - A(0,0)*A(1,3)*A(2,2)*A(3,1) -
    A(0,2)*A(1,0)*A(2,3)*A(3,1) + A(0,0)*A(1,2)*A(2,3)*A(3,1) +
    A(0,3)*A(1,1)*A(2,0)*A(3,2) - A(0,1)*A(1,3)*A(2,0)*A(3,2) -
    A(0,3)*A(1,0)*A(2,1)*A(3,2) + A(0,0)*A(1,3)*A(2,1)*A(3,2) +
    A(0,1)*A(1,0)*A(2,3)*A(3,2) - A(0,0)*A(1,1)*A(2,3)*A(3,2) -
    A(0,2)*A(1,1)*A(2,0)*A(3,3) + A(0,1)*A(1,2)*A(2,0)*A(3,3) +
    A(0,2)*A(1,0)*A(2,1)*A(3,3) - A(0,0)*A(1,2)*A(2,1)*A(3,3) -
    A(0,1)*A(1,0)*A(2,2)*A(3,3) + A(0,0)*A(1,1)*A(2,2)*A(3,3);

}

__device__
float gpuMatrix4x4::norm( void ) const {
  return sqrt( entries[0].norm2() +
      entries[1].norm2() +
      entries[2].norm2() +
      entries[3].norm2());
}

__device__
gpuMatrix4x4 gpuMatrix4x4::operator-( void ) const {

  // returns -A (Negation).
  const gpuMatrix4x4& A( *this );
  gpuMatrix4x4 B;

  B(0,0) = -A(0,0); B(0,1) = -A(0,1); B(0,2) = -A(0,2); B(0,3) = -A(0,3);
  B(1,0) = -A(1,0); B(1,1) = -A(1,1); B(1,2) = -A(1,2); B(1,3) = -A(1,3);
  B(2,0) = -A(2,0); B(2,1) = -A(2,1); B(2,2) = -A(2,2); B(2,3) = -A(2,3);
  B(3,0) = -A(3,0); B(3,1) = -A(3,1); B(3,2) = -A(3,2); B(3,3) = -A(3,3);

  return B;
}

__device__
void gpuMatrix4x4::operator+=( const gpuMatrix4x4& B ) {

  gpuMatrix4x4& A( *this );
  float* Aij = (float*) &A;
  const float* Bij = (const float*) &B;

  // Add the 16 contigous vector packed float values.
  *Aij++ += *Bij++;//0
  *Aij++ += *Bij++;
  *Aij++ += *Bij++;
  *Aij++ += *Bij++;
  *Aij++ += *Bij++;//4
  *Aij++ += *Bij++;
  *Aij++ += *Bij++;
  *Aij++ += *Bij++;
  *Aij++ += *Bij++;//8
  *Aij++ += *Bij++;
  *Aij++ += *Bij++;
  *Aij++ += *Bij++;
  *Aij++ += *Bij++;//12
  *Aij++ += *Bij++;
  *Aij++ += *Bij++;
  *Aij++ += *Bij++;//15
  //16.

}

__device__
gpuMatrix4x4 gpuMatrix4x4::operator-( const gpuMatrix4x4& B ) const {
  const gpuMatrix4x4& A( *this );
  gpuMatrix4x4 C;

  for( int i = 0; i < 4; i++ )
    for( int j = 0; j < 4; j++ )
    {
      C(i,j) = A(i,j) - B(i,j);
    }

  return C;
}

__device__
gpuMatrix4x4 gpuMatrix4x4::operator*( float c ) const {
  const gpuMatrix4x4& A( *this );
  gpuMatrix4x4 B;

  for( int i = 0; i < 4; i++ )
    for( int j = 0; j < 4; j++ )
    {
      B(i,j) = c*A(i,j);
    }

  return B;
}

// Returns c*A.
__device__
gpuMatrix4x4 operator*( float c, const gpuMatrix4x4& A ) {

  gpuMatrix4x4 cA;
  const float* Aij = (const float*) &A;
  float* cAij = (float*) &cA;

  *cAij++ = c * (*Aij++);//0
  *cAij++ = c * (*Aij++);
  *cAij++ = c * (*Aij++);
  *cAij++ = c * (*Aij++);
  *cAij++ = c * (*Aij++);//4
  *cAij++ = c * (*Aij++);
  *cAij++ = c * (*Aij++);
  *cAij++ = c * (*Aij++);
  *cAij++ = c * (*Aij++);//8
  *cAij++ = c * (*Aij++);
  *cAij++ = c * (*Aij++);
  *cAij++ = c * (*Aij++);
  *cAij++ = c * (*Aij++);//12
  *cAij++ = c * (*Aij++);
  *cAij++ = c * (*Aij++);
  *cAij++ = c * (*Aij++);//15
  //16
  return cA;
}

// Tradiational Grade School Multiplication. N^3
__device__
gpuMatrix4x4 gpuMatrix4x4::operator*( const gpuMatrix4x4& B ) const {
  const gpuMatrix4x4& A( *this );
  gpuMatrix4x4 C;

  for( int i = 0; i < 4; i++ )
    for( int j = 0; j < 4; j++ )
    {
      C(i,j) = 0.;

      for( int k = 0; k < 4; k++ )
      {
        C(i,j) += A(i,k)*B(k,j);
      }
    }

  return C;
}


__device__
gpuVector4D gpuMatrix4x4::operator*( const gpuVector4D& x ) const {
  return x[0]*entries[0] + // Add up products for each matrix column.
    x[1]*entries[1] +
    x[2]*entries[2] +
    x[3]*entries[3];
}

// Naive Transposition.
__device__
gpuMatrix4x4 gpuMatrix4x4::T( void ) const {
  const gpuMatrix4x4& A( *this );
  gpuMatrix4x4 B;

  for( int i = 0; i < 4; i++ )
    for( int j = 0; j < 4; j++ )
    {
      B(i,j) = A(j,i);
    }

  return B;
}

__device__
gpuMatrix4x4 gpuMatrix4x4::inv( void ) const {
  const gpuMatrix4x4& A( *this );
  gpuMatrix4x4 B;

  // Hardcoded in Fully Symbolic computation.

  B(0,0) = A(1,2)*A(2,3)*A(3,1) - A(1,3)*A(2,2)*A(3,1) + A(1,3)*A(2,1)*A(3,2) - A(1,1)*A(2,3)*A(3,2) - A(1,2)*A(2,1)*A(3,3) + A(1,1)*A(2,2)*A(3,3);
  B(0,1) = A(0,3)*A(2,2)*A(3,1) - A(0,2)*A(2,3)*A(3,1) - A(0,3)*A(2,1)*A(3,2) + A(0,1)*A(2,3)*A(3,2) + A(0,2)*A(2,1)*A(3,3) - A(0,1)*A(2,2)*A(3,3);
  B(0,2) = A(0,2)*A(1,3)*A(3,1) - A(0,3)*A(1,2)*A(3,1) + A(0,3)*A(1,1)*A(3,2) - A(0,1)*A(1,3)*A(3,2) - A(0,2)*A(1,1)*A(3,3) + A(0,1)*A(1,2)*A(3,3);
  B(0,3) = A(0,3)*A(1,2)*A(2,1) - A(0,2)*A(1,3)*A(2,1) - A(0,3)*A(1,1)*A(2,2) + A(0,1)*A(1,3)*A(2,2) + A(0,2)*A(1,1)*A(2,3) - A(0,1)*A(1,2)*A(2,3);
  B(1,0) = A(1,3)*A(2,2)*A(3,0) - A(1,2)*A(2,3)*A(3,0) - A(1,3)*A(2,0)*A(3,2) + A(1,0)*A(2,3)*A(3,2) + A(1,2)*A(2,0)*A(3,3) - A(1,0)*A(2,2)*A(3,3);
  B(1,1) = A(0,2)*A(2,3)*A(3,0) - A(0,3)*A(2,2)*A(3,0) + A(0,3)*A(2,0)*A(3,2) - A(0,0)*A(2,3)*A(3,2) - A(0,2)*A(2,0)*A(3,3) + A(0,0)*A(2,2)*A(3,3);
  B(1,2) = A(0,3)*A(1,2)*A(3,0) - A(0,2)*A(1,3)*A(3,0) - A(0,3)*A(1,0)*A(3,2) + A(0,0)*A(1,3)*A(3,2) + A(0,2)*A(1,0)*A(3,3) - A(0,0)*A(1,2)*A(3,3);
  B(1,3) = A(0,2)*A(1,3)*A(2,0) - A(0,3)*A(1,2)*A(2,0) + A(0,3)*A(1,0)*A(2,2) - A(0,0)*A(1,3)*A(2,2) - A(0,2)*A(1,0)*A(2,3) + A(0,0)*A(1,2)*A(2,3);
  B(2,0) = A(1,1)*A(2,3)*A(3,0) - A(1,3)*A(2,1)*A(3,0) + A(1,3)*A(2,0)*A(3,1) - A(1,0)*A(2,3)*A(3,1) - A(1,1)*A(2,0)*A(3,3) + A(1,0)*A(2,1)*A(3,3);
  B(2,1) = A(0,3)*A(2,1)*A(3,0) - A(0,1)*A(2,3)*A(3,0) - A(0,3)*A(2,0)*A(3,1) + A(0,0)*A(2,3)*A(3,1) + A(0,1)*A(2,0)*A(3,3) - A(0,0)*A(2,1)*A(3,3);
  B(2,2) = A(0,1)*A(1,3)*A(3,0) - A(0,3)*A(1,1)*A(3,0) + A(0,3)*A(1,0)*A(3,1) - A(0,0)*A(1,3)*A(3,1) - A(0,1)*A(1,0)*A(3,3) + A(0,0)*A(1,1)*A(3,3);
  B(2,3) = A(0,3)*A(1,1)*A(2,0) - A(0,1)*A(1,3)*A(2,0) - A(0,3)*A(1,0)*A(2,1) + A(0,0)*A(1,3)*A(2,1) + A(0,1)*A(1,0)*A(2,3) - A(0,0)*A(1,1)*A(2,3);
  B(3,0) = A(1,2)*A(2,1)*A(3,0) - A(1,1)*A(2,2)*A(3,0) - A(1,2)*A(2,0)*A(3,1) + A(1,0)*A(2,2)*A(3,1) + A(1,1)*A(2,0)*A(3,2) - A(1,0)*A(2,1)*A(3,2);
  B(3,1) = A(0,1)*A(2,2)*A(3,0) - A(0,2)*A(2,1)*A(3,0) + A(0,2)*A(2,0)*A(3,1) - A(0,0)*A(2,2)*A(3,1) - A(0,1)*A(2,0)*A(3,2) + A(0,0)*A(2,1)*A(3,2);
  B(3,2) = A(0,2)*A(1,1)*A(3,0) - A(0,1)*A(1,2)*A(3,0) - A(0,2)*A(1,0)*A(3,1) + A(0,0)*A(1,2)*A(3,1) + A(0,1)*A(1,0)*A(3,2) - A(0,0)*A(1,1)*A(3,2);
  B(3,3) = A(0,1)*A(1,2)*A(2,0) - A(0,2)*A(1,1)*A(2,0) + A(0,2)*A(1,0)*A(2,1) - A(0,0)*A(1,2)*A(2,1) - A(0,1)*A(1,0)*A(2,2) + A(0,0)*A(1,1)*A(2,2);

  // Invertable iff the determinant is not equal to zero.
  B /= det();

  return B;
}

__device__
void gpuMatrix4x4::operator/=( float x ) {
  gpuMatrix4x4& A( *this );
  float rx = 1./x;

  for( int i = 0; i < 4; i++ )
    for( int j = 0; j < 4; j++ )
    {
      A( i, j ) *= rx;
    }
}

__device__
gpuMatrix4x4 gpuMatrix4x4::identity( void ) {
  gpuMatrix4x4 B;

  B(0,0) = 1.; B(0,1) = 0.; B(0,2) = 0.; B(0,3) = 0.;
  B(1,0) = 0.; B(1,1) = 1.; B(1,2) = 0.; B(1,3) = 0.;
  B(2,0) = 0.; B(2,1) = 0.; B(2,2) = 1.; B(2,3) = 0.;
  B(3,0) = 0.; B(3,1) = 0.; B(3,2) = 0.; B(3,3) = 1.;

  return B;
}

__device__
gpuMatrix4x4 outer( const gpuVector4D& u, const gpuVector4D& v ) {
  gpuMatrix4x4 B;

  // Opposite of an inner product.
  for( int i = 0; i < 4; i++ )
    for( int j = 0; j < 4; j++ )
    {
      B( i, j ) = u[i]*v[j];
    }

  return B;
}

	.text
	.file	"gpuMatrix4x4.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
