#include <cuda_runtime.h>
#include <iostream>

// CUDA kernel for bilinear resize
__global__ void bilinear_resize_kernel(
    const float* input, float* output,
    int in_width, int in_height,
    int out_width, int out_height
) {
    int out_x = blockIdx.x * blockDim.x + threadIdx.x;
    int out_y = blockIdx.y * blockDim.y + threadIdx.y;

    if (out_x >= out_width || out_y >= out_height) return;

    // Normalize output pixel coordinates to input image coordinates
    float x_ratio = (float)(in_width - 1) / (out_width - 1);
    float y_ratio = (float)(in_height - 1) / (out_height - 1);

    float in_x = out_x * x_ratio;
    float in_y = out_y * y_ratio;

    int x_low = floorf(in_x);
    int y_low = floorf(in_y);
    int x_high = min(x_low + 1, in_width - 1);
    int y_high = min(y_low + 1, in_height - 1);

    float x_lerp = in_x - x_low;
    float y_lerp = in_y - y_low;

    // Fetch 4 neighbors
    float top_left = input[y_low * in_width + x_low];
    float top_right = input[y_low * in_width + x_high];
    float bottom_left = input[y_high * in_width + x_low];
    float bottom_right = input[y_high * in_width + x_high];

    // Interpolate
    float top = top_left + (top_right - top_left) * x_lerp;
    float bottom = bottom_left + (bottom_right - bottom_left) * x_lerp;
    float value = top + (bottom - top) * y_lerp;
    value = roundf(value * 1e5f) / 1e5f;

    output[out_y * out_width + out_x] = value;
}

void bilinear_resize(
    const float* h_input, float* h_output,
    int in_width, int in_height,
    int out_width, int out_height
) {
    float *d_input, *d_output;
    size_t in_size = in_width * in_height * sizeof(float);
    size_t out_size = out_width * out_height * sizeof(float);

    cudaMalloc(&d_input, in_size);
    cudaMalloc(&d_output, out_size);

    cudaMemcpy(d_input, h_input, in_size, cudaMemcpyHostToDevice);

    dim3 block(16, 16);
    dim3 grid((out_width + block.x - 1) / block.x, (out_height + block.y - 1) / block.y);

    bilinear_resize_kernel<<<grid, block>>>(d_input, d_output, in_width, in_height, out_width, out_height);

    cudaMemcpy(h_output, d_output, out_size, cudaMemcpyDeviceToHost);

    cudaFree(d_input);
    cudaFree(d_output);
}

int main() {
    const int in_width = 4, in_height = 4;
    const int out_width = 8, out_height = 8;

    float h_input[in_width * in_height] = {
        1, 2, 3, 4,
        5, 6, 7, 8,
        9,10,11,12,
        13,14,15,16
    };

    float h_output[out_width * out_height];

    bilinear_resize(h_input, h_output, in_width, in_height, out_width, out_height);

    for (int y = 0; y < out_height; ++y) {
        for (int x = 0; x < out_width; ++x) {
            std::cout << h_output[y * out_width + x] << " ";
        }
        std::cout << "\n";
    }

    return 0;
}

	.file	"tmpxft_00047e31_00000000-6_26.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3378:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3378:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z49__device_stub__Z22bilinear_resize_kernelPKfPfiiiiPKfPfiiii
	.type	_Z49__device_stub__Z22bilinear_resize_kernelPKfPfiiiiPKfPfiiii, @function
_Z49__device_stub__Z22bilinear_resize_kernelPKfPfiiiiPKfPfiiii:
.LFB3400:
	.cfi_startproc
	endbr64
	subq	$168, %rsp
	.cfi_def_cfa_offset 176
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 8(%rsp)
	leaq	48(%rsp), %rcx
	movl	%r8d, 4(%rsp)
	movl	%r9d, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 136(%rsp)
	movq	%rsp, %rax
	movq	%rax, 144(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	48(%rsp)
	.cfi_def_cfa_offset 184
	leaq	_Z22bilinear_resize_kernelPKfPfiiii(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 192
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 184
	popq	%rdx
	.cfi_def_cfa_offset 176
.L2:
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$168, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3400:
	.size	_Z49__device_stub__Z22bilinear_resize_kernelPKfPfiiiiPKfPfiiii, .-_Z49__device_stub__Z22bilinear_resize_kernelPKfPfiiiiPKfPfiiii
	.globl	_Z22bilinear_resize_kernelPKfPfiiii
	.type	_Z22bilinear_resize_kernelPKfPfiiii, @function
_Z22bilinear_resize_kernelPKfPfiiii:
.LFB3401:
	.cfi_startproc
	endbr64
	jmp	_Z49__device_stub__Z22bilinear_resize_kernelPKfPfiiiiPKfPfiiii
	.cfi_endproc
.LFE3401:
	.size	_Z22bilinear_resize_kernelPKfPfiiii, .-_Z22bilinear_resize_kernelPKfPfiiii
	.globl	_Z15bilinear_resizePKfPfiiii
	.type	_Z15bilinear_resizePKfPfiiii, @function
_Z15bilinear_resizePKfPfiiii:
.LFB3374:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rsi, %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movl	%edx, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movl	%edx, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	imull	%ecx, %r13d
	movl	%r8d, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	imull	%r9d, %r12d
	movl	%r9d, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%r8d, %ebx
	movslq	%r13d, %r13
	salq	$2, %r13
	movslq	%r12d, %r12
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	salq	$2, %r12
	movq	%r13, %rsi
	movq	%rdi, 8(%rsp)
	leaq	16(%rsp), %rdi
	movq	%fs:40, %rax
	movq	%rax, 56(%rsp)
	xorl	%eax, %eax
	movl	%ecx, 4(%rsp)
	call	cudaMalloc@PLT
	leaq	24(%rsp), %rdi
	movq	%r12, %rsi
	call	cudaMalloc@PLT
	movq	8(%rsp), %rsi
	movq	16(%rsp), %rdi
	movq	%r13, %rdx
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	leal	15(%rbx), %eax
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	shrl	$4, %eax
	movl	$1, %ecx
	movl	$1, %esi
	movabsq	$68719476752, %rdx
	movl	%eax, 44(%rsp)
	leal	15(%rbp), %eax
	shrl	$4, %eax
	movl	%eax, 48(%rsp)
	movq	44(%rsp), %rdi
	call	__cudaPushCallConfiguration@PLT
	movl	4(%rsp), %r10d
	testl	%eax, %eax
	jne	.L9
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	%ebp, %r9d
	movl	%ebx, %r8d
	movl	%r10d, %ecx
	movl	%r14d, %edx
	call	_Z49__device_stub__Z22bilinear_resize_kernelPKfPfiiiiPKfPfiiii
.L9:
	movq	24(%rsp), %rsi
	movl	$2, %ecx
	movq	%r12, %rdx
	movq	%r15, %rdi
	call	cudaMemcpy@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$72, %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
.LFE3374:
	.size	_Z15bilinear_resizePKfPfiiii, .-_Z15bilinear_resizePKfPfiiii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	" "
.LC2:
	.string	"\n"
	.section	.rodata
	.align 32
.LC0:
	.long	1065353216
	.long	1073741824
	.long	1077936128
	.long	1082130432
	.long	1084227584
	.long	1086324736
	.long	1088421888
	.long	1090519040
	.long	1091567616
	.long	1092616192
	.long	1093664768
	.long	1094713344
	.long	1095761920
	.long	1096810496
	.long	1097859072
	.long	1098907648
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3375:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movl	$16, %ecx
	leaq	.LC0(%rip), %rsi
	movl	$8, %r9d
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movl	$8, %r8d
	movl	$4, %edx
	xorl	%r12d, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	leaq	_ZSt4cout(%rip), %r13
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	subq	$344, %rsp
	.cfi_def_cfa_offset 384
	movq	%fs:40, %rax
	movq	%rax, 328(%rsp)
	xorl	%eax, %eax
	leaq	72(%rsp), %rbx
	leaq	8(%rsp), %rdi
	rep movsl
	leaq	8(%rsp), %rdi
	movl	$4, %ecx
	movq	%rbx, %rsi
	call	_Z15bilinear_resizePKfPfiiii
.L13:
	xorl	%ebp, %ebp
.L14:
	movq	%r13, %rdi
	cvtss2sd	(%rbx,%rbp,4), %xmm0
	incq	%rbp
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movl	$1, %edx
	leaq	.LC1(%rip), %rsi
	movq	%rax, %rdi
	call	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l@PLT
	cmpq	$8, %rbp
	jne	.L14
	movl	$1, %edx
	movq	%r13, %rdi
	addl	$8, %r12d
	addq	$32, %rbx
	leaq	.LC2(%rip), %rsi
	call	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l@PLT
	cmpl	$64, %r12d
	jne	.L13
	movq	328(%rsp), %rax
	subq	%fs:40, %rax
	je	.L16
	call	__stack_chk_fail@PLT
.L16:
	addq	$344, %rsp
	.cfi_def_cfa_offset 40
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3375:
	.size	main, .-main
	.section	.rodata.str1.1
.LC3:
	.string	"_Z22bilinear_resize_kernelPKfPfiiii"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3403:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 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	.LC3(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z22bilinear_resize_kernelPKfPfiiii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	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	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE3403:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.text.startup
	.type	_GLOBAL__sub_I__Z15bilinear_resizePKfPfiiii, @function
_GLOBAL__sub_I__Z15bilinear_resizePKfPfiiii:
.LFB4005:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	leaq	_ZStL8__ioinit(%rip), %rbp
	movq	%rbp, %rdi
	call	_ZNSt8ios_base4InitC1Ev@PLT
	movq	_ZNSt8ios_base4InitD1Ev@GOTPCREL(%rip), %rdi
	movq	%rbp, %rsi
	popq	%rbp
	.cfi_def_cfa_offset 8
	leaq	__dso_handle(%rip), %rdx
	jmp	__cxa_atexit@PLT
	.cfi_endproc
.LFE4005:
	.size	_GLOBAL__sub_I__Z15bilinear_resizePKfPfiiii, .-_GLOBAL__sub_I__Z15bilinear_resizePKfPfiiii
	.section	.init_array
	.align 8
	.quad	_GLOBAL__sub_I__Z15bilinear_resizePKfPfiiii
	.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
	.local	_ZStL8__ioinit
	.comm	_ZStL8__ioinit,1,1
	.hidden	__dso_handle
	.ident	"GCC: (Ubuntu 11.4.0-1ubuntu1~22.04) 11.4.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

#include <hip/hip_runtime.h>
#include <iostream>

// HIP kernel for bilinear resize
__global__ void bilinear_resize_kernel(
    const float* input, float* output,
    int in_width, int in_height,
    int out_width, int out_height
) {
    int out_x = blockIdx.x * blockDim.x + threadIdx.x;
    int out_y = blockIdx.y * blockDim.y + threadIdx.y;

    if (out_x >= out_width || out_y >= out_height) return;

    double x_ratio = (double)(in_width - 1) / (out_width - 1);
    double y_ratio = (double)(in_height - 1) / (out_height - 1);

    double in_x = x_ratio * out_x;
    double in_y = y_ratio * out_y;

    int x_low = (int)floor(in_x);
    int y_low = (int)floor(in_y);
    int x_high = min(x_low + 1, in_width - 1);
    int y_high = min(y_low + 1, in_height - 1);

    double x_lerp = in_x - x_low;
    double y_lerp = in_y - y_low;

    double top_left = (double)input[y_low * in_width + x_low];
    double top_right = (double)input[y_low * in_width + x_high];
    double bottom_left = (double)input[y_high * in_width + x_low];
    double bottom_right = (double)input[y_high * in_width + x_high];

    double top = top_left + (top_right - top_left) * x_lerp;
    double bottom = bottom_left + (bottom_right - bottom_left) * x_lerp;
    double value = top + (bottom - top) * y_lerp;

    output[out_y * out_width + out_x] = (float)value;
}




void bilinear_resize(
    const float* h_input, float* h_output,
    int in_width, int in_height,
    int out_width, int out_height
) {
    float *d_input, *d_output;
    size_t in_size = in_width * in_height * sizeof(float);
    size_t out_size = out_width * out_height * sizeof(float);

    hipMalloc(&d_input, in_size);
    hipMalloc(&d_output, out_size);

    hipMemcpy(d_input, h_input, in_size, hipMemcpyHostToDevice);

    dim3 block(16, 16);
    dim3 grid((out_width + block.x - 1) / block.x, (out_height + block.y - 1) / block.y);

    hipLaunchKernelGGL(bilinear_resize_kernel, grid, block, 0, 0, d_input, d_output, in_width, in_height, out_width, out_height);

    hipMemcpy(h_output, d_output, out_size, hipMemcpyDeviceToHost);

    hipFree(d_input);
    hipFree(d_output);
}

int main() {
    const int in_width = 4, in_height = 4;
    const int out_width = 8, out_height = 8;

    float h_input[in_width * in_height] = {
        1, 2, 3, 4,
        5, 6, 7, 8,
        9,10,11,12,
        13,14,15,16
    };

    float h_output[out_width * out_height];

    bilinear_resize(h_input, h_output, in_width, in_height, out_width, out_height);

    for (int y = 0; y < out_height; ++y) {
        for (int x = 0; x < out_width; ++x) {
            std::cout << h_output[y * out_width + x] << " ";
        }
        std::cout << "\n";
    }

    return 0;
}

	.text
	.file	"26.cpp"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.globl	_Z37__device_stub__bilinear_resize_kernelPKfPfiiii # -- Begin function _Z37__device_stub__bilinear_resize_kernelPKfPfiiii
	.p2align	4, 0x90
	.type	_Z37__device_stub__bilinear_resize_kernelPKfPfiiii,@function
_Z37__device_stub__bilinear_resize_kernelPKfPfiiii: # @_Z37__device_stub__bilinear_resize_kernelPKfPfiiii
	.cfi_startproc
# %bb.0:
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 72(%rsp)
	movq	%rsi, 64(%rsp)
	movl	%edx, 12(%rsp)
	movl	%ecx, 8(%rsp)
	movl	%r8d, 4(%rsp)
	movl	%r9d, (%rsp)
	leaq	72(%rsp), %rax
	movq	%rax, 80(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 96(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 112(%rsp)
	movq	%rsp, %rax
	movq	%rax, 120(%rsp)
	leaq	48(%rsp), %rdi
	leaq	32(%rsp), %rsi
	leaq	24(%rsp), %rdx
	leaq	16(%rsp), %rcx
	callq	__hipPopCallConfiguration
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	movq	32(%rsp), %rcx
	movl	40(%rsp), %r8d
	leaq	80(%rsp), %r9
	movl	$_Z22bilinear_resize_kernelPKfPfiiii, %edi
	pushq	16(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	32(%rsp)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$152, %rsp
	.cfi_adjust_cfa_offset -152
	retq
.Lfunc_end0:
	.size	_Z37__device_stub__bilinear_resize_kernelPKfPfiiii, .Lfunc_end0-_Z37__device_stub__bilinear_resize_kernelPKfPfiiii
	.cfi_endproc
                                        # -- End function
	.globl	_Z15bilinear_resizePKfPfiiii    # -- Begin function _Z15bilinear_resizePKfPfiiii
	.p2align	4, 0x90
	.type	_Z15bilinear_resizePKfPfiiii,@function
_Z15bilinear_resizePKfPfiiii:           # @_Z15bilinear_resizePKfPfiiii
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$168, %rsp
	.cfi_def_cfa_offset 224
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%r9d, %r15d
	movl	%r8d, %r12d
	movl	%edx, %r13d
	movq	%rsi, 40(%rsp)                  # 8-byte Spill
	movq	%rdi, %rbx
	movl	%ecx, 20(%rsp)                  # 4-byte Spill
	movl	%ecx, %eax
	imull	%edx, %eax
	movslq	%eax, %rbp
	shlq	$2, %rbp
	movl	%r9d, %eax
	imull	%r8d, %eax
	movslq	%eax, %r14
	shlq	$2, %r14
	leaq	8(%rsp), %rdi
	movq	%rbp, %rsi
	callq	hipMalloc
	movq	%rsp, %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	movq	8(%rsp), %rdi
	movq	%rbx, %rsi
	movq	%rbp, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leal	15(%r12), %eax
	shrl	$4, %eax
	leal	15(%r15), %edi
	shrl	$4, %edi
	shlq	$32, %rdi
	orq	%rax, %rdi
	movabsq	$68719476752, %rdx              # imm = 0x1000000010
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_2
# %bb.1:
	movq	8(%rsp), %rax
	movq	(%rsp), %rcx
	movq	%rax, 104(%rsp)
	movq	%rcx, 96(%rsp)
	movl	%r13d, 36(%rsp)
	movl	20(%rsp), %eax                  # 4-byte Reload
	movl	%eax, 32(%rsp)
	movl	%r12d, 28(%rsp)
	movl	%r15d, 24(%rsp)
	leaq	104(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	96(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	36(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	80(%rsp), %rdi
	leaq	64(%rsp), %rsi
	leaq	56(%rsp), %rdx
	leaq	48(%rsp), %rcx
	callq	__hipPopCallConfiguration
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	movq	64(%rsp), %rcx
	movl	72(%rsp), %r8d
	leaq	112(%rsp), %r9
	movl	$_Z22bilinear_resize_kernelPKfPfiiii, %edi
	pushq	48(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	64(%rsp)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
.LBB1_2:
	movq	(%rsp), %rsi
	movq	40(%rsp), %rdi                  # 8-byte Reload
	movq	%r14, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	(%rsp), %rdi
	callq	hipFree
	addq	$168, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z15bilinear_resizePKfPfiiii, .Lfunc_end1-_Z15bilinear_resizePKfPfiiii
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function main
.LCPI2_0:
	.long	0x3f800000                      # float 1
	.long	0x40000000                      # float 2
	.long	0x40400000                      # float 3
	.long	0x40800000                      # float 4
.LCPI2_1:
	.long	0x40a00000                      # float 5
	.long	0x40c00000                      # float 6
	.long	0x40e00000                      # float 7
	.long	0x41000000                      # float 8
.LCPI2_2:
	.long	0x41100000                      # float 9
	.long	0x41200000                      # float 10
	.long	0x41300000                      # float 11
	.long	0x41400000                      # float 12
.LCPI2_3:
	.long	0x41500000                      # float 13
	.long	0x41600000                      # float 14
	.long	0x41700000                      # float 15
	.long	0x41800000                      # float 16
	.text
	.globl	main
	.p2align	4, 0x90
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	subq	$320, %rsp                      # imm = 0x140
	.cfi_def_cfa_offset 352
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movaps	.LCPI2_0(%rip), %xmm0           # xmm0 = [1.0E+0,2.0E+0,3.0E+0,4.0E+0]
	movaps	%xmm0, (%rsp)
	movaps	.LCPI2_1(%rip), %xmm0           # xmm0 = [5.0E+0,6.0E+0,7.0E+0,8.0E+0]
	movaps	%xmm0, 16(%rsp)
	movaps	.LCPI2_2(%rip), %xmm0           # xmm0 = [9.0E+0,1.0E+1,1.1E+1,1.2E+1]
	movaps	%xmm0, 32(%rsp)
	movaps	.LCPI2_3(%rip), %xmm0           # xmm0 = [1.3E+1,1.4E+1,1.5E+1,1.6E+1]
	movaps	%xmm0, 48(%rsp)
	movq	%rsp, %rdi
	leaq	64(%rsp), %rbx
	movq	%rbx, %rsi
	movl	$4, %edx
	movl	$4, %ecx
	movl	$8, %r8d
	movl	$8, %r9d
	callq	_Z15bilinear_resizePKfPfiiii
	xorl	%r14d, %r14d
	.p2align	4, 0x90
.LBB2_1:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_2 Depth 2
	xorl	%r15d, %r15d
	.p2align	4, 0x90
.LBB2_2:                                #   Parent Loop BB2_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movss	(%rbx,%r15,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
	cvtss2sd	%xmm0, %xmm0
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movl	$.L.str, %esi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	incq	%r15
	cmpq	$8, %r15
	jne	.LBB2_2
# %bb.3:                                #   in Loop: Header=BB2_1 Depth=1
	movl	$_ZSt4cout, %edi
	movl	$.L.str.1, %esi
	movl	$1, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	incq	%r14
	addq	$32, %rbx
	cmpq	$8, %r14
	jne	.LBB2_1
# %bb.4:
	xorl	%eax, %eax
	addq	$320, %rsp                      # imm = 0x140
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	main, .Lfunc_end2-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	cmpq	$0, __hip_gpubin_handle(%rip)
	jne	.LBB3_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB3_2:
	movq	__hip_gpubin_handle(%rip), %rdi
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z22bilinear_resize_kernelPKfPfiiii, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end3:
	.size	__hip_module_ctor, .Lfunc_end3-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB4_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB4_2:
	retq
.Lfunc_end4:
	.size	__hip_module_dtor, .Lfunc_end4-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z22bilinear_resize_kernelPKfPfiiii,@object # @_Z22bilinear_resize_kernelPKfPfiiii
	.section	.rodata,"a",@progbits
	.globl	_Z22bilinear_resize_kernelPKfPfiiii
	.p2align	3, 0x0
_Z22bilinear_resize_kernelPKfPfiiii:
	.quad	_Z37__device_stub__bilinear_resize_kernelPKfPfiiii
	.size	_Z22bilinear_resize_kernelPKfPfiiii, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	" "
	.size	.L.str, 2

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"\n"
	.size	.L.str.1, 2

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z22bilinear_resize_kernelPKfPfiiii"
	.size	.L__unnamed_1, 36

	.type	__hip_fatbin_wrapper,@object    # @__hip_fatbin_wrapper
	.section	.hipFatBinSegment,"a",@progbits
	.p2align	3, 0x0
__hip_fatbin_wrapper:
	.long	1212764230                      # 0x48495046
	.long	1                               # 0x1
	.quad	__hip_fatbin
	.quad	0
	.size	__hip_fatbin_wrapper, 24

	.type	__hip_gpubin_handle,@object     # @__hip_gpubin_handle
	.local	__hip_gpubin_handle
	.comm	__hip_gpubin_handle,8,8
	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	__hip_module_ctor
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.section	".linker-options","e",@llvm_linker_options
	.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 _Z37__device_stub__bilinear_resize_kernelPKfPfiiii
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z22bilinear_resize_kernelPKfPfiiii
	.addrsig_sym _ZSt4cout
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z22bilinear_resize_kernelPKfPfiiii ; -- Begin function _Z22bilinear_resize_kernelPKfPfiiii
	.globl	_Z22bilinear_resize_kernelPKfPfiiii
	.p2align	8
	.type	_Z22bilinear_resize_kernelPKfPfiiii,@function
_Z22bilinear_resize_kernelPKfPfiiii:    ; @_Z22bilinear_resize_kernelPKfPfiiii
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x2c
	s_load_b64 s[4:5], s[0:1], 0x18
	v_and_b32_e32 v2, 0x3ff, v0
	v_bfe_u32 v3, v0, 10, 10
	s_waitcnt lgkmcnt(0)
	s_and_b32 s3, s2, 0xffff
	s_lshr_b32 s2, s2, 16
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mad_u64_u32 v[0:1], null, s14, s3, v[2:3]
	v_mad_u64_u32 v[1:2], null, s15, s2, v[3:4]
	v_cmp_gt_i32_e32 vcc_lo, s4, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_gt_i32_e64 s2, s5, v1
	s_and_b32 s2, vcc_lo, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s3, s2
	s_cbranch_execz .LBB0_2
; %bb.1:
	s_load_b64 s[6:7], s[0:1], 0x10
	s_add_i32 s2, s4, -1
	s_add_i32 s3, s5, -1
	v_cvt_f64_i32_e32 v[2:3], s2
	v_cvt_f64_i32_e32 v[6:7], s3
	s_waitcnt lgkmcnt(0)
	s_add_i32 s5, s6, -1
	s_add_i32 s7, s7, -1
	v_cvt_f64_i32_e32 v[4:5], s5
	v_cvt_f64_i32_e32 v[8:9], s7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_div_scale_f64 v[10:11], null, v[2:3], v[2:3], v[4:5]
	v_div_scale_f64 v[12:13], null, v[6:7], v[6:7], v[8:9]
	v_div_scale_f64 v[22:23], vcc_lo, v[4:5], v[2:3], v[4:5]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_rcp_f64_e32 v[14:15], v[10:11]
	v_rcp_f64_e32 v[16:17], v[12:13]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[18:19], -v[10:11], v[14:15], 1.0
	v_fma_f64 v[20:21], -v[12:13], v[16:17], 1.0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[14:15], v[14:15], v[18:19], v[14:15]
	v_fma_f64 v[16:17], v[16:17], v[20:21], v[16:17]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[18:19], -v[10:11], v[14:15], 1.0
	v_fma_f64 v[20:21], -v[12:13], v[16:17], 1.0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_fma_f64 v[14:15], v[14:15], v[18:19], v[14:15]
	v_div_scale_f64 v[18:19], s2, v[8:9], v[6:7], v[8:9]
	v_fma_f64 v[16:17], v[16:17], v[20:21], v[16:17]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f64 v[20:21], v[22:23], v[14:15]
	v_mul_f64 v[24:25], v[18:19], v[16:17]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[10:11], -v[10:11], v[20:21], v[22:23]
	v_fma_f64 v[12:13], -v[12:13], v[24:25], v[18:19]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_div_fmas_f64 v[10:11], v[10:11], v[14:15], v[20:21]
	s_mov_b32 vcc_lo, s2
	v_cvt_f64_i32_e32 v[14:15], v1
	v_div_fmas_f64 v[12:13], v[12:13], v[16:17], v[24:25]
	s_load_b128 s[0:3], s[0:1], 0x0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_div_fixup_f64 v[2:3], v[10:11], v[2:3], v[4:5]
	v_div_fixup_f64 v[6:7], v[12:13], v[6:7], v[8:9]
	v_cvt_f64_i32_e32 v[8:9], v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f64 v[4:5], v[6:7], v[14:15]
	v_mul_f64 v[10:11], v[2:3], v[8:9]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_floor_f64_e32 v[4:5], v[4:5]
	v_floor_f64_e32 v[10:11], v[10:11]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cvt_i32_f64_e32 v20, v[4:5]
	v_cvt_i32_f64_e32 v18, v[10:11]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v4, 1, v20
	v_mul_lo_u32 v5, v20, s6
	v_add_nc_u32_e32 v10, 1, v18
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_min_i32_e32 v11, s7, v4
	v_min_i32_e32 v13, s5, v10
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v4, v5, v18
	v_mul_lo_u32 v16, v11, s6
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v10, v5, v13
	v_ashrrev_i32_e32 v5, 31, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_ashrrev_i32_e32 v11, 31, v10
	v_add_nc_u32_e32 v12, v16, v18
	v_add_nc_u32_e32 v16, v16, v13
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_lshlrev_b64 v[4:5], 2, v[4:5]
	v_lshlrev_b64 v[10:11], 2, v[10:11]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_ashrrev_i32_e32 v13, 31, v12
	v_ashrrev_i32_e32 v17, 31, v16
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v4, vcc_lo, s0, v4
	v_lshlrev_b64 v[12:13], 2, v[12:13]
	v_add_co_ci_u32_e32 v5, vcc_lo, s1, v5, vcc_lo
	v_add_co_u32 v10, vcc_lo, s0, v10
	v_lshlrev_b64 v[16:17], 2, v[16:17]
	v_add_co_ci_u32_e32 v11, vcc_lo, s1, v11, vcc_lo
	v_add_co_u32 v12, vcc_lo, s0, v12
	v_add_co_ci_u32_e32 v13, vcc_lo, s1, v13, vcc_lo
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_u32 v16, vcc_lo, s0, v16
	v_add_co_ci_u32_e32 v17, vcc_lo, s1, v17, vcc_lo
	s_clause 0x3
	global_load_b32 v19, v[4:5], off
	global_load_b32 v21, v[10:11], off
	global_load_b32 v22, v[12:13], off
	global_load_b32 v23, v[16:17], off
	v_cvt_f64_i32_e32 v[4:5], v18
	s_delay_alu instid0(VALU_DEP_1)
	v_fma_f64 v[2:3], v[2:3], v[8:9], -v[4:5]
	s_waitcnt vmcnt(3)
	v_cvt_f64_f32_e32 v[10:11], v19
	s_waitcnt vmcnt(2)
	v_cvt_f64_f32_e32 v[12:13], v21
	s_waitcnt vmcnt(1)
	v_cvt_f64_f32_e32 v[16:17], v22
	s_waitcnt vmcnt(0)
	v_cvt_f64_f32_e32 v[18:19], v23
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_f64 v[4:5], v[12:13], -v[10:11]
	v_cvt_f64_i32_e32 v[12:13], v20
	v_add_f64 v[8:9], v[18:19], -v[16:17]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[4:5], v[2:3], v[4:5], v[10:11]
	v_fma_f64 v[6:7], v[6:7], v[14:15], -v[12:13]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[2:3], v[2:3], v[8:9], v[16:17]
	v_add_f64 v[2:3], v[2:3], -v[4:5]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[2:3], v[6:7], v[2:3], v[4:5]
	v_cvt_f32_f64_e32 v4, v[2:3]
	v_mad_u64_u32 v[2:3], null, v1, s4, v[0:1]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	v_lshlrev_b64 v[0:1], 2, v[2:3]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s2, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s3, v1, vcc_lo
	global_store_b32 v[0:1], v4, off
.LBB0_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z22bilinear_resize_kernelPKfPfiiii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.amdhsa_user_sgpr_count 14
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 26
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z22bilinear_resize_kernelPKfPfiiii, .Lfunc_end0-_Z22bilinear_resize_kernelPKfPfiiii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 828
; NumSgprs: 18
; NumVgprs: 26
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 3
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 26
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.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_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         20
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         28
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z22bilinear_resize_kernelPKfPfiiii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z22bilinear_resize_kernelPKfPfiiii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     26
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
