#include "depthconv_cuda_kernel.h"

#include <cstdio>

#define CUDA_KERNEL_LOOP(i, n)                                                 \
  for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < (n);                 \
       i += blockDim.x * gridDim.x)

const int CUDA_NUM_THREADS = 1024;

inline int GET_BLOCKS(const int N) {
  return (N + CUDA_NUM_THREADS - 1) / CUDA_NUM_THREADS;
}

template <typename DType>
__device__ DType get_gradient_weight(DType argmax_h, DType argmax_w,
                                     const int h, const int w, const int height,
                                     const int width) {

  if (argmax_h < 0 || argmax_h > height || argmax_w < 0 || argmax_w > width) {
    // empty
    return 0;
  }

  argmax_h = max(argmax_h, (DType)0.0f);
  argmax_w = max(argmax_w, (DType)0.0f);

  int argmax_h_low = (int)argmax_h;
  int argmax_w_low = (int)argmax_w;
  int argmax_h_high;
  int argmax_w_high;
  if (argmax_h_low >= height - 1) {
    argmax_h_high = argmax_h_low = height - 1;
    argmax_h = (DType)argmax_h_low;
  } else {
    argmax_h_high = argmax_h_low + 1;
  }
  if (argmax_w_low >= width - 1) {
    argmax_w_high = argmax_w_low = width - 1;
    argmax_w = (DType)argmax_w_low;
  } else {
    argmax_w_high = argmax_w_low + 1;
  }
  DType weight = 0;
  if (h == argmax_h_low) {
    if (w == argmax_w_low) {
      weight = (h + 1 - argmax_h) * (w + 1 - argmax_w);
    } else if (w == argmax_w_high) {
      weight = (h + 1 - argmax_h) * (argmax_w + 1 - w);
    }
  } else if (h == argmax_h_high) {
    if (w == argmax_w_low) {
      weight = (argmax_h + 1 - h) * (w + 1 - argmax_w);
    } else if (w == argmax_w_high) {
      weight = (argmax_h + 1 - h) * (argmax_w + 1 - w);
    }
  }
  return weight;
}

template <typename DType>
__global__ void depthconv_im2col_gpu_kernel(
    const int n, const DType *data_im, const DType *data_depth,
    const int height, const int width, const int kernel_h, const int kernel_w,
    const int pad_h, const int pad_w, const int stride_h, const int stride_w,
    const int dilation_h, const int dilation_w, const int height_col,
    const int width_col, DType *data_col) {
  // CxHxW --> (khxkw)x(CxHxW) 
  CUDA_KERNEL_LOOP(index, n) {
    const int w_col = index % width_col;
    const int h_col = (index / width_col) % height_col;
    const int c_im = (index / width_col) / height_col;
    const int c_col = c_im * kernel_h * kernel_w;


    const int h_in = h_col * stride_h - pad_h;
    const int w_in = w_col * stride_w - pad_w;
    DType *data_col_ptr = data_col + (c_col * height_col + h_col) * width_col + w_col;
    const DType *data_im_ptr = data_im + (c_im * height + h_in) * width + w_in;
    const DType *data_depth_ptr = data_depth + h_in * width + w_in;
    DType Di = 0.;
    bool valid = true;
    if ((h_in + dilation_h * (kernel_h - 1) / 2)>=0 &&
         w_in  + dilation_w * (kernel_w - 1) / 2 >= 0 &&
         (h_in + dilation_h * (kernel_h - 1) / 2) < height &&
         w_in  + dilation_w * (kernel_w - 1) / 2 < width)
        Di = data_depth[(h_in + dilation_h * (kernel_h - 1) / 2) * width + w_in  + dilation_w * (kernel_w - 1) / 2];
    else
        valid = false;
    //const DType Di = data_depth[(h_in + (kernel_h - 1) / 2 + dilation_h - 1) * width + (w_in + (kernel_w - 1) / 2 + dilation_w - 1)]; 

    for (int i = 0; i < kernel_h; ++i) {
      for (int j = 0; j < kernel_w; ++j) {
        DType val = static_cast<DType>(0);
        DType Dval = static_cast<DType>(0);
        const int h_im = h_in + i * dilation_h;
        const int w_im = w_in + j * dilation_w;
        if (h_im >= 0 && w_im >= 0 && h_im < height && w_im < width) {
          const int map_h = i * dilation_h;
          const int map_w = j * dilation_w;
          val = data_im_ptr[map_h * width + map_w];
          if (valid)
            Dval = data_depth_ptr[map_h * width + map_w];
	  //printf("%f,%d\n",Dval,h_in * width + w_in+map_h * width + map_w - ((h_in + (kernel_h - 1) / 2 + dilation_h - 1) * width + (w_in + (kernel_w - 1) / 2 + dilation_w - 1)));
          // printf("Di-Dval: %f, %f\n", Di, Dval);
	  // if (exp(-abs(Di - Dval))<0.2)
	  //	printf("Di-Dval: %f\n", exp(-abs(Di - Dval)));
          val *= exp(-abs(Di - Dval));
        }
        *data_col_ptr = val;
        data_col_ptr += height_col * width_col;
      }
    }
  }
}

template <typename DType>
void depthconv_im2col(cudaStream_t stream, const DType *data_im, const DType *data_depth, const int channels,
                       const int height, const int width, const int ksize_h,
                       const int ksize_w, const int pad_h, const int pad_w,
                       const int stride_h, const int stride_w,
                       const int dilation_h, const int dilation_w, DType *data_col) {
  // We are going to launch channels * height_col * width_col kernels, each
  // kernel responsible for copying a single-channel grid.
  int height_col =
      (height + 2 * pad_h - (dilation_h * (ksize_h - 1) + 1)) / stride_h + 1;
  int width_col =
      (width + 2 * pad_w - (dilation_w * (ksize_w - 1) + 1)) / stride_w + 1;
  int num_kernels = channels * height_col * width_col;
  // Launch
  depthconv_im2col_gpu_kernel<<<GET_BLOCKS(num_kernels), CUDA_NUM_THREADS, 0,
                                 stream>>>(
      num_kernels, data_im, data_depth, height, width, ksize_h, ksize_w, pad_h, pad_w, stride_h, stride_w, dilation_h, dilation_w, height_col, width_col, data_col);

  cudaError_t err = cudaGetLastError();
  if (err != cudaSuccess) {
    printf("error in depthconv_im2col: %s\n", cudaGetErrorString(err));
    // TODO(BZ) panic
  }
}

template void depthconv_im2col<float>(
    cudaStream_t stream, const float *data_im, const float *data_depth, 
    const int channels, const int height, const int width, const int ksize_h,
    const int ksize_w, const int pad_h, const int pad_w, const int stride_h,
    const int stride_w, const int dilation_h, const int dilation_w, float *data_col);

/*template void depthconv_im2col<double>(
    cudaStream_t stream, const double *data_im, const double *data_depth,
    const int channels, const int height, const int width, const int ksize_h,
    const int ksize_w, const int pad_h, const int pad_w, const int stride_h,
    const int stride_w, const int dilation_h, const int dilation_w, double *data_col);*/

template <typename DType>
__global__ void depthconv_col2im_gpu_kernel(
    const int n, const DType *data_col, const DType *data_depth,
    const int channels, const int height, const int width, const int kernel_h,
    const int kernel_w, const int pad_h, const int pad_w, const int stride_h,
    const int stride_w, const int dilation_h, const int dilation_w, const int height_col,
    const int width_col, DType *grad_im) {
  CUDA_KERNEL_LOOP(index, n) {
    for (int ii = 0; ii < kernel_h * kernel_w; ii++){
      int ii_index = ii + index * kernel_h * kernel_w;
      const int j = (ii_index / width_col / height_col) % kernel_w;
      const int i = (ii_index / width_col / height_col / kernel_w) % kernel_h;
      const int c = ii_index / width_col / height_col / kernel_w / kernel_h;
      // compute the start and end of the output

      int w_out = ii_index % width_col;
      int h_out = (ii_index / width_col) % height_col;
      int w_in = w_out * stride_w - pad_w;
      int h_in = h_out * stride_h - pad_h;

      //const DType cur_inv_h_data = h_in + i * dilation_h;
      //const DType cur_inv_w_data = w_in + j * dilation_w;

      const DType cur_top_grad = data_col[ii_index];
      const int cur_h = h_in + i * dilation_h;//(int)cur_inv_h_data;
      const int cur_w = w_in + j * dilation_w;//(int)cur_inv_w_data;

       DType Di = 0.;
      bool valid = true;
      if ((h_in + dilation_h * (kernel_h - 1) / 2)>=0 &&
             w_in  + dilation_w * (kernel_w - 1) / 2 >= 0 &&
             (h_in + dilation_h * (kernel_h - 1) / 2) < height &&
             w_in  + dilation_w * (kernel_w - 1) / 2 < width)
            Di = data_depth[(h_in + dilation_h * (kernel_h - 1) / 2) * width + w_in  + dilation_w * (kernel_w - 1) / 2];
      else
            valid = false;
//      const DType Di = data_depth[(h_in + dilation_h * (kernel_h - 1) / 2) * width + w_in  + dilation_w * (kernel_w - 1) / 2];
      //const DType Di = data_depth[(h_in + (kernel_h - 1) / 2 + dilation_h - 1) * width + w_in  + (kernel_w - 1) / 2 + dilation_w - 1];
      //printf("%d\n",(h_in + dilation_h * (kernel_h - 1) / 2) * width + w_in  + dilation_w * (kernel_w - 1) / 2);
      //data_depth[cur_h * width + cur_w];
      // data_depth[(h_in + (kernel_h - 1) / 2 + dilation_h - 1) * width + w_in  + (kernel_w - 1) / 2 + dilation_w - 1];
      

      int cur_bottom_grad_pos =
          (c * height + cur_h) * width + cur_w;
      int cur_bottom_depth_pos=
          (cur_h) * width + cur_w;
      //printf("%d,%d,%d,%d\n",i,j,((h_in + dilation_h * (kernel_h - 1) / 2) * width + w_in  + dilation_w * (kernel_w - 1) / 2-cur_bottom_depth_pos),dilation_h);
      //printf("%d\n",((h_in + dilation_h * (kernel_h - 1) / 2) * width + w_in  + dilation_w * (kernel_w - 1) / 2-cur_bottom_depth_pos));
      DType Dval = 0.;
      if (valid)
        Dval = data_depth[cur_bottom_depth_pos];
      if (cur_h >= 0 && cur_h < height && cur_w  >= 0 &&
              cur_w  < width)
        atomicAdd(grad_im + cur_bottom_grad_pos, cur_top_grad * exp(-abs(Di - Dval)));
      
    }
  }
}

template <typename DType>
void depthconv_col2im(cudaStream_t stream, const DType *data_col,
                       const DType *data_depth, const int channels,
                       const int height, const int width, const int ksize_h,
                       const int ksize_w, const int pad_h, const int pad_w,
                       const int stride_h, const int stride_w,
                       const int dilation_h, const int dilation_w, DType *grad_im) {

  int height_col =
      (height + 2 * pad_h - (dilation_h * (ksize_h - 1) + 1)) / stride_h + 1;
  int width_col =
      (width + 2 * pad_w - (dilation_w * (ksize_w - 1) + 1)) / stride_w + 1;
  int num_kernels = channels * height_col * width_col;
  // int channel_per_depthconv_group = channels / depthconv_group;
  // To avoid involving atomic operations, we will launch one kernel per
  // bottom dimension, and then in the kernel add up the top dimensions.
  depthconv_col2im_gpu_kernel<<<GET_BLOCKS(num_kernels), CUDA_NUM_THREADS, 0,
                                 stream>>>(
      num_kernels, data_col, data_depth, channels, height, width, ksize_h,
      ksize_w, pad_h, pad_w, stride_h, stride_w, dilation_h, dilation_w,
      height_col, width_col, grad_im);

  cudaError_t err = cudaGetLastError();
  if (err != cudaSuccess) {
    printf("error in depthconv_col2im: %s\n", cudaGetErrorString(err));
    // TODO(BZ) panic
  }
}

template void depthconv_col2im<float>(
    cudaStream_t stream, const float *data_col, const float *data_depth,
    const int channels, const int height, const int width, const int ksize_h,
    const int ksize_w, const int pad_h, const int pad_w, const int stride_h,
    const int stride_w, const int dilation_h, const int dilation_w, float *grad_im);

/*template void depthconv_col2im<double>(
    cudaStream_t stream, const double *data_col, const double *data_depth,
    const int channels, const int height, const int width, const int ksize_h,
    const int ksize_w, const int pad_h, const int pad_w, const int stride_h,
    const int stride_w, const int dilation_h, const int dilation_w, double *grad_im);*/

	.file	"tmpxft_0021642b_00000000-6_depthconv_cuda_kernel.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2037:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2037:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_ZL74__device_stub__Z27depthconv_col2im_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiiPS0_iPKfS0_iiiiiiiiiiiiiPf, @function
_ZL74__device_stub__Z27depthconv_col2im_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiiPS0_iPKfS0_iiiiiiiiiiiiiPf:
.LFB2061:
	.cfi_startproc
	subq	$264, %rsp
	.cfi_def_cfa_offset 272
	movq	352(%rsp), %rax
	movl	%edi, 44(%rsp)
	leaq	64(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	48(%rsp), %rdx
	movl	%ecx, 40(%rsp)
	leaq	56(%rsp), %rcx
	movl	%r8d, 20(%rsp)
	movl	%r9d, 16(%rsp)
	movq	%rax, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 248(%rsp)
	xorl	%eax, %eax
	leaq	44(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 112(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	272(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	280(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	288(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	296(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	304(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	312(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	320(%rsp), %rax
	movq	%rax, 208(%rsp)
	leaq	328(%rsp), %rax
	movq	%rax, 216(%rsp)
	leaq	336(%rsp), %rax
	movq	%rax, 224(%rsp)
	leaq	344(%rsp), %rax
	movq	%rax, 232(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 240(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	56(%rsp)
	.cfi_def_cfa_offset 280
	leaq	_Z27depthconv_col2im_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiiPS0_(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 288
	movq	92(%rsp), %rcx
	movl	100(%rsp), %r8d
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	leaq	128(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 280
	popq	%rdx
	.cfi_def_cfa_offset 272
.L2:
	movq	248(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$264, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2061:
	.size	_ZL74__device_stub__Z27depthconv_col2im_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiiPS0_iPKfS0_iiiiiiiiiiiiiPf, .-_ZL74__device_stub__Z27depthconv_col2im_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiiPS0_iPKfS0_iiiiiiiiiiiiiPf
	.type	_ZL73__device_stub__Z27depthconv_im2col_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiPS0_iPKfS0_iiiiiiiiiiiiPf, @function
_ZL73__device_stub__Z27depthconv_im2col_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiPS0_iPKfS0_iiiiiiiiiiiiPf:
.LFB2059:
	.cfi_startproc
	subq	$264, %rsp
	.cfi_def_cfa_offset 272
	movq	344(%rsp), %rax
	movl	%edi, 44(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	56(%rsp), %rdx
	movl	%ecx, 40(%rsp)
	leaq	64(%rsp), %rcx
	movl	%r8d, 20(%rsp)
	movl	%r9d, 16(%rsp)
	movq	%rax, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 248(%rsp)
	xorl	%eax, %eax
	leaq	44(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 120(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	272(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	280(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	288(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	296(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	304(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	312(%rsp), %rax
	movq	%rax, 208(%rsp)
	leaq	320(%rsp), %rax
	movq	%rax, 216(%rsp)
	leaq	328(%rsp), %rax
	movq	%rax, 224(%rsp)
	leaq	336(%rsp), %rax
	movq	%rax, 232(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 240(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L7
	pushq	64(%rsp)
	.cfi_def_cfa_offset 280
	leaq	_Z27depthconv_im2col_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiPS0_(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 288
	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 280
	popq	%rdx
	.cfi_def_cfa_offset 272
.L7:
	movq	248(%rsp), %rax
	subq	%fs:40, %rax
	je	.L9
	call	__stack_chk_fail@PLT
.L9:
	addq	$264, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2059:
	.size	_ZL73__device_stub__Z27depthconv_im2col_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiPS0_iPKfS0_iiiiiiiiiiiiPf, .-_ZL73__device_stub__Z27depthconv_im2col_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiPS0_iPKfS0_iiiiiiiiiiiiPf
	.section	.text._Z27depthconv_im2col_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiPS0_,"axG",@progbits,_Z27depthconv_im2col_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiPS0_,comdat
	.weak	_Z27depthconv_im2col_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiPS0_
	.type	_Z27depthconv_im2col_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiPS0_, @function
_Z27depthconv_im2col_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiPS0_:
.LFB2112:
	.cfi_startproc
	endbr64
	jmp	_ZL73__device_stub__Z27depthconv_im2col_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiPS0_iPKfS0_iiiiiiiiiiiiPf
	.cfi_endproc
.LFE2112:
	.size	_Z27depthconv_im2col_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiPS0_, .-_Z27depthconv_im2col_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiPS0_
	.section	.text._Z27depthconv_col2im_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiiPS0_,"axG",@progbits,_Z27depthconv_col2im_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiiPS0_,comdat
	.weak	_Z27depthconv_col2im_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiiPS0_
	.type	_Z27depthconv_col2im_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiiPS0_, @function
_Z27depthconv_col2im_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiiPS0_:
.LFB2114:
	.cfi_startproc
	endbr64
	jmp	_ZL74__device_stub__Z27depthconv_col2im_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiiPS0_iPKfS0_iiiiiiiiiiiiiPf
	.cfi_endproc
.LFE2114:
	.size	_Z27depthconv_col2im_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiiPS0_, .-_Z27depthconv_col2im_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiiPS0_
	.section	.rodata._Z16depthconv_im2colIfEvP11CUstream_stPKT_S4_iiiiiiiiiiiPS2_.str1.1,"aMS",@progbits,1
.LC0:
	.string	"error in depthconv_im2col: %s\n"
	.section	.text._Z16depthconv_im2colIfEvP11CUstream_stPKT_S4_iiiiiiiiiiiPS2_,"axG",@progbits,_Z16depthconv_im2colIfEvP11CUstream_stPKT_S4_iiiiiiiiiiiPS2_,comdat
	.weak	_Z16depthconv_im2colIfEvP11CUstream_stPKT_S4_iiiiiiiiiiiPS2_
	.type	_Z16depthconv_im2colIfEvP11CUstream_stPKT_S4_iiiiiiiiiiiPS2_, @function
_Z16depthconv_im2colIfEvP11CUstream_stPKT_S4_iiiiiiiiiiiPS2_:
.LFB2110:
	.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	$104, %rsp
	.cfi_def_cfa_offset 160
	movl	160(%rsp), %eax
	movl	208(%rsp), %r10d
	movq	%rdx, 40(%rsp)
	movl	176(%rsp), %r14d
	movl	216(%rsp), %r11d
	movq	%rsi, 32(%rsp)
	movl	%eax, 20(%rsp)
	movl	192(%rsp), %eax
	movl	20(%rsp), %ebx
	movl	%r8d, 12(%rsp)
	movl	%eax, 24(%rsp)
	movl	200(%rsp), %eax
	leal	-1(%rbx), %edx
	movl	%r9d, 16(%rsp)
	movl	168(%rsp), %r13d
	imull	%r10d, %edx
	movl	%eax, 28(%rsp)
	movq	224(%rsp), %rax
	movl	184(%rsp), %r15d
	movq	%rax, 48(%rsp)
	leal	(%r8,%r14,2), %eax
	xorl	%r8d, %r8d
	incl	%edx
	movl	%r10d, 60(%rsp)
	subl	%edx, %eax
	movl	%r11d, 56(%rsp)
	cltd
	idivl	24(%rsp)
	leal	-1(%r13), %edx
	imull	%r11d, %edx
	incl	%edx
	leal	1(%rax), %ebx
	leal	(%r9,%r15,2), %eax
	movq	%rdi, %r9
	subl	%edx, %eax
	imull	%ebx, %ecx
	cltd
	idivl	28(%rsp)
	movl	%ecx, %r12d
	movl	$1024, %ecx
	leal	1(%rax), %ebp
	imull	%ebp, %r12d
	leal	1023(%r12), %eax
	cltd
	idivl	%ecx
	movl	$4194305, %edx
	movl	$1, %ecx
	salq	$10, %rdx
	movl	%eax, 72(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 76(%rsp)
	movq	72(%rsp), %rdi
	movl	80(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	movl	56(%rsp), %r11d
	movl	60(%rsp), %r10d
	testl	%eax, %eax
	jne	.L14
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	movl	%r12d, %edi
	pushq	%rbp
	.cfi_def_cfa_offset 176
	pushq	%rbx
	.cfi_def_cfa_offset 184
	pushq	%r11
	.cfi_def_cfa_offset 192
	pushq	%r10
	.cfi_def_cfa_offset 200
	movl	68(%rsp), %eax
	pushq	%rax
	.cfi_def_cfa_offset 208
	movl	72(%rsp), %eax
	pushq	%rax
	.cfi_def_cfa_offset 216
	pushq	%r15
	.cfi_def_cfa_offset 224
	pushq	%r14
	.cfi_def_cfa_offset 232
	pushq	%r13
	.cfi_def_cfa_offset 240
	movl	100(%rsp), %r9d
	movl	96(%rsp), %r8d
	movl	92(%rsp), %ecx
	movq	120(%rsp), %rdx
	movq	112(%rsp), %rsi
	call	_ZL73__device_stub__Z27depthconv_im2col_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiPS0_iPKfS0_iiiiiiiiiiiiPf
	addq	$80, %rsp
	.cfi_def_cfa_offset 160
.L14:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L13
	call	cudaGetErrorString@PLT
	addq	$104, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 56
	movl	$2, %edi
	leaq	.LC0(%rip), %rsi
	popq	%rbx
	.cfi_def_cfa_offset 48
	movq	%rax, %rdx
	popq	%rbp
	.cfi_def_cfa_offset 40
	xorl	%eax, %eax
	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
	jmp	__printf_chk@PLT
.L13:
	.cfi_restore_state
	addq	$104, %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
.LFE2110:
	.size	_Z16depthconv_im2colIfEvP11CUstream_stPKT_S4_iiiiiiiiiiiPS2_, .-_Z16depthconv_im2colIfEvP11CUstream_stPKT_S4_iiiiiiiiiiiPS2_
	.section	.rodata._Z16depthconv_col2imIfEvP11CUstream_stPKT_S4_iiiiiiiiiiiPS2_.str1.1,"aMS",@progbits,1
.LC1:
	.string	"error in depthconv_col2im: %s\n"
	.section	.text._Z16depthconv_col2imIfEvP11CUstream_stPKT_S4_iiiiiiiiiiiPS2_,"axG",@progbits,_Z16depthconv_col2imIfEvP11CUstream_stPKT_S4_iiiiiiiiiiiPS2_,comdat
	.weak	_Z16depthconv_col2imIfEvP11CUstream_stPKT_S4_iiiiiiiiiiiPS2_
	.type	_Z16depthconv_col2imIfEvP11CUstream_stPKT_S4_iiiiiiiiiiiPS2_, @function
_Z16depthconv_col2imIfEvP11CUstream_stPKT_S4_iiiiiiiiiiiPS2_:
.LFB2111:
	.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
	movl	%ecx, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$104, %rsp
	.cfi_def_cfa_offset 160
	movl	192(%rsp), %eax
	movl	176(%rsp), %ebx
	movl	%ecx, 12(%rsp)
	movl	$1024, %ecx
	movl	160(%rsp), %r13d
	movl	208(%rsp), %r10d
	movq	%rdx, 40(%rsp)
	movl	%eax, 24(%rsp)
	movl	200(%rsp), %eax
	leal	-1(%r13), %edx
	movl	216(%rsp), %r11d
	movl	%r10d, 60(%rsp)
	imull	%r10d, %edx
	movl	%eax, 28(%rsp)
	movq	224(%rsp), %rax
	movq	%rsi, 32(%rsp)
	movl	168(%rsp), %r14d
	movq	%rax, 48(%rsp)
	leal	(%r8,%rbx,2), %eax
	movl	184(%rsp), %r15d
	incl	%edx
	movl	%r8d, 16(%rsp)
	xorl	%r8d, %r8d
	subl	%edx, %eax
	movl	%r9d, 20(%rsp)
	cltd
	idivl	24(%rsp)
	movl	%r11d, 56(%rsp)
	leal	-1(%r14), %edx
	imull	%r11d, %edx
	incl	%edx
	incl	%eax
	movl	%eax, 8(%rsp)
	leal	(%r9,%r15,2), %eax
	movq	%rdi, %r9
	subl	%edx, %eax
	cltd
	idivl	28(%rsp)
	leal	1(%rax), %ebp
	movl	8(%rsp), %eax
	imull	%eax, %r12d
	imull	%ebp, %r12d
	leal	1023(%r12), %eax
	cltd
	idivl	%ecx
	movl	$4194305, %edx
	movl	$1, %ecx
	salq	$10, %rdx
	movl	%eax, 72(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 76(%rsp)
	movq	72(%rsp), %rdi
	movl	80(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	movl	56(%rsp), %r11d
	movl	60(%rsp), %r10d
	testl	%eax, %eax
	jne	.L18
	pushq	%rax
	.cfi_def_cfa_offset 168
	movl	%r12d, %edi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 176
	pushq	%rbp
	.cfi_def_cfa_offset 184
	movl	32(%rsp), %eax
	pushq	%rax
	.cfi_def_cfa_offset 192
	pushq	%r11
	.cfi_def_cfa_offset 200
	pushq	%r10
	.cfi_def_cfa_offset 208
	movl	76(%rsp), %eax
	pushq	%rax
	.cfi_def_cfa_offset 216
	movl	80(%rsp), %eax
	pushq	%rax
	.cfi_def_cfa_offset 224
	pushq	%r15
	.cfi_def_cfa_offset 232
	pushq	%rbx
	.cfi_def_cfa_offset 240
	pushq	%r14
	.cfi_def_cfa_offset 248
	pushq	%r13
	.cfi_def_cfa_offset 256
	movl	116(%rsp), %r9d
	movl	112(%rsp), %r8d
	movl	108(%rsp), %ecx
	movq	136(%rsp), %rdx
	movq	128(%rsp), %rsi
	call	_ZL74__device_stub__Z27depthconv_col2im_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiiPS0_iPKfS0_iiiiiiiiiiiiiPf
	addq	$96, %rsp
	.cfi_def_cfa_offset 160
.L18:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L17
	call	cudaGetErrorString@PLT
	addq	$104, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 56
	movl	$2, %edi
	leaq	.LC1(%rip), %rsi
	popq	%rbx
	.cfi_def_cfa_offset 48
	movq	%rax, %rdx
	popq	%rbp
	.cfi_def_cfa_offset 40
	xorl	%eax, %eax
	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
	jmp	__printf_chk@PLT
.L17:
	.cfi_restore_state
	addq	$104, %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
.LFE2111:
	.size	_Z16depthconv_col2imIfEvP11CUstream_stPKT_S4_iiiiiiiiiiiPS2_, .-_Z16depthconv_col2imIfEvP11CUstream_stPKT_S4_iiiiiiiiiiiPS2_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC2:
	.string	"_Z27depthconv_col2im_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiiPS0_"
.LC3:
	.string	"_Z27depthconv_im2col_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiPS0_"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2064:
	.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	_Z27depthconv_col2im_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiiPS0_(%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	_Z27depthconv_im2col_gpu_kernelIfEviPKT_S2_iiiiiiiiiiiiPS0_(%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
.LFE2064:
	.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>
template <typename DType>
void depthconv_im2col(hipStream_t stream, const DType *data_im,
                       const DType *data_depth, const int channels,
                       const int height, const int width, const int ksize_h,
                       const int ksize_w, const int pad_h, const int pad_w,
                       const int stride_h, const int stride_w,
                       const int dilation_h, const int dilation_w, DType *data_col);

template <typename DType>
void depthconv_col2im(hipStream_t stream, const DType *data_col,
                       const DType *data_depth, const int channels,
                       const int height, const int width, const int ksize_h,
                       const int ksize_w, const int pad_h, const int pad_w,
                       const int stride_h, const int stride_w,
                       const int dilation_h, const int dilation_w, DType *grad_im);


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