
#include "dcnv2.cuh"
#include <cublas_v2.h>
#include "common.cuh"
 
template<typename scalar_t>
static __device__ scalar_t dcn_im2col_bilinear(
    const scalar_t* input, const int width_step, 
    const int width, const int height,
    scalar_t y, scalar_t x
){
    int y_low = floor(y);
    int x_low = floor(x);
    int y_high = y_low + 1;
    int x_high = x_low + 1;

    scalar_t ly = y - y_low;
    scalar_t lx = x - x_low;
    scalar_t hy = 1 - ly;
    scalar_t hx = 1 - lx;

    /*
    p0(x_low, y_low, hx*hy)             p1(x_high, y_low, lx*hy)
            anchor(x, y)
 
    p2(x_low, y_high, hx*ly)            p3(x_high, y_high, lx*ly)
     */

    scalar_t v0 = 0;
    if(x_low >= 0 && y_low >= 0)
        v0 = input[y_low * width_step + x_low];
    
    scalar_t v1 = 0;
    if(x_high < width && y_low >= 0)
        v1 = input[y_low * width_step + x_high];

    scalar_t v2 = 0;
    if(x_low >= 0 && y_high < height)
        v2 = input[y_high * width_step + x_low];

    scalar_t v3 = 0;
    if(x_high < width && y_high < height)
        v3 = input[y_high * width_step + x_high];

    // printf("%f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %d, %d, %d, %d\n",
    //     v0, v1, v2, v3,
    //     lx, ly, hx, hy, x, y, x_low, y_low, x_high, y_high
    // );
    return  v0 * (hx * hy) + 
            v1 * (lx * hy) + 
            v2 * (hx * ly) + 
            v3 * (lx * ly);
}

// 循环的次数，是batch,  hoxwo,  channels
template<typename scalar_t>
static __global__ void deformable_conv_im2col_kernel(
    const int edge,
    const scalar_t* input_data, const scalar_t* coord_offset,const scalar_t* coord_weight,
    const int batch, const int channels, const int height, const int width, 
    const int output_height, const int output_width, const int output_area,
    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 deformable_groups, const int channel_per_deformable_groups,
    scalar_t* column_data
){
    const int position = blockIdx.x * blockDim.x + threadIdx.x;
    if(position >= edge) return;

    // input_data    b c h w
    // coord_offset  b (deformable_groups * 2 * kh * kw) h w
    // coord_weight  b (deformable_groups * 1 * kh * kw) h w
    // deformable_groups 为 把输入分为多少个group。每个group内使用一组offset, weight
    // channel_per_deformable_groups = channels / deformable_groups

    // 每次进行一个核的赋值
    const int col_channel   = position % channels;
    const int col_y         = (position / channels) % output_area;
    const int col_batch     = (position / channels / output_area) % batch;

    // column坐标映射到输入的起点坐标
    const int output_x      = col_y % output_width;
    const int output_y      = col_y / output_width;
    const int input_x       = output_x * stride_w - pad_w;
    const int input_y       = output_y * stride_h - pad_h;
    const int input_channel = col_channel;
    const int input_batch   = col_batch;
    const int deformable_groups_index = input_channel / channel_per_deformable_groups;

    const int input_position = (input_batch * channels + input_channel) * height * width;
    scalar_t* column_data_ptr  = column_data + position * kernel_h * kernel_w;
    const scalar_t* input_data_ptr   = input_data + input_position;

    for(int i = 0; i < kernel_h; ++i){
        for(int j = 0; j < kernel_w; ++j){
            const scalar_t offset_y = coord_offset[(((((input_batch * deformable_groups + deformable_groups_index) * 2 + 0) * kernel_h + i) * kernel_w + j) * height + input_y) * width + input_x];
            const scalar_t offset_x = coord_offset[(((((input_batch * deformable_groups + deformable_groups_index) * 2 + 1) * kernel_h + i) * kernel_w + j) * height + input_y) * width + input_x];
            const scalar_t weight   = coord_weight[(((((input_batch * deformable_groups + deformable_groups_index) * 1 + 0) * kernel_h + i) * kernel_w + j) * height + input_y) * width + input_x];
            const scalar_t input_y_new_position = input_y + i * dilation_h + offset_y;
            const scalar_t input_x_new_position = input_x + j * dilation_w + offset_x;
            scalar_t value = 0;

            if(input_y_new_position > -1 && input_x_new_position > -1 && input_y_new_position < height && input_x_new_position < width){
                // 如果坐标在范围内，才需要进行插值。否则不需要
                value = dcn_im2col_bilinear<scalar_t>(input_data_ptr, width, width, height, input_y_new_position, input_x_new_position) * weight;
            }
            *column_data_ptr++ = value;
        }
    }
}

template<typename scalar_t>
void deformable_conv_im2col(
    const scalar_t* input_data, const scalar_t* coord_offset, const scalar_t* coord_weight,
    const int batch, const int channels, const int height, const int width, 
    const int output_height, const int output_width, const int output_area,
    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 deformable_groups,
    scalar_t* column_data, cudaStream_t stream
){
    const int num_kernels = batch * output_area * channels;
    int threads = 512;
    int blocks = ceil(num_kernels / (scalar_t)threads);
    const int channel_per_deformable_groups = channels / deformable_groups;

    deformable_conv_im2col_kernel<<<blocks, threads, 0, stream>>>(
        num_kernels, 
        input_data, coord_offset, coord_weight,
        batch, channels, height, width,
        output_height, output_width, output_area, 
        kernel_h, kernel_w,
        pad_h, pad_w,
        stride_h, stride_w,
        dilation_h, dilation_w,
        deformable_groups, channel_per_deformable_groups,
        column_data
    );

    CHECK_KERNEL_AND_PRINT_ERROR;
}


// 循环的次数，是batch,  hoxwo,  channels
template<typename scalar_t>
static __global__ void deformable_conv_col2im_kernel(
    const int edge,
    const scalar_t* dcolumn,
    const scalar_t* input, const scalar_t* coord_offset,const scalar_t* coord_weight,
    const int batch, const int channels, const int height, const int width, 
    const int output_height, const int output_width, const int output_area,
    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 deformable_groups, const int channel_per_deformable_groups,
    scalar_t* dinput, scalar_t* dcoord_offset, scalar_t* dcoord_weight
){
    const int position = blockIdx.x * blockDim.x + threadIdx.x;
    if(position >= edge) return;

    // 每次进行一个核的赋值
    const int col_channel   = position % channels;
    const int col_y         = (position / channels) % output_area;
    const int col_batch     = (position / channels / output_area) % batch;

    // column坐标映射到输入的起点坐标
    const int output_x      = col_y % output_width;
    const int output_y      = col_y / output_width;
    const int input_x       = output_x * stride_w - pad_w;
    const int input_y       = output_y * stride_h - pad_h;
    const int input_channel = col_channel;
    const int input_batch   = col_batch;
    const int deformable_groups_index = input_channel / channel_per_deformable_groups;

    // dcolumn is batch x channels*kernel_area x output_area
    const int input_position        = (input_batch * channels + input_channel) * height * width;
    scalar_t* dinput_data_ptr       = dinput + input_position;
    const scalar_t* input_data_ptr  = input  + input_position;
    for(int i = 0; i < kernel_h; ++i){
        for(int j = 0; j < kernel_w; ++j){
            int offset_y_position = (((((input_batch * deformable_groups + deformable_groups_index) * 2 + 0) * kernel_h + i) * kernel_w + j) * height + input_y) * width + input_x;
            int offset_x_position = (((((input_batch * deformable_groups + deformable_groups_index) * 2 + 1) * kernel_h + i) * kernel_w + j) * height + input_y) * width + input_x;
            int weight_position   = (((((input_batch * deformable_groups + deformable_groups_index) * 1 + 0) * kernel_h + i) * kernel_w + j) * height + input_y) * width + input_x;
            int column_grad_position = ((((col_batch * channels + col_channel) * kernel_h + i) * kernel_w + j) * output_height + output_y) * output_width + output_x;
            const scalar_t column_grad = dcolumn[column_grad_position];
            const scalar_t offset_y = coord_offset[offset_y_position];
            const scalar_t offset_x = coord_offset[offset_x_position];
            const scalar_t weight   = coord_weight[weight_position];
            const scalar_t input_y_new_position = input_y + i * dilation_h + offset_y;
            const scalar_t input_x_new_position = input_x + j * dilation_w + offset_x;
            scalar_t dvalue = column_grad * weight;

            if(input_y_new_position > -1 && input_x_new_position > -1 && input_y_new_position < height && input_x_new_position < width){
                // 如果坐标在范围内，才需要进行插值。否则不需要
                //value = dcn_im2col_bilinear<scalar_t>(input_data_ptr, width, width, height, input_y_new_position, input_x_new_position);

                int y_low = floor(input_y_new_position);
                int x_low = floor(input_x_new_position);
                int y_high = y_low + 1;
                int x_high = x_low + 1;
                int width_step = width;

                scalar_t ly = input_y_new_position - y_low;
                scalar_t lx = input_x_new_position - x_low;
                scalar_t hy = 1 - ly;
                scalar_t hx = 1 - lx;

                /*
                p0(x_low, y_low, hx*hy)             p1(x_high, y_low, lx*hy)
                        anchor(x, y)
            
                p2(x_low, y_high, hx*ly)            p3(x_high, y_high, lx*ly)
                */

                scalar_t v0 = 0;
                if(x_low >= 0 && y_low >= 0){
                    v0 = input_data_ptr[y_low * width_step + x_low];
                    atomicAdd(dinput_data_ptr + y_low * width_step + x_low, dvalue * hx * hy);
                }

                scalar_t v1 = 0;
                if(x_high < width && y_low >= 0){
                    v1 = input_data_ptr[y_low * width_step + x_high];
                    atomicAdd(dinput_data_ptr + y_low * width_step + x_high, dvalue * lx * hy);
                }

                scalar_t v2 = 0;
                if(x_low >= 0 && y_high < height){
                    v2 = input_data_ptr[y_high * width_step + x_low];
                    atomicAdd(dinput_data_ptr + y_high * width_step + x_low, dvalue * hx * ly);
                }

                scalar_t v3 = 0;
                if(x_high < width && y_high < height){
                    v3 = input_data_ptr[y_high * width_step + x_high];
                    atomicAdd(dinput_data_ptr + y_high * width_step + x_high, dvalue * lx * ly);
                }
                
                atomicAdd(dcoord_offset + offset_y_position, (v0 * hx * (-1) + v1 * lx * (-1) + v2 * hx * (+1) + v3 * lx * (+1)) * dvalue);
                atomicAdd(dcoord_offset + offset_x_position, (v0 * hy * (-1) + v1 * hy * (+1) + v2 * ly * (-1) + v3 * ly * (+1)) * dvalue);

                scalar_t ret_value = 
                        v0 * (hx * hy) + 
                        v1 * (lx * hy) + 
                        v2 * (hx * ly) + 
                        v3 * (lx * ly);
                atomicAdd(dcoord_weight + weight_position, column_grad * ret_value);
            }
        } 
    }
} 

template<typename scalar_t>
void deformable_conv_col2im(
    const scalar_t* dcolumn, const scalar_t* column,
    const scalar_t* input, const scalar_t* coord_offset, const scalar_t* coord_weight,
    const int batch, const int channels, const int height, const int width, 
    const int output_height, const int output_width, const int output_area,
    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 deformable_groups,
    scalar_t* dinput, scalar_t* dcoord_offset, scalar_t* dcoord_weight,
    cudaStream_t stream
){
    // channels*kernel_area x output_area 
    const int num_kernels = batch * output_area * channels;
    int threads = 256;
    int blocks = ceil(num_kernels / (float)threads);
    const int channel_per_deformable_groups = channels / deformable_groups;
  
    deformable_conv_col2im_kernel<<<blocks, threads, 0, stream>>>(
        num_kernels, 
        dcolumn,
        input, coord_offset, coord_weight,
        batch, channels, height, width,
        output_height, output_width, output_area, 
        kernel_h, kernel_w,
        pad_h, pad_w,
        stride_h, stride_w,
        dilation_h, dilation_w,
        deformable_groups, channel_per_deformable_groups,
        dinput, dcoord_offset, dcoord_weight
    );

    CHECK_KERNEL_AND_PRINT_ERROR;
}



#define DefineTemplate(name)                        \
static auto name##_float = name<float>;             \
static auto name##_double = name<double>;

DefineTemplate(deformable_conv_im2col);
DefineTemplate(deformable_conv_col2im);
	.file	"tmpxft_00268524_00000000-6_dcnv2.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2841:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2841:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_Z29deformable_conv_col2im_kernelIdEviPKT_S2_S2_S2_iiiiiiiiiiiiiiiiiPS0_S3_S3_, @function
_Z29deformable_conv_col2im_kernelIdEviPKT_S2_S2_S2_iiiiiiiiiiiiiiiiiPS0_S3_S3_:
.LFB2926:
	.cfi_startproc
	endbr64
	subq	$408, %rsp
	.cfi_def_cfa_offset 416
	movq	%fs:40, %rax
	movq	%rax, 392(%rsp)
	xorl	%eax, %eax
	movl	416(%rsp), %eax
	movl	%edi, (%rsp)
	leaq	144(%rsp), %rdi
	movq	%rsi, 72(%rsp)
	leaq	156(%rsp), %rsi
	movl	%eax, 8(%rsp)
	movl	424(%rsp), %eax
	movq	%rdx, 80(%rsp)
	leaq	128(%rsp), %rdx
	movl	%eax, 12(%rsp)
	movl	432(%rsp), %eax
	movq	%rcx, 88(%rsp)
	leaq	136(%rsp), %rcx
	movl	%eax, 16(%rsp)
	movl	440(%rsp), %eax
	movq	%r8, 96(%rsp)
	movl	%eax, 20(%rsp)
	movl	448(%rsp), %eax
	movl	%r9d, 4(%rsp)
	movl	%eax, 24(%rsp)
	movl	456(%rsp), %eax
	movl	%eax, 28(%rsp)
	movl	464(%rsp), %eax
	movl	%eax, 32(%rsp)
	movl	472(%rsp), %eax
	movl	%eax, 36(%rsp)
	movl	480(%rsp), %eax
	movl	%eax, 40(%rsp)
	movl	488(%rsp), %eax
	movl	%eax, 44(%rsp)
	movl	496(%rsp), %eax
	movl	%eax, 48(%rsp)
	movl	504(%rsp), %eax
	movl	%eax, 52(%rsp)
	movl	512(%rsp), %eax
	movl	%eax, 56(%rsp)
	movl	520(%rsp), %eax
	movl	%eax, 60(%rsp)
	movl	528(%rsp), %eax
	movl	%eax, 64(%rsp)
	movl	536(%rsp), %eax
	movl	%eax, 68(%rsp)
	movq	544(%rsp), %rax
	movq	%rax, 104(%rsp)
	movq	552(%rsp), %rax
	movq	%rax, 112(%rsp)
	movq	560(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	%rsp, %rax
	movq	%rax, 192(%rsp)
	leaq	72(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	80(%rsp), %rax
	movq	%rax, 208(%rsp)
	leaq	88(%rsp), %rax
	movq	%rax, 216(%rsp)
	leaq	96(%rsp), %rax
	movq	%rax, 224(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 232(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 240(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 248(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 256(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 264(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 272(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 280(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 288(%rsp)
	leaq	36(%rsp), %rax
	movq	%rax, 296(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 304(%rsp)
	leaq	44(%rsp), %rax
	movq	%rax, 312(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 320(%rsp)
	leaq	52(%rsp), %rax
	movq	%rax, 328(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 336(%rsp)
	leaq	60(%rsp), %rax
	movq	%rax, 344(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 352(%rsp)
	leaq	68(%rsp), %rax
	movq	%rax, 360(%rsp)
	leaq	104(%rsp), %rax
	movq	%rax, 368(%rsp)
	leaq	112(%rsp), %rax
	movq	%rax, 376(%rsp)
	leaq	120(%rsp), %rax
	movq	%rax, 384(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 144(%rsp)
	movl	$1, 152(%rsp)
	movq	%rax, 156(%rsp)
	movl	$1, 164(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	136(%rsp)
	.cfi_def_cfa_offset 424
	leaq	_Z29deformable_conv_col2im_kernelIdEviPKT_S2_S2_S2_iiiiiiiiiiiiiiiiiPS0_S3_S3_(%rip), %rdi
	pushq	136(%rsp)
	.cfi_def_cfa_offset 432
	movq	172(%rsp), %rcx
	movl	180(%rsp), %r8d
	movq	160(%rsp), %rsi
	movl	168(%rsp), %edx
	leaq	208(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 424
	popq	%rdx
	.cfi_def_cfa_offset 416
.L2:
	movq	392(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$408, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2926:
	.size	_Z29deformable_conv_col2im_kernelIdEviPKT_S2_S2_S2_iiiiiiiiiiiiiiiiiPS0_S3_S3_, .-_Z29deformable_conv_col2im_kernelIdEviPKT_S2_S2_S2_iiiiiiiiiiiiiiiiiPS0_S3_S3_
	.type	_Z29deformable_conv_im2col_kernelIfEviPKT_S2_S2_iiiiiiiiiiiiiiiiiPS0_, @function
_Z29deformable_conv_im2col_kernelIfEviPKT_S2_S2_iiiiiiiiiiiiiiiiiPS0_:
.LFB2922:
	.cfi_startproc
	endbr64
	subq	$360, %rsp
	.cfi_def_cfa_offset 368
	movq	%fs:40, %rax
	movq	%rax, 344(%rsp)
	xorl	%eax, %eax
	movl	368(%rsp), %eax
	movl	%edi, (%rsp)
	leaq	120(%rsp), %rdi
	movq	%rsi, 72(%rsp)
	leaq	132(%rsp), %rsi
	movl	%eax, 12(%rsp)
	movl	376(%rsp), %eax
	movq	%rdx, 80(%rsp)
	leaq	104(%rsp), %rdx
	movl	%eax, 16(%rsp)
	movl	384(%rsp), %eax
	movq	%rcx, 88(%rsp)
	leaq	112(%rsp), %rcx
	movl	%eax, 20(%rsp)
	movl	392(%rsp), %eax
	movl	%r8d, 4(%rsp)
	movl	%eax, 24(%rsp)
	movl	400(%rsp), %eax
	movl	%r9d, 8(%rsp)
	movl	%eax, 28(%rsp)
	movl	408(%rsp), %eax
	movl	%eax, 32(%rsp)
	movl	416(%rsp), %eax
	movl	%eax, 36(%rsp)
	movl	424(%rsp), %eax
	movl	%eax, 40(%rsp)
	movl	432(%rsp), %eax
	movl	%eax, 44(%rsp)
	movl	440(%rsp), %eax
	movl	%eax, 48(%rsp)
	movl	448(%rsp), %eax
	movl	%eax, 52(%rsp)
	movl	456(%rsp), %eax
	movl	%eax, 56(%rsp)
	movl	464(%rsp), %eax
	movl	%eax, 60(%rsp)
	movl	472(%rsp), %eax
	movl	%eax, 64(%rsp)
	movl	480(%rsp), %eax
	movl	%eax, 68(%rsp)
	movq	488(%rsp), %rax
	movq	%rax, 96(%rsp)
	movq	%rsp, %rax
	movq	%rax, 168(%rsp)
	leaq	72(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	80(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	88(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 208(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 216(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 224(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 232(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 240(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 248(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 256(%rsp)
	leaq	36(%rsp), %rax
	movq	%rax, 264(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 272(%rsp)
	leaq	44(%rsp), %rax
	movq	%rax, 280(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 288(%rsp)
	leaq	52(%rsp), %rax
	movq	%rax, 296(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 304(%rsp)
	leaq	60(%rsp), %rax
	movq	%rax, 312(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 320(%rsp)
	leaq	68(%rsp), %rax
	movq	%rax, 328(%rsp)
	leaq	96(%rsp), %rax
	movq	%rax, 336(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 120(%rsp)
	movl	$1, 128(%rsp)
	movq	%rax, 132(%rsp)
	movl	$1, 140(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L7
	pushq	112(%rsp)
	.cfi_def_cfa_offset 376
	leaq	_Z29deformable_conv_im2col_kernelIfEviPKT_S2_S2_iiiiiiiiiiiiiiiiiPS0_(%rip), %rdi
	pushq	112(%rsp)
	.cfi_def_cfa_offset 384
	movq	148(%rsp), %rcx
	movl	156(%rsp), %r8d
	movq	136(%rsp), %rsi
	movl	144(%rsp), %edx
	leaq	184(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 376
	popq	%rdx
	.cfi_def_cfa_offset 368
.L7:
	movq	344(%rsp), %rax
	subq	%fs:40, %rax
	je	.L9
	call	__stack_chk_fail@PLT
.L9:
	addq	$360, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2922:
	.size	_Z29deformable_conv_im2col_kernelIfEviPKT_S2_S2_iiiiiiiiiiiiiiiiiPS0_, .-_Z29deformable_conv_im2col_kernelIfEviPKT_S2_S2_iiiiiiiiiiiiiiiiiPS0_
	.type	_Z29deformable_conv_im2col_kernelIdEviPKT_S2_S2_iiiiiiiiiiiiiiiiiPS0_, @function
_Z29deformable_conv_im2col_kernelIdEviPKT_S2_S2_iiiiiiiiiiiiiiiiiPS0_:
.LFB2924:
	.cfi_startproc
	endbr64
	subq	$360, %rsp
	.cfi_def_cfa_offset 368
	movq	%fs:40, %rax
	movq	%rax, 344(%rsp)
	xorl	%eax, %eax
	movl	368(%rsp), %eax
	movl	%edi, (%rsp)
	leaq	120(%rsp), %rdi
	movq	%rsi, 72(%rsp)
	leaq	132(%rsp), %rsi
	movl	%eax, 12(%rsp)
	movl	376(%rsp), %eax
	movq	%rdx, 80(%rsp)
	leaq	104(%rsp), %rdx
	movl	%eax, 16(%rsp)
	movl	384(%rsp), %eax
	movq	%rcx, 88(%rsp)
	leaq	112(%rsp), %rcx
	movl	%eax, 20(%rsp)
	movl	392(%rsp), %eax
	movl	%r8d, 4(%rsp)
	movl	%eax, 24(%rsp)
	movl	400(%rsp), %eax
	movl	%r9d, 8(%rsp)
	movl	%eax, 28(%rsp)
	movl	408(%rsp), %eax
	movl	%eax, 32(%rsp)
	movl	416(%rsp), %eax
	movl	%eax, 36(%rsp)
	movl	424(%rsp), %eax
	movl	%eax, 40(%rsp)
	movl	432(%rsp), %eax
	movl	%eax, 44(%rsp)
	movl	440(%rsp), %eax
	movl	%eax, 48(%rsp)
	movl	448(%rsp), %eax
	movl	%eax, 52(%rsp)
	movl	456(%rsp), %eax
	movl	%eax, 56(%rsp)
	movl	464(%rsp), %eax
	movl	%eax, 60(%rsp)
	movl	472(%rsp), %eax
	movl	%eax, 64(%rsp)
	movl	480(%rsp), %eax
	movl	%eax, 68(%rsp)
	movq	488(%rsp), %rax
	movq	%rax, 96(%rsp)
	movq	%rsp, %rax
	movq	%rax, 168(%rsp)
	leaq	72(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	80(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	88(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 208(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 216(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 224(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 232(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 240(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 248(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 256(%rsp)
	leaq	36(%rsp), %rax
	movq	%rax, 264(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 272(%rsp)
	leaq	44(%rsp), %rax
	movq	%rax, 280(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 288(%rsp)
	leaq	52(%rsp), %rax
	movq	%rax, 296(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 304(%rsp)
	leaq	60(%rsp), %rax
	movq	%rax, 312(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 320(%rsp)
	leaq	68(%rsp), %rax
	movq	%rax, 328(%rsp)
	leaq	96(%rsp), %rax
	movq	%rax, 336(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 120(%rsp)
	movl	$1, 128(%rsp)
	movq	%rax, 132(%rsp)
	movl	$1, 140(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L11
	pushq	112(%rsp)
	.cfi_def_cfa_offset 376
	leaq	_Z29deformable_conv_im2col_kernelIdEviPKT_S2_S2_iiiiiiiiiiiiiiiiiPS0_(%rip), %rdi
	pushq	112(%rsp)
	.cfi_def_cfa_offset 384
	movq	148(%rsp), %rcx
	movl	156(%rsp), %r8d
	movq	136(%rsp), %rsi
	movl	144(%rsp), %edx
	leaq	184(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 376
	popq	%rdx
	.cfi_def_cfa_offset 368
.L11:
	movq	344(%rsp), %rax
	subq	%fs:40, %rax
	je	.L13
	call	__stack_chk_fail@PLT
.L13:
	addq	$360, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2924:
	.size	_Z29deformable_conv_im2col_kernelIdEviPKT_S2_S2_iiiiiiiiiiiiiiiiiPS0_, .-_Z29deformable_conv_im2col_kernelIdEviPKT_S2_S2_iiiiiiiiiiiiiiiiiPS0_
	.type	_Z29deformable_conv_col2im_kernelIfEviPKT_S2_S2_S2_iiiiiiiiiiiiiiiiiPS0_S3_S3_, @function
_Z29deformable_conv_col2im_kernelIfEviPKT_S2_S2_S2_iiiiiiiiiiiiiiiiiPS0_S3_S3_:
.LFB2925:
	.cfi_startproc
	endbr64
	subq	$408, %rsp
	.cfi_def_cfa_offset 416
	movq	%fs:40, %rax
	movq	%rax, 392(%rsp)
	xorl	%eax, %eax
	movl	416(%rsp), %eax
	movl	%edi, (%rsp)
	leaq	144(%rsp), %rdi
	movq	%rsi, 72(%rsp)
	leaq	156(%rsp), %rsi
	movl	%eax, 8(%rsp)
	movl	424(%rsp), %eax
	movq	%rdx, 80(%rsp)
	leaq	128(%rsp), %rdx
	movl	%eax, 12(%rsp)
	movl	432(%rsp), %eax
	movq	%rcx, 88(%rsp)
	leaq	136(%rsp), %rcx
	movl	%eax, 16(%rsp)
	movl	440(%rsp), %eax
	movq	%r8, 96(%rsp)
	movl	%eax, 20(%rsp)
	movl	448(%rsp), %eax
	movl	%r9d, 4(%rsp)
	movl	%eax, 24(%rsp)
	movl	456(%rsp), %eax
	movl	%eax, 28(%rsp)
	movl	464(%rsp), %eax
	movl	%eax, 32(%rsp)
	movl	472(%rsp), %eax
	movl	%eax, 36(%rsp)
	movl	480(%rsp), %eax
	movl	%eax, 40(%rsp)
	movl	488(%rsp), %eax
	movl	%eax, 44(%rsp)
	movl	496(%rsp), %eax
	movl	%eax, 48(%rsp)
	movl	504(%rsp), %eax
	movl	%eax, 52(%rsp)
	movl	512(%rsp), %eax
	movl	%eax, 56(%rsp)
	movl	520(%rsp), %eax
	movl	%eax, 60(%rsp)
	movl	528(%rsp), %eax
	movl	%eax, 64(%rsp)
	movl	536(%rsp), %eax
	movl	%eax, 68(%rsp)
	movq	544(%rsp), %rax
	movq	%rax, 104(%rsp)
	movq	552(%rsp), %rax
	movq	%rax, 112(%rsp)
	movq	560(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	%rsp, %rax
	movq	%rax, 192(%rsp)
	leaq	72(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	80(%rsp), %rax
	movq	%rax, 208(%rsp)
	leaq	88(%rsp), %rax
	movq	%rax, 216(%rsp)
	leaq	96(%rsp), %rax
	movq	%rax, 224(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 232(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 240(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 248(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 256(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 264(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 272(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 280(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 288(%rsp)
	leaq	36(%rsp), %rax
	movq	%rax, 296(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 304(%rsp)
	leaq	44(%rsp), %rax
	movq	%rax, 312(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 320(%rsp)
	leaq	52(%rsp), %rax
	movq	%rax, 328(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 336(%rsp)
	leaq	60(%rsp), %rax
	movq	%rax, 344(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 352(%rsp)
	leaq	68(%rsp), %rax
	movq	%rax, 360(%rsp)
	leaq	104(%rsp), %rax
	movq	%rax, 368(%rsp)
	leaq	112(%rsp), %rax
	movq	%rax, 376(%rsp)
	leaq	120(%rsp), %rax
	movq	%rax, 384(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 144(%rsp)
	movl	$1, 152(%rsp)
	movq	%rax, 156(%rsp)
	movl	$1, 164(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L15
	pushq	136(%rsp)
	.cfi_def_cfa_offset 424
	leaq	_Z29deformable_conv_col2im_kernelIfEviPKT_S2_S2_S2_iiiiiiiiiiiiiiiiiPS0_S3_S3_(%rip), %rdi
	pushq	136(%rsp)
	.cfi_def_cfa_offset 432
	movq	172(%rsp), %rcx
	movl	180(%rsp), %r8d
	movq	160(%rsp), %rsi
	movl	168(%rsp), %edx
	leaq	208(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 424
	popq	%rdx
	.cfi_def_cfa_offset 416
.L15:
	movq	392(%rsp), %rax
	subq	%fs:40, %rax
	je	.L17
	call	__stack_chk_fail@PLT
.L17:
	addq	$408, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2925:
	.size	_Z29deformable_conv_col2im_kernelIfEviPKT_S2_S2_S2_iiiiiiiiiiiiiiiiiPS0_S3_S3_, .-_Z29deformable_conv_col2im_kernelIfEviPKT_S2_S2_S2_iiiiiiiiiiiiiiiiiPS0_S3_S3_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z29deformable_conv_col2im_kernelIdEviPKT_S2_S2_S2_iiiiiiiiiiiiiiiiiPS0_S3_S3_"
.LC1:
	.string	"_Z29deformable_conv_col2im_kernelIfEviPKT_S2_S2_S2_iiiiiiiiiiiiiiiiiPS0_S3_S3_"
.LC2:
	.string	"_Z29deformable_conv_im2col_kernelIdEviPKT_S2_S2_iiiiiiiiiiiiiiiiiPS0_"
.LC3:
	.string	"_Z29deformable_conv_im2col_kernelIfEviPKT_S2_S2_iiiiiiiiiiiiiiiiiPS0_"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2872:
	.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	.LC0(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z29deformable_conv_col2im_kernelIdEviPKT_S2_S2_S2_iiiiiiiiiiiiiiiiiPS0_S3_S3_(%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	.LC1(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z29deformable_conv_col2im_kernelIfEviPKT_S2_S2_S2_iiiiiiiiiiiiiiiiiPS0_S3_S3_(%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
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC2(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z29deformable_conv_im2col_kernelIdEviPKT_S2_S2_iiiiiiiiiiiiiiiiiPS0_(%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
	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	_Z29deformable_conv_im2col_kernelIfEviPKT_S2_S2_iiiiiiiiiiiiiiiiiPS0_(%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
.LFE2872:
	.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:



#ifndef DCNV2_CUH
#define DCNV2_CUH

#include <hip/hip_runtime.h>

template<typename scalar_t>
void deformable_conv_im2col(
    const scalar_t* input_data, const scalar_t* coord_offset, const scalar_t* coord_weight,
    const int batch, const int channels, const int height, const int width, 
    const int output_height, const int output_width, const int output_area,
    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 deformable_group,
    scalar_t* column_data, hipStream_t stream
);

template<typename scalar_t>
void deformable_conv_col2im(
    const scalar_t* dcolumn, const scalar_t* column,
    const scalar_t* input, const scalar_t* coord_offset, const scalar_t* coord_weight,
    const int batch, const int channels, const int height, const int width, 
    const int output_height, const int output_width, const int output_area,
    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 deformable_group,
    scalar_t* dinput, scalar_t* dcoord_offset, scalar_t* dcoord_weight,
    hipStream_t stream
);

#endif //DCNV2_CUH
	.text
	.file	"dcnv2.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
