#include "Utils.cuh"
#include "cp4Conv2d.cuh"
#include "cp4Conv2dBackwardData.cuh"
#include "cp4Conv2dBackwardFilter.cuh"
#include "cp4Conv2dForward.cuh"
#include <iostream>
#include <stdlib.h>

using namespace std;

/*******************************************************************************
 * Unified memory Tensorized call of Convolution in GPU
 ******************************************************************************/
Tensor CP::Conv2dForward(Tensor const Input,
                         Tensor const FT,
                         Tensor const FC,
                         Tensor const FY,
                         Tensor const FX,
                         unsigned     pad) {

  tensor_shape params;
  params.N    = Input.shape[0];
  params.C    = Input.shape[1];
  params.H    = Input.shape[2];
  params.W    = Input.shape[3];
  params.pad  = pad;
  params.Rank = FT.shape[1];
  params.T    = FT.shape[0];
  params.Y    = FY.shape[0];
  params.X    = FX.shape[0];

  Tensor Out{ params.N, params.T, params.H, params.W };

  cp4_conv2d_forward_gpu(params,
                         Input.m_data,
                         FT.m_data,
                         FC.m_data,
                         FY.m_data,
                         FX.m_data,
                         Out.m_data);

  return Out;
}

/*******************************************************************************
 * Unified memory Tensorized call of Convolution Backward Data in GPU
 ******************************************************************************/
/* Tensor CP::Conv2dBackwardData(Tensor const Upstream, */
/*                               Tensor const FT, */
/*                               Tensor const FC, */
/*                               Tensor const FY, */
/*                               Tensor const FX, */
/*                               unsigned     pad) { */

/*   tensor_shape params; */
/*   params.N    = Upstream.shape[0]; */
/*   params.T    = Upstream.shape[1]; */
/*   params.H    = Upstream.shape[2]; */
/*   params.W    = Upstream.shape[3]; */
/*   params.pad  = pad; */
/*   params.Rank = FT.shape[1]; */
/*   params.C    = FC.shape[0]; */
/*   params.Y    = FY.shape[0]; */
/*   params.X    = FX.shape[0]; */

/*   Tensor Out{ params.N, params.C, params.H, params.W }; */

/*   cp4_conv2d_backward_data_gpu(params, */
/*                                Upstream.m_data, */
/*                                FT.m_data, */
/*                                FC.m_data, */
/*                                FY.m_data, */
/*                                FX.m_data, */
/*                                Out.m_data); */

/*   return Out; */
/* } */


/*******************************************************************************
 * Unified memory Tensorized call of Convolution Backward Filter in GPU
 ******************************************************************************/
/* Tensor CP::Conv2dBackwardFilter(Tensor const dLdO, */
/*                                 Tensor const In, */
/*                                 Tensor const FT, */
/*                                 Tensor const FC, */
/*                                 Tensor const FY, */
/*                                 Tensor const FX, */
/*                                 unsigned     pad) { */

/*   tensor_shape s; */
/*   s.N    = dLdO.shape[0]; */
/*   s.T    = dLdO.shape[1]; */
/*   s.H    = dLdO.shape[2]; */
/*   s.W    = dLdO.shape[3]; */
/*   s.pad  = pad; */
/*   s.Rank = FT.shape[1]; */
/*   s.C    = FC.shape[0]; */
/*   s.Y    = FY.shape[0]; */
/*   s.X    = FX.shape[0]; */

/*   Tensor dFT{ s.T, s.Rank }; */
/*   Tensor dFC{ s.C, s.Rank }; */
/*   Tensor dFY{ s.Y, s.Rank }; */
/*   Tensor dFX{ s.X, s.Rank }; */
/*   /1* Tensor dFF{ s.T, s.C, s.Y, s.X }; *1/ */


/*   /1* cp4_conv2d_backward_filter_full_gpu( *1/ */
/*   /1*     s, dFF.m_data, In.m_data, dLdO.m_data); *1/ */

/*   /1* return dFF; *1/ */

/*   cp4_conv2d_backward_filter_t_gpu(s, */
/*                                    dFT.m_data, */
/*                                    In.m_data, */
/*                                    dLdO.m_data, */
/*                                    FC.m_data, */
/*                                    FY.m_data, */
/*                                    FX.m_data); */

/*   /1* cout << dFT.m_data[0] << endl; *1/ */
/*   /1* cout << dFC.m_data[0] << endl; *1/ */
/*   /1* cout << dFY.m_data[0] << endl; *1/ */
/*   /1* cout << dFX.m_data[0] << endl; *1/ */

/*   cp4_conv2d_backward_filter_c_gpu(s, */
/*                                    dFC.m_data, */
/*                                    In.m_data, */
/*                                    dLdO.m_data, */
/*                                    FT.m_data, */
/*                                    FY.m_data, */
/*                                    FX.m_data); */

/*   cp4_conv2d_backward_filter_y_gpu(s, */
/*                                    dFY.m_data, */
/*                                    In.m_data, */
/*                                    dLdO.m_data, */
/*                                    FT.m_data, */
/*                                    FC.m_data, */
/*                                    FX.m_data); */


/*   cp4_conv2d_backward_filter_x_gpu(s, */
/*                                    dFX.m_data, */
/*                                    In.m_data, */
/*                                    dLdO.m_data, */
/*                                    FT.m_data, */
/*                                    FC.m_data, */
/*                                    FY.m_data); */

/*   return cp4recom(dFT, dFC, dFY, dFX); */
/* } */

/*******************************************************************************
 * Run_convolution operation with a profile count loop
 ******************************************************************************/
std::pair<float, unsigned>
CP::run_convolution(tensor_shape p, unsigned PROFCOUNT) {

  float* In;
  float* Out;
  float* FT;
  float* FC;
  float* FX;
  float* FY;

  cudaMalloc(&In, p.N * p.C * p.H * p.W * sizeof(float));
  cudaMalloc(&FT, p.T * p.Rank * sizeof(float));
  cudaMalloc(&FC, p.C * p.Rank * sizeof(float));
  cudaMalloc(&FY, p.Y * p.Rank * sizeof(float));
  cudaMalloc(&FX, p.X * p.Rank * sizeof(float));
  cudaMalloc(&Out, p.N * p.T * p.H * p.W * sizeof(float));


  unsigned mem =
      sizeof(float) * ((p.N * p.T * p.H * p.W) + (p.N * p.C * p.H * p.W)
                       + p.Rank * (p.T + p.C + p.X + p.Y));

  float us = cp4_conv2d_forward_gpu(p, In, FT, FC, FY, FX, Out, PROFCOUNT);

  cudaFree(In);
  cudaFree(FT);
  cudaFree(FC);
  cudaFree(FY);
  cudaFree(FX);
  cudaFree(Out);

  return std::make_pair(us,mem);
}


/*******************************************************************************
 * Main function. call 1 instance of kernel execution
 ******************************************************************************/
int main(int argc, char** argv) {

  unsigned N    = 5;
  unsigned C    = 32;
  unsigned H    = 1024;
  unsigned W    = 1024;
  unsigned pad  = 1;
  unsigned T    = 32;
  unsigned Y    = 3;
  unsigned X    = 3;
  unsigned Rank = 8;

  if (argc != 11) {
    cerr << "Using Default shape" << endl;
    cudaSetDevice(0);
  } else {
    N    = atoi(argv[1]);
    C    = atoi(argv[2]);
    H    = atoi(argv[3]);
    W    = atoi(argv[4]);
    pad  = atoi(argv[5]);
    T    = atoi(argv[6]);
    Y    = atoi(argv[7]);
    X    = atoi(argv[8]);
    Rank = atoi(argv[9]);
    cudaSetDevice(atoi(argv[10]));
  }

  tensor_shape params;
  params.N    = N;
  params.C    = C;
  params.H    = H;
  params.W    = W;
  params.pad  = pad;
  params.Rank = Rank;
  params.T    = T;
  params.Y    = Y;
  params.X    = X;

  CP::run_convolution(params, 1);
}

	.file	"tmpxft_002a1fbe_00000000-6_cp4Conv2d.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4692:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4692:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_ZN2CP13Conv2dForwardE6TensorS0_S0_S0_S0_j
	.type	_ZN2CP13Conv2dForwardE6TensorS0_S0_S0_S0_j, @function
_ZN2CP13Conv2dForwardE6TensorS0_S0_S0_S0_j:
.LFB4660:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4660
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register 6
	pushq	%r15
	pushq	%r14
	pushq	%r13
	.cfi_offset 15, -24
	.cfi_offset 14, -32
	.cfi_offset 13, -40
	movq	%rdx, %r13
	pushq	%r12
	.cfi_offset 12, -48
	movq	%rsi, %r12
	pushq	%rbx
	.cfi_offset 3, -56
	movq	%rdi, %rbx
	subq	$104, %rsp
	movq	%rcx, -144(%rbp)
	movq	%r9, -120(%rbp)
	movq	%fs:40, %rdx
	movq	%rdx, -56(%rbp)
	movq	8(%rsi), %rdx
	movq	%r8, -136(%rbp)
	movl	4(%rdx), %ecx
	movl	(%rdx), %r9d
	movl	12(%rdx), %r15d
	movl	%ecx, -88(%rbp)
	movl	8(%rdx), %ecx
	movl	16(%rbp), %edx
	movl	%edx, -96(%rbp)
	movq	8(%r13), %rdx
	movl	4(%rdx), %esi
	movl	(%rdx), %r14d
	movq	8(%r8), %rdx
	movl	%esi, -76(%rbp)
	leaq	-72(%rbp), %rsi
	movl	(%rdx), %edx
	movq	-120(%rbp), %rax
	movl	%r9d, -128(%rbp)
	movl	%ecx, -124(%rbp)
	movl	%edx, -84(%rbp)
	movq	8(%rax), %rdx
	movl	(%rdx), %edx
	movl	%r9d, -72(%rbp)
	movl	%ecx, -64(%rbp)
	movl	%edx, -80(%rbp)
	movl	$4, %edx
	movl	%r14d, -68(%rbp)
	movl	%r15d, -60(%rbp)
.LEHB0:
	call	_ZN6TensorC1ESt16initializer_listIjE@PLT
.LEHE0:
	movl	-128(%rbp), %r9d
	movl	-124(%rbp), %ecx
	movl	%r15d, -100(%rbp)
	movq	-136(%rbp), %r8
	movq	(%r12), %r11
	movl	%r14d, -92(%rbp)
	leaq	-108(%rbp), %rsi
	movq	-144(%rbp), %rax
	movl	%r9d, -108(%rbp)
	movq	(%r8), %r10
	movl	%ecx, -104(%rbp)
	movl	$9, %ecx
	movq	(%rax), %rdx
	movq	0(%r13), %rax
	pushq	$1
	subq	$40, %rsp
	movq	%rsp, %rdi
	rep movsl
	movq	%r10, %rcx
	movq	%rax, %rsi
	movq	-120(%rbp), %rdi
	movq	(%rbx), %r9
	movq	(%rdi), %r8
	movq	%r11, %rdi
.LEHB1:
	.cfi_escape 0x2e,0x30
	call	_Z22cp4_conv2d_forward_gpu12tensor_shapePKfS1_S1_S1_S1_Pfj@PLT
.LEHE1:
	addq	$48, %rsp
	movq	-56(%rbp), %rax
	subq	%fs:40, %rax
	je	.L5
	jmp	.L7
.L6:
	endbr64
	movq	%rax, %r14
.L3:
	movq	%rbx, %rdi
	call	_ZN6TensorD1Ev@PLT
	movq	-56(%rbp), %rax
	subq	%fs:40, %rax
	jne	.L7
	movq	%r14, %rdi
.LEHB2:
	call	_Unwind_Resume@PLT
.LEHE2:
.L7:
	call	__stack_chk_fail@PLT
.L5:
	leaq	-40(%rbp), %rsp
	movq	%rbx, %rax
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa 7, 8
	ret
	.cfi_endproc
.LFE4660:
	.globl	__gxx_personality_v0
	.section	.gcc_except_table,"a",@progbits
.LLSDA4660:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4660-.LLSDACSB4660
.LLSDACSB4660:
	.uleb128 .LEHB0-.LFB4660
	.uleb128 .LEHE0-.LEHB0
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB1-.LFB4660
	.uleb128 .LEHE1-.LEHB1
	.uleb128 .L6-.LFB4660
	.uleb128 0
	.uleb128 .LEHB2-.LFB4660
	.uleb128 .LEHE2-.LEHB2
	.uleb128 0
	.uleb128 0
.LLSDACSE4660:
	.text
	.size	_ZN2CP13Conv2dForwardE6TensorS0_S0_S0_S0_j, .-_ZN2CP13Conv2dForwardE6TensorS0_S0_S0_S0_j
	.globl	_ZN2CP15run_convolutionE12tensor_shapej
	.type	_ZN2CP15run_convolutionE12tensor_shapej, @function
_ZN2CP15run_convolutionE12tensor_shapej:
.LFB4678:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$88, %rsp
	.cfi_def_cfa_offset 144
	movl	148(%rsp), %r13d
	movl	164(%rsp), %ebp
	movl	%edi, 12(%rsp)
	leaq	24(%rsp), %rdi
	imull	144(%rsp), %r13d
	movl	176(%rsp), %r12d
	imull	152(%rsp), %r13d
	movl	160(%rsp), %r15d
	movl	168(%rsp), %r14d
	movl	172(%rsp), %ebx
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	movl	%r13d, %esi
	imull	%ebp, %esi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	movl	%r12d, %esi
	leaq	40(%rsp), %rdi
	imull	%r15d, %esi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	movl	%r12d, %esi
	leaq	48(%rsp), %rdi
	imull	%ebp, %esi
	addl	%r15d, %ebp
	salq	$2, %rsi
	call	cudaMalloc@PLT
	movl	%r12d, %esi
	leaq	64(%rsp), %rdi
	imull	%r14d, %esi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	movl	%r12d, %esi
	leaq	56(%rsp), %rdi
	imull	%ebx, %esi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	movl	%r13d, %esi
	leaq	32(%rsp), %rdi
	imull	%r15d, %esi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	movl	12(%rsp), %edx
	movl	$9, %ecx
	pushq	%rdx
	.cfi_def_cfa_offset 152
	subq	$40, %rsp
	.cfi_def_cfa_offset 192
	movq	%rsp, %rdi
	leaq	192(%rsp), %rsi
	rep movsl
	movq	80(%rsp), %r9
	movq	104(%rsp), %r8
	movq	112(%rsp), %rcx
	movq	96(%rsp), %rdx
	movq	88(%rsp), %rsi
	movq	72(%rsp), %rdi
	call	_Z22cp4_conv2d_forward_gpu12tensor_shapePKfS1_S1_S1_S1_Pfj@PLT
	movss	%xmm0, 60(%rsp)
	movq	72(%rsp), %rdi
	addq	$48, %rsp
	.cfi_def_cfa_offset 144
	call	cudaFree@PLT
	movq	40(%rsp), %rdi
	call	cudaFree@PLT
	movq	48(%rsp), %rdi
	call	cudaFree@PLT
	movq	64(%rsp), %rdi
	call	cudaFree@PLT
	movq	56(%rsp), %rdi
	call	cudaFree@PLT
	movq	32(%rsp), %rdi
	call	cudaFree@PLT
	addl	%r14d, %ebx
	movss	12(%rsp), %xmm0
	addl	%ebp, %ebx
	imull	%r13d, %ebp
	imull	%r12d, %ebx
	movd	%xmm0, %eax
	addl	%ebp, %ebx
	salq	$34, %rbx
	orq	%rbx, %rax
	movq	72(%rsp), %rdx
	subq	%fs:40, %rdx
	je	.L11
	call	__stack_chk_fail@PLT
.L11:
	addq	$88, %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
.LFE4678:
	.size	_ZN2CP15run_convolutionE12tensor_shapej, .-_ZN2CP15run_convolutionE12tensor_shapej
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"Using Default shape"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB4689:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	cmpl	$11, %edi
	je	.L14
	leaq	.LC0(%rip), %rsi
	movl	$8, %r14d
	movl	$3, %ebp
	movl	$3, %r12d
	leaq	_ZSt4cerr(%rip), %rdi
	movl	$32, %r13d
	movl	$1, %r15d
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	xorl	%edi, %edi
	call	cudaSetDevice@PLT
	movl	$1024, %edx
	movl	$1024, %ecx
	movl	$32, %esi
	movl	$5, %r8d
	jmp	.L15
.L14:
	movq	8(%rsi), %rdi
	movq	%rsi, %rbx
	call	atoi@PLT
	movq	16(%rbx), %rdi
	movl	%eax, 12(%rsp)
	call	atoi@PLT
	movq	24(%rbx), %rdi
	movl	%eax, 8(%rsp)
	call	atoi@PLT
	movq	32(%rbx), %rdi
	movl	%eax, 4(%rsp)
	call	atoi@PLT
	movq	40(%rbx), %rdi
	movl	%eax, (%rsp)
	call	atoi@PLT
	movq	48(%rbx), %rdi
	movl	%eax, %r15d
	call	atoi@PLT
	movq	56(%rbx), %rdi
	movl	%eax, %r13d
	call	atoi@PLT
	movq	64(%rbx), %rdi
	movl	%eax, %r12d
	call	atoi@PLT
	movq	72(%rbx), %rdi
	movl	%eax, %ebp
	call	atoi@PLT
	movq	80(%rbx), %rdi
	movl	%eax, %r14d
	call	atoi@PLT
	movl	%eax, %edi
	call	cudaSetDevice@PLT
	movl	12(%rsp), %r8d
	movl	8(%rsp), %esi
	movl	4(%rsp), %ecx
	movl	(%rsp), %edx
.L15:
	movl	%ecx, 32(%rsp)
	subq	$48, %rsp
	.cfi_def_cfa_offset 176
	movl	$9, %ecx
	movl	%r15d, 88(%rsp)
	movl	%r14d, 108(%rsp)
	movl	%r13d, 92(%rsp)
	movl	%r12d, 100(%rsp)
	movl	%ebp, 104(%rsp)
	movl	%r8d, 76(%rsp)
	movl	%esi, 96(%rsp)
	movl	%edx, 84(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rsp, %rdi
	rep movsl
	movl	$1, %edi
	call	_ZN2CP15run_convolutionE12tensor_shapej
	addq	$120, %rsp
	.cfi_def_cfa_offset 56
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4689:
	.size	main, .-main
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4715:
	.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
.LFE4715:
	.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
	.hidden	DW.ref.__gxx_personality_v0
	.weak	DW.ref.__gxx_personality_v0
	.section	.data.rel.local.DW.ref.__gxx_personality_v0,"awG",@progbits,DW.ref.__gxx_personality_v0,comdat
	.align 8
	.type	DW.ref.__gxx_personality_v0, @object
	.size	DW.ref.__gxx_personality_v0, 8
DW.ref.__gxx_personality_v0:
	.quad	__gxx_personality_v0
	.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:

#ifndef CP4CONV2D_H
#define CP4CONV2D_H

#include "Tensor.cuh"

namespace CP {

  Tensor Conv2dForward(Tensor const,
                       Tensor const,
                       Tensor const,
                       Tensor const,
                       Tensor const,
                       unsigned pad = 1);

  /* Tensor Conv2dBackwardData(Tensor const, */
  /*                           Tensor const, */
  /*                           Tensor const, */
  /*                           Tensor const, */
  /*                           Tensor const, */
  /*                           unsigned pad = 1); */

  /* Tensor Conv2dBackwardFilter(Tensor const, */
  /*                             Tensor const, */
  /*                             Tensor const, */
  /*                             Tensor const, */
  /*                             Tensor const, */
  /*                             Tensor const, */
  /*                             unsigned pad = 1); */

  std::pair<float, unsigned>
  run_convolution(tensor_shape, unsigned PROFCOUNT = 1);

} // namespace CP


#endif /* CP4CONV2D_H */

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