/*
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.
*/

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

#define OFFSET_BANK(idx) ({ __typeof__ (idx) _idx = idx; ((_idx) + ((_idx) / 32)); })

__global__ void blockreduce_argmax_kernel(
    uint32_t block_dim,
    uint32_t num_blocks,
    const float *xs,
    float *x_max,
    uint32_t *x_argmax)
{
  __shared__ float cache[1024 + 32];
  __shared__ uint32_t cache_idx[1024 + 32];
  uint32_t tid = threadIdx.x;
  uint32_t block = blockIdx.x;
  uint32_t i = tid + block * block_dim;
  if (tid < block_dim && block < num_blocks) {
    cache[OFFSET_BANK(tid)] = xs[i];
  } else {
    cache[OFFSET_BANK(tid)] = -CUDART_INF_F;
  }
  cache_idx[OFFSET_BANK(tid)] = tid;
  __syncthreads();
  for (uint32_t s = 1; s < blockDim.x; s *= 2) {
    if (tid < block_dim && block < num_blocks) {
      if (tid % (2*s) == 0 && (tid + s) < block_dim && cache[OFFSET_BANK(tid)] < cache[OFFSET_BANK(tid + s)]) {
        cache[OFFSET_BANK(tid)]     = cache[OFFSET_BANK(tid + s)];
        cache_idx[OFFSET_BANK(tid)] = cache_idx[OFFSET_BANK(tid + s)];
      }
    }
    __syncthreads();
  }
  if (tid < block_dim && block < num_blocks) {
    if (tid == 0) {
      x_max[block] = cache[0];
      if (NULL != x_argmax) {
        x_argmax[block] = cache_idx[0];
      }
    }
  }
}

extern "C" void arraydiff_cuda_kernel_blockreduce_max_argmax_f32(
    size_t block_dim,
    size_t num_blocks,
    const float *xs,
    float *xs_max,
    uint32_t *xs_argmax,
    cudaStream_t stream)
{
  // XXX: assert(block_dim <= 1024);
  // FIXME(20151022): could make more efficient use of blocks but w/e.
  blockreduce_argmax_kernel<<<num_blocks, 1024, 0, stream>>>(
      block_dim, num_blocks, xs, xs_max, xs_argmax);
}

/*__global__ void blockreduce_sum_kernel(
    uint32_t block_dim,
    uint32_t num_blocks,
    const float *xs,
    float *xs_sum)
{
  __shared__ float cache[1024 + 32];
  uint32_t tid = threadIdx.x;
  uint32_t block = blockIdx.x;
  uint32_t i = tid + block * block_dim;
  if (tid < block_dim && block < num_blocks) {
    cache[OFFSET_BANK(tid)] = xs[i];
  } else {
    cache[OFFSET_BANK(tid)] = 0.0f;
  }
  __syncthreads();
  for (uint32_t s = 1; s < blockDim.x; s *= 2) {
    if (tid < block_dim && block < num_blocks) {
      if (tid % (2*s) == 0 && (tid + s) < block_dim) {
        cache[OFFSET_BANK(tid)] += cache[OFFSET_BANK(tid + s)];
      }
    }
    __syncthreads();
  }
  if (tid < block_dim && block < num_blocks) {
    if (tid == 0) {
      xs_sum[block] = cache[0];
    }
  }
}*/

__global__ void blockreduce_sum_kernel(
    uint32_t block_dim,
    uint32_t num_blocks,
    const float *xs,
    float *xs_sum)
{
  __shared__ float cache[1024];
  uint32_t tid = threadIdx.x;
  uint32_t block = blockIdx.x;
  uint32_t idx = tid + block_dim * block;
  if (tid < block_dim && block < num_blocks) {
    cache[tid] = xs[idx];
  } else {
    cache[tid] = 0.0f;
  }
  __syncthreads();
  threadblock1024_reduce_sum_f32(cache);
  if (tid < block_dim && block < num_blocks) {
    if (tid == 0) {
      xs_sum[block] = cache[0];
    }
  }
}

extern "C" void arraydiff_cuda_kernel_blockreduce_sum_f32(
    size_t block_dim,
    size_t num_blocks,
    const float *xs,
    float *xs_sum,
    cudaStream_t stream)
{
  // XXX: assert(block_dim <= 1024);
  // FIXME(20151022): could make more efficient use of blocks but w/e.
  blockreduce_sum_kernel<<<num_blocks, 1024, 0, stream>>>(
      block_dim, num_blocks, xs, xs_sum);
}

__global__ void reduce_index_fwd_f32_kernel(
    uint32_t dim,
    uint32_t batch_sz,
    const float *x,
    const uint32_t *index,
    float *y)
{
  uint32_t idx = threadIdx.x + blockDim.x * blockIdx.x;
  if (idx < batch_sz) {
    y[idx] = x[index[idx] + dim * idx];
  }
}

extern "C" void arraydiff_cuda_kernel_reduce_index_fwd_f32(
    size_t dim,
    size_t batch_sz,
    const float *x,
    const uint32_t *index,
    float *y,
    cudaStream_t stream)
{
  size_t n = batch_sz;
  reduce_index_fwd_f32_kernel<<<(n+1024-1)/1024, 1024, 0, stream>>>(
      dim, batch_sz, x, index, y);
}

__global__ void reduce_index_bwd_f32_kernel(
    uint32_t dim,
    uint32_t batch_sz,
    const float *dy,
    const uint32_t *index,
    float *dx)
{
  uint32_t idx = threadIdx.x + blockDim.x * blockIdx.x;
  if (idx < batch_sz) {
    dx[index[idx] + dim * idx] += dy[idx];
  }
}

extern "C" void arraydiff_cuda_kernel_reduce_index_bwd_f32(
    size_t dim,
    size_t batch_sz,
    const float *dy,
    const uint32_t *index,
    float *dx,
    cudaStream_t stream)
{
  size_t n = batch_sz;
  reduce_index_bwd_f32_kernel<<<(n+1024-1)/1024, 1024, 0, stream>>>(
      dim, batch_sz, dy, index, dx);
}

	.file	"tmpxft_00236b7a_00000000-6_reduce.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	_Z52__device_stub__Z25blockreduce_argmax_kerneljjPKfPfPjjjPKfPfPj
	.type	_Z52__device_stub__Z25blockreduce_argmax_kerneljjPKfPfPjjjPKfPfPj, @function
_Z52__device_stub__Z25blockreduce_argmax_kerneljjPKfPfPjjjPKfPfPj:
.LFB2038:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movl	%edi, 28(%rsp)
	leaq	48(%rsp), %rdi
	movl	%esi, 24(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 16(%rsp)
	leaq	32(%rsp), %rdx
	movq	%rcx, 8(%rsp)
	leaq	40(%rsp), %rcx
	movq	%r8, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	%rsp, %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z25blockreduce_argmax_kerneljjPKfPfPj(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 176
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L2:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2038:
	.size	_Z52__device_stub__Z25blockreduce_argmax_kerneljjPKfPfPjjjPKfPfPj, .-_Z52__device_stub__Z25blockreduce_argmax_kerneljjPKfPfPjjjPKfPfPj
	.globl	_Z25blockreduce_argmax_kerneljjPKfPfPj
	.type	_Z25blockreduce_argmax_kerneljjPKfPfPj, @function
_Z25blockreduce_argmax_kerneljjPKfPfPj:
.LFB2039:
	.cfi_startproc
	endbr64
	jmp	_Z52__device_stub__Z25blockreduce_argmax_kerneljjPKfPfPjjjPKfPfPj
	.cfi_endproc
.LFE2039:
	.size	_Z25blockreduce_argmax_kerneljjPKfPfPj, .-_Z25blockreduce_argmax_kerneljjPKfPfPj
	.globl	arraydiff_cuda_kernel_blockreduce_max_argmax_f32
	.type	arraydiff_cuda_kernel_blockreduce_max_argmax_f32, @function
arraydiff_cuda_kernel_blockreduce_max_argmax_f32:
.LFB2010:
	.cfi_startproc
	endbr64
	movabsq	$4294967297, %rax
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movl	%esi, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	movq	%r8, %r13
	xorl	%r8d, %r8d
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	movq	%rcx, %r12
	movl	$1, %ecx
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	movq	%rdx, %rbp
	movl	$4194305, %edx
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	movq	%rdi, %rbx
	salq	$10, %rdx
	subq	$32, %rsp
	.cfi_def_cfa_offset 80
	movl	%esi, 8(%rsp)
	movq	%rax, 12(%rsp)
	movq	8(%rsp), %rdi
	movl	16(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L8
	addq	$32, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 48
	movq	%r13, %r8
	movq	%r12, %rcx
	movq	%rbp, %rdx
	movl	%r14d, %esi
	movl	%ebx, %edi
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	jmp	_Z52__device_stub__Z25blockreduce_argmax_kerneljjPKfPfPjjjPKfPfPj
.L8:
	.cfi_restore_state
	addq	$32, %rsp
	.cfi_def_cfa_offset 48
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2010:
	.size	arraydiff_cuda_kernel_blockreduce_max_argmax_f32, .-arraydiff_cuda_kernel_blockreduce_max_argmax_f32
	.globl	_Z47__device_stub__Z22blockreduce_sum_kerneljjPKfPfjjPKfPf
	.type	_Z47__device_stub__Z22blockreduce_sum_kerneljjPKfPfjjPKfPf, @function
_Z47__device_stub__Z22blockreduce_sum_kerneljjPKfPfjjPKfPf:
.LFB2040:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movl	%edi, 28(%rsp)
	leaq	56(%rsp), %rdi
	movl	%esi, 24(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 16(%rsp)
	leaq	40(%rsp), %rdx
	movq	%rcx, 8(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L11
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z22blockreduce_sum_kerneljjPKfPf(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 176
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L11:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L13
	call	__stack_chk_fail@PLT
.L13:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2040:
	.size	_Z47__device_stub__Z22blockreduce_sum_kerneljjPKfPfjjPKfPf, .-_Z47__device_stub__Z22blockreduce_sum_kerneljjPKfPfjjPKfPf
	.globl	_Z22blockreduce_sum_kerneljjPKfPf
	.type	_Z22blockreduce_sum_kerneljjPKfPf, @function
_Z22blockreduce_sum_kerneljjPKfPf:
.LFB2041:
	.cfi_startproc
	endbr64
	jmp	_Z47__device_stub__Z22blockreduce_sum_kerneljjPKfPfjjPKfPf
	.cfi_endproc
.LFE2041:
	.size	_Z22blockreduce_sum_kerneljjPKfPf, .-_Z22blockreduce_sum_kerneljjPKfPf
	.globl	arraydiff_cuda_kernel_blockreduce_sum_f32
	.type	arraydiff_cuda_kernel_blockreduce_sum_f32, @function
arraydiff_cuda_kernel_blockreduce_sum_f32:
.LFB2011:
	.cfi_startproc
	endbr64
	movabsq	$4294967297, %rax
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movl	%esi, %r13d
	movq	%r8, %r9
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	xorl	%r8d, %r8d
	movq	%rcx, %r12
	movl	$1, %ecx
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movq	%rdx, %rbp
	movl	$4194305, %edx
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movq	%rdi, %rbx
	salq	$10, %rdx
	subq	$40, %rsp
	.cfi_def_cfa_offset 80
	movl	%esi, 8(%rsp)
	movq	%rax, 12(%rsp)
	movq	8(%rsp), %rdi
	movl	16(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L16
	addq	$40, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 40
	movq	%r12, %rcx
	movq	%rbp, %rdx
	movl	%r13d, %esi
	movl	%ebx, %edi
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	jmp	_Z47__device_stub__Z22blockreduce_sum_kerneljjPKfPfjjPKfPf
.L16:
	.cfi_restore_state
	addq	$40, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2011:
	.size	arraydiff_cuda_kernel_blockreduce_sum_f32, .-arraydiff_cuda_kernel_blockreduce_sum_f32
	.globl	_Z55__device_stub__Z27reduce_index_fwd_f32_kerneljjPKfPKjPfjjPKfPKjPf
	.type	_Z55__device_stub__Z27reduce_index_fwd_f32_kerneljjPKfPKjPfjjPKfPKjPf, @function
_Z55__device_stub__Z27reduce_index_fwd_f32_kerneljjPKfPKjPfjjPKfPKjPf:
.LFB2042:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movl	%edi, 28(%rsp)
	leaq	48(%rsp), %rdi
	movl	%esi, 24(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 16(%rsp)
	leaq	32(%rsp), %rdx
	movq	%rcx, 8(%rsp)
	leaq	40(%rsp), %rcx
	movq	%r8, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	%rsp, %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L19
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z27reduce_index_fwd_f32_kerneljjPKfPKjPf(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 176
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L19:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L21
	call	__stack_chk_fail@PLT
.L21:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2042:
	.size	_Z55__device_stub__Z27reduce_index_fwd_f32_kerneljjPKfPKjPfjjPKfPKjPf, .-_Z55__device_stub__Z27reduce_index_fwd_f32_kerneljjPKfPKjPfjjPKfPKjPf
	.globl	_Z27reduce_index_fwd_f32_kerneljjPKfPKjPf
	.type	_Z27reduce_index_fwd_f32_kerneljjPKfPKjPf, @function
_Z27reduce_index_fwd_f32_kerneljjPKfPKjPf:
.LFB2043:
	.cfi_startproc
	endbr64
	jmp	_Z55__device_stub__Z27reduce_index_fwd_f32_kerneljjPKfPKjPfjjPKfPKjPf
	.cfi_endproc
.LFE2043:
	.size	_Z27reduce_index_fwd_f32_kerneljjPKfPKjPf, .-_Z27reduce_index_fwd_f32_kerneljjPKfPKjPf
	.globl	arraydiff_cuda_kernel_reduce_index_fwd_f32
	.type	arraydiff_cuda_kernel_reduce_index_fwd_f32, @function
arraydiff_cuda_kernel_reduce_index_fwd_f32:
.LFB2012:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	leaq	1023(%rsi), %rax
	movq	%r8, %r14
	xorl	%r8d, %r8d
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	shrq	$10, %rax
	movq	%rcx, %r13
	movl	$1, %ecx
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	movq	%rdx, %r12
	movl	$4194305, %edx
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	movq	%rdi, %rbp
	salq	$10, %rdx
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	movq	%rsi, %rbx
	subq	$32, %rsp
	.cfi_def_cfa_offset 80
	movl	%eax, 8(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 12(%rsp)
	movq	8(%rsp), %rdi
	movl	16(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L24
	addq	$32, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 48
	movq	%r14, %r8
	movq	%r13, %rcx
	movq	%r12, %rdx
	movl	%ebx, %esi
	movl	%ebp, %edi
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	jmp	_Z55__device_stub__Z27reduce_index_fwd_f32_kerneljjPKfPKjPfjjPKfPKjPf
.L24:
	.cfi_restore_state
	addq	$32, %rsp
	.cfi_def_cfa_offset 48
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2012:
	.size	arraydiff_cuda_kernel_reduce_index_fwd_f32, .-arraydiff_cuda_kernel_reduce_index_fwd_f32
	.globl	_Z55__device_stub__Z27reduce_index_bwd_f32_kerneljjPKfPKjPfjjPKfPKjPf
	.type	_Z55__device_stub__Z27reduce_index_bwd_f32_kerneljjPKfPKjPfjjPKfPKjPf, @function
_Z55__device_stub__Z27reduce_index_bwd_f32_kerneljjPKfPKjPfjjPKfPKjPf:
.LFB2044:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movl	%edi, 28(%rsp)
	leaq	48(%rsp), %rdi
	movl	%esi, 24(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 16(%rsp)
	leaq	32(%rsp), %rdx
	movq	%rcx, 8(%rsp)
	leaq	40(%rsp), %rcx
	movq	%r8, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	%rsp, %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L27
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z27reduce_index_bwd_f32_kerneljjPKfPKjPf(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 176
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L27:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L29
	call	__stack_chk_fail@PLT
.L29:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2044:
	.size	_Z55__device_stub__Z27reduce_index_bwd_f32_kerneljjPKfPKjPfjjPKfPKjPf, .-_Z55__device_stub__Z27reduce_index_bwd_f32_kerneljjPKfPKjPfjjPKfPKjPf
	.globl	_Z27reduce_index_bwd_f32_kerneljjPKfPKjPf
	.type	_Z27reduce_index_bwd_f32_kerneljjPKfPKjPf, @function
_Z27reduce_index_bwd_f32_kerneljjPKfPKjPf:
.LFB2045:
	.cfi_startproc
	endbr64
	jmp	_Z55__device_stub__Z27reduce_index_bwd_f32_kerneljjPKfPKjPfjjPKfPKjPf
	.cfi_endproc
.LFE2045:
	.size	_Z27reduce_index_bwd_f32_kerneljjPKfPKjPf, .-_Z27reduce_index_bwd_f32_kerneljjPKfPKjPf
	.globl	arraydiff_cuda_kernel_reduce_index_bwd_f32
	.type	arraydiff_cuda_kernel_reduce_index_bwd_f32, @function
arraydiff_cuda_kernel_reduce_index_bwd_f32:
.LFB2013:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	leaq	1023(%rsi), %rax
	movq	%r8, %r14
	xorl	%r8d, %r8d
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	shrq	$10, %rax
	movq	%rcx, %r13
	movl	$1, %ecx
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	movq	%rdx, %r12
	movl	$4194305, %edx
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	movq	%rdi, %rbp
	salq	$10, %rdx
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	movq	%rsi, %rbx
	subq	$32, %rsp
	.cfi_def_cfa_offset 80
	movl	%eax, 8(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 12(%rsp)
	movq	8(%rsp), %rdi
	movl	16(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L32
	addq	$32, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 48
	movq	%r14, %r8
	movq	%r13, %rcx
	movq	%r12, %rdx
	movl	%ebx, %esi
	movl	%ebp, %edi
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	jmp	_Z55__device_stub__Z27reduce_index_bwd_f32_kerneljjPKfPKjPfjjPKfPKjPf
.L32:
	.cfi_restore_state
	addq	$32, %rsp
	.cfi_def_cfa_offset 48
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2013:
	.size	arraydiff_cuda_kernel_reduce_index_bwd_f32, .-arraydiff_cuda_kernel_reduce_index_bwd_f32
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z27reduce_index_bwd_f32_kerneljjPKfPKjPf"
.LC1:
	.string	"_Z27reduce_index_fwd_f32_kerneljjPKfPKjPf"
.LC2:
	.string	"_Z22blockreduce_sum_kerneljjPKfPf"
.LC3:
	.string	"_Z25blockreduce_argmax_kerneljjPKfPfPj"
	.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	_Z27reduce_index_bwd_f32_kerneljjPKfPKjPf(%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	_Z27reduce_index_fwd_f32_kerneljjPKfPKjPf(%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	_Z22blockreduce_sum_kerneljjPKfPf(%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	_Z25blockreduce_argmax_kerneljjPKfPfPj(%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.
*/


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

#define OFFSET_BANK(idx) ({ __typeof__ (idx) _idx = idx; ((_idx) + ((_idx) / 32)); })

__global__ void blockreduce_argmax_kernel(
    uint32_t block_dim,
    uint32_t num_blocks,
    const float *xs,
    float *x_max,
    uint32_t *x_argmax)
{
  __shared__ float cache[1024 + 32];
  __shared__ uint32_t cache_idx[1024 + 32];
  uint32_t tid = threadIdx.x;
  uint32_t block = blockIdx.x;
  uint32_t i = tid + block * block_dim;
  if (tid < block_dim && block < num_blocks) {
    cache[OFFSET_BANK(tid)] = xs[i];
  } else {
    cache[OFFSET_BANK(tid)] = -HIP_INF_F;
  }
  cache_idx[OFFSET_BANK(tid)] = tid;
  __syncthreads();
  for (uint32_t s = 1; s < blockDim.x; s *= 2) {
    if (tid < block_dim && block < num_blocks) {
      if (tid % (2*s) == 0 && (tid + s) < block_dim && cache[OFFSET_BANK(tid)] < cache[OFFSET_BANK(tid + s)]) {
        cache[OFFSET_BANK(tid)]     = cache[OFFSET_BANK(tid + s)];
        cache_idx[OFFSET_BANK(tid)] = cache_idx[OFFSET_BANK(tid + s)];
      }
    }
    __syncthreads();
  }
  if (tid < block_dim && block < num_blocks) {
    if (tid == 0) {
      x_max[block] = cache[0];
      if (NULL != x_argmax) {
        x_argmax[block] = cache_idx[0];
      }
    }
  }
}

extern "C" void arraydiff_cuda_kernel_blockreduce_max_argmax_f32(
    size_t block_dim,
    size_t num_blocks,
    const float *xs,
    float *xs_max,
    uint32_t *xs_argmax,
    hipStream_t stream)
{
  // XXX: assert(block_dim <= 1024);
  // FIXME(20151022): could make more efficient use of blocks but w/e.
  blockreduce_argmax_kernel<<<num_blocks, 1024, 0, stream>>>(
      block_dim, num_blocks, xs, xs_max, xs_argmax);
}

/*__global__ void blockreduce_sum_kernel(
    uint32_t block_dim,
    uint32_t num_blocks,
    const float *xs,
    float *xs_sum)
{
  __shared__ float cache[1024 + 32];
  uint32_t tid = threadIdx.x;
  uint32_t block = blockIdx.x;
  uint32_t i = tid + block * block_dim;
  if (tid < block_dim && block < num_blocks) {
    cache[OFFSET_BANK(tid)] = xs[i];
  } else {
    cache[OFFSET_BANK(tid)] = 0.0f;
  }
  __syncthreads();
  for (uint32_t s = 1; s < blockDim.x; s *= 2) {
    if (tid < block_dim && block < num_blocks) {
      if (tid % (2*s) == 0 && (tid + s) < block_dim) {
        cache[OFFSET_BANK(tid)] += cache[OFFSET_BANK(tid + s)];
      }
    }
    __syncthreads();
  }
  if (tid < block_dim && block < num_blocks) {
    if (tid == 0) {
      xs_sum[block] = cache[0];
    }
  }
}*/

__global__ void blockreduce_sum_kernel(
    uint32_t block_dim,
    uint32_t num_blocks,
    const float *xs,
    float *xs_sum)
{
  __shared__ float cache[1024];
  uint32_t tid = threadIdx.x;
  uint32_t block = blockIdx.x;
  uint32_t idx = tid + block_dim * block;
  if (tid < block_dim && block < num_blocks) {
    cache[tid] = xs[idx];
  } else {
    cache[tid] = 0.0f;
  }
  __syncthreads();
  threadblock1024_reduce_sum_f32(cache);
  if (tid < block_dim && block < num_blocks) {
    if (tid == 0) {
      xs_sum[block] = cache[0];
    }
  }
}

extern "C" void arraydiff_cuda_kernel_blockreduce_sum_f32(
    size_t block_dim,
    size_t num_blocks,
    const float *xs,
    float *xs_sum,
    hipStream_t stream)
{
  // XXX: assert(block_dim <= 1024);
  // FIXME(20151022): could make more efficient use of blocks but w/e.
  blockreduce_sum_kernel<<<num_blocks, 1024, 0, stream>>>(
      block_dim, num_blocks, xs, xs_sum);
}

__global__ void reduce_index_fwd_f32_kernel(
    uint32_t dim,
    uint32_t batch_sz,
    const float *x,
    const uint32_t *index,
    float *y)
{
  uint32_t idx = threadIdx.x + blockDim.x * blockIdx.x;
  if (idx < batch_sz) {
    y[idx] = x[index[idx] + dim * idx];
  }
}

extern "C" void arraydiff_cuda_kernel_reduce_index_fwd_f32(
    size_t dim,
    size_t batch_sz,
    const float *x,
    const uint32_t *index,
    float *y,
    hipStream_t stream)
{
  size_t n = batch_sz;
  reduce_index_fwd_f32_kernel<<<(n+1024-1)/1024, 1024, 0, stream>>>(
      dim, batch_sz, x, index, y);
}

__global__ void reduce_index_bwd_f32_kernel(
    uint32_t dim,
    uint32_t batch_sz,
    const float *dy,
    const uint32_t *index,
    float *dx)
{
  uint32_t idx = threadIdx.x + blockDim.x * blockIdx.x;
  if (idx < batch_sz) {
    dx[index[idx] + dim * idx] += dy[idx];
  }
}

extern "C" void arraydiff_cuda_kernel_reduce_index_bwd_f32(
    size_t dim,
    size_t batch_sz,
    const float *dy,
    const uint32_t *index,
    float *dx,
    hipStream_t stream)
{
  size_t n = batch_sz;
  reduce_index_bwd_f32_kernel<<<(n+1024-1)/1024, 1024, 0, stream>>>(
      dim, batch_sz, dy, index, dx);
}

	.text
	.file	"reduce.hip"
	.globl	_Z40__device_stub__blockreduce_argmax_kerneljjPKfPfPj # -- Begin function _Z40__device_stub__blockreduce_argmax_kerneljjPKfPfPj
	.type	_Z40__device_stub__blockreduce_argmax_kerneljjPKfPfPj,@function
_Z40__device_stub__blockreduce_argmax_kerneljjPKfPfPj: # @_Z40__device_stub__blockreduce_argmax_kerneljjPKfPfPj
	.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	$128, %rsp
	.cfi_def_cfa_offset 176
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	4(%rsp), %rax
	movl	%edi, (%rax)
	movq	%rsp, %rdi
	movl	%esi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	24(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%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	$_Z25blockreduce_argmax_kerneljjPKfPfPj, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$144, %rsp
	.cfi_adjust_cfa_offset -144
	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	_Z40__device_stub__blockreduce_argmax_kerneljjPKfPfPj, .Lfunc_end0-_Z40__device_stub__blockreduce_argmax_kerneljjPKfPfPj
	.cfi_endproc
                                        # -- End function
	.globl	arraydiff_cuda_kernel_blockreduce_max_argmax_f32 # -- Begin function arraydiff_cuda_kernel_blockreduce_max_argmax_f32
	.type	arraydiff_cuda_kernel_blockreduce_max_argmax_f32,@function
arraydiff_cuda_kernel_blockreduce_max_argmax_f32: # @arraydiff_cuda_kernel_blockreduce_max_argmax_f32
	.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
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%r8, %rbx
	movq	%rcx, %r14
	movq	%rdx, %r15
	movq	%rsi, %r12
	movq	%rdi, %r13
	movl	%r12d, %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_2
# %bb.1:
	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
.LBB1_2:
	.cfi_def_cfa_offset 48
	movl	%r13d, %edi
	movl	%r12d, %esi
	movq	%r15, %rdx
	movq	%r14, %rcx
	movq	%rbx, %r8
	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
	jmp	_Z40__device_stub__blockreduce_argmax_kerneljjPKfPfPj # TAILCALL
.Lfunc_end1:
	.size	arraydiff_cuda_kernel_blockreduce_max_argmax_f32, .Lfunc_end1-arraydiff_cuda_kernel_blockreduce_max_argmax_f32
	.cfi_endproc
                                        # -- End function
	.globl	_Z37__device_stub__blockreduce_sum_kerneljjPKfPf # -- Begin function _Z37__device_stub__blockreduce_sum_kerneljjPKfPf
	.type	_Z37__device_stub__blockreduce_sum_kerneljjPKfPf,@function
_Z37__device_stub__blockreduce_sum_kerneljjPKfPf: # @_Z37__device_stub__blockreduce_sum_kerneljjPKfPf
	.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	$112, %rsp
	.cfi_def_cfa_offset 160
	.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	8(%rsp), %rdi
	movl	%esi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%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	$_Z22blockreduce_sum_kerneljjPKfPf, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	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	_Z37__device_stub__blockreduce_sum_kerneljjPKfPf, .Lfunc_end2-_Z37__device_stub__blockreduce_sum_kerneljjPKfPf
	.cfi_endproc
                                        # -- End function
	.globl	arraydiff_cuda_kernel_blockreduce_sum_f32 # -- Begin function arraydiff_cuda_kernel_blockreduce_sum_f32
	.type	arraydiff_cuda_kernel_blockreduce_sum_f32,@function
arraydiff_cuda_kernel_blockreduce_sum_f32: # @arraydiff_cuda_kernel_blockreduce_sum_f32
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%r8, %r9
	movq	%rcx, %rbx
	movq	%rdx, %r14
	movq	%rsi, %r15
	movq	%rdi, %r12
	movl	%r15d, %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_2
# %bb.1:
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.LBB3_2:
	.cfi_def_cfa_offset 48
	movl	%r12d, %edi
	movl	%r15d, %esi
	movq	%r14, %rdx
	movq	%rbx, %rcx
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	jmp	_Z37__device_stub__blockreduce_sum_kerneljjPKfPf # TAILCALL
.Lfunc_end3:
	.size	arraydiff_cuda_kernel_blockreduce_sum_f32, .Lfunc_end3-arraydiff_cuda_kernel_blockreduce_sum_f32
	.cfi_endproc
                                        # -- End function
	.globl	_Z42__device_stub__reduce_index_fwd_f32_kerneljjPKfPKjPf # -- Begin function _Z42__device_stub__reduce_index_fwd_f32_kerneljjPKfPKjPf
	.type	_Z42__device_stub__reduce_index_fwd_f32_kerneljjPKfPKjPf,@function
_Z42__device_stub__reduce_index_fwd_f32_kerneljjPKfPKjPf: # @_Z42__device_stub__reduce_index_fwd_f32_kerneljjPKfPKjPf
	.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	$128, %rsp
	.cfi_def_cfa_offset 176
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	4(%rsp), %rax
	movl	%edi, (%rax)
	movq	%rsp, %rdi
	movl	%esi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	24(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%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	$_Z27reduce_index_fwd_f32_kerneljjPKfPKjPf, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$144, %rsp
	.cfi_adjust_cfa_offset -144
	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	_Z42__device_stub__reduce_index_fwd_f32_kerneljjPKfPKjPf, .Lfunc_end4-_Z42__device_stub__reduce_index_fwd_f32_kerneljjPKfPKjPf
	.cfi_endproc
                                        # -- End function
	.globl	arraydiff_cuda_kernel_reduce_index_fwd_f32 # -- Begin function arraydiff_cuda_kernel_reduce_index_fwd_f32
	.type	arraydiff_cuda_kernel_reduce_index_fwd_f32,@function
arraydiff_cuda_kernel_reduce_index_fwd_f32: # @arraydiff_cuda_kernel_reduce_index_fwd_f32
	.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
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%r8, %rbx
	movq	%rcx, %r14
	movq	%rdx, %r15
	movq	%rsi, %r12
	movq	%rdi, %r13
	leaq	1023(%rsi), %rax
	shrq	$10, %rax
	movl	%eax, %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_2
# %bb.1:
	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
.LBB5_2:
	.cfi_def_cfa_offset 48
	movl	%r13d, %edi
	movl	%r12d, %esi
	movq	%r15, %rdx
	movq	%r14, %rcx
	movq	%rbx, %r8
	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
	jmp	_Z42__device_stub__reduce_index_fwd_f32_kerneljjPKfPKjPf # TAILCALL
.Lfunc_end5:
	.size	arraydiff_cuda_kernel_reduce_index_fwd_f32, .Lfunc_end5-arraydiff_cuda_kernel_reduce_index_fwd_f32
	.cfi_endproc
                                        # -- End function
	.globl	_Z42__device_stub__reduce_index_bwd_f32_kerneljjPKfPKjPf # -- Begin function _Z42__device_stub__reduce_index_bwd_f32_kerneljjPKfPKjPf
	.type	_Z42__device_stub__reduce_index_bwd_f32_kerneljjPKfPKjPf,@function
_Z42__device_stub__reduce_index_bwd_f32_kerneljjPKfPKjPf: # @_Z42__device_stub__reduce_index_bwd_f32_kerneljjPKfPKjPf
	.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	$128, %rsp
	.cfi_def_cfa_offset 176
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	4(%rsp), %rax
	movl	%edi, (%rax)
	movq	%rsp, %rdi
	movl	%esi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	24(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%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	$_Z27reduce_index_bwd_f32_kerneljjPKfPKjPf, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$144, %rsp
	.cfi_adjust_cfa_offset -144
	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	_Z42__device_stub__reduce_index_bwd_f32_kerneljjPKfPKjPf, .Lfunc_end6-_Z42__device_stub__reduce_index_bwd_f32_kerneljjPKfPKjPf
	.cfi_endproc
                                        # -- End function
	.globl	arraydiff_cuda_kernel_reduce_index_bwd_f32 # -- Begin function arraydiff_cuda_kernel_reduce_index_bwd_f32
	.type	arraydiff_cuda_kernel_reduce_index_bwd_f32,@function
arraydiff_cuda_kernel_reduce_index_bwd_f32: # @arraydiff_cuda_kernel_reduce_index_bwd_f32
	.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
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%r8, %rbx
	movq	%rcx, %r14
	movq	%rdx, %r15
	movq	%rsi, %r12
	movq	%rdi, %r13
	leaq	1023(%rsi), %rax
	shrq	$10, %rax
	movl	%eax, %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_2
# %bb.1:
	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
.LBB7_2:
	.cfi_def_cfa_offset 48
	movl	%r13d, %edi
	movl	%r12d, %esi
	movq	%r15, %rdx
	movq	%r14, %rcx
	movq	%rbx, %r8
	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
	jmp	_Z42__device_stub__reduce_index_bwd_f32_kerneljjPKfPKjPf # TAILCALL
.Lfunc_end7:
	.size	arraydiff_cuda_kernel_reduce_index_bwd_f32, .Lfunc_end7-arraydiff_cuda_kernel_reduce_index_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	$_Z25blockreduce_argmax_kerneljjPKfPfPj, %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	$_Z22blockreduce_sum_kerneljjPKfPf, %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	$_Z27reduce_index_fwd_f32_kerneljjPKfPKjPf, %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	$_Z27reduce_index_bwd_f32_kerneljjPKfPKjPf, %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	_Z25blockreduce_argmax_kerneljjPKfPfPj,@object # @_Z25blockreduce_argmax_kerneljjPKfPfPj
	.section	.rodata,"a",@progbits
	.globl	_Z25blockreduce_argmax_kerneljjPKfPfPj
	.p2align	3, 0x0
_Z25blockreduce_argmax_kerneljjPKfPfPj:
	.quad	_Z40__device_stub__blockreduce_argmax_kerneljjPKfPfPj
	.size	_Z25blockreduce_argmax_kerneljjPKfPfPj, 8

	.type	_Z22blockreduce_sum_kerneljjPKfPf,@object # @_Z22blockreduce_sum_kerneljjPKfPf
	.globl	_Z22blockreduce_sum_kerneljjPKfPf
	.p2align	3, 0x0
_Z22blockreduce_sum_kerneljjPKfPf:
	.quad	_Z37__device_stub__blockreduce_sum_kerneljjPKfPf
	.size	_Z22blockreduce_sum_kerneljjPKfPf, 8

	.type	_Z27reduce_index_fwd_f32_kerneljjPKfPKjPf,@object # @_Z27reduce_index_fwd_f32_kerneljjPKfPKjPf
	.globl	_Z27reduce_index_fwd_f32_kerneljjPKfPKjPf
	.p2align	3, 0x0
_Z27reduce_index_fwd_f32_kerneljjPKfPKjPf:
	.quad	_Z42__device_stub__reduce_index_fwd_f32_kerneljjPKfPKjPf
	.size	_Z27reduce_index_fwd_f32_kerneljjPKfPKjPf, 8

	.type	_Z27reduce_index_bwd_f32_kerneljjPKfPKjPf,@object # @_Z27reduce_index_bwd_f32_kerneljjPKfPKjPf
	.globl	_Z27reduce_index_bwd_f32_kerneljjPKfPKjPf
	.p2align	3, 0x0
_Z27reduce_index_bwd_f32_kerneljjPKfPKjPf:
	.quad	_Z42__device_stub__reduce_index_bwd_f32_kerneljjPKfPKjPf
	.size	_Z27reduce_index_bwd_f32_kerneljjPKfPKjPf, 8

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

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z22blockreduce_sum_kerneljjPKfPf"
	.size	.L__unnamed_2, 34

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"_Z27reduce_index_fwd_f32_kerneljjPKfPKjPf"
	.size	.L__unnamed_3, 42

	.type	.L__unnamed_4,@object           # @3
.L__unnamed_4:
	.asciz	"_Z27reduce_index_bwd_f32_kerneljjPKfPKjPf"
	.size	.L__unnamed_4, 42

	.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 _Z40__device_stub__blockreduce_argmax_kerneljjPKfPfPj
	.addrsig_sym _Z37__device_stub__blockreduce_sum_kerneljjPKfPf
	.addrsig_sym _Z42__device_stub__reduce_index_fwd_f32_kerneljjPKfPKjPf
	.addrsig_sym _Z42__device_stub__reduce_index_bwd_f32_kerneljjPKfPKjPf
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z25blockreduce_argmax_kerneljjPKfPfPj
	.addrsig_sym _Z22blockreduce_sum_kerneljjPKfPf
	.addrsig_sym _Z27reduce_index_fwd_f32_kerneljjPKfPKjPf
	.addrsig_sym _Z27reduce_index_bwd_f32_kerneljjPKfPKjPf
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z25blockreduce_argmax_kerneljjPKfPfPj ; -- Begin function _Z25blockreduce_argmax_kerneljjPKfPfPj
	.globl	_Z25blockreduce_argmax_kerneljjPKfPfPj
	.p2align	8
	.type	_Z25blockreduce_argmax_kerneljjPKfPfPj,@function
_Z25blockreduce_argmax_kerneljjPKfPfPj: ; @_Z25blockreduce_argmax_kerneljjPKfPfPj
; %bb.0:
	s_clause 0x2
	s_load_b64 s[12:13], s[0:1], 0x0
	s_load_b128 s[4:7], s[0:1], 0x8
	s_load_b64 s[8:9], s[0:1], 0x18
	v_mov_b32_e32 v2, 0xff800000
	s_mov_b32 s10, s15
	s_waitcnt lgkmcnt(0)
	v_cmp_gt_u32_e32 vcc_lo, s12, v0
	s_cmp_lt_u32 s15, s13
	s_cselect_b32 s3, -1, 0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s11, vcc_lo, s3
	s_and_saveexec_b32 s13, s11
	s_cbranch_execz .LBB0_2
; %bb.1:
	v_mad_u64_u32 v[1:2], null, s10, s12, v[0:1]
	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, s2, s4, v1
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v2, s2, s5, v2, s2
	global_load_b32 v2, v[1:2], off
.LBB0_2:
	s_or_b32 exec_lo, exec_lo, s13
	v_lshrrev_b32_e32 v1, 5, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_add_lshl_u32 v1, v1, v0, 2
	s_waitcnt vmcnt(0)
	ds_store_b32 v1, v2
	ds_store_b32 v1, v0 offset:4224
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_load_b32 s0, s[0:1], 0x2c
	s_waitcnt lgkmcnt(0)
	s_and_b32 s1, s0, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_lt_u32 s1, 2
	s_cbranch_scc1 .LBB0_10
; %bb.3:                                ; %.lr.ph
	v_add_nc_u32_e32 v2, 0x1080, v1
	s_mov_b32 s2, 1
.LBB0_4:                                ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 s5, s2
	s_lshl_b32 s2, s2, 1
	s_and_saveexec_b32 s4, s11
	s_cbranch_execz .LBB0_9
; %bb.5:                                ;   in Loop: Header=BB0_4 Depth=1
	s_add_i32 s0, s2, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v3, s0, v0
	v_cmp_eq_u32_e64 s0, 0, v3
	s_delay_alu instid0(VALU_DEP_1)
	s_and_b32 exec_lo, exec_lo, s0
	s_cbranch_execz .LBB0_9
; %bb.6:                                ;   in Loop: Header=BB0_4 Depth=1
	v_add_nc_u32_e32 v3, s5, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_gt_u32_e64 s0, s12, v3
	s_and_b32 exec_lo, exec_lo, s0
	s_cbranch_execz .LBB0_9
; %bb.7:                                ;   in Loop: Header=BB0_4 Depth=1
	v_lshrrev_b32_e32 v4, 5, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v3, v4, v3
	v_lshlrev_b32_e32 v3, 2, v3
	ds_load_b32 v5, v1
	ds_load_b32 v4, v3
	s_waitcnt lgkmcnt(0)
	v_cmp_lt_f32_e64 s0, v5, v4
	s_delay_alu instid0(VALU_DEP_1)
	s_and_b32 exec_lo, exec_lo, s0
	s_cbranch_execz .LBB0_9
; %bb.8:                                ;   in Loop: Header=BB0_4 Depth=1
	ds_load_b32 v3, v3 offset:4224
	ds_store_b32 v1, v4
	s_waitcnt lgkmcnt(1)
	ds_store_b32 v2, v3
.LBB0_9:                                ; %._crit_edge84
                                        ;   in Loop: Header=BB0_4 Depth=1
	s_or_b32 exec_lo, exec_lo, s4
	s_cmp_ge_u32 s2, s1
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_cbranch_scc0 .LBB0_4
.LBB0_10:                               ; %._crit_edge
	v_cmp_eq_u32_e64 s0, 0, v0
	s_mov_b32 s11, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s0, s0, s3
	s_and_b32 s0, vcc_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_13
; %bb.11:
	v_mov_b32_e32 v0, 0
	s_lshl_b64 s[0:1], s[10:11], 2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s2, s6, s0
	s_addc_u32 s3, s7, s1
	ds_load_b32 v1, v0
	s_cmp_eq_u64 s[8:9], 0
	s_waitcnt lgkmcnt(0)
	global_store_b32 v0, v1, s[2:3]
	s_cbranch_scc1 .LBB0_13
; %bb.12:
	ds_load_b32 v1, v0 offset:4224
	s_add_u32 s0, s8, s0
	s_addc_u32 s1, s9, s1
	s_waitcnt lgkmcnt(0)
	global_store_b32 v0, v1, s[0:1]
.LBB0_13:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z25blockreduce_argmax_kerneljjPKfPfPj
		.amdhsa_group_segment_fixed_size 8448
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.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 6
		.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	_Z25blockreduce_argmax_kerneljjPKfPfPj, .Lfunc_end0-_Z25blockreduce_argmax_kerneljjPKfPfPj
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 536
; NumSgprs: 18
; NumVgprs: 6
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 8448 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 6
; 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	_Z22blockreduce_sum_kerneljjPKfPf ; -- Begin function _Z22blockreduce_sum_kerneljjPKfPf
	.globl	_Z22blockreduce_sum_kerneljjPKfPf
	.p2align	8
	.type	_Z22blockreduce_sum_kerneljjPKfPf,@function
_Z22blockreduce_sum_kerneljjPKfPf:      ; @_Z22blockreduce_sum_kerneljjPKfPf
; %bb.0:
	s_clause 0x1
	s_load_b64 s[6:7], s[0:1], 0x0
	s_load_b128 s[0:3], s[0:1], 0x8
	v_mov_b32_e32 v2, 0
	s_mov_b32 s4, s15
	s_waitcnt lgkmcnt(0)
	v_cmp_gt_u32_e32 vcc_lo, s6, v0
	s_cmp_lt_u32 s15, s7
	s_cselect_b32 s5, -1, 0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s8, vcc_lo, s5
	s_and_saveexec_b32 s7, s8
	s_cbranch_execz .LBB1_2
; %bb.1:
	v_mad_u64_u32 v[1:2], null, s4, s6, v[0:1]
	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, s0, s0, v1
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v2, s0, s1, v2, s0
	global_load_b32 v2, v[1:2], off
.LBB1_2:
	s_or_b32 exec_lo, exec_lo, s7
	v_lshlrev_b32_e32 v1, 2, v0
	s_movk_i32 s1, 0x200
	s_waitcnt vmcnt(0)
	ds_store_b32 v1, v2
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
.LBB1_3:                                ; =>This Inner Loop Header: Depth=1
	v_cmp_gt_u32_e64 s0, s1, v0
	s_delay_alu instid0(VALU_DEP_1)
	s_and_saveexec_b32 s6, s0
	s_cbranch_execz .LBB1_5
; %bb.4:                                ;   in Loop: Header=BB1_3 Depth=1
	v_lshl_add_u32 v2, s1, 2, v1
	ds_load_b32 v2, v2
	ds_load_b32 v3, v1
	s_waitcnt lgkmcnt(0)
	v_add_f32_e32 v2, v2, v3
	ds_store_b32 v1, v2
.LBB1_5:                                ;   in Loop: Header=BB1_3 Depth=1
	s_or_b32 exec_lo, exec_lo, s6
	s_lshr_b32 s0, s1, 1
	s_cmp_gt_u32 s1, 1
	s_mov_b32 s1, s0
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_cbranch_scc1 .LBB1_3
; %bb.6:                                ; %_Z30threadblock1024_reduce_sum_f32Pf.exit
	v_cmp_eq_u32_e64 s0, 0, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	s_and_b32 s0, s0, s5
	s_mov_b32 s5, 0
	s_and_b32 s0, vcc_lo, s0
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB1_8
; %bb.7:
	v_mov_b32_e32 v0, 0
	s_lshl_b64 s[0:1], s[4:5], 2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s0, s2, s0
	s_addc_u32 s1, s3, s1
	ds_load_b32 v1, v0
	s_waitcnt lgkmcnt(0)
	global_store_b32 v0, v1, s[0:1]
.LBB1_8:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z22blockreduce_sum_kerneljjPKfPf
		.amdhsa_group_segment_fixed_size 4096
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 24
		.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 4
		.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	_Z22blockreduce_sum_kerneljjPKfPf, .Lfunc_end1-_Z22blockreduce_sum_kerneljjPKfPf
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 332
; NumSgprs: 18
; NumVgprs: 4
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 4096 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 4
; 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	_Z27reduce_index_fwd_f32_kerneljjPKfPKjPf ; -- Begin function _Z27reduce_index_fwd_f32_kerneljjPKfPKjPf
	.globl	_Z27reduce_index_fwd_f32_kerneljjPKfPKjPf
	.p2align	8
	.type	_Z27reduce_index_fwd_f32_kerneljjPKfPKjPf,@function
_Z27reduce_index_fwd_f32_kerneljjPKfPKjPf: ; @_Z27reduce_index_fwd_f32_kerneljjPKfPKjPf
; %bb.0:
	s_clause 0x1
	s_load_b32 s4, s[0:1], 0x2c
	s_load_b64 s[2:3], s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s4, s4, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s4, v[0:1]
	v_cmp_gt_u32_e32 vcc_lo, s3, v1
	s_and_saveexec_b32 s3, vcc_lo
	s_cbranch_execz .LBB2_2
; %bb.1:
	s_load_b128 s[4:7], s[0:1], 0x8
	v_mov_b32_e32 v2, 0
	s_load_b64 s[0:1], s[0:1], 0x18
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[3:4], 2, v[1:2]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v5, vcc_lo, s6, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e32 v6, vcc_lo, s7, v4, vcc_lo
	global_load_b32 v0, v[5:6], off
	s_waitcnt vmcnt(0)
	v_mad_u64_u32 v[5:6], null, v1, s2, v[0:1]
	v_mov_b32_e32 v6, v2
	v_lshlrev_b64 v[0:1], 2, v[5:6]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s4, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s5, v1, vcc_lo
	global_load_b32 v2, v[0:1], off
	v_add_co_u32 v0, vcc_lo, s0, v3
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v4, vcc_lo
	s_waitcnt vmcnt(0)
	global_store_b32 v[0:1], v2, off
.LBB2_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z27reduce_index_fwd_f32_kerneljjPKfPKjPf
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.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 7
		.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	_Z27reduce_index_fwd_f32_kerneljjPKfPKjPf, .Lfunc_end2-_Z27reduce_index_fwd_f32_kerneljjPKfPKjPf
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 200
; NumSgprs: 18
; NumVgprs: 7
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 7
; Occupancy: 16
; WaveLimiterHint : 1
; 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	_Z27reduce_index_bwd_f32_kerneljjPKfPKjPf ; -- Begin function _Z27reduce_index_bwd_f32_kerneljjPKfPKjPf
	.globl	_Z27reduce_index_bwd_f32_kerneljjPKfPKjPf
	.p2align	8
	.type	_Z27reduce_index_bwd_f32_kerneljjPKfPKjPf,@function
_Z27reduce_index_bwd_f32_kerneljjPKfPKjPf: ; @_Z27reduce_index_bwd_f32_kerneljjPKfPKjPf
; %bb.0:
	s_clause 0x1
	s_load_b32 s4, s[0:1], 0x2c
	s_load_b64 s[2:3], s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s4, s4, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s4, v[0:1]
	v_cmp_gt_u32_e32 vcc_lo, s3, v1
	s_and_saveexec_b32 s3, vcc_lo
	s_cbranch_execz .LBB3_2
; %bb.1:
	s_load_b128 s[4:7], s[0:1], 0x8
	v_mov_b32_e32 v2, 0
	s_load_b64 s[0:1], s[0:1], 0x18
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[3:4], 2, v[1:2]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v5, vcc_lo, s6, v3
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v6, vcc_lo, s7, v4, vcc_lo
	global_load_b32 v0, v[5:6], off
	s_waitcnt vmcnt(0)
	v_mad_u64_u32 v[5:6], null, v1, s2, v[0:1]
	v_mov_b32_e32 v6, v2
	v_add_co_u32 v2, vcc_lo, s4, v3
	v_add_co_ci_u32_e32 v3, vcc_lo, s5, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[5:6]
	v_add_co_u32 v0, vcc_lo, s0, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_load_b32 v2, v[2:3], off
	global_load_b32 v3, v[0:1], off
	s_waitcnt vmcnt(0)
	v_add_f32_e32 v2, v2, v3
	global_store_b32 v[0:1], v2, off
.LBB3_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z27reduce_index_bwd_f32_kerneljjPKfPKjPf
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.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 7
		.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_end3:
	.size	_Z27reduce_index_bwd_f32_kerneljjPKfPKjPf, .Lfunc_end3-_Z27reduce_index_bwd_f32_kerneljjPKfPKjPf
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 216
; NumSgprs: 18
; NumVgprs: 7
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 7
; Occupancy: 16
; WaveLimiterHint : 1
; 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
      - .offset:         4
        .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
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 8448
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z25blockreduce_argmax_kerneljjPKfPfPj
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z25blockreduce_argmax_kerneljjPKfPfPj.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     6
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .offset:         4
        .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
    .group_segment_fixed_size: 4096
    .kernarg_segment_align: 8
    .kernarg_segment_size: 24
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z22blockreduce_sum_kerneljjPKfPf
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z22blockreduce_sum_kerneljjPKfPf.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     4
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .offset:         4
        .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
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z27reduce_index_fwd_f32_kerneljjPKfPKjPf
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z27reduce_index_fwd_f32_kerneljjPKfPKjPf.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     7
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .offset:         4
        .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
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z27reduce_index_bwd_f32_kerneljjPKfPKjPf
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z27reduce_index_bwd_f32_kerneljjPKfPKjPf.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     7
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
