#include "Utils.cuh"
#include "iostream"

#include <curand.h>

using namespace std;


// Simple cuda error checking macro
#define ErrChk(ans) \
  { CudaAssert((ans), __FILE__, __LINE__); }
inline void
CudaAssert(cudaError_t code, const char* file, int line, bool abort = true) {
  if (code != cudaSuccess) {
    fprintf(
        stderr, "CudaAssert: %s %s %d\n", cudaGetErrorString(code), file, line);
    if (abort) exit(code);
  }
}


Tensor random_fill(std::initializer_list<unsigned> lst) {

  curandGenerator_t gen;
  Tensor            A(lst);
  curandCreateGenerator(&gen, CURAND_RNG_PSEUDO_DEFAULT);
  curandSetPseudoRandomGeneratorSeed(gen, 1234ULL);
  curandGenerateUniform(gen, A.m_data, A.size());
  cudaDeviceSynchronize();
  curandDestroyGenerator(gen);

  return A;
};


__global__ void cp_recompose(float* __restrict__ Out,
                             const float* __restrict__ FilterK,
                             const float* __restrict__ FilterC,
                             const float* __restrict__ FilterH,
                             const float* __restrict__ FilterW,
                             const unsigned rank,
                             const unsigned FK,
                             const unsigned FC,
                             const unsigned FH,
                             const unsigned FW) {

  // clang-format off
  for (unsigned fc = threadIdx.z + blockIdx.z * blockDim.z; fc < FC; fc += blockDim.z*gridDim.z)
  for (unsigned fh = threadIdx.y + blockIdx.y * blockDim.y; fh < FH; fh += blockDim.y*gridDim.y)
  for (unsigned fw = threadIdx.x + blockIdx.x * blockDim.x; fw < FW; fw += blockDim.x*gridDim.x)
  for (unsigned fk = 0; fk < FK; ++fk) {

    float pixel = 0.0f;

    for (unsigned rr = 0; rr < rank; ++rr) {
      pixel += FilterK[fk * rank + rr]
             * FilterC[fc * rank + rr]
             * FilterH[fh * rank + rr]
             * FilterW[fw * rank + rr];
    }

    Out[fk*FC*FH*FW + fc*FH*FW + fh*FW + fw] = pixel;
  }
  // clang-format on
}

Tensor
cp4recom(Tensor FilterK, Tensor FilterC, Tensor FilterH, Tensor FilterW) {

  const unsigned rank = FilterK.shape[1];
  const unsigned FK   = FilterK.shape[0];
  const unsigned FC   = FilterC.shape[0];
  const unsigned FH   = FilterH.shape[0];
  const unsigned FW   = FilterW.shape[0];
  Tensor         Out  = { FK, FC, FH, FW };

  const unsigned W_dim = (FW / 8) + ((FW % 8) != 0);
  const unsigned H_dim = (FH / 8) + ((FH % 8) != 0);
  const unsigned C_dim = (FC / 8) + ((FC % 8) != 0);
  const dim3     Gshp(W_dim, H_dim, C_dim);
  const dim3     Bshp(8, 8, 8);

  cp_recompose<<<Gshp, Bshp>>>(Out.m_data,
                               FilterK.m_data,
                               FilterC.m_data,
                               FilterH.m_data,
                               FilterW.m_data,
                               rank,
                               FK,
                               FC,
                               FH,
                               FW);

  ErrChk(cudaPeekAtLastError());
  ErrChk(cudaDeviceSynchronize());

  return Out;
}

/* template <unsigned BlockSize> */
/* __device__ void warpMax(volatile float* sdata, unsigned tid) { */
/*   if (BlockSize >= 64) sdata[tid] = fmaxf(sdata[tid], sdata[tid + 32]); */
/*   if (BlockSize >= 32) sdata[tid] = fmaxf(sdata[tid], sdata[tid + 16]); */
/*   if (BlockSize >= 16) sdata[tid] = fmaxf(sdata[tid], sdata[tid + 8]); */
/*   if (BlockSize >= 8) sdata[tid]  = fmaxf(sdata[tid], sdata[tid + 4]); */
/*   if (BlockSize >= 4) sdata[tid]  = fmaxf(sdata[tid], sdata[tid + 2]); */
/*   if (BlockSize >= 2) sdata[tid]  = fmaxf(sdata[tid], sdata[tid + 1]); */
/* } */

/* template <unsigned BlockSize> */
/* __global__ void    maxCuda(bool* __restrict__ Over, */
/*                         const float* __restrict__ delta, */
/*                         const float    tol, */
/*                         const unsigned n) { */
/*   extern __shared__ float sdata[]; */

/*   unsigned tid    = threadIdx.x; */
/*   unsigned stride = 2 * BlockSize * gridDim.x; */
/*   sdata[tid]      = 0; */

/*   for (unsigned i = blockIdx.x * (2 * BlockSize) + tid; i < n; i += stride)
 */
/*     sdata[tid]    = fmaxf(delta[i], delta[i + BlockSize]); */
/*   __syncthreads(); */

/*   if (BlockSize >= 512) { */
/*     if (tid < 256) { sdata[tid] = fmaxf(sdata[tid], sdata[tid + 256]); } */
/*     __syncthreads(); */
/*   } */
/*   if (BlockSize >= 256) { */
/*     if (tid < 128) { sdata[tid] = fmaxf(sdata[tid], sdata[tid + 128]); } */
/*     __syncthreads(); */
/*   } */
/*   if (BlockSize >= 128) { */
/*     if (tid < 64) { sdata[tid] = fmaxf(sdata[tid], sdata[tid + 64]); } */
/*     __syncthreads(); */
/*   } */

/*   if (tid < 32) warpMax(sdata, tid); */
/*   if (tid == 0) Over[blockIdx.x] = sdata[0]; */
/* } */

__global__ void arrayRelativeDelta(int*        Exceeds,
                                   const float tol,
                                   const int   size,
                                   const float* __restrict__ A,
                                   const float* __restrict__ B) {
  unsigned tid = threadIdx.x + blockIdx.x * blockDim.x;
  if (tid >= size) return;

  float a = A[tid];
  float b = B[tid];
  float d = fabsf(a - b);
  float err = (a == 0 || b == 0) ? d : fmaxf(d / fabsf(a), d / fabsf(b));
  if (err > tol) atomicAdd(Exceeds, 1);
}


bool AllClose(Tensor A, Tensor B, float tolerance) {
  if (A.size() != B.size()) { return false; }

  int* Exceeds;
  ErrChk(cudaMalloc(&Exceeds, sizeof(int)));
  ErrChk(cudaMemset(Exceeds, 0, sizeof(int)));

  unsigned BlockSize = A.size() < 512 ? A.size() : 512;
  unsigned GridSize = (A.size() / 512) + ((A.size() % 512) != 0);

  arrayRelativeDelta<<<GridSize, BlockSize>>>(
      Exceeds, tolerance, A.size(), A.m_data, B.m_data);

  ErrChk(cudaPeekAtLastError());
  ErrChk(cudaDeviceSynchronize());

  int host_Exceeds;
  ErrChk(
      cudaMemcpy(&host_Exceeds, Exceeds, sizeof(int), cudaMemcpyDeviceToHost));

  return host_Exceeds == 0;
}

	.file	"tmpxft_002de21a_00000000-6_Utils.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4678:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4678:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.rodata._Z10CudaAssert9cudaErrorPKcib.str1.1,"aMS",@progbits,1
.LC0:
	.string	"CudaAssert: %s %s %d\n"
	.section	.text._Z10CudaAssert9cudaErrorPKcib,"axG",@progbits,_Z10CudaAssert9cudaErrorPKcib,comdat
	.weak	_Z10CudaAssert9cudaErrorPKcib
	.type	_Z10CudaAssert9cudaErrorPKcib, @function
_Z10CudaAssert9cudaErrorPKcib:
.LFB4660:
	.cfi_startproc
	endbr64
	testl	%edi, %edi
	je	.L8
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movl	%ecx, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movl	%edi, %ebx
	subq	$24, %rsp
	.cfi_def_cfa_offset 48
	movl	%edx, 12(%rsp)
	movq	%rsi, (%rsp)
	call	cudaGetErrorString@PLT
	movl	12(%rsp), %r9d
	movq	(%rsp), %r8
	leaq	.LC0(%rip), %rdx
	movq	stderr(%rip), %rdi
	movq	%rax, %rcx
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	testb	%bpl, %bpl
	je	.L2
	movl	%ebx, %edi
	call	exit@PLT
.L2:
	addq	$24, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
.L8:
	.cfi_restore 3
	.cfi_restore 6
	ret
	.cfi_endproc
.LFE4660:
	.size	_Z10CudaAssert9cudaErrorPKcib, .-_Z10CudaAssert9cudaErrorPKcib
	.text
	.globl	_Z11random_fillSt16initializer_listIjE
	.type	_Z11random_fillSt16initializer_listIjE, @function
_Z11random_fillSt16initializer_listIjE:
.LFB4661:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4661
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rdi, %rbx
	subq	$24, %rsp
	.cfi_def_cfa_offset 48
	movq	%fs:40, %rax
	movq	%rax, 8(%rsp)
	xorl	%eax, %eax
.LEHB0:
	call	_ZN6TensorC1ESt16initializer_listIjE@PLT
.LEHE0:
	movq	%rsp, %rdi
	movl	$100, %esi
.LEHB1:
	call	curandCreateGenerator@PLT
	movq	(%rsp), %rdi
	movl	$1234, %esi
	call	curandSetPseudoRandomGeneratorSeed@PLT
	movq	%rbx, %rdi
	call	_ZNK6Tensor4sizeEv@PLT
	movq	(%rbx), %rsi
	movq	(%rsp), %rdi
	movq	%rax, %rdx
	call	curandGenerateUniform@PLT
	call	cudaDeviceSynchronize@PLT
	movq	(%rsp), %rdi
	call	curandDestroyGenerator@PLT
.LEHE1:
	jmp	.L20
.L16:
	endbr64
	movq	%rax, %rbp
.L13:
	movq	%rbx, %rdi
	call	_ZN6TensorD1Ev@PLT
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L19
	movq	%rbp, %rdi
.LEHB2:
	call	_Unwind_Resume@PLT
.LEHE2:
.L20:
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	je	.L15
.L19:
	call	__stack_chk_fail@PLT
.L15:
	addq	$24, %rsp
	.cfi_def_cfa_offset 24
	movq	%rbx, %rax
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4661:
	.globl	__gxx_personality_v0
	.section	.gcc_except_table,"a",@progbits
.LLSDA4661:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4661-.LLSDACSB4661
.LLSDACSB4661:
	.uleb128 .LEHB0-.LFB4661
	.uleb128 .LEHE0-.LEHB0
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB1-.LFB4661
	.uleb128 .LEHE1-.LEHB1
	.uleb128 .L16-.LFB4661
	.uleb128 0
	.uleb128 .LEHB2-.LFB4661
	.uleb128 .LEHE2-.LEHB2
	.uleb128 0
	.uleb128 0
.LLSDACSE4661:
	.text
	.size	_Z11random_fillSt16initializer_listIjE, .-_Z11random_fillSt16initializer_listIjE
	.globl	_Z49__device_stub__Z12cp_recomposePfPKfS1_S1_S1_jjjjjPfPKfS1_S1_S1_jjjjj
	.type	_Z49__device_stub__Z12cp_recomposePfPKfS1_S1_S1_jjjjjPfPKfS1_S1_S1_jjjjj, @function
_Z49__device_stub__Z12cp_recomposePfPKfS1_S1_S1_jjjjjPfPKfS1_S1_S1_jjjjj:
.LFB4700:
	.cfi_startproc
	endbr64
	subq	$216, %rsp
	.cfi_def_cfa_offset 224
	movl	%r9d, 12(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 200(%rsp)
	xorl	%eax, %eax
	leaq	16(%rsp), %rax
	movq	%rdi, 16(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rax, 120(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	224(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	232(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	240(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	248(%rsp), %rax
	movq	%rsi, 24(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rdx, 32(%rsp)
	leaq	56(%rsp), %rdx
	movq	%rcx, 40(%rsp)
	leaq	64(%rsp), %rcx
	movq	%rax, 192(%rsp)
	movabsq	$4294967297, %rax
	movq	%r8, 48(%rsp)
	movq	%rax, 72(%rsp)
	movl	$1, 80(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L21
	pushq	64(%rsp)
	.cfi_def_cfa_offset 232
	leaq	_Z12cp_recomposePfPKfS1_S1_S1_jjjjj(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 240
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	136(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 232
	popq	%rdx
	.cfi_def_cfa_offset 224
.L21:
	movq	200(%rsp), %rax
	subq	%fs:40, %rax
	je	.L23
	call	__stack_chk_fail@PLT
.L23:
	addq	$216, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4700:
	.size	_Z49__device_stub__Z12cp_recomposePfPKfS1_S1_S1_jjjjjPfPKfS1_S1_S1_jjjjj, .-_Z49__device_stub__Z12cp_recomposePfPKfS1_S1_S1_jjjjjPfPKfS1_S1_S1_jjjjj
	.globl	_Z12cp_recomposePfPKfS1_S1_S1_jjjjj
	.type	_Z12cp_recomposePfPKfS1_S1_S1_jjjjj, @function
_Z12cp_recomposePfPKfS1_S1_S1_jjjjj:
.LFB4701:
	.cfi_startproc
	endbr64
	jmp	_Z49__device_stub__Z12cp_recomposePfPKfS1_S1_S1_jjjjjPfPKfS1_S1_S1_jjjjj
	.cfi_endproc
.LFE4701:
	.size	_Z12cp_recomposePfPKfS1_S1_S1_jjjjj, .-_Z12cp_recomposePfPKfS1_S1_S1_jjjjj
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"/home/ubuntu/Datasets/stackv2/train-structured/Areustle/ParallelTNNLayers/master/Cuda_kernels/Utils.cu"
	.text
	.globl	_Z8cp4recom6TensorS_S_S_
	.type	_Z8cp4recom6TensorS_S_S_, @function
_Z8cp4recom6TensorS_S_S_:
.LFB4674:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4674
	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
	pushq	%r12
	.cfi_offset 15, -24
	.cfi_offset 14, -32
	.cfi_offset 13, -40
	.cfi_offset 12, -48
	movq	%r8, %r12
	pushq	%rbx
	.cfi_offset 3, -56
	movq	%rdx, %rbx
	subq	$88, %rsp
	movq	%rdi, -104(%rbp)
	movq	%rcx, -128(%rbp)
	movq	%rsi, -120(%rbp)
	movq	%fs:40, %rax
	movq	%rax, -56(%rbp)
	movq	%rsi, %rax
	leaq	-72(%rbp), %rsi
	movq	8(%rax), %rax
	movl	4(%rax), %ecx
	movl	(%rax), %eax
	movl	%eax, -108(%rbp)
	movq	8(%rdx), %rax
	movl	$4, %edx
	movl	%ecx, -112(%rbp)
	movl	(%rax), %r13d
	movq	-128(%rbp), %rax
	movq	8(%rax), %rax
	movl	(%rax), %r14d
	movq	8(%r8), %rax
	movl	(%rax), %r15d
	movl	-108(%rbp), %eax
	movl	%eax, -72(%rbp)
	movl	%r13d, -68(%rbp)
	movl	%r14d, -64(%rbp)
	movl	%r15d, -60(%rbp)
.LEHB3:
	call	_ZN6TensorC1ESt16initializer_listIjE@PLT
.LEHE3:
	xorl	%eax, %eax
	movl	%r15d, %edx
	movl	$8, %ecx
	testl	$7, %r15d
	setne	%al
	shrl	$3, %edx
	addl	%edx, %eax
	movl	%r14d, %edx
	movl	%eax, -96(%rbp)
	xorl	%eax, %eax
	testl	$7, %r14d
	setne	%al
	shrl	$3, %edx
	xorl	%esi, %esi
	addl	%edx, %eax
	testl	$7, %r13d
	movabsq	$34359738376, %rdx
	movl	%eax, -92(%rbp)
	movl	%r13d, %eax
	setne	%sil
	movq	-96(%rbp), %rdi
	shrl	$3, %eax
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	addl	%eax, %esi
.LEHB4:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	je	.L27
.L29:
	call	cudaPeekAtLastError@PLT
	leaq	.LC1(%rip), %rbx
	movl	%eax, %edi
	movl	$1, %ecx
	movl	$94, %edx
	movq	%rbx, %rsi
	call	_Z10CudaAssert9cudaErrorPKcib
	jmp	.L38
.L27:
	movl	-108(%rbp), %eax
	movl	-112(%rbp), %r9d
	pushq	%r15
	movq	(%r12), %r8
	movq	(%rbx), %rdx
	pushq	%r14
	pushq	%r13
	pushq	%rax
	movq	-128(%rbp), %rax
	movq	(%rax), %rcx
	movq	-120(%rbp), %rax
	movq	(%rax), %rsi
	movq	-104(%rbp), %rax
	movq	(%rax), %rdi
	.cfi_escape 0x2e,0x20
	call	_Z49__device_stub__Z12cp_recomposePfPKfS1_S1_S1_jjjjjPfPKfS1_S1_S1_jjjjj
	addq	$32, %rsp
	jmp	.L29
.L38:
	.cfi_escape 0x2e,0
	call	cudaDeviceSynchronize@PLT
	movl	%eax, %edi
	movl	$1, %ecx
	movl	$95, %edx
	movq	%rbx, %rsi
	call	_Z10CudaAssert9cudaErrorPKcib
.LEHE4:
	jmp	.L39
.L33:
	endbr64
	movq	%rax, %rbx
.L30:
	movq	-104(%rbp), %rdi
	call	_ZN6TensorD1Ev@PLT
	movq	-56(%rbp), %rax
	subq	%fs:40, %rax
	jne	.L37
	movq	%rbx, %rdi
.LEHB5:
	call	_Unwind_Resume@PLT
.LEHE5:
.L39:
	movq	-56(%rbp), %rax
	subq	%fs:40, %rax
	je	.L32
.L37:
	call	__stack_chk_fail@PLT
.L32:
	movq	-104(%rbp), %rax
	leaq	-40(%rbp), %rsp
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa 7, 8
	ret
	.cfi_endproc
.LFE4674:
	.section	.gcc_except_table
.LLSDA4674:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4674-.LLSDACSB4674
.LLSDACSB4674:
	.uleb128 .LEHB3-.LFB4674
	.uleb128 .LEHE3-.LEHB3
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB4-.LFB4674
	.uleb128 .LEHE4-.LEHB4
	.uleb128 .L33-.LFB4674
	.uleb128 0
	.uleb128 .LEHB5-.LFB4674
	.uleb128 .LEHE5-.LEHB5
	.uleb128 0
	.uleb128 0
.LLSDACSE4674:
	.text
	.size	_Z8cp4recom6TensorS_S_S_, .-_Z8cp4recom6TensorS_S_S_
	.globl	_Z46__device_stub__Z18arrayRelativeDeltaPifiPKfS1_PifiPKfS1_
	.type	_Z46__device_stub__Z18arrayRelativeDeltaPifiPKfS1_PifiPKfS1_, @function
_Z46__device_stub__Z18arrayRelativeDeltaPifiPKfS1_PifiPKfS1_:
.LFB4702:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 8(%rsp)
	leaq	48(%rsp), %rdi
	movl	%esi, (%rsp)
	leaq	60(%rsp), %rsi
	movss	%xmm0, 4(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movq	%rdx, 16(%rsp)
	leaq	32(%rsp), %rdx
	movq	%rax, 96(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 104(%rsp)
	movq	%rsp, %rax
	movq	%rax, 112(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	24(%rsp), %rax
	movq	%rcx, 24(%rsp)
	leaq	40(%rsp), %rcx
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movl	$1, 56(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L40
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z18arrayRelativeDeltaPifiPKfS1_(%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
.L40:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L42
	call	__stack_chk_fail@PLT
.L42:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4702:
	.size	_Z46__device_stub__Z18arrayRelativeDeltaPifiPKfS1_PifiPKfS1_, .-_Z46__device_stub__Z18arrayRelativeDeltaPifiPKfS1_PifiPKfS1_
	.globl	_Z18arrayRelativeDeltaPifiPKfS1_
	.type	_Z18arrayRelativeDeltaPifiPKfS1_, @function
_Z18arrayRelativeDeltaPifiPKfS1_:
.LFB4703:
	.cfi_startproc
	endbr64
	jmp	_Z46__device_stub__Z18arrayRelativeDeltaPifiPKfS1_PifiPKfS1_
	.cfi_endproc
.LFE4703:
	.size	_Z18arrayRelativeDeltaPifiPKfS1_, .-_Z18arrayRelativeDeltaPifiPKfS1_
	.globl	_Z8AllClose6TensorS_f
	.type	_Z8AllClose6TensorS_f, @function
_Z8AllClose6TensorS_f:
.LFB4675:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movd	%xmm0, %r15d
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movq	%rsi, %r13
	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
	movq	%rdi, %rbx
	subq	$56, %rsp
	.cfi_def_cfa_offset 112
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	call	_ZNK6Tensor4sizeEv@PLT
	movq	%r13, %rdi
	movq	%rax, %rbp
	call	_ZNK6Tensor4sizeEv@PLT
	movq	%rax, %rdx
	xorl	%eax, %eax
	cmpq	%rdx, %rbp
	jne	.L45
	leaq	8(%rsp), %rdi
	movl	$4, %esi
	leaq	.LC1(%rip), %rbp
	movl	$512, %r14d
	call	cudaMalloc@PLT
	movl	$1, %ecx
	movl	$163, %edx
	movq	%rbp, %rsi
	movl	%eax, %edi
	call	_Z10CudaAssert9cudaErrorPKcib
	movq	8(%rsp), %rdi
	movl	$4, %edx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	movl	$1, %ecx
	movl	$164, %edx
	movq	%rbp, %rsi
	movl	%eax, %edi
	call	_Z10CudaAssert9cudaErrorPKcib
	movq	%rbx, %rdi
	call	_ZNK6Tensor4sizeEv@PLT
	cmpq	$511, %rax
	ja	.L47
	movq	%rbx, %rdi
	call	_ZNK6Tensor4sizeEv@PLT
	movl	%eax, %r14d
.L47:
	movq	%rbx, %rdi
	call	_ZNK6Tensor4sizeEv@PLT
	movq	%rbx, %rdi
	movq	%rax, %r12
	call	_ZNK6Tensor4sizeEv@PLT
	movl	%r14d, 28(%rsp)
	movabsq	$4294967297, %rdx
	testl	$511, %eax
	movq	%rdx, 20(%rsp)
	movl	24(%rsp), %esi
	setne	%al
	shrq	$9, %r12
	movq	%rdx, 32(%rsp)
	movl	36(%rsp), %ecx
	movzbl	%al, %eax
	movq	28(%rsp), %rdx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	addl	%eax, %r12d
	movl	%r12d, 16(%rsp)
	movq	16(%rsp), %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L48
	movq	0(%r13), %r13
	movq	(%rbx), %r12
	movq	%rbx, %rdi
	call	_ZNK6Tensor4sizeEv@PLT
	movq	8(%rsp), %rdi
	movd	%r15d, %xmm0
	movq	%r13, %rcx
	movq	%r12, %rdx
	movl	%eax, %esi
	call	_Z46__device_stub__Z18arrayRelativeDeltaPifiPKfS1_PifiPKfS1_
.L48:
	call	cudaPeekAtLastError@PLT
	movl	$1, %ecx
	movl	$172, %edx
	movq	%rbp, %rsi
	movl	%eax, %edi
	call	_Z10CudaAssert9cudaErrorPKcib
	call	cudaDeviceSynchronize@PLT
	movl	$1, %ecx
	movl	$173, %edx
	movq	%rbp, %rsi
	movl	%eax, %edi
	call	_Z10CudaAssert9cudaErrorPKcib
	movq	8(%rsp), %rsi
	leaq	28(%rsp), %rdi
	movl	$2, %ecx
	movl	$4, %edx
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$176, %edx
	movq	%rbp, %rsi
	movl	%eax, %edi
	call	_Z10CudaAssert9cudaErrorPKcib
	cmpl	$0, 28(%rsp)
	sete	%al
.L45:
	movq	40(%rsp), %rdx
	subq	%fs:40, %rdx
	je	.L49
	call	__stack_chk_fail@PLT
.L49:
	addq	$56, %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
.LFE4675:
	.size	_Z8AllClose6TensorS_f, .-_Z8AllClose6TensorS_f
	.section	.rodata.str1.1
.LC2:
	.string	"_Z18arrayRelativeDeltaPifiPKfS1_"
.LC3:
	.string	"_Z12cp_recomposePfPKfS1_S1_S1_jjjjj"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4705:
	.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	.LC2(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z18arrayRelativeDeltaPifiPKfS1_(%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	.LC3(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z12cp_recomposePfPKfS1_S1_S1_jjjjj(%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
.LFE4705:
	.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 CONV_UTILS_H
#define CONV_UTILS_H

#include "Tensor.cuh"

#include <cstddef>
#include <initializer_list>

Tensor random_fill(std::initializer_list<unsigned> lst);

Tensor cp4recom(Tensor A, Tensor B, Tensor C, Tensor D);

bool AllClose(Tensor, Tensor, float tolerance=1e-3);

#endif /* CONV_UTILS_H */

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