/*
Copyright 2017 the arraydiff authors

Licensed under the Apache License, Version 2.0 (the "License");
you may not use this file except in compliance with the License.
You may obtain a copy of the License at

    http://www.apache.org/licenses/LICENSE-2.0

Unless required by applicable law or agreed to in writing, software
distributed under the License is distributed on an "AS IS" BASIS,
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and
limitations under the License.

Caffe copyright:

COPYRIGHT

All contributions by the University of California:
Copyright (c) 2014-2017 The Regents of the University of California (Regents)
All rights reserved.

All other contributions:
Copyright (c) 2014-2017, the respective contributors
All rights reserved.

Caffe uses a shared copyright model: each contributor holds copyright over
their contributions to Caffe. The project versioning records all such
contribution and copyright details. If a contributor wants to further mark
their specific copyright on a particular contribution, they should indicate
their copyright solely in the commit message of the change when it is
committed.

LICENSE

Redistribution and use in source and binary forms, with or without
modification, are permitted provided that the following conditions are met: 

1. Redistributions of source code must retain the above copyright notice, this
   list of conditions and the following disclaimer. 
2. Redistributions in binary form must reproduce the above copyright notice,
   this list of conditions and the following disclaimer in the documentation
   and/or other materials provided with the distribution. 

THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR
ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
(INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*/

#include "common.cuh"
#include <cuda_runtime_api.h>
#include <math_constants.h>
#include <stdint.h>

__global__ void max_pool_fwd_f32_kernel(
    const uint32_t nthreads,
    const float* const bottom_data,
    const int num,
    const int channels,
    const int height,
    const int width,
    const int pooled_height,
    const int pooled_width,
    const int kernel_h,
    const int kernel_w,
    const int stride_h,
    const int stride_w,
    const int pad_h,
    const int pad_w,
    float* const top_data,
    int32_t* const mask)
{
  uint32_t index = threadIdx.x + blockDim.x * blockIdx.x;
  if (index < nthreads) {
    const int pw = index % pooled_width;
    const int ph = (index / pooled_width) % pooled_height;
    const int c = (index / pooled_width / pooled_height) % channels;
    const int n = index / pooled_width / pooled_height / channels;
    int hstart = ph * stride_h - pad_h;
    int wstart = pw * stride_w - pad_w;
    const int hend = min(hstart + kernel_h, height);
    const int wend = min(wstart + kernel_w, width);
    hstart = max(hstart, 0);
    wstart = max(wstart, 0);
    float maxval = -CUDART_INF_F;
    int maxidx = -1;
    const float* const bottom_slice =
        bottom_data + (n * channels + c) * height * width;
    for (int h = hstart; h < hend; ++h) {
      for (int w = wstart; w < wend; ++w) {
        if (bottom_slice[h * width + w] > maxval) {
          maxidx = h * width + w;
          maxval = bottom_slice[maxidx];
        }
      }
    }
    if (NULL != top_data) {
      top_data[index] = maxval;
    }
    if (NULL != mask) {
      mask[index] = maxidx;
    }
  }
}

extern "C" void arraydiff_cuda_kernel_max_pool_fwd_f32(
    size_t x_w, size_t x_h, size_t chan_dim, size_t batch_sz,
    size_t y_w, size_t y_h,
    size_t kernel_w, size_t kernel_h,
    size_t stride_w, size_t stride_h,
    size_t pad_w, size_t pad_h,
    const float *x,
    float *maybe_y,
    int32_t *maybe_mask,
    cudaStream_t stream)
{
  uint32_t n = y_w * y_h * chan_dim * batch_sz;
  max_pool_fwd_f32_kernel<<<(n+1024-1)/1024, 1024, 0, stream>>>(
      n, x,
      batch_sz, chan_dim, x_h, x_w,
      y_h, y_w,
      kernel_h, kernel_w,
      stride_h, stride_w,
      pad_h, pad_w,
      maybe_y, maybe_mask);
}

__global__ void max_pool_bwd_f32_kernel(
    const uint32_t nthreads,
    const float* const top_diff,
    const int32_t* const mask,
    const int num,
    const int channels,
    const int height,
    const int width,
    const int pooled_height,
    const int pooled_width,
    const int kernel_h,
    const int kernel_w,
    const int stride_h,
    const int stride_w,
    const int pad_h,
    const int pad_w,
    float* const bottom_diff)
{
  uint32_t index = threadIdx.x + blockDim.x * blockIdx.x;
  if (index < nthreads) {
    // find out the local index
    // find out the local offset
    const int w = index % width;
    const int h = (index / width) % height;
    const int c = (index / width / height) % channels;
    const int n = index / width / height / channels;
    const int phstart =
         (h + pad_h < kernel_h) ? 0 : (h + pad_h - kernel_h) / stride_h + 1;
    const int phend = min((h + pad_h) / stride_h + 1, pooled_height);
    const int pwstart =
         (w + pad_w < kernel_w) ? 0 : (w + pad_w - kernel_w) / stride_w + 1;
    const int pwend = min((w + pad_w) / stride_w + 1, pooled_width);
    float gradient = 0.0f;
    const int offset = (n * channels + c) * pooled_height * pooled_width;
    const float* const top_diff_slice = top_diff + offset;
    const int* const mask_slice = mask + offset;
    for (int ph = phstart; ph < phend; ++ph) {
      for (int pw = pwstart; pw < pwend; ++pw) {
        if (mask_slice[ph * pooled_width + pw] == h * width + w) {
          gradient += top_diff_slice[ph * pooled_width + pw];
        }
      }
    }
    bottom_diff[index] += gradient;
  }
}

extern "C" void arraydiff_cuda_kernel_max_pool_bwd_f32(
    size_t x_w, size_t x_h, size_t chan_dim, size_t batch_sz,
    size_t y_w, size_t y_h,
    size_t kernel_w, size_t kernel_h,
    size_t stride_w, size_t stride_h,
    size_t pad_w, size_t pad_h,
    const float *dy,
    const int32_t *mask,
    float *dx,
    cudaStream_t stream)
{
  uint32_t n = x_w * x_h * chan_dim * batch_sz;
  max_pool_bwd_f32_kernel<<<(n+1024-1)/1024, 1024, 0, stream>>>(
      n, dy, mask,
      batch_sz, chan_dim, x_h, x_w,
      y_h, y_w,
      kernel_h, kernel_w,
      stride_h, stride_w,
      pad_h, pad_w,
      dx);
}

__global__ void avg_pool_fwd_f32_kernel(
    const uint32_t nthreads,
    const float* const bottom_data,
    const int num,
    const int channels,
    const int height,
    const int width,
    const int pooled_height,
    const int pooled_width,
    const int kernel_h,
    const int kernel_w,
    const int stride_h,
    const int stride_w,
    const int pad_h,
    const int pad_w,
    float* const top_data)
{
  uint32_t index = threadIdx.x + blockDim.x * blockIdx.x;
  if (index < nthreads) {
    const int pw = index % pooled_width;
    const int ph = (index / pooled_width) % pooled_height;
    const int c = (index / pooled_width / pooled_height) % channels;
    const int n = index / pooled_width / pooled_height / channels;
    int hstart = ph * stride_h - pad_h;
    int wstart = pw * stride_w - pad_w;
    int hend = min(hstart + kernel_h, height + pad_h);
    int wend = min(wstart + kernel_w, width + pad_w);
    const int pool_size = (hend - hstart) * (wend - wstart);
    hstart = max(hstart, 0);
    wstart = max(wstart, 0);
    hend = min(hend, height);
    wend = min(wend, width);
    float aveval = 0.0f;
    const float* const bottom_slice =
        bottom_data + (n * channels + c) * height * width;
    for (int h = hstart; h < hend; ++h) {
      for (int w = wstart; w < wend; ++w) {
        aveval += bottom_slice[h * width + w];
      }
    }
    top_data[index] = aveval / pool_size;
  }
}

extern "C" void arraydiff_cuda_kernel_avg_pool_fwd_f32(
    size_t x_w, size_t x_h, size_t chan_dim, size_t batch_sz,
    size_t y_w, size_t y_h,
    size_t kernel_w, size_t kernel_h,
    size_t stride_w, size_t stride_h,
    size_t pad_w, size_t pad_h,
    const float *x,
    float *y,
    cudaStream_t stream)
{
  uint32_t n = y_w * y_h * chan_dim * batch_sz;
  avg_pool_fwd_f32_kernel<<<(n+1024-1)/1024, 1024, 0, stream>>>(
      n, x,
      batch_sz, chan_dim, x_h, x_w,
      y_h, y_w,
      kernel_h, kernel_w,
      stride_h, stride_w,
      pad_h, pad_w,
      y);
}

__global__ void avg_pool_bwd_f32_kernel(
    const uint32_t nthreads,
    const float* const top_diff,
    const int num,
    const int channels,
    const int height,
    const int width,
    const int pooled_height,
    const int pooled_width,
    const int kernel_h,
    const int kernel_w,
    const int stride_h,
    const int stride_w,
    const int pad_h,
    const int pad_w,
    float* const bottom_diff)
{
  uint32_t index = threadIdx.x + blockDim.x * blockIdx.x;
  if (index < nthreads) {
    // find out the local index
    // find out the local offset
    const int w = index % width + pad_w;
    const int h = (index / width) % height + pad_h;
    const int c = (index / width / height) % channels;
    const int n = index / width / height / channels;
    const int phstart = (h < kernel_h) ? 0 : (h - kernel_h) / stride_h + 1;
    const int phend = min(h / stride_h + 1, pooled_height);
    const int pwstart = (w < kernel_w) ? 0 : (w - kernel_w) / stride_w + 1;
    const int pwend = min(w / stride_w + 1, pooled_width);
    float gradient = 0.0f;
    const float* const top_diff_slice =
        top_diff + (n * channels + c) * pooled_height * pooled_width;
    for (int ph = phstart; ph < phend; ++ph) {
      for (int pw = pwstart; pw < pwend; ++pw) {
        // figure out the pooling size
        int hstart = ph * stride_h - pad_h;
        int wstart = pw * stride_w - pad_w;
        int hend = min(hstart + kernel_h, height + pad_h);
        int wend = min(wstart + kernel_w, width + pad_w);
        int pool_size = (hend - hstart) * (wend - wstart);
        gradient += top_diff_slice[ph * pooled_width + pw] / pool_size;
      }
    }
    bottom_diff[index] += gradient;
  }
}

extern "C" void arraydiff_cuda_kernel_avg_pool_bwd_f32(
    size_t x_w, size_t x_h, size_t chan_dim, size_t batch_sz,
    size_t y_w, size_t y_h,
    size_t kernel_w, size_t kernel_h,
    size_t stride_w, size_t stride_h,
    size_t pad_w, size_t pad_h,
    const float *dy,
    float *dx,
    cudaStream_t stream)
{
  uint32_t n = x_w * x_h * chan_dim * batch_sz;
  avg_pool_bwd_f32_kernel<<<(n+1024-1)/1024, 1024, 0, stream>>>(
      n, dy,
      batch_sz, chan_dim, x_h, x_w,
      y_h, y_w,
      kernel_h, kernel_w,
      stride_h, stride_w,
      pad_h, pad_w,
      dx);
}

	.file	"tmpxft_003bb10b_00000000-6_pool.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2016:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2016:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z61__device_stub__Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPijPKfiiiiiiiiiiiiPfPi
	.type	_Z61__device_stub__Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPijPKfiiiiiiiiiiiiPfPi, @function
_Z61__device_stub__Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPijPKfiiiiiiiiiiiiPfPi:
.LFB2038:
	.cfi_startproc
	endbr64
	subq	$264, %rsp
	.cfi_def_cfa_offset 272
	movq	336(%rsp), %rax
	movl	%edi, 44(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rax, 8(%rsp)
	movq	344(%rsp), %rax
	movl	%edx, 40(%rsp)
	leaq	56(%rsp), %rdx
	movl	%ecx, 28(%rsp)
	leaq	64(%rsp), %rcx
	movl	%r8d, 24(%rsp)
	movl	%r9d, 20(%rsp)
	movq	%rax, (%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	40(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	20(%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	8(%rsp), %rax
	movq	%rax, 232(%rsp)
	movq	%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	.L2
	pushq	64(%rsp)
	.cfi_def_cfa_offset 280
	leaq	_Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi(%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
.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
.LFE2038:
	.size	_Z61__device_stub__Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPijPKfiiiiiiiiiiiiPfPi, .-_Z61__device_stub__Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPijPKfiiiiiiiiiiiiPfPi
	.globl	_Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi
	.type	_Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi, @function
_Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi:
.LFB2039:
	.cfi_startproc
	endbr64
	jmp	_Z61__device_stub__Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPijPKfiiiiiiiiiiiiPfPi
	.cfi_endproc
.LFE2039:
	.size	_Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi, .-_Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi
	.globl	arraydiff_cuda_kernel_max_pool_fwd_f32
	.type	arraydiff_cuda_kernel_max_pool_fwd_f32, @function
arraydiff_cuda_kernel_max_pool_fwd_f32:
.LFB2010:
	.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
	movq	%r8, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movq	%r9, %r13
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdx, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rcx, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%edx, %ebx
	movl	$4194305, %edx
	imull	%ecx, %ebx
	salq	$10, %rdx
	movl	$1, %ecx
	subq	$120, %rsp
	.cfi_def_cfa_offset 176
	movq	176(%rsp), %rax
	imull	%r9d, %ebx
	movq	%rdi, (%rsp)
	movq	240(%rsp), %r11
	movq	%rsi, 8(%rsp)
	movq	%rax, 16(%rsp)
	movq	192(%rsp), %rax
	imull	%r8d, %ebx
	movq	184(%rsp), %r15
	xorl	%r8d, %r8d
	movq	%rax, 24(%rsp)
	movq	200(%rsp), %rax
	movq	%rax, 32(%rsp)
	movq	208(%rsp), %rax
	movq	%rax, 40(%rsp)
	movq	216(%rsp), %rax
	movq	%rax, 48(%rsp)
	movq	224(%rsp), %rax
	movq	%rax, 56(%rsp)
	movq	232(%rsp), %rax
	movq	%rax, 64(%rsp)
	leal	1023(%rbx), %eax
	shrl	$10, %eax
	movq	%r11, 72(%rsp)
	movq	248(%rsp), %r9
	movl	%eax, 88(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 92(%rsp)
	movq	88(%rsp), %rdi
	movl	96(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	movq	72(%rsp), %r11
	testl	%eax, %eax
	jne	.L8
	movl	%r14d, 184(%rsp)
	movl	(%rsp), %r9d
	movl	%ebp, %edx
	movl	%ebx, %edi
	movq	64(%rsp), %rax
	movl	8(%rsp), %r8d
	movl	%r12d, %ecx
	movl	%r13d, 176(%rsp)
	movq	%r11, 248(%rsp)
	movq	56(%rsp), %rsi
	movq	%rax, 240(%rsp)
	movl	40(%rsp), %eax
	movl	%r15d, 192(%rsp)
	movl	%eax, 232(%rsp)
	movl	48(%rsp), %eax
	movl	%eax, 224(%rsp)
	movl	24(%rsp), %eax
	movl	%eax, 216(%rsp)
	movl	32(%rsp), %eax
	movl	%eax, 208(%rsp)
	movl	16(%rsp), %eax
	movl	%eax, 200(%rsp)
	addq	$120, %rsp
	.cfi_remember_state
	.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
	jmp	_Z61__device_stub__Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPijPKfiiiiiiiiiiiiPfPi
.L8:
	.cfi_restore_state
	addq	$120, %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
.LFE2010:
	.size	arraydiff_cuda_kernel_max_pool_fwd_f32, .-arraydiff_cuda_kernel_max_pool_fwd_f32
	.globl	_Z62__device_stub__Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPfjPKfPKiiiiiiiiiiiiiPf
	.type	_Z62__device_stub__Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPfjPKfPKiiiiiiiiiiiiiPf, @function
_Z62__device_stub__Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPfjPKfPKiiiiiiiiiiiiiPf:
.LFB2040:
	.cfi_startproc
	endbr64
	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	.L11
	pushq	64(%rsp)
	.cfi_def_cfa_offset 280
	leaq	_Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf(%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
.L11:
	movq	248(%rsp), %rax
	subq	%fs:40, %rax
	je	.L13
	call	__stack_chk_fail@PLT
.L13:
	addq	$264, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2040:
	.size	_Z62__device_stub__Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPfjPKfPKiiiiiiiiiiiiiPf, .-_Z62__device_stub__Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPfjPKfPKiiiiiiiiiiiiiPf
	.globl	_Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf
	.type	_Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf, @function
_Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf:
.LFB2041:
	.cfi_startproc
	endbr64
	jmp	_Z62__device_stub__Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPfjPKfPKiiiiiiiiiiiiiPf
	.cfi_endproc
.LFE2041:
	.size	_Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf, .-_Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf
	.globl	arraydiff_cuda_kernel_max_pool_bwd_f32
	.type	arraydiff_cuda_kernel_max_pool_bwd_f32, @function
arraydiff_cuda_kernel_max_pool_bwd_f32:
.LFB2011:
	.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
	movq	%rdi, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movq	%rsi, %r13
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdx, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rcx, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%edx, %ebx
	movl	$4194305, %edx
	imull	%ecx, %ebx
	salq	$10, %rdx
	movl	$1, %ecx
	subq	$120, %rsp
	.cfi_def_cfa_offset 176
	movq	176(%rsp), %rax
	imull	%esi, %ebx
	movq	%r8, (%rsp)
	xorl	%r8d, %r8d
	movq	208(%rsp), %r10
	movq	240(%rsp), %r11
	movq	%r9, 8(%rsp)
	movq	%rax, 16(%rsp)
	movq	184(%rsp), %rax
	imull	%edi, %ebx
	movq	%r10, 72(%rsp)
	movq	216(%rsp), %r15
	movq	%rax, 24(%rsp)
	movq	192(%rsp), %rax
	movq	%rax, 32(%rsp)
	movq	200(%rsp), %rax
	movq	%rax, 40(%rsp)
	movq	224(%rsp), %rax
	movq	%rax, 48(%rsp)
	movq	232(%rsp), %rax
	movq	%rax, 56(%rsp)
	leal	1023(%rbx), %eax
	shrl	$10, %eax
	movq	%r11, 64(%rsp)
	movq	248(%rsp), %r9
	movl	%eax, 88(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 92(%rsp)
	movq	88(%rsp), %rdi
	movl	96(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	movq	64(%rsp), %r11
	movq	72(%rsp), %r10
	testl	%eax, %eax
	jne	.L16
	movl	32(%rsp), %eax
	movl	%r15d, 232(%rsp)
	movl	%ebp, %ecx
	movl	%ebx, %edi
	movq	56(%rsp), %rdx
	movq	48(%rsp), %rsi
	movl	%r13d, %r9d
	movl	%r12d, %r8d
	movl	%eax, 224(%rsp)
	movl	40(%rsp), %eax
	movl	%r14d, 176(%rsp)
	movl	%eax, 216(%rsp)
	movl	16(%rsp), %eax
	movq	%r11, 248(%rsp)
	movl	%eax, 208(%rsp)
	movl	24(%rsp), %eax
	movl	%r10d, 240(%rsp)
	movl	%eax, 200(%rsp)
	movl	(%rsp), %eax
	movl	%eax, 192(%rsp)
	movl	8(%rsp), %eax
	movl	%eax, 184(%rsp)
	addq	$120, %rsp
	.cfi_remember_state
	.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
	jmp	_Z62__device_stub__Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPfjPKfPKiiiiiiiiiiiiiPf
.L16:
	.cfi_restore_state
	addq	$120, %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
.LFE2011:
	.size	arraydiff_cuda_kernel_max_pool_bwd_f32, .-arraydiff_cuda_kernel_max_pool_bwd_f32
	.globl	_Z59__device_stub__Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfjPKfiiiiiiiiiiiiPf
	.type	_Z59__device_stub__Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfjPKfiiiiiiiiiiiiPf, @function
_Z59__device_stub__Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfjPKfiiiiiiiiiiiiPf:
.LFB2042:
	.cfi_startproc
	endbr64
	subq	$248, %rsp
	.cfi_def_cfa_offset 256
	movq	320(%rsp), %rax
	movl	%edi, 44(%rsp)
	leaq	64(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	76(%rsp), %rsi
	movl	%edx, 40(%rsp)
	leaq	48(%rsp), %rdx
	movl	%ecx, 28(%rsp)
	leaq	56(%rsp), %rcx
	movl	%r8d, 24(%rsp)
	movl	%r9d, 20(%rsp)
	movq	%rax, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 232(%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	40(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	256(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	264(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	272(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	280(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	288(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	296(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	304(%rsp), %rax
	movq	%rax, 208(%rsp)
	leaq	312(%rsp), %rax
	movq	%rax, 216(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 224(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L19
	pushq	56(%rsp)
	.cfi_def_cfa_offset 264
	leaq	_Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 272
	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 264
	popq	%rdx
	.cfi_def_cfa_offset 256
.L19:
	movq	232(%rsp), %rax
	subq	%fs:40, %rax
	je	.L21
	call	__stack_chk_fail@PLT
.L21:
	addq	$248, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2042:
	.size	_Z59__device_stub__Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfjPKfiiiiiiiiiiiiPf, .-_Z59__device_stub__Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfjPKfiiiiiiiiiiiiPf
	.globl	_Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf
	.type	_Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf, @function
_Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf:
.LFB2043:
	.cfi_startproc
	endbr64
	jmp	_Z59__device_stub__Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfjPKfiiiiiiiiiiiiPf
	.cfi_endproc
.LFE2043:
	.size	_Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf, .-_Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf
	.globl	arraydiff_cuda_kernel_avg_pool_fwd_f32
	.type	arraydiff_cuda_kernel_avg_pool_fwd_f32, @function
arraydiff_cuda_kernel_avg_pool_fwd_f32:
.LFB2012:
	.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
	movq	%r8, %r13
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%r9, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rdx, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rcx, %rbx
	subq	$120, %rsp
	.cfi_def_cfa_offset 176
	movq	176(%rsp), %rax
	movq	208(%rsp), %r10
	movq	%rdi, 8(%rsp)
	movq	232(%rsp), %r11
	movq	%rsi, 16(%rsp)
	movq	%rax, 24(%rsp)
	movq	192(%rsp), %rax
	movq	%r10, 72(%rsp)
	movq	184(%rsp), %r15
	movq	%rax, 32(%rsp)
	movq	200(%rsp), %rax
	movq	%r11, 64(%rsp)
	movq	%rax, 40(%rsp)
	movq	216(%rsp), %rax
	movq	%rax, 48(%rsp)
	movq	224(%rsp), %rax
	movq	%rax, 56(%rsp)
	movl	%edx, %eax
	movl	$4194305, %edx
	imull	%ecx, %eax
	salq	$10, %rdx
	movl	$1, %ecx
	imull	%r9d, %eax
	imull	%r8d, %eax
	xorl	%r8d, %r8d
	movl	%eax, %r14d
	addl	$1023, %eax
	shrl	$10, %eax
	movl	%eax, 88(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 92(%rsp)
	movq	240(%rsp), %r9
	movq	88(%rsp), %rdi
	movl	96(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	movq	64(%rsp), %r11
	movq	72(%rsp), %r10
	testl	%eax, %eax
	jne	.L24
	movl	48(%rsp), %eax
	movl	%r15d, 192(%rsp)
	movl	%ebp, %ecx
	movl	%ebx, %edx
	movl	%r13d, 184(%rsp)
	movl	8(%rsp), %r9d
	movl	%r14d, %edi
	movl	%eax, 224(%rsp)
	movl	32(%rsp), %eax
	movl	%r12d, 176(%rsp)
	movl	16(%rsp), %r8d
	movl	%eax, 216(%rsp)
	movl	40(%rsp), %eax
	movq	%r11, 240(%rsp)
	movq	56(%rsp), %rsi
	movl	%eax, 208(%rsp)
	movl	24(%rsp), %eax
	movl	%r10d, 232(%rsp)
	movl	%eax, 200(%rsp)
	addq	$120, %rsp
	.cfi_remember_state
	.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
	jmp	_Z59__device_stub__Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfjPKfiiiiiiiiiiiiPf
.L24:
	.cfi_restore_state
	addq	$120, %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
.LFE2012:
	.size	arraydiff_cuda_kernel_avg_pool_fwd_f32, .-arraydiff_cuda_kernel_avg_pool_fwd_f32
	.globl	_Z59__device_stub__Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPfjPKfiiiiiiiiiiiiPf
	.type	_Z59__device_stub__Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPfjPKfiiiiiiiiiiiiPf, @function
_Z59__device_stub__Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPfjPKfiiiiiiiiiiiiPf:
.LFB2044:
	.cfi_startproc
	endbr64
	subq	$248, %rsp
	.cfi_def_cfa_offset 256
	movq	320(%rsp), %rax
	movl	%edi, 44(%rsp)
	leaq	64(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	76(%rsp), %rsi
	movl	%edx, 40(%rsp)
	leaq	48(%rsp), %rdx
	movl	%ecx, 28(%rsp)
	leaq	56(%rsp), %rcx
	movl	%r8d, 24(%rsp)
	movl	%r9d, 20(%rsp)
	movq	%rax, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 232(%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	40(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	256(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	264(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	272(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	280(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	288(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	296(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	304(%rsp), %rax
	movq	%rax, 208(%rsp)
	leaq	312(%rsp), %rax
	movq	%rax, 216(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 224(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L27
	pushq	56(%rsp)
	.cfi_def_cfa_offset 264
	leaq	_Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 272
	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 264
	popq	%rdx
	.cfi_def_cfa_offset 256
.L27:
	movq	232(%rsp), %rax
	subq	%fs:40, %rax
	je	.L29
	call	__stack_chk_fail@PLT
.L29:
	addq	$248, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2044:
	.size	_Z59__device_stub__Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPfjPKfiiiiiiiiiiiiPf, .-_Z59__device_stub__Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPfjPKfiiiiiiiiiiiiPf
	.globl	_Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf
	.type	_Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf, @function
_Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf:
.LFB2045:
	.cfi_startproc
	endbr64
	jmp	_Z59__device_stub__Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPfjPKfiiiiiiiiiiiiPf
	.cfi_endproc
.LFE2045:
	.size	_Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf, .-_Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf
	.globl	arraydiff_cuda_kernel_avg_pool_bwd_f32
	.type	arraydiff_cuda_kernel_avg_pool_bwd_f32, @function
arraydiff_cuda_kernel_avg_pool_bwd_f32:
.LFB2013:
	.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
	movq	%rdi, %r13
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rsi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rdx, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rcx, %rbx
	subq	$120, %rsp
	.cfi_def_cfa_offset 176
	movq	176(%rsp), %rax
	movq	208(%rsp), %r10
	movq	%r8, 16(%rsp)
	xorl	%r8d, %r8d
	movq	232(%rsp), %r11
	movq	%r9, 24(%rsp)
	movq	%rax, 32(%rsp)
	movq	184(%rsp), %rax
	movq	%r10, 72(%rsp)
	movq	192(%rsp), %r14
	movq	%rax, 40(%rsp)
	movq	200(%rsp), %rax
	movq	%r11, 64(%rsp)
	movq	216(%rsp), %r15
	movq	%rax, 48(%rsp)
	movq	224(%rsp), %rax
	movq	%rax, 56(%rsp)
	movl	%edx, %eax
	movl	$4194305, %edx
	imull	%ecx, %eax
	salq	$10, %rdx
	movl	$1, %ecx
	imull	%esi, %eax
	imull	%edi, %eax
	movl	%eax, 12(%rsp)
	addl	$1023, %eax
	shrl	$10, %eax
	movl	%eax, 88(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 92(%rsp)
	movq	240(%rsp), %r9
	movq	88(%rsp), %rdi
	movl	96(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	movq	64(%rsp), %r11
	movq	72(%rsp), %r10
	testl	%eax, %eax
	jne	.L32
	movl	48(%rsp), %eax
	movl	%r15d, 224(%rsp)
	movl	%ebp, %ecx
	movl	%ebx, %edx
	movq	56(%rsp), %rsi
	movl	12(%rsp), %edi
	movl	%r13d, %r9d
	movl	%r12d, %r8d
	movl	%eax, 208(%rsp)
	movl	32(%rsp), %eax
	movl	%r14d, 216(%rsp)
	movl	%eax, 200(%rsp)
	movl	40(%rsp), %eax
	movq	%r11, 240(%rsp)
	movl	%eax, 192(%rsp)
	movl	16(%rsp), %eax
	movl	%r10d, 232(%rsp)
	movl	%eax, 184(%rsp)
	movl	24(%rsp), %eax
	movl	%eax, 176(%rsp)
	addq	$120, %rsp
	.cfi_remember_state
	.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
	jmp	_Z59__device_stub__Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPfjPKfiiiiiiiiiiiiPf
.L32:
	.cfi_restore_state
	addq	$120, %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
.LFE2013:
	.size	arraydiff_cuda_kernel_avg_pool_bwd_f32, .-arraydiff_cuda_kernel_avg_pool_bwd_f32
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf"
.LC1:
	.string	"_Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf"
.LC2:
	.string	"_Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf"
.LC3:
	.string	"_Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2047:
	.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	_Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf(%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	_Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf(%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	_Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf(%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	_Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi(%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
.LFE2047:
	.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:

/*
Copyright 2017 the arraydiff authors

Licensed under the Apache License, Version 2.0 (the "License");
you may not use this file except in compliance with the License.
You may obtain a copy of the License at

    http://www.apache.org/licenses/LICENSE-2.0

Unless required by applicable law or agreed to in writing, software
distributed under the License is distributed on an "AS IS" BASIS,
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and
limitations under the License.

Caffe copyright:

COPYRIGHT

All contributions by the University of California:
Copyright (c) 2014-2017 The Regents of the University of California (Regents)
All rights reserved.

All other contributions:
Copyright (c) 2014-2017, the respective contributors
All rights reserved.

Caffe uses a shared copyright model: each contributor holds copyright over
their contributions to Caffe. The project versioning records all such
contribution and copyright details. If a contributor wants to further mark
their specific copyright on a particular contribution, they should indicate
their copyright solely in the commit message of the change when it is
committed.

LICENSE

Redistribution and use in source and binary forms, with or without
modification, are permitted provided that the following conditions are met: 

1. Redistributions of source code must retain the above copyright notice, this
   list of conditions and the following disclaimer. 
2. Redistributions in binary form must reproduce the above copyright notice,
   this list of conditions and the following disclaimer in the documentation
   and/or other materials provided with the distribution. 

THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR
ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
(INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*/


#include <hip/hip_runtime.h>
#include "common.cuh"
#include <hip/hip_runtime_api.h>
#include <hip/hip_math_constants.h>
#include <stdint.h>

__global__ void max_pool_fwd_f32_kernel(
    const uint32_t nthreads,
    const float* const bottom_data,
    const int num,
    const int channels,
    const int height,
    const int width,
    const int pooled_height,
    const int pooled_width,
    const int kernel_h,
    const int kernel_w,
    const int stride_h,
    const int stride_w,
    const int pad_h,
    const int pad_w,
    float* const top_data,
    int32_t* const mask)
{
  uint32_t index = threadIdx.x + blockDim.x * blockIdx.x;
  if (index < nthreads) {
    const int pw = index % pooled_width;
    const int ph = (index / pooled_width) % pooled_height;
    const int c = (index / pooled_width / pooled_height) % channels;
    const int n = index / pooled_width / pooled_height / channels;
    int hstart = ph * stride_h - pad_h;
    int wstart = pw * stride_w - pad_w;
    const int hend = min(hstart + kernel_h, height);
    const int wend = min(wstart + kernel_w, width);
    hstart = max(hstart, 0);
    wstart = max(wstart, 0);
    float maxval = -HIP_INF_F;
    int maxidx = -1;
    const float* const bottom_slice =
        bottom_data + (n * channels + c) * height * width;
    for (int h = hstart; h < hend; ++h) {
      for (int w = wstart; w < wend; ++w) {
        if (bottom_slice[h * width + w] > maxval) {
          maxidx = h * width + w;
          maxval = bottom_slice[maxidx];
        }
      }
    }
    if (NULL != top_data) {
      top_data[index] = maxval;
    }
    if (NULL != mask) {
      mask[index] = maxidx;
    }
  }
}

extern "C" void arraydiff_cuda_kernel_max_pool_fwd_f32(
    size_t x_w, size_t x_h, size_t chan_dim, size_t batch_sz,
    size_t y_w, size_t y_h,
    size_t kernel_w, size_t kernel_h,
    size_t stride_w, size_t stride_h,
    size_t pad_w, size_t pad_h,
    const float *x,
    float *maybe_y,
    int32_t *maybe_mask,
    hipStream_t stream)
{
  uint32_t n = y_w * y_h * chan_dim * batch_sz;
  max_pool_fwd_f32_kernel<<<(n+1024-1)/1024, 1024, 0, stream>>>(
      n, x,
      batch_sz, chan_dim, x_h, x_w,
      y_h, y_w,
      kernel_h, kernel_w,
      stride_h, stride_w,
      pad_h, pad_w,
      maybe_y, maybe_mask);
}

__global__ void max_pool_bwd_f32_kernel(
    const uint32_t nthreads,
    const float* const top_diff,
    const int32_t* const mask,
    const int num,
    const int channels,
    const int height,
    const int width,
    const int pooled_height,
    const int pooled_width,
    const int kernel_h,
    const int kernel_w,
    const int stride_h,
    const int stride_w,
    const int pad_h,
    const int pad_w,
    float* const bottom_diff)
{
  uint32_t index = threadIdx.x + blockDim.x * blockIdx.x;
  if (index < nthreads) {
    // find out the local index
    // find out the local offset
    const int w = index % width;
    const int h = (index / width) % height;
    const int c = (index / width / height) % channels;
    const int n = index / width / height / channels;
    const int phstart =
         (h + pad_h < kernel_h) ? 0 : (h + pad_h - kernel_h) / stride_h + 1;
    const int phend = min((h + pad_h) / stride_h + 1, pooled_height);
    const int pwstart =
         (w + pad_w < kernel_w) ? 0 : (w + pad_w - kernel_w) / stride_w + 1;
    const int pwend = min((w + pad_w) / stride_w + 1, pooled_width);
    float gradient = 0.0f;
    const int offset = (n * channels + c) * pooled_height * pooled_width;
    const float* const top_diff_slice = top_diff + offset;
    const int* const mask_slice = mask + offset;
    for (int ph = phstart; ph < phend; ++ph) {
      for (int pw = pwstart; pw < pwend; ++pw) {
        if (mask_slice[ph * pooled_width + pw] == h * width + w) {
          gradient += top_diff_slice[ph * pooled_width + pw];
        }
      }
    }
    bottom_diff[index] += gradient;
  }
}

extern "C" void arraydiff_cuda_kernel_max_pool_bwd_f32(
    size_t x_w, size_t x_h, size_t chan_dim, size_t batch_sz,
    size_t y_w, size_t y_h,
    size_t kernel_w, size_t kernel_h,
    size_t stride_w, size_t stride_h,
    size_t pad_w, size_t pad_h,
    const float *dy,
    const int32_t *mask,
    float *dx,
    hipStream_t stream)
{
  uint32_t n = x_w * x_h * chan_dim * batch_sz;
  max_pool_bwd_f32_kernel<<<(n+1024-1)/1024, 1024, 0, stream>>>(
      n, dy, mask,
      batch_sz, chan_dim, x_h, x_w,
      y_h, y_w,
      kernel_h, kernel_w,
      stride_h, stride_w,
      pad_h, pad_w,
      dx);
}

__global__ void avg_pool_fwd_f32_kernel(
    const uint32_t nthreads,
    const float* const bottom_data,
    const int num,
    const int channels,
    const int height,
    const int width,
    const int pooled_height,
    const int pooled_width,
    const int kernel_h,
    const int kernel_w,
    const int stride_h,
    const int stride_w,
    const int pad_h,
    const int pad_w,
    float* const top_data)
{
  uint32_t index = threadIdx.x + blockDim.x * blockIdx.x;
  if (index < nthreads) {
    const int pw = index % pooled_width;
    const int ph = (index / pooled_width) % pooled_height;
    const int c = (index / pooled_width / pooled_height) % channels;
    const int n = index / pooled_width / pooled_height / channels;
    int hstart = ph * stride_h - pad_h;
    int wstart = pw * stride_w - pad_w;
    int hend = min(hstart + kernel_h, height + pad_h);
    int wend = min(wstart + kernel_w, width + pad_w);
    const int pool_size = (hend - hstart) * (wend - wstart);
    hstart = max(hstart, 0);
    wstart = max(wstart, 0);
    hend = min(hend, height);
    wend = min(wend, width);
    float aveval = 0.0f;
    const float* const bottom_slice =
        bottom_data + (n * channels + c) * height * width;
    for (int h = hstart; h < hend; ++h) {
      for (int w = wstart; w < wend; ++w) {
        aveval += bottom_slice[h * width + w];
      }
    }
    top_data[index] = aveval / pool_size;
  }
}

extern "C" void arraydiff_cuda_kernel_avg_pool_fwd_f32(
    size_t x_w, size_t x_h, size_t chan_dim, size_t batch_sz,
    size_t y_w, size_t y_h,
    size_t kernel_w, size_t kernel_h,
    size_t stride_w, size_t stride_h,
    size_t pad_w, size_t pad_h,
    const float *x,
    float *y,
    hipStream_t stream)
{
  uint32_t n = y_w * y_h * chan_dim * batch_sz;
  avg_pool_fwd_f32_kernel<<<(n+1024-1)/1024, 1024, 0, stream>>>(
      n, x,
      batch_sz, chan_dim, x_h, x_w,
      y_h, y_w,
      kernel_h, kernel_w,
      stride_h, stride_w,
      pad_h, pad_w,
      y);
}

__global__ void avg_pool_bwd_f32_kernel(
    const uint32_t nthreads,
    const float* const top_diff,
    const int num,
    const int channels,
    const int height,
    const int width,
    const int pooled_height,
    const int pooled_width,
    const int kernel_h,
    const int kernel_w,
    const int stride_h,
    const int stride_w,
    const int pad_h,
    const int pad_w,
    float* const bottom_diff)
{
  uint32_t index = threadIdx.x + blockDim.x * blockIdx.x;
  if (index < nthreads) {
    // find out the local index
    // find out the local offset
    const int w = index % width + pad_w;
    const int h = (index / width) % height + pad_h;
    const int c = (index / width / height) % channels;
    const int n = index / width / height / channels;
    const int phstart = (h < kernel_h) ? 0 : (h - kernel_h) / stride_h + 1;
    const int phend = min(h / stride_h + 1, pooled_height);
    const int pwstart = (w < kernel_w) ? 0 : (w - kernel_w) / stride_w + 1;
    const int pwend = min(w / stride_w + 1, pooled_width);
    float gradient = 0.0f;
    const float* const top_diff_slice =
        top_diff + (n * channels + c) * pooled_height * pooled_width;
    for (int ph = phstart; ph < phend; ++ph) {
      for (int pw = pwstart; pw < pwend; ++pw) {
        // figure out the pooling size
        int hstart = ph * stride_h - pad_h;
        int wstart = pw * stride_w - pad_w;
        int hend = min(hstart + kernel_h, height + pad_h);
        int wend = min(wstart + kernel_w, width + pad_w);
        int pool_size = (hend - hstart) * (wend - wstart);
        gradient += top_diff_slice[ph * pooled_width + pw] / pool_size;
      }
    }
    bottom_diff[index] += gradient;
  }
}

extern "C" void arraydiff_cuda_kernel_avg_pool_bwd_f32(
    size_t x_w, size_t x_h, size_t chan_dim, size_t batch_sz,
    size_t y_w, size_t y_h,
    size_t kernel_w, size_t kernel_h,
    size_t stride_w, size_t stride_h,
    size_t pad_w, size_t pad_h,
    const float *dy,
    float *dx,
    hipStream_t stream)
{
  uint32_t n = x_w * x_h * chan_dim * batch_sz;
  avg_pool_bwd_f32_kernel<<<(n+1024-1)/1024, 1024, 0, stream>>>(
      n, dy,
      batch_sz, chan_dim, x_h, x_w,
      y_h, y_w,
      kernel_h, kernel_w,
      stride_h, stride_w,
      pad_h, pad_w,
      dx);
}

	.text
	.file	"pool.hip"
	.globl	_Z38__device_stub__max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi # -- Begin function _Z38__device_stub__max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi
	.type	_Z38__device_stub__max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi,@function
_Z38__device_stub__max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi: # @_Z38__device_stub__max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$208, %rsp
	.cfi_def_cfa_offset 256
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	20(%rsp), %rax
	movl	%edi, (%rax)
	leaq	40(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	16(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	12(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	8(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	4(%rsp), %r8
	movl	%r9d, (%r8)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	256(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	264(%rsp), %rax
	movq	%rax, 56(%rbx)
	leaq	272(%rsp), %rax
	movq	%rax, 64(%rbx)
	leaq	280(%rsp), %rax
	movq	%rax, 72(%rbx)
	leaq	288(%rsp), %rax
	movq	%rax, 80(%rbx)
	leaq	296(%rsp), %rax
	movq	%rax, 88(%rbx)
	leaq	304(%rsp), %rax
	movq	%rax, 96(%rbx)
	leaq	312(%rsp), %rax
	movq	%rax, 104(%rbx)
	leaq	320(%rsp), %rax
	movq	%rax, 112(%rbx)
	leaq	328(%rsp), %rax
	movq	%rax, 120(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	32(%rsp), %r12
	leaq	24(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$224, %rsp
	.cfi_adjust_cfa_offset -224
	popq	%rbx
	.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
	retq
.Lfunc_end0:
	.size	_Z38__device_stub__max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi, .Lfunc_end0-_Z38__device_stub__max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi
	.cfi_endproc
                                        # -- End function
	.globl	arraydiff_cuda_kernel_max_pool_fwd_f32 # -- Begin function arraydiff_cuda_kernel_max_pool_fwd_f32
	.type	arraydiff_cuda_kernel_max_pool_fwd_f32,@function
arraydiff_cuda_kernel_max_pool_fwd_f32: # @arraydiff_cuda_kernel_max_pool_fwd_f32
	.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
	pushq	%rax
	.cfi_def_cfa_offset 64
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%r9, %rbx
	movq	%r8, %r15
	movq	%rcx, %rbp
	movq	%rdx, %r13
	movq	%rsi, %r12
	movq	%rdi, (%rsp)                    # 8-byte Spill
	movq	136(%rsp), %r9
	movq	%rcx, %rax
	imulq	%rdx, %rax
	movq	%r8, %r14
	imulq	%rbx, %r14
	imulq	%rax, %r14
	leal	1023(%r14), %edi
	shrl	$10, %edi
	btsq	$32, %rdi
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$1024, %rdx                     # imm = 0x400
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	je	.LBB1_1
# %bb.2:
	addq	$8, %rsp
	.cfi_def_cfa_offset 56
	jmp	.LBB1_3
.LBB1_1:
	.cfi_def_cfa_offset 64
	movq	112(%rsp), %rsi
	movq	104(%rsp), %rax
	movq	96(%rsp), %r11
	movq	80(%rsp), %r10
	movl	%r14d, %edi
	movl	%ebp, %edx
	movl	%r13d, %ecx
	movl	%r12d, %r8d
	movq	(%rsp), %r9                     # 8-byte Reload
                                        # kill: def $r9d killed $r9d killed $r9
	pushq	128(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	128(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	%r11
	.cfi_adjust_cfa_offset 8
	pushq	%rax
	.cfi_adjust_cfa_offset 8
	pushq	%r10
	.cfi_adjust_cfa_offset 8
	pushq	128(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	112(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	128(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	%r15
	.cfi_adjust_cfa_offset 8
	pushq	%rbx
	.cfi_adjust_cfa_offset 8
	callq	_Z38__device_stub__max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi
	addq	$88, %rsp
	.cfi_adjust_cfa_offset -88
.LBB1_3:
	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	arraydiff_cuda_kernel_max_pool_fwd_f32, .Lfunc_end1-arraydiff_cuda_kernel_max_pool_fwd_f32
	.cfi_endproc
                                        # -- End function
	.globl	_Z38__device_stub__max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf # -- Begin function _Z38__device_stub__max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf
	.type	_Z38__device_stub__max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf,@function
_Z38__device_stub__max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf: # @_Z38__device_stub__max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$208, %rsp
	.cfi_def_cfa_offset 256
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	12(%rsp), %rax
	movl	%edi, (%rax)
	leaq	40(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	32(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	8(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	4(%rsp), %rcx
	movl	%r8d, (%rcx)
	movq	%rsp, %r8
	movl	%r9d, (%r8)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	256(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	264(%rsp), %rax
	movq	%rax, 56(%rbx)
	leaq	272(%rsp), %rax
	movq	%rax, 64(%rbx)
	leaq	280(%rsp), %rax
	movq	%rax, 72(%rbx)
	leaq	288(%rsp), %rax
	movq	%rax, 80(%rbx)
	leaq	296(%rsp), %rax
	movq	%rax, 88(%rbx)
	leaq	304(%rsp), %rax
	movq	%rax, 96(%rbx)
	leaq	312(%rsp), %rax
	movq	%rax, 104(%rbx)
	leaq	320(%rsp), %rax
	movq	%rax, 112(%rbx)
	leaq	328(%rsp), %rax
	movq	%rax, 120(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$224, %rsp
	.cfi_adjust_cfa_offset -224
	popq	%rbx
	.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
	retq
.Lfunc_end2:
	.size	_Z38__device_stub__max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf, .Lfunc_end2-_Z38__device_stub__max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf
	.cfi_endproc
                                        # -- End function
	.globl	arraydiff_cuda_kernel_max_pool_bwd_f32 # -- Begin function arraydiff_cuda_kernel_max_pool_bwd_f32
	.type	arraydiff_cuda_kernel_max_pool_bwd_f32,@function
arraydiff_cuda_kernel_max_pool_bwd_f32: # @arraydiff_cuda_kernel_max_pool_bwd_f32
	.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
	pushq	%rax
	.cfi_def_cfa_offset 64
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%r9, (%rsp)                     # 8-byte Spill
	movq	%r8, %r15
	movq	%rcx, %rbp
	movq	%rdx, %rbx
	movq	%rsi, %r12
	movq	%rdi, %r14
	movq	136(%rsp), %r9
	movq	%rsi, %rax
	imulq	%rdi, %rax
	movq	%rdx, %r13
	imulq	%rcx, %r13
	imulq	%rax, %r13
	leal	1023(%r13), %edi
	shrl	$10, %edi
	btsq	$32, %rdi
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$1024, %rdx                     # imm = 0x400
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	je	.LBB3_1
# %bb.2:
	addq	$8, %rsp
	.cfi_def_cfa_offset 56
	jmp	.LBB3_3
.LBB3_1:
	.cfi_def_cfa_offset 64
	movq	120(%rsp), %rdx
	movq	112(%rsp), %rsi
	movq	104(%rsp), %rax
	movq	96(%rsp), %r11
	movq	80(%rsp), %r10
	movl	%r13d, %edi
	movl	%ebp, %ecx
	movl	%ebx, %r8d
	movl	%r12d, %r9d
	pushq	128(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	%r11
	.cfi_adjust_cfa_offset 8
	pushq	%rax
	.cfi_adjust_cfa_offset 8
	pushq	%r10
	.cfi_adjust_cfa_offset 8
	pushq	120(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	104(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	120(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	%r15
	.cfi_adjust_cfa_offset 8
	pushq	64(%rsp)                        # 8-byte Folded Reload
	.cfi_adjust_cfa_offset 8
	pushq	%r14
	.cfi_adjust_cfa_offset 8
	callq	_Z38__device_stub__max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf
	addq	$88, %rsp
	.cfi_adjust_cfa_offset -88
.LBB3_3:
	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_end3:
	.size	arraydiff_cuda_kernel_max_pool_bwd_f32, .Lfunc_end3-arraydiff_cuda_kernel_max_pool_bwd_f32
	.cfi_endproc
                                        # -- End function
	.globl	_Z38__device_stub__avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf # -- Begin function _Z38__device_stub__avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf
	.type	_Z38__device_stub__avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf,@function
_Z38__device_stub__avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf: # @_Z38__device_stub__avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$208, %rsp
	.cfi_def_cfa_offset 256
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	20(%rsp), %rax
	movl	%edi, (%rax)
	leaq	40(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	16(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	12(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	8(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	4(%rsp), %r8
	movl	%r9d, (%r8)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	256(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	264(%rsp), %rax
	movq	%rax, 56(%rbx)
	leaq	272(%rsp), %rax
	movq	%rax, 64(%rbx)
	leaq	280(%rsp), %rax
	movq	%rax, 72(%rbx)
	leaq	288(%rsp), %rax
	movq	%rax, 80(%rbx)
	leaq	296(%rsp), %rax
	movq	%rax, 88(%rbx)
	leaq	304(%rsp), %rax
	movq	%rax, 96(%rbx)
	leaq	312(%rsp), %rax
	movq	%rax, 104(%rbx)
	leaq	320(%rsp), %rax
	movq	%rax, 112(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	32(%rsp), %r12
	leaq	24(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$224, %rsp
	.cfi_adjust_cfa_offset -224
	popq	%rbx
	.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
	retq
.Lfunc_end4:
	.size	_Z38__device_stub__avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf, .Lfunc_end4-_Z38__device_stub__avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf
	.cfi_endproc
                                        # -- End function
	.globl	arraydiff_cuda_kernel_avg_pool_fwd_f32 # -- Begin function arraydiff_cuda_kernel_avg_pool_fwd_f32
	.type	arraydiff_cuda_kernel_avg_pool_fwd_f32,@function
arraydiff_cuda_kernel_avg_pool_fwd_f32: # @arraydiff_cuda_kernel_avg_pool_fwd_f32
	.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
	pushq	%rax
	.cfi_def_cfa_offset 64
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%r9, %rbx
	movq	%r8, %r15
	movq	%rcx, %rbp
	movq	%rdx, %r13
	movq	%rsi, %r12
	movq	%rdi, (%rsp)                    # 8-byte Spill
	movq	128(%rsp), %r9
	movq	%rcx, %rax
	imulq	%rdx, %rax
	movq	%r8, %r14
	imulq	%rbx, %r14
	imulq	%rax, %r14
	leal	1023(%r14), %edi
	shrl	$10, %edi
	btsq	$32, %rdi
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$1024, %rdx                     # imm = 0x400
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	je	.LBB5_1
# %bb.2:
	addq	$8, %rsp
	.cfi_def_cfa_offset 56
	jmp	.LBB5_3
.LBB5_1:
	.cfi_def_cfa_offset 64
	movq	112(%rsp), %rsi
	movq	104(%rsp), %rax
	movq	96(%rsp), %r11
	movq	80(%rsp), %r10
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	movl	%r14d, %edi
	movl	%ebp, %edx
	movl	%r13d, %ecx
	movl	%r12d, %r8d
	movq	8(%rsp), %r9                    # 8-byte Reload
                                        # kill: def $r9d killed $r9d killed $r9
	pushq	128(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	%r11
	.cfi_adjust_cfa_offset 8
	pushq	%rax
	.cfi_adjust_cfa_offset 8
	pushq	%r10
	.cfi_adjust_cfa_offset 8
	pushq	128(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	112(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	128(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	%r15
	.cfi_adjust_cfa_offset 8
	pushq	%rbx
	.cfi_adjust_cfa_offset 8
	callq	_Z38__device_stub__avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf
	addq	$88, %rsp
	.cfi_adjust_cfa_offset -88
.LBB5_3:
	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_end5:
	.size	arraydiff_cuda_kernel_avg_pool_fwd_f32, .Lfunc_end5-arraydiff_cuda_kernel_avg_pool_fwd_f32
	.cfi_endproc
                                        # -- End function
	.globl	_Z38__device_stub__avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf # -- Begin function _Z38__device_stub__avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf
	.type	_Z38__device_stub__avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf,@function
_Z38__device_stub__avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf: # @_Z38__device_stub__avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$208, %rsp
	.cfi_def_cfa_offset 256
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	20(%rsp), %rax
	movl	%edi, (%rax)
	leaq	40(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	16(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	12(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	8(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	4(%rsp), %r8
	movl	%r9d, (%r8)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	256(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	264(%rsp), %rax
	movq	%rax, 56(%rbx)
	leaq	272(%rsp), %rax
	movq	%rax, 64(%rbx)
	leaq	280(%rsp), %rax
	movq	%rax, 72(%rbx)
	leaq	288(%rsp), %rax
	movq	%rax, 80(%rbx)
	leaq	296(%rsp), %rax
	movq	%rax, 88(%rbx)
	leaq	304(%rsp), %rax
	movq	%rax, 96(%rbx)
	leaq	312(%rsp), %rax
	movq	%rax, 104(%rbx)
	leaq	320(%rsp), %rax
	movq	%rax, 112(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	32(%rsp), %r12
	leaq	24(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$224, %rsp
	.cfi_adjust_cfa_offset -224
	popq	%rbx
	.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
	retq
.Lfunc_end6:
	.size	_Z38__device_stub__avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf, .Lfunc_end6-_Z38__device_stub__avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf
	.cfi_endproc
                                        # -- End function
	.globl	arraydiff_cuda_kernel_avg_pool_bwd_f32 # -- Begin function arraydiff_cuda_kernel_avg_pool_bwd_f32
	.type	arraydiff_cuda_kernel_avg_pool_bwd_f32,@function
arraydiff_cuda_kernel_avg_pool_bwd_f32: # @arraydiff_cuda_kernel_avg_pool_bwd_f32
	.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
	pushq	%rax
	.cfi_def_cfa_offset 64
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%r9, (%rsp)                     # 8-byte Spill
	movq	%r8, %r14
	movq	%rcx, %rbp
	movq	%rdx, %rbx
	movq	%rsi, %r12
	movq	%rdi, %r15
	movq	128(%rsp), %r9
	movq	%rsi, %rax
	imulq	%rdi, %rax
	movq	%rdx, %r13
	imulq	%rcx, %r13
	imulq	%rax, %r13
	leal	1023(%r13), %edi
	shrl	$10, %edi
	btsq	$32, %rdi
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$1024, %rdx                     # imm = 0x400
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	je	.LBB7_1
# %bb.2:
	addq	$8, %rsp
	.cfi_def_cfa_offset 56
	jmp	.LBB7_3
.LBB7_1:
	.cfi_def_cfa_offset 64
	movq	112(%rsp), %rsi
	movq	104(%rsp), %rax
	movq	96(%rsp), %r11
	movq	80(%rsp), %r10
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	movl	%r13d, %edi
	movl	%ebp, %edx
	movl	%ebx, %ecx
	movl	%r12d, %r8d
	movl	%r15d, %r9d
	pushq	128(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	%r11
	.cfi_adjust_cfa_offset 8
	pushq	%rax
	.cfi_adjust_cfa_offset 8
	pushq	%r10
	.cfi_adjust_cfa_offset 8
	pushq	128(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	112(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	128(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	%r14
	.cfi_adjust_cfa_offset 8
	pushq	72(%rsp)                        # 8-byte Folded Reload
	.cfi_adjust_cfa_offset 8
	callq	_Z38__device_stub__avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf
	addq	$88, %rsp
	.cfi_adjust_cfa_offset -88
.LBB7_3:
	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_end7:
	.size	arraydiff_cuda_kernel_avg_pool_bwd_f32, .Lfunc_end7-arraydiff_cuda_kernel_avg_pool_bwd_f32
	.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:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB8_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB8_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf, %esi
	movl	$.L__unnamed_4, %edx
	movl	$.L__unnamed_4, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end8:
	.size	__hip_module_ctor, .Lfunc_end8-__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	.LBB9_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
.LBB9_2:
	retq
.Lfunc_end9:
	.size	__hip_module_dtor, .Lfunc_end9-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi,@object # @_Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi
	.section	.rodata,"a",@progbits
	.globl	_Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi
	.p2align	3, 0x0
_Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi:
	.quad	_Z38__device_stub__max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi
	.size	_Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi, 8

	.type	_Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf,@object # @_Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf
	.globl	_Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf
	.p2align	3, 0x0
_Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf:
	.quad	_Z38__device_stub__max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf
	.size	_Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf, 8

	.type	_Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf,@object # @_Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf
	.globl	_Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf
	.p2align	3, 0x0
_Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf:
	.quad	_Z38__device_stub__avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf
	.size	_Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf, 8

	.type	_Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf,@object # @_Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf
	.globl	_Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf
	.p2align	3, 0x0
_Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf:
	.quad	_Z38__device_stub__avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf
	.size	_Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf, 8

	.type	.L__unnamed_1,@object           # @0
	.section	.rodata.str1.1,"aMS",@progbits,1
.L__unnamed_1:
	.asciz	"_Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi"
	.size	.L__unnamed_1, 48

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf"
	.size	.L__unnamed_2, 49

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"_Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf"
	.size	.L__unnamed_3, 46

	.type	.L__unnamed_4,@object           # @3
.L__unnamed_4:
	.asciz	"_Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf"
	.size	.L__unnamed_4, 46

	.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

	.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 _Z38__device_stub__max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi
	.addrsig_sym _Z38__device_stub__max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf
	.addrsig_sym _Z38__device_stub__avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf
	.addrsig_sym _Z38__device_stub__avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi
	.addrsig_sym _Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf
	.addrsig_sym _Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf
	.addrsig_sym _Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi ; -- Begin function _Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi
	.globl	_Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi
	.p2align	8
	.type	_Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi,@function
_Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi: ; @_Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x5c
	s_load_b32 s3, s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_u32_e64 s3, v1
	s_cbranch_execz .LBB0_13
; %bb.1:
	s_load_b256 s[4:11], s[0:1], 0x18
	v_mov_b32_e32 v7, 0xff800000
	s_waitcnt lgkmcnt(0)
	v_cvt_f32_u32_e32 v0, s7
	s_sub_i32 s2, 0, s7
	v_cvt_f32_u32_e32 v3, s6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v0, v0
	v_rcp_iflag_f32_e32 v3, v3
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v0, 0x4f7ffffe, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v0, v0
	v_mul_lo_u32 v2, s2, v0
	s_sub_i32 s2, 0, s6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v2, v0, v2
	v_add_nc_u32_e32 v0, v0, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v0, v1, v0
	v_dual_mul_f32 v2, 0x4f7ffffe, v3 :: v_dual_add_nc_u32 v5, 1, v0
	v_mul_lo_u32 v3, v0, s7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v2, v2
	v_mul_lo_u32 v4, s2, v2
	s_load_b64 s[2:3], s[0:1], 0x38
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v3, v1, v3
	v_subrev_nc_u32_e32 v6, s7, v3
	v_cmp_le_u32_e32 vcc_lo, s7, v3
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_hi_u32 v4, v2, v4
	v_dual_cndmask_b32 v0, v0, v5 :: v_dual_cndmask_b32 v3, v3, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v4, v2, v4
	v_add_nc_u32_e32 v5, 1, v0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_le_u32_e32 vcc_lo, s7, v3
	v_cndmask_b32_e32 v2, v0, v5, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v0, v2, v4
	v_mul_lo_u32 v3, v0, s6
	v_add_nc_u32_e32 v4, 1, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v3, v2, v3
	v_subrev_nc_u32_e32 v5, s6, v3
	v_cmp_le_u32_e32 vcc_lo, s6, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v0, v0, v4 :: v_dual_cndmask_b32 v3, v3, v5
	v_add_nc_u32_e32 v4, 1, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_le_u32_e32 vcc_lo, s6, v3
	v_cndmask_b32_e32 v3, v0, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v0, v3, s6
	s_mov_b32 s6, exec_lo
	v_sub_nc_u32_e32 v0, v2, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v0, v0, s10
	s_waitcnt lgkmcnt(0)
	v_subrev_nc_u32_e32 v0, s2, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v4, s8, v0
	v_max_i32_e32 v5, 0, v0
	v_mov_b32_e32 v0, -1
	s_mov_b32 s8, 0
	v_min_i32_e32 v6, s4, v4
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_lt_i32_e64 v5, v6
	s_cbranch_execz .LBB0_9
; %bb.2:                                ; %.preheader.lr.ph
	v_mul_lo_u32 v0, v2, s7
	s_mul_i32 s2, s5, s4
	s_load_b64 s[12:13], s[0:1], 0x8
	v_mul_lo_u32 v2, s2, v3
	v_mul_lo_u32 v8, s5, v5
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_sub_nc_u32_e32 v0, v1, v0
	v_ashrrev_i32_e32 v3, 31, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v0, v0, s11
	v_lshlrev_b64 v[2:3], 2, v[2:3]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v4, s3, v0
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v11, vcc_lo, s12, v2
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v12, vcc_lo, s13, v3, vcc_lo
	v_dual_mov_b32 v0, -1 :: v_dual_add_nc_u32 v7, s9, v4
	v_max_i32_e32 v9, 0, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_min_i32_e32 v10, s5, v7
	v_dual_mov_b32 v7, 0xff800000 :: v_dual_add_nc_u32 v2, v9, v8
	s_delay_alu instid0(VALU_DEP_2)
	v_cmp_lt_i32_e32 vcc_lo, v9, v10
.LBB0_3:                                ; %.preheader
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_5 Depth 2
	s_and_saveexec_b32 s7, vcc_lo
	s_cbranch_execz .LBB0_7
; %bb.4:                                ; %.lr.ph
                                        ;   in Loop: Header=BB0_3 Depth=1
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v3, 31, v2
	v_mov_b32_e32 v13, v9
	s_mov_b32 s9, 0
	v_lshlrev_b64 v[3:4], 2, v[2:3]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v3, s2, v11, v3
	v_add_co_ci_u32_e64 v4, s2, v12, v4, s2
.LBB0_5:                                ;   Parent Loop BB0_3 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b32 v14, v[3:4], off
	v_add_nc_u32_e32 v15, v8, v13
	v_add_nc_u32_e32 v13, 1, v13
	v_add_co_u32 v3, s2, v3, 4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e64 v4, s2, 0, v4, s2
	v_cmp_ge_i32_e64 s4, v13, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	s_or_b32 s9, s4, s9
	s_waitcnt vmcnt(0)
	v_cmp_gt_f32_e64 s3, v14, v7
	v_cndmask_b32_e64 v7, v7, v14, s3
	v_cndmask_b32_e64 v0, v0, v15, s3
	s_and_not1_b32 exec_lo, exec_lo, s9
	s_cbranch_execnz .LBB0_5
; %bb.6:                                ; %Flow128
                                        ;   in Loop: Header=BB0_3 Depth=1
	s_or_b32 exec_lo, exec_lo, s9
.LBB0_7:                                ; %Flow129
                                        ;   in Loop: Header=BB0_3 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_3) | instid1(VALU_DEP_3)
	s_or_b32 exec_lo, exec_lo, s7
	v_add_nc_u32_e32 v5, 1, v5
	v_add_nc_u32_e32 v8, s5, v8
	v_add_nc_u32_e32 v2, s5, v2
	v_cmp_ge_i32_e64 s2, v5, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s8, s2, s8
	s_and_not1_b32 exec_lo, exec_lo, s8
	s_cbranch_execnz .LBB0_3
; %bb.8:                                ; %Flow130
	s_or_b32 exec_lo, exec_lo, s8
.LBB0_9:                                ; %Flow131
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s6
	s_load_b128 s[0:3], s[0:1], 0x40
	s_waitcnt lgkmcnt(0)
	s_cmp_lg_u64 s[0:1], 0
	s_cbranch_scc0 .LBB0_11
; %bb.10:
	v_mov_b32_e32 v2, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[2:3], 2, v[1:2]
	v_add_co_u32 v2, vcc_lo, s0, v2
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v3, vcc_lo, s1, v3, vcc_lo
	global_store_b32 v[2:3], v7, off
.LBB0_11:
	s_cmp_eq_u64 s[2:3], 0
	s_cbranch_scc1 .LBB0_13
; %bb.12:
	v_mov_b32_e32 v2, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[1:2], 2, v[1:2]
	v_add_co_u32 v1, vcc_lo, s2, v1
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v2, vcc_lo, s3, v2, vcc_lo
	global_store_b32 v[1:2], v0, off
.LBB0_13:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 336
		.amdhsa_user_sgpr_count 15
		.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 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 16
		.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	_Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi, .Lfunc_end0-_Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 872
; NumSgprs: 18
; NumVgprs: 16
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 16
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.protected	_Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf ; -- Begin function _Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf
	.globl	_Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf
	.p2align	8
	.type	_Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf,@function
_Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf: ; @_Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x5c
	s_load_b32 s3, s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_u32_e64 s3, v1
	s_cbranch_execz .LBB1_16
; %bb.1:
	s_load_b256 s[4:11], s[0:1], 0x20
	v_dual_mov_b32 v9, 0 :: v_dual_mov_b32 v10, 0
	s_waitcnt lgkmcnt(0)
	v_cvt_f32_u32_e32 v0, s5
	s_sub_i32 s2, 0, s5
	v_cvt_f32_u32_e32 v3, s4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v0, v0
	v_rcp_iflag_f32_e32 v3, v3
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v0, 0x4f7ffffe, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v0, v0
	v_mul_lo_u32 v2, s2, v0
	s_sub_i32 s2, 0, s4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v2, v0, v2
	v_add_nc_u32_e32 v0, v0, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v0, v1, v0
	v_dual_mul_f32 v2, 0x4f7ffffe, v3 :: v_dual_add_nc_u32 v5, 1, v0
	v_mul_lo_u32 v3, v0, s5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v2, v2
	v_mul_lo_u32 v4, s2, v2
	s_load_b64 s[2:3], s[0:1], 0x40
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v3, v1, v3
	v_subrev_nc_u32_e32 v6, s5, v3
	v_cmp_le_u32_e32 vcc_lo, s5, v3
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_hi_u32 v4, v2, v4
	v_dual_cndmask_b32 v0, v0, v5 :: v_dual_cndmask_b32 v3, v3, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v2, v2, v4
	v_add_nc_u32_e32 v5, 1, v0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_le_u32_e32 vcc_lo, s5, v3
	v_cndmask_b32_e32 v0, v0, v5, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v2, v0, v2
	v_mul_lo_u32 v3, v2, s4
	v_add_nc_u32_e32 v4, 1, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v3, v0, v3
	v_subrev_nc_u32_e32 v5, s4, v3
	v_cmp_le_u32_e32 vcc_lo, s4, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v2, v2, v4 :: v_dual_cndmask_b32 v3, v3, v5
	v_add_nc_u32_e32 v4, 1, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_le_u32_e32 vcc_lo, s4, v3
	v_cndmask_b32_e32 v3, v2, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v2, v3, s4
	v_sub_nc_u32_e32 v5, v0, v2
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v4, s2, v5
	s_mov_b32 s2, exec_lo
	v_cmpx_le_i32_e64 s8, v4
	s_cbranch_execz .LBB1_3
; %bb.2:
	s_ashr_i32 s4, s10, 31
	v_subrev_nc_u32_e32 v7, s8, v4
	s_add_i32 s12, s10, s4
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_xor_b32 s12, s12, s4
	v_ashrrev_i32_e32 v8, 31, v7
	v_cvt_f32_u32_e32 v2, s12
	s_sub_i32 s13, 0, s12
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v7, v7, v8
	v_rcp_iflag_f32_e32 v2, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v7, v7, v8
	v_xor_b32_e32 v8, s4, v8
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v2, 0x4f7ffffe, v2
	v_cvt_u32_f32_e32 v2, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v6, s13, v2
	v_mul_hi_u32 v6, v2, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v2, v2, v6
	v_mul_hi_u32 v2, v7, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v6, v2, s12
	v_sub_nc_u32_e32 v6, v7, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v9, s12, v6
	v_cmp_le_u32_e32 vcc_lo, s12, v6
	v_dual_cndmask_b32 v6, v6, v9 :: v_dual_add_nc_u32 v7, 1, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v2, v2, v7, vcc_lo
	v_cmp_le_u32_e32 vcc_lo, s12, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v7, 1, v2
	v_cndmask_b32_e32 v2, v2, v7, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v2, v2, v8
	v_sub_nc_u32_e32 v2, v2, v8
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v9, 1, v2
.LBB1_3:
	s_or_b32 exec_lo, exec_lo, s2
	v_mul_lo_u32 v0, v0, s5
	s_mov_b32 s2, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v2, v1, v0
	v_mov_b32_e32 v0, 0
	v_add_nc_u32_e32 v6, s3, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_le_i32_e64 s9, v6
	s_cbranch_execz .LBB1_5
; %bb.4:
	s_ashr_i32 s3, s11, 31
	v_subrev_nc_u32_e32 v8, s9, v6
	s_add_i32 s4, s11, s3
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_xor_b32 s4, s4, s3
	v_ashrrev_i32_e32 v11, 31, v8
	v_cvt_f32_u32_e32 v0, s4
	s_sub_i32 s8, 0, s4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v8, v8, v11
	v_rcp_iflag_f32_e32 v0, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v8, v8, v11
	v_xor_b32_e32 v11, s3, v11
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v0, 0x4f7ffffe, v0
	v_cvt_u32_f32_e32 v0, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v7, s8, v0
	v_mul_hi_u32 v7, v0, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v0, v0, v7
	v_mul_hi_u32 v0, v8, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v7, v0, s4
	v_sub_nc_u32_e32 v7, v8, v7
	v_add_nc_u32_e32 v8, 1, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v12, s4, v7
	v_cmp_le_u32_e32 vcc_lo, s4, v7
	v_cndmask_b32_e32 v7, v7, v12, vcc_lo
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v0, v0, v8, vcc_lo
	v_cmp_le_u32_e32 vcc_lo, s4, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v8, 1, v0
	v_cndmask_b32_e32 v0, v0, v8, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v0, v0, v11
	v_sub_nc_u32_e32 v0, v0, v11
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v0, 1, v0
.LBB1_5:
	s_or_b32 exec_lo, exec_lo, s2
	s_ashr_i32 s2, s10, 31
	v_ashrrev_i32_e32 v11, 31, v4
	s_add_i32 s3, s10, s2
	s_mov_b32 s8, 0
	s_xor_b32 s3, s3, s2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v4, v4, v11
	v_cvt_f32_u32_e32 v7, s3
	s_sub_i32 s4, 0, s3
	v_xor_b32_e32 v4, v4, v11
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v7, v7
	v_xor_b32_e32 v11, s2, v11
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v7, 0x4f7ffffe, v7
	v_cvt_u32_f32_e32 v7, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v8, s4, v7
	s_mov_b32 s4, exec_lo
	v_mul_hi_u32 v8, v7, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v7, v7, v8
	v_mul_hi_u32 v7, v4, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v8, v7, s3
	v_sub_nc_u32_e32 v4, v4, v8
	v_add_nc_u32_e32 v8, 1, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v12, s3, v4
	v_cmp_le_u32_e32 vcc_lo, s3, v4
	v_cndmask_b32_e32 v4, v4, v12, vcc_lo
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v7, v7, v8, vcc_lo
	v_cmp_le_u32_e32 vcc_lo, s3, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v8, 1, v7
	v_cndmask_b32_e32 v4, v7, v8, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v4, v4, v11
	v_sub_nc_u32_e32 v4, v4, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v4, 1, v4
	v_min_i32_e32 v11, s6, v4
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_lt_i32_e64 v9, v11
	s_cbranch_execz .LBB1_15
; %bb.6:                                ; %.preheader.lr.ph
	s_ashr_i32 s2, s11, 31
	v_ashrrev_i32_e32 v8, 31, v6
	s_add_i32 s3, s11, s2
	s_mul_i32 s6, s7, s6
	s_xor_b32 s3, s3, s2
	s_load_b128 s[12:15], s[0:1], 0x8
	v_cvt_f32_u32_e32 v4, s3
	s_sub_i32 s9, 0, s3
	v_add_nc_u32_e32 v6, v6, v8
	v_mul_lo_u32 v3, s6, v3
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_rcp_iflag_f32_e32 v4, v4
	v_xor_b32_e32 v6, v6, v8
	v_xor_b32_e32 v8, s2, v8
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v4, 0x4f7ffffe, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v4, v4
	v_mul_lo_u32 v7, s9, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v7, v4, v7
	v_add_nc_u32_e32 v4, v4, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v4, v6, v4
	v_mul_lo_u32 v7, v4, s3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v6, v6, v7
	v_subrev_nc_u32_e32 v10, s3, v6
	v_cmp_le_u32_e32 vcc_lo, s3, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v6, v6, v10 :: v_dual_add_nc_u32 v7, 1, v4
	v_cndmask_b32_e32 v4, v4, v7, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_le_u32_e32 vcc_lo, s3, v6
	v_add_nc_u32_e32 v7, 1, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v6, v4, v7, vcc_lo
	v_ashrrev_i32_e32 v4, 31, v3
	v_xor_b32_e32 v10, v6, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_lshlrev_b64 v[6:7], 2, v[3:4]
	v_mad_u64_u32 v[3:4], null, v5, s5, v[2:3]
	v_mad_u64_u32 v[4:5], null, v9, s7, v[0:1]
	v_sub_nc_u32_e32 v8, v10, v8
	v_mov_b32_e32 v10, 0
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s14, v6
	v_add_co_ci_u32_e32 v12, vcc_lo, s15, v7, vcc_lo
	v_add_nc_u32_e32 v5, 1, v8
	v_add_co_u32 v14, vcc_lo, s12, v6
	v_add_co_ci_u32_e32 v15, vcc_lo, s13, v7, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_min_i32_e32 v13, s7, v5
	v_cmp_lt_i32_e32 vcc_lo, v0, v13
.LBB1_7:                                ; %.preheader
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB1_9 Depth 2
	s_and_saveexec_b32 s5, vcc_lo
	s_cbranch_execz .LBB1_13
; %bb.8:                                ; %.lr.ph
                                        ;   in Loop: Header=BB1_7 Depth=1
	v_ashrrev_i32_e32 v5, 31, v4
	v_mov_b32_e32 v16, v0
	s_mov_b32 s6, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[7:8], 2, v[4:5]
	v_add_co_u32 v5, s2, v2, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e64 v6, s2, v12, v8, s2
	v_add_co_u32 v7, s2, v14, v7
	v_add_co_ci_u32_e64 v8, s2, v15, v8, s2
.LBB1_9:                                ;   Parent Loop BB1_7 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b32 v17, v[5:6], off
	s_mov_b32 s3, exec_lo
	s_waitcnt vmcnt(0)
	v_cmpx_eq_u32_e64 v17, v3
	s_cbranch_execz .LBB1_11
; %bb.10:                               ;   in Loop: Header=BB1_9 Depth=2
	global_load_b32 v17, v[7:8], off
	s_waitcnt vmcnt(0)
	v_add_f32_e32 v10, v10, v17
.LBB1_11:                               ;   in Loop: Header=BB1_9 Depth=2
	s_or_b32 exec_lo, exec_lo, s3
	v_add_nc_u32_e32 v16, 1, v16
	v_add_co_u32 v5, s2, v5, 4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e64 v6, s2, 0, v6, s2
	v_cmp_ge_i32_e64 s2, v16, v13
	v_add_co_u32 v7, s3, v7, 4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e64 v8, s3, 0, v8, s3
	s_or_b32 s6, s2, s6
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s6
	s_cbranch_execnz .LBB1_9
; %bb.12:                               ; %Flow
                                        ;   in Loop: Header=BB1_7 Depth=1
	s_or_b32 exec_lo, exec_lo, s6
.LBB1_13:                               ; %Flow131
                                        ;   in Loop: Header=BB1_7 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	s_or_b32 exec_lo, exec_lo, s5
	v_add_nc_u32_e32 v9, 1, v9
	v_add_nc_u32_e32 v4, s7, v4
	v_cmp_ge_i32_e64 s2, v9, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s8, s2, s8
	s_and_not1_b32 exec_lo, exec_lo, s8
	s_cbranch_execnz .LBB1_7
; %bb.14:                               ; %Flow132
	s_or_b32 exec_lo, exec_lo, s8
.LBB1_15:                               ; %Flow133
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	s_or_b32 exec_lo, exec_lo, s4
	s_load_b64 s[0:1], s[0:1], 0x48
	v_mov_b32_e32 v2, 0
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s0, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_load_b32 v2, v[0:1], off
	s_waitcnt vmcnt(0)
	v_add_f32_e32 v2, v10, v2
	global_store_b32 v[0:1], v2, off
.LBB1_16:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 336
		.amdhsa_user_sgpr_count 15
		.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 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 18
		.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_end1:
	.size	_Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf, .Lfunc_end1-_Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 1600
; NumSgprs: 18
; NumVgprs: 18
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 2
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 18
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.protected	_Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf ; -- Begin function _Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf
	.globl	_Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf
	.p2align	8
	.type	_Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf,@function
_Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf: ; @_Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x54
	s_load_b32 s3, s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_u32_e64 s3, v1
	s_cbranch_execz .LBB2_10
; %bb.1:
	s_load_b256 s[4:11], s[0:1], 0x18
	v_mov_b32_e32 v8, 0
	s_waitcnt lgkmcnt(0)
	v_cvt_f32_u32_e32 v0, s7
	s_sub_i32 s2, 0, s7
	v_cvt_f32_u32_e32 v3, s6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v0, v0
	v_rcp_iflag_f32_e32 v3, v3
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v0, 0x4f7ffffe, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v0, v0
	v_mul_lo_u32 v2, s2, v0
	s_sub_i32 s2, 0, s6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v2, v0, v2
	v_add_nc_u32_e32 v0, v0, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v0, v1, v0
	v_dual_mul_f32 v2, 0x4f7ffffe, v3 :: v_dual_add_nc_u32 v5, 1, v0
	v_mul_lo_u32 v3, v0, s7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v2, v2
	v_mul_lo_u32 v4, s2, v2
	s_load_b64 s[2:3], s[0:1], 0x38
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v3, v1, v3
	v_subrev_nc_u32_e32 v6, s7, v3
	v_cmp_le_u32_e32 vcc_lo, s7, v3
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_hi_u32 v4, v2, v4
	v_dual_cndmask_b32 v0, v0, v5 :: v_dual_cndmask_b32 v3, v3, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v5, 1, v0
	v_cmp_le_u32_e32 vcc_lo, s7, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v3, v0, v5 :: v_dual_add_nc_u32 v2, v2, v4
	v_mul_hi_u32 v0, v3, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v2, v0, s6
	v_add_nc_u32_e32 v4, 1, v0
	v_sub_nc_u32_e32 v2, v3, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_subrev_nc_u32_e32 v5, s6, v2
	v_cmp_le_u32_e32 vcc_lo, s6, v2
	v_cndmask_b32_e32 v0, v0, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v2, v2, v5, vcc_lo
	v_add_nc_u32_e32 v4, 1, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_le_u32_e32 vcc_lo, s6, v2
	v_cndmask_b32_e32 v0, v0, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v2, v0, s6
	s_mov_b32 s6, exec_lo
	v_sub_nc_u32_e32 v2, v3, v2
	v_mul_lo_u32 v3, v3, s7
	s_mov_b32 s7, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v2, v2, s10
	v_sub_nc_u32_e32 v3, v1, v3
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v5, s2, v2
	v_mul_lo_u32 v2, v3, s11
	s_add_i32 s2, s2, s4
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v3, s8, v5
	v_max_i32_e32 v10, 0, v5
	v_subrev_nc_u32_e32 v6, s3, v2
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_min_i32_e32 v7, s2, v3
	s_add_i32 s2, s3, s5
	v_add_nc_u32_e32 v2, s9, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_min_i32_e32 v11, s4, v7
	v_min_i32_e32 v9, s2, v2
	s_delay_alu instid0(VALU_DEP_2)
	v_cmpx_lt_i32_e64 v10, v11
	s_cbranch_execz .LBB2_9
; %bb.2:                                ; %.preheader.lr.ph
	s_mul_i32 s4, s5, s4
	s_load_b64 s[2:3], s[0:1], 0x8
	v_mul_lo_u32 v2, s4, v0
	v_max_i32_e32 v0, 0, v6
	v_min_i32_e32 v12, s5, v9
	v_mov_b32_e32 v8, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_lt_i32_e32 vcc_lo, v0, v12
	v_ashrrev_i32_e32 v3, 31, v2
	v_lshlrev_b64 v[13:14], 2, v[2:3]
	v_mad_u64_u32 v[2:3], null, s5, v10, v[0:1]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v13, s2, s2, v13
	v_add_co_ci_u32_e64 v14, s2, s3, v14, s2
.LBB2_3:                                ; %.preheader
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB2_5 Depth 2
	s_and_saveexec_b32 s4, vcc_lo
	s_cbranch_execz .LBB2_7
; %bb.4:                                ; %.lr.ph
                                        ;   in Loop: Header=BB2_3 Depth=1
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v3, 31, v2
	v_mov_b32_e32 v15, v0
	s_mov_b32 s8, 0
	v_lshlrev_b64 v[3:4], 2, v[2:3]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v3, s2, v13, v3
	v_add_co_ci_u32_e64 v4, s2, v14, v4, s2
.LBB2_5:                                ;   Parent Loop BB2_3 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b32 v16, v[3:4], off
	v_add_nc_u32_e32 v15, 1, v15
	v_add_co_u32 v3, s2, v3, 4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e64 v4, s2, 0, v4, s2
	s_waitcnt vmcnt(0)
	v_add_f32_e32 v8, v8, v16
	v_cmp_ge_i32_e64 s3, v15, v12
	s_or_b32 s8, s3, s8
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s8
	s_cbranch_execnz .LBB2_5
; %bb.6:                                ; %Flow
                                        ;   in Loop: Header=BB2_3 Depth=1
	s_or_b32 exec_lo, exec_lo, s8
.LBB2_7:                                ; %Flow111
                                        ;   in Loop: Header=BB2_3 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	s_or_b32 exec_lo, exec_lo, s4
	v_add_nc_u32_e32 v10, 1, v10
	v_add_nc_u32_e32 v2, s5, v2
	v_cmp_ge_i32_e64 s2, v10, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s7, s2, s7
	s_and_not1_b32 exec_lo, exec_lo, s7
	s_cbranch_execnz .LBB2_3
; %bb.8:                                ; %Flow112
	s_or_b32 exec_lo, exec_lo, s7
.LBB2_9:                                ; %Flow113
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	s_or_b32 exec_lo, exec_lo, s6
	v_sub_nc_u32_e32 v0, v7, v5
	v_sub_nc_u32_e32 v2, v9, v6
	s_load_b64 s[0:1], s[0:1], 0x40
	v_mul_lo_u32 v0, v0, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_f32_i32_e32 v3, v0
	v_div_scale_f32 v0, null, v3, v3, v8
	v_div_scale_f32 v5, vcc_lo, v8, v3, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f32_e32 v4, v0
	s_waitcnt_depctr 0xfff
	v_fma_f32 v2, -v0, v4, 1.0
	v_fmac_f32_e32 v4, v2, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v6, v5, v4
	v_fma_f32 v2, -v0, v6, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v6, v2, v4
	v_mov_b32_e32 v2, 0
	v_fma_f32 v0, -v0, v6, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_div_fmas_f32 v4, v0, v4, v6
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_fixup_f32 v2, v4, v3, v8
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v0, vcc_lo, s0, v0
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_store_b32 v[0:1], v2, off
.LBB2_10:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 328
		.amdhsa_user_sgpr_count 15
		.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 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 17
		.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_end2:
	.size	_Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf, .Lfunc_end2-_Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 900
; NumSgprs: 18
; NumVgprs: 17
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 2
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 17
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.protected	_Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf ; -- Begin function _Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf
	.globl	_Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf
	.p2align	8
	.type	_Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf,@function
_Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf: ; @_Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x54
	s_load_b32 s3, s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_u32_e64 s3, v1
	s_cbranch_execz .LBB3_14
; %bb.1:
	s_clause 0x1
	s_load_b256 s[4:11], s[0:1], 0x18
	s_load_b64 s[12:13], s[0:1], 0x38
	s_waitcnt lgkmcnt(0)
	v_cvt_f32_u32_e32 v0, s5
	s_sub_i32 s2, 0, s5
	v_cvt_f32_u32_e32 v3, s4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v0, v0
	v_rcp_iflag_f32_e32 v3, v3
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v0, 0x4f7ffffe, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v0, v0
	v_mul_lo_u32 v2, s2, v0
	s_sub_i32 s2, 0, s4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v2, v0, v2
	v_add_nc_u32_e32 v0, v0, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v0, v1, v0
	v_dual_mul_f32 v2, 0x4f7ffffe, v3 :: v_dual_add_nc_u32 v5, 1, v0
	v_mul_lo_u32 v3, v0, s5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v2, v2
	v_mul_lo_u32 v4, s2, v2
	s_mov_b32 s2, exec_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v3, v1, v3
	v_subrev_nc_u32_e32 v6, s5, v3
	v_cmp_le_u32_e32 vcc_lo, s5, v3
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_hi_u32 v4, v2, v4
	v_dual_cndmask_b32 v0, v0, v5 :: v_dual_cndmask_b32 v3, v3, v6
	v_mov_b32_e32 v6, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v2, v2, v4
	v_add_nc_u32_e32 v5, 1, v0
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_le_u32_e32 vcc_lo, s5, v3
	v_cndmask_b32_e32 v0, v0, v5, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v2, v0, v2
	v_mul_lo_u32 v3, v2, s4
	v_add_nc_u32_e32 v4, 1, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v3, v0, v3
	v_subrev_nc_u32_e32 v5, s4, v3
	v_cmp_le_u32_e32 vcc_lo, s4, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v2, v2, v4 :: v_dual_cndmask_b32 v3, v3, v5
	v_dual_mov_b32 v5, 0 :: v_dual_add_nc_u32 v4, 1, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_le_u32_e32 vcc_lo, s4, v3
	v_cndmask_b32_e32 v2, v2, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v3, v2, s4
	v_sub_nc_u32_e32 v3, v0, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v3, s12, v3
	v_cmpx_le_i32_e64 s8, v3
	s_cbranch_execz .LBB3_3
; %bb.2:
	s_ashr_i32 s3, s10, 31
	v_subrev_nc_u32_e32 v7, s8, v3
	s_add_i32 s14, s10, s3
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_xor_b32 s14, s14, s3
	v_ashrrev_i32_e32 v8, 31, v7
	v_cvt_f32_u32_e32 v4, s14
	s_sub_i32 s15, 0, s14
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v7, v7, v8
	v_rcp_iflag_f32_e32 v4, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v7, v7, v8
	v_xor_b32_e32 v8, s3, v8
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v4, 0x4f7ffffe, v4
	v_cvt_u32_f32_e32 v4, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v5, s15, v4
	v_mul_hi_u32 v5, v4, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v4, v4, v5
	v_mul_hi_u32 v4, v7, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v5, v4, s14
	v_sub_nc_u32_e32 v5, v7, v5
	v_add_nc_u32_e32 v7, 1, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v9, s14, v5
	v_cmp_le_u32_e32 vcc_lo, s14, v5
	v_dual_cndmask_b32 v5, v5, v9 :: v_dual_cndmask_b32 v4, v4, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_le_u32_e32 vcc_lo, s14, v5
	v_add_nc_u32_e32 v7, 1, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v4, v4, v7, vcc_lo
	v_xor_b32_e32 v4, v4, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v4, v4, v8
	v_add_nc_u32_e32 v5, 1, v4
.LBB3_3:
	s_or_b32 exec_lo, exec_lo, s2
	v_mul_lo_u32 v0, v0, s5
	s_mov_b32 s2, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v0, v1, v0
	v_add_nc_u32_e32 v4, s13, v0
	v_mov_b32_e32 v0, 0
	s_delay_alu instid0(VALU_DEP_2)
	v_cmpx_le_i32_e64 s9, v4
	s_cbranch_execz .LBB3_5
; %bb.4:
	s_ashr_i32 s3, s11, 31
	v_subrev_nc_u32_e32 v8, s9, v4
	s_add_i32 s14, s11, s3
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_xor_b32 s14, s14, s3
	v_ashrrev_i32_e32 v9, 31, v8
	v_cvt_f32_u32_e32 v0, s14
	s_sub_i32 s15, 0, s14
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v8, v8, v9
	v_rcp_iflag_f32_e32 v0, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v8, v8, v9
	v_xor_b32_e32 v9, s3, v9
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v0, 0x4f7ffffe, v0
	v_cvt_u32_f32_e32 v0, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v7, s15, v0
	v_mul_hi_u32 v7, v0, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v0, v0, v7
	v_mul_hi_u32 v0, v8, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v7, v0, s14
	v_sub_nc_u32_e32 v7, v8, v7
	v_add_nc_u32_e32 v8, 1, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v10, s14, v7
	v_cmp_le_u32_e32 vcc_lo, s14, v7
	v_dual_cndmask_b32 v7, v7, v10 :: v_dual_cndmask_b32 v0, v0, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_le_u32_e32 vcc_lo, s14, v7
	v_add_nc_u32_e32 v8, 1, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v0, v0, v8, vcc_lo
	v_xor_b32_e32 v0, v0, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v0, v0, v9
	v_add_nc_u32_e32 v0, 1, v0
.LBB3_5:
	s_or_b32 exec_lo, exec_lo, s2
	s_ashr_i32 s2, s10, 31
	v_ashrrev_i32_e32 v9, 31, v3
	s_add_i32 s3, s10, s2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_xor_b32 s3, s3, s2
	v_add_nc_u32_e32 v3, v3, v9
	v_cvt_f32_u32_e32 v7, s3
	s_sub_i32 s14, 0, s3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v3, v3, v9
	v_rcp_iflag_f32_e32 v7, v7
	v_xor_b32_e32 v9, s2, v9
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v7, 0x4f7ffffe, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v7, v7
	v_mul_lo_u32 v8, s14, v7
	s_mov_b32 s14, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v8, v7, v8
	v_add_nc_u32_e32 v7, v7, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v7, v3, v7
	v_mul_lo_u32 v8, v7, s3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v3, v3, v8
	v_subrev_nc_u32_e32 v10, s3, v3
	v_cmp_le_u32_e32 vcc_lo, s3, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v3, v3, v10 :: v_dual_add_nc_u32 v8, 1, v7
	v_cndmask_b32_e32 v7, v7, v8, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_le_u32_e32 vcc_lo, s3, v3
	v_add_nc_u32_e32 v8, 1, v7
	s_mov_b32 s3, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v3, v7, v8, vcc_lo
	v_xor_b32_e32 v3, v3, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v3, v3, v9
	v_add_nc_u32_e32 v3, 1, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_min_i32_e32 v7, s6, v3
	v_cmpx_lt_i32_e64 v5, v7
	s_cbranch_execz .LBB3_13
; %bb.6:                                ; %.preheader.lr.ph
	s_ashr_i32 s2, s11, 31
	v_ashrrev_i32_e32 v8, 31, v4
	s_add_i32 s15, s11, s2
	s_mul_i32 s6, s7, s6
	s_xor_b32 s15, s15, s2
	v_mul_lo_u32 v2, s6, v2
	v_cvt_f32_u32_e32 v3, s15
	s_sub_i32 s16, 0, s15
	s_add_i32 s4, s12, s4
	s_add_i32 s5, s13, s5
	v_add_nc_u32_e32 v4, v4, v8
	v_rcp_iflag_f32_e32 v3, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v4, v4, v8
	v_xor_b32_e32 v8, s2, v8
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v3, 0x4f7ffffe, v3
	v_cvt_u32_f32_e32 v3, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v6, s16, v3
	s_load_b64 s[16:17], s[0:1], 0x8
	v_mul_hi_u32 v6, v3, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v3, v3, v6
	v_mul_hi_u32 v3, v4, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v6, v3, s15
	v_sub_nc_u32_e32 v4, v4, v6
	v_add_nc_u32_e32 v6, 1, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v9, s15, v4
	v_cmp_le_u32_e32 vcc_lo, s15, v4
	v_dual_cndmask_b32 v4, v4, v9 :: v_dual_cndmask_b32 v3, v3, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_le_u32_e32 vcc_lo, s15, v4
	v_add_nc_u32_e32 v6, 1, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v3, v3, v6, vcc_lo
	v_mul_lo_u32 v6, v0, s11
	v_xor_b32_e32 v4, v3, v8
	v_ashrrev_i32_e32 v3, 31, v2
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v9, s9, v6
	v_sub_nc_u32_e32 v4, v4, v8
	s_delay_alu instid0(VALU_DEP_3)
	v_lshlrev_b64 v[11:12], 2, v[2:3]
	v_mad_u64_u32 v[2:3], null, v5, s7, v[0:1]
	v_sub_nc_u32_e32 v8, s13, v6
	v_mov_b32_e32 v6, 0
	v_add_nc_u32_e32 v4, 1, v4
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v11, vcc_lo, s16, v11
	v_subrev_nc_u32_e32 v9, s13, v9
	v_add_co_ci_u32_e32 v12, vcc_lo, s17, v12, vcc_lo
	v_min_i32_e32 v10, s7, v4
	s_delay_alu instid0(VALU_DEP_1)
	v_cmp_lt_i32_e64 s2, v0, v10
.LBB3_7:                                ; %.preheader
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB3_9 Depth 2
	s_delay_alu instid0(VALU_DEP_1)
	s_and_saveexec_b32 s6, s2
	s_cbranch_execz .LBB3_11
; %bb.8:                                ; %.lr.ph
                                        ;   in Loop: Header=BB3_7 Depth=1
	v_mul_lo_u32 v3, v5, s10
	v_mov_b32_e32 v14, v9
	s_mov_b32 s9, 0
	v_mov_b32_e32 v13, v8
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v15, s12, v3
	v_ashrrev_i32_e32 v3, 31, v2
	v_add_nc_u32_e32 v16, s8, v15
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[3:4], 2, v[2:3]
	v_min_i32_e32 v16, s4, v16
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v3, vcc_lo, v11, v3
	v_add_co_ci_u32_e32 v4, vcc_lo, v12, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_3)
	v_sub_nc_u32_e32 v15, v16, v15
	v_mov_b32_e32 v16, v0
.LBB3_9:                                ;   Parent Loop BB3_7 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b32 v17, v[3:4], off
	v_min_i32_e32 v18, s5, v14
	v_add_nc_u32_e32 v16, 1, v16
	v_add_nc_u32_e32 v14, s11, v14
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v18, v18, v13
	v_subrev_nc_u32_e32 v13, s11, v13
	v_mul_lo_u32 v18, v18, v15
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cvt_f32_i32_e32 v18, v18
	s_waitcnt vmcnt(0)
	v_div_scale_f32 v19, null, v18, v18, v17
	v_div_scale_f32 v22, vcc_lo, v17, v18, v17
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f32_e32 v20, v19
	s_waitcnt_depctr 0xfff
	v_fma_f32 v21, -v19, v20, 1.0
	v_fmac_f32_e32 v20, v21, v20
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v21, v22, v20
	v_fma_f32 v23, -v19, v21, v22
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v21, v23, v20
	v_fma_f32 v19, -v19, v21, v22
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_div_fmas_f32 v19, v19, v20, v21
	v_add_co_u32 v3, vcc_lo, v3, 4
	v_add_co_ci_u32_e32 v4, vcc_lo, 0, v4, vcc_lo
	v_div_fixup_f32 v17, v19, v18, v17
	v_cmp_ge_i32_e32 vcc_lo, v16, v10
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_add_f32_e32 v6, v6, v17
	s_or_b32 s9, vcc_lo, s9
	s_and_not1_b32 exec_lo, exec_lo, s9
	s_cbranch_execnz .LBB3_9
; %bb.10:                               ; %Flow
                                        ;   in Loop: Header=BB3_7 Depth=1
	s_or_b32 exec_lo, exec_lo, s9
.LBB3_11:                               ; %Flow131
                                        ;   in Loop: Header=BB3_7 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	s_or_b32 exec_lo, exec_lo, s6
	v_add_nc_u32_e32 v5, 1, v5
	v_add_nc_u32_e32 v2, s7, v2
	v_cmp_ge_i32_e32 vcc_lo, v5, v7
	s_or_b32 s14, vcc_lo, s14
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s14
	s_cbranch_execnz .LBB3_7
; %bb.12:                               ; %Flow132
	s_or_b32 exec_lo, exec_lo, s14
.LBB3_13:                               ; %Flow133
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	s_or_b32 exec_lo, exec_lo, s3
	s_load_b64 s[0:1], s[0:1], 0x40
	v_mov_b32_e32 v2, 0
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s0, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_load_b32 v2, v[0:1], off
	s_waitcnt vmcnt(0)
	v_add_f32_e32 v2, v6, v2
	global_store_b32 v[0:1], v2, off
.LBB3_14:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 328
		.amdhsa_user_sgpr_count 15
		.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 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 24
		.amdhsa_next_free_sgpr 18
		.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_end3:
	.size	_Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf, .Lfunc_end3-_Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 1688
; NumSgprs: 20
; NumVgprs: 24
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 2
; NumSGPRsForWavesPerEU: 20
; NumVGPRsForWavesPerEU: 24
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.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:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .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:     by_value
      - .offset:         36
        .size:           4
        .value_kind:     by_value
      - .offset:         40
        .size:           4
        .value_kind:     by_value
      - .offset:         44
        .size:           4
        .value_kind:     by_value
      - .offset:         48
        .size:           4
        .value_kind:     by_value
      - .offset:         52
        .size:           4
        .value_kind:     by_value
      - .offset:         56
        .size:           4
        .value_kind:     by_value
      - .offset:         60
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         64
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         72
        .size:           8
        .value_kind:     global_buffer
      - .offset:         80
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         84
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         88
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         92
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         94
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         96
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         98
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         100
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         102
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         120
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         128
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         136
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         144
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 336
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z23max_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPfPi.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     16
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         28
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     by_value
      - .offset:         36
        .size:           4
        .value_kind:     by_value
      - .offset:         40
        .size:           4
        .value_kind:     by_value
      - .offset:         44
        .size:           4
        .value_kind:     by_value
      - .offset:         48
        .size:           4
        .value_kind:     by_value
      - .offset:         52
        .size:           4
        .value_kind:     by_value
      - .offset:         56
        .size:           4
        .value_kind:     by_value
      - .offset:         60
        .size:           4
        .value_kind:     by_value
      - .offset:         64
        .size:           4
        .value_kind:     by_value
      - .offset:         68
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         72
        .size:           8
        .value_kind:     global_buffer
      - .offset:         80
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         84
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         88
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         92
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         94
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         96
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         98
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         100
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         102
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         120
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         128
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         136
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         144
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 336
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z23max_pool_bwd_f32_kerneljPKfPKiiiiiiiiiiiiiPf.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     18
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .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:     by_value
      - .offset:         36
        .size:           4
        .value_kind:     by_value
      - .offset:         40
        .size:           4
        .value_kind:     by_value
      - .offset:         44
        .size:           4
        .value_kind:     by_value
      - .offset:         48
        .size:           4
        .value_kind:     by_value
      - .offset:         52
        .size:           4
        .value_kind:     by_value
      - .offset:         56
        .size:           4
        .value_kind:     by_value
      - .offset:         60
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         64
        .size:           8
        .value_kind:     global_buffer
      - .offset:         72
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         76
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         80
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         84
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         86
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         88
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         90
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         92
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         94
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         112
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         120
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         128
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         136
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 328
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z23avg_pool_fwd_f32_kerneljPKfiiiiiiiiiiiiPf.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     17
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .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:     by_value
      - .offset:         36
        .size:           4
        .value_kind:     by_value
      - .offset:         40
        .size:           4
        .value_kind:     by_value
      - .offset:         44
        .size:           4
        .value_kind:     by_value
      - .offset:         48
        .size:           4
        .value_kind:     by_value
      - .offset:         52
        .size:           4
        .value_kind:     by_value
      - .offset:         56
        .size:           4
        .value_kind:     by_value
      - .offset:         60
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         64
        .size:           8
        .value_kind:     global_buffer
      - .offset:         72
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         76
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         80
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         84
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         86
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         88
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         90
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         92
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         94
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         112
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         120
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         128
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         136
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 328
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf
    .private_segment_fixed_size: 0
    .sgpr_count:     20
    .sgpr_spill_count: 0
    .symbol:         _Z23avg_pool_bwd_f32_kerneljPKfiiiiiiiiiiiiPf.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     24
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
