/*-----------
 *
 * matrixMulGlobal.cu
 *
 * This is the source file for matrix multiplication with global memory only.
 *
 * This kernel is from NVIDIA CUDA samples. reduction_kernel.cu.
 *
 * streamsOptBenchmark/reduction_kernel.cu
 *
 * By Hao Li
 *
 *------------
 */

/*
    Parallel reduction kernels
*/

#include <stdio.h>
// #include "structs.h"
// #include "functions.h"

// Utility class used to avoid linker errors with extern
// unsized shared memory arrays with templated type
struct SharedMemory
{
    __device__ inline operator       float *()
    {
        extern __shared__ float __smem[];
        return (float *)__smem;
    }

    __device__ inline operator const float *() const
    {
        extern __shared__ float __smem[];
        return (float *)__smem;
    }
};

// // specialize for double to avoid unaligned memory
// // access compile errors
// template<>
// struct SharedMemory<double>
// {
//     __device__ inline operator       double *()
//     {
//         extern __shared__ double __smem_d[];
//         return (double *)__smem_d;
//     }

//     __device__ inline operator const double *() const
//     {
//         extern __shared__ double __smem_d[];
//         return (double *)__smem_d;
//     }
// };

/*
    Parallel sum reduction using shared memory
    - takes log(n) steps for n input elements
    - uses n threads
    - only works for power-of-2 arrays
*/

/* This reduction interleaves which threads are active by using the modulo
   operator.  This operator is very expensive on GPUs, and the interleaved
   inactivity means that no whole warps are active, which is also very
   inefficient */
__global__ void reduce0(float *g_idata, float *g_odata, unsigned int n)
{
    for(int l = 0; l < 100000; l++)
    {
    float *sdata = SharedMemory();

    // load shared mem
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;

    sdata[tid] = (i < n) ? g_idata[i] : 0;

    __syncthreads();

    // do reduction in shared mem
    for (unsigned int s=1; s < blockDim.x; s *= 2)
    {
        // modulo arithmetic is slow!
        if ((tid % (2*s)) == 0)
        {
            sdata[tid] += sdata[tid + s];
        }

        __syncthreads();
    }

    // write result for this block to global mem
    if (tid == 0) g_odata[blockIdx.x] = sdata[0];
    }
}

/* This version uses contiguous threads, but its interleaved
   addressing results in many shared memory bank conflicts.
*/
__global__ void reduce1(float *g_idata, float *g_odata, unsigned int n)
{
    for(int l = 0; l < 100000; l++)
    {
    float *sdata = SharedMemory();

    // load shared mem
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;

    sdata[tid] = (i < n) ? g_idata[i] : 0;

    __syncthreads();

    // do reduction in shared mem
    for (unsigned int s=1; s < blockDim.x; s *= 2)
    {
        int index = 2 * s * tid;

        if (index < blockDim.x)
        {
            sdata[index] += sdata[index + s];
        }

        __syncthreads();
    }

    // write result for this block to global mem
    if (tid == 0) g_odata[blockIdx.x] = sdata[0];
    }
}

/*
    This version uses sequential addressing -- no divergence or bank conflicts.
*/
__global__ void reduce2(float *g_idata, float *g_odata, unsigned int n)
{
    for(int l = 0; l < 100000; l++)
    {
    float *sdata = SharedMemory();

    // load shared mem
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;

    sdata[tid] = (i < n) ? g_idata[i] : 0;

    __syncthreads();

    // do reduction in shared mem
    for (unsigned int s=blockDim.x/2; s>0; s>>=1)
    {
        if (tid < s)
        {
            sdata[tid] += sdata[tid + s];
        }

        __syncthreads();
    }

    // write result for this block to global mem
    if (tid == 0) g_odata[blockIdx.x] = sdata[0];
    }
}

/*
    This version uses n/2 threads --
    it performs the first level of reduction when reading from global memory.
*/
__global__ void reduce3(float *g_idata, float *g_odata, unsigned int n)
{
    for(int l = 0; l < 100000; l++)
    {
    float *sdata = SharedMemory();

    // perform first level of reduction,
    // reading from global memory, writing to shared memory
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;

    int mySum = (i < n) ? g_idata[i] : 0;

    if (i + blockDim.x < n)
        mySum += g_idata[i+blockDim.x];

    sdata[tid] = mySum;
    __syncthreads();

    // do reduction in shared mem
    for (unsigned int s=blockDim.x/2; s>0; s>>=1)
    {
        if (tid < s)
        {
            sdata[tid] = mySum = mySum + sdata[tid + s];
        }

        __syncthreads();
    }

    // write result for this block to global mem
    if (tid == 0) g_odata[blockIdx.x] = mySum;
    }
}

////////////////////////////////////////////////////////////////////////////////
// Wrapper function for kernel launch
////////////////////////////////////////////////////////////////////////////////
// void reduce(int size, int threads, int blocks,
//        int whichKernel, int *d_idata, int *d_odata)
// {
//     dim3 dimBlock(threads, 1, 1);
//     dim3 dimGrid(blocks, 1, 1);

//     // when there is only one warp per block, we need to allocate two warps
//     // worth of shared memory so that we don't index shared memory out of bounds
//     int smemSize = (threads <= 32) ? 2 * threads * sizeof(int) : threads * sizeof(int);

//     // choose which of the optimized versions of reduction to launch
//     switch (whichKernel)
//     {
//         case 0:
//             reduce0<<< dimGrid, dimBlock, smemSize >>>(d_idata, d_odata, size);
//             break;

//         case 1:
//             reduce1<<< dimGrid, dimBlock, smemSize >>>(d_idata, d_odata, size);
//             break;

//         case 2:
//             reduce2<<< dimGrid, dimBlock, smemSize >>>(d_idata, d_odata, size);
//             break;

//         case 3:
//             reduce3<<< dimGrid, dimBlock, smemSize >>>(d_idata, d_odata, size);
//             break;
//     }
// }

// int main(int argc, char **argv){
//  int matrixDataSize = sizeof(int) * MATRIX_SIZE * MATRIX_SIZE;

//  Matrix h_A, h_C;
//  Matrix d_A, d_C;

//  initMatrix(h_A, matrixDataSize, onHOST);
//  initMatrix(h_C, matrixDataSize, onHOST);
//  initMatrix(d_A, matrixDataSize, onDEVICE);
//  initMatrix(d_C, matrixDataSize, onDEVICE);

//  cudaMemcpy(d_A.elements, h_A.elements, matrixDataSize, cudaMemcpyHostToDevice);

//  // Invoke kernel
//  // dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
//  // dim3 dimGrid(h_B.width / dimBlock.x, h_A.height / dimBlock.y);

//  // execute the kernel
//  for(int i =0; i < 4; i++){
//  reduce(matrixDataSize, h_A.width / BLOCK_SIZE * h_A.height / BLOCK_SIZE, 
//     BLOCK_SIZE*BLOCK_SIZE, i, d_A.elements, d_C.elements);
// }

//  cudaMemcpy(h_C.elements, d_C.elements, matrixDataSize, cudaMemcpyDeviceToHost);

//  free(h_A.elements);
//  free(h_C.elements);
//  cudaFree(d_A.elements);
//  cudaFree(d_C.elements);

//  return 0;
// }

	.file	"tmpxft_0021c677_00000000-6_reduction_kernel.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2031:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2031:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z29__device_stub__Z7reduce0PfS_jPfS_j
	.type	_Z29__device_stub__Z7reduce0PfS_jPfS_j, @function
_Z29__device_stub__Z7reduce0PfS_jPfS_j:
.LFB2053:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%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 152
	leaq	_Z7reduce0PfS_j(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	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 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L2:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2053:
	.size	_Z29__device_stub__Z7reduce0PfS_jPfS_j, .-_Z29__device_stub__Z7reduce0PfS_jPfS_j
	.globl	_Z7reduce0PfS_j
	.type	_Z7reduce0PfS_j, @function
_Z7reduce0PfS_j:
.LFB2054:
	.cfi_startproc
	endbr64
	jmp	_Z29__device_stub__Z7reduce0PfS_jPfS_j
	.cfi_endproc
.LFE2054:
	.size	_Z7reduce0PfS_j, .-_Z7reduce0PfS_j
	.globl	_Z29__device_stub__Z7reduce1PfS_jPfS_j
	.type	_Z29__device_stub__Z7reduce1PfS_jPfS_j, @function
_Z29__device_stub__Z7reduce1PfS_jPfS_j:
.LFB2055:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L8
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z7reduce1PfS_j(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	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 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L8:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2055:
	.size	_Z29__device_stub__Z7reduce1PfS_jPfS_j, .-_Z29__device_stub__Z7reduce1PfS_jPfS_j
	.globl	_Z7reduce1PfS_j
	.type	_Z7reduce1PfS_j, @function
_Z7reduce1PfS_j:
.LFB2056:
	.cfi_startproc
	endbr64
	jmp	_Z29__device_stub__Z7reduce1PfS_jPfS_j
	.cfi_endproc
.LFE2056:
	.size	_Z7reduce1PfS_j, .-_Z7reduce1PfS_j
	.globl	_Z29__device_stub__Z7reduce2PfS_jPfS_j
	.type	_Z29__device_stub__Z7reduce2PfS_jPfS_j, @function
_Z29__device_stub__Z7reduce2PfS_jPfS_j:
.LFB2057:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L13
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z7reduce2PfS_j(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	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 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L13:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L15
	call	__stack_chk_fail@PLT
.L15:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2057:
	.size	_Z29__device_stub__Z7reduce2PfS_jPfS_j, .-_Z29__device_stub__Z7reduce2PfS_jPfS_j
	.globl	_Z7reduce2PfS_j
	.type	_Z7reduce2PfS_j, @function
_Z7reduce2PfS_j:
.LFB2058:
	.cfi_startproc
	endbr64
	jmp	_Z29__device_stub__Z7reduce2PfS_jPfS_j
	.cfi_endproc
.LFE2058:
	.size	_Z7reduce2PfS_j, .-_Z7reduce2PfS_j
	.globl	_Z29__device_stub__Z7reduce3PfS_jPfS_j
	.type	_Z29__device_stub__Z7reduce3PfS_jPfS_j, @function
_Z29__device_stub__Z7reduce3PfS_jPfS_j:
.LFB2059:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L18
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z7reduce3PfS_j(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	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 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L18:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L20
	call	__stack_chk_fail@PLT
.L20:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2059:
	.size	_Z29__device_stub__Z7reduce3PfS_jPfS_j, .-_Z29__device_stub__Z7reduce3PfS_jPfS_j
	.globl	_Z7reduce3PfS_j
	.type	_Z7reduce3PfS_j, @function
_Z7reduce3PfS_j:
.LFB2060:
	.cfi_startproc
	endbr64
	jmp	_Z29__device_stub__Z7reduce3PfS_jPfS_j
	.cfi_endproc
.LFE2060:
	.size	_Z7reduce3PfS_j, .-_Z7reduce3PfS_j
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z7reduce3PfS_j"
.LC1:
	.string	"_Z7reduce2PfS_j"
.LC2:
	.string	"_Z7reduce1PfS_j"
.LC3:
	.string	"_Z7reduce0PfS_j"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2062:
	.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	_Z7reduce3PfS_j(%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	_Z7reduce2PfS_j(%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	_Z7reduce1PfS_j(%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	_Z7reduce0PfS_j(%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
.LFE2062:
	.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:

/*-----------
 *
 * matrixMulGlobal.cu
 *
 * This is the source file for matrix multiplication with global memory only.
 *
 * This kernel is from NVIDIA CUDA samples. reduction_kernel.cu.
 *
 * streamsOptBenchmark/reduction_kernel.cu
 *
 * By Hao Li
 *
 *------------
 */

/*
    Parallel reduction kernels
*/


#include <hip/hip_runtime.h>
#include <stdio.h>
// #include "structs.h"
// #include "functions.h"

// Utility class used to avoid linker errors with extern
// unsized shared memory arrays with templated type
struct SharedMemory
{
    __device__ inline operator       float *()
    {
        extern __shared__ float __smem[];
        return (float *)__smem;
    }

    __device__ inline operator const float *() const
    {
        extern __shared__ float __smem[];
        return (float *)__smem;
    }
};

// // specialize for double to avoid unaligned memory
// // access compile errors
// template<>
// struct SharedMemory<double>
// {
//     __device__ inline operator       double *()
//     {
//         extern __shared__ double __smem_d[];
//         return (double *)__smem_d;
//     }

//     __device__ inline operator const double *() const
//     {
//         extern __shared__ double __smem_d[];
//         return (double *)__smem_d;
//     }
// };

/*
    Parallel sum reduction using shared memory
    - takes log(n) steps for n input elements
    - uses n threads
    - only works for power-of-2 arrays
*/

/* This reduction interleaves which threads are active by using the modulo
   operator.  This operator is very expensive on GPUs, and the interleaved
   inactivity means that no whole warps are active, which is also very
   inefficient */
__global__ void reduce0(float *g_idata, float *g_odata, unsigned int n)
{
    for(int l = 0; l < 100000; l++)
    {
    float *sdata = SharedMemory();

    // load shared mem
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;

    sdata[tid] = (i < n) ? g_idata[i] : 0;

    __syncthreads();

    // do reduction in shared mem
    for (unsigned int s=1; s < blockDim.x; s *= 2)
    {
        // modulo arithmetic is slow!
        if ((tid % (2*s)) == 0)
        {
            sdata[tid] += sdata[tid + s];
        }

        __syncthreads();
    }

    // write result for this block to global mem
    if (tid == 0) g_odata[blockIdx.x] = sdata[0];
    }
}

/* This version uses contiguous threads, but its interleaved
   addressing results in many shared memory bank conflicts.
*/
__global__ void reduce1(float *g_idata, float *g_odata, unsigned int n)
{
    for(int l = 0; l < 100000; l++)
    {
    float *sdata = SharedMemory();

    // load shared mem
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;

    sdata[tid] = (i < n) ? g_idata[i] : 0;

    __syncthreads();

    // do reduction in shared mem
    for (unsigned int s=1; s < blockDim.x; s *= 2)
    {
        int index = 2 * s * tid;

        if (index < blockDim.x)
        {
            sdata[index] += sdata[index + s];
        }

        __syncthreads();
    }

    // write result for this block to global mem
    if (tid == 0) g_odata[blockIdx.x] = sdata[0];
    }
}

/*
    This version uses sequential addressing -- no divergence or bank conflicts.
*/
__global__ void reduce2(float *g_idata, float *g_odata, unsigned int n)
{
    for(int l = 0; l < 100000; l++)
    {
    float *sdata = SharedMemory();

    // load shared mem
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;

    sdata[tid] = (i < n) ? g_idata[i] : 0;

    __syncthreads();

    // do reduction in shared mem
    for (unsigned int s=blockDim.x/2; s>0; s>>=1)
    {
        if (tid < s)
        {
            sdata[tid] += sdata[tid + s];
        }

        __syncthreads();
    }

    // write result for this block to global mem
    if (tid == 0) g_odata[blockIdx.x] = sdata[0];
    }
}

/*
    This version uses n/2 threads --
    it performs the first level of reduction when reading from global memory.
*/
__global__ void reduce3(float *g_idata, float *g_odata, unsigned int n)
{
    for(int l = 0; l < 100000; l++)
    {
    float *sdata = SharedMemory();

    // perform first level of reduction,
    // reading from global memory, writing to shared memory
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;

    int mySum = (i < n) ? g_idata[i] : 0;

    if (i + blockDim.x < n)
        mySum += g_idata[i+blockDim.x];

    sdata[tid] = mySum;
    __syncthreads();

    // do reduction in shared mem
    for (unsigned int s=blockDim.x/2; s>0; s>>=1)
    {
        if (tid < s)
        {
            sdata[tid] = mySum = mySum + sdata[tid + s];
        }

        __syncthreads();
    }

    // write result for this block to global mem
    if (tid == 0) g_odata[blockIdx.x] = mySum;
    }
}

////////////////////////////////////////////////////////////////////////////////
// Wrapper function for kernel launch
////////////////////////////////////////////////////////////////////////////////
// void reduce(int size, int threads, int blocks,
//        int whichKernel, int *d_idata, int *d_odata)
// {
//     dim3 dimBlock(threads, 1, 1);
//     dim3 dimGrid(blocks, 1, 1);

//     // when there is only one warp per block, we need to allocate two warps
//     // worth of shared memory so that we don't index shared memory out of bounds
//     int smemSize = (threads <= 32) ? 2 * threads * sizeof(int) : threads * sizeof(int);

//     // choose which of the optimized versions of reduction to launch
//     switch (whichKernel)
//     {
//         case 0:
//             reduce0<<< dimGrid, dimBlock, smemSize >>>(d_idata, d_odata, size);
//             break;

//         case 1:
//             reduce1<<< dimGrid, dimBlock, smemSize >>>(d_idata, d_odata, size);
//             break;

//         case 2:
//             reduce2<<< dimGrid, dimBlock, smemSize >>>(d_idata, d_odata, size);
//             break;

//         case 3:
//             reduce3<<< dimGrid, dimBlock, smemSize >>>(d_idata, d_odata, size);
//             break;
//     }
// }

// int main(int argc, char **argv){
//  int matrixDataSize = sizeof(int) * MATRIX_SIZE * MATRIX_SIZE;

//  Matrix h_A, h_C;
//  Matrix d_A, d_C;

//  initMatrix(h_A, matrixDataSize, onHOST);
//  initMatrix(h_C, matrixDataSize, onHOST);
//  initMatrix(d_A, matrixDataSize, onDEVICE);
//  initMatrix(d_C, matrixDataSize, onDEVICE);

//  cudaMemcpy(d_A.elements, h_A.elements, matrixDataSize, cudaMemcpyHostToDevice);

//  // Invoke kernel
//  // dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
//  // dim3 dimGrid(h_B.width / dimBlock.x, h_A.height / dimBlock.y);

//  // execute the kernel
//  for(int i =0; i < 4; i++){
//  reduce(matrixDataSize, h_A.width / BLOCK_SIZE * h_A.height / BLOCK_SIZE, 
//     BLOCK_SIZE*BLOCK_SIZE, i, d_A.elements, d_C.elements);
// }

//  cudaMemcpy(h_C.elements, d_C.elements, matrixDataSize, cudaMemcpyDeviceToHost);

//  free(h_A.elements);
//  free(h_C.elements);
//  cudaFree(d_A.elements);
//  cudaFree(d_C.elements);

//  return 0;
// }

	.text
	.file	"reduction_kernel.hip"
	.globl	_Z22__device_stub__reduce0PfS_j # -- Begin function _Z22__device_stub__reduce0PfS_j
	.type	_Z22__device_stub__reduce0PfS_j,@function
_Z22__device_stub__reduce0PfS_j:        # @_Z22__device_stub__reduce0PfS_j
	.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	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%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	$_Z7reduce0PfS_j, %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_end0:
	.size	_Z22__device_stub__reduce0PfS_j, .Lfunc_end0-_Z22__device_stub__reduce0PfS_j
	.cfi_endproc
                                        # -- End function
	.globl	_Z22__device_stub__reduce1PfS_j # -- Begin function _Z22__device_stub__reduce1PfS_j
	.type	_Z22__device_stub__reduce1PfS_j,@function
_Z22__device_stub__reduce1PfS_j:        # @_Z22__device_stub__reduce1PfS_j
	.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	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%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	$_Z7reduce1PfS_j, %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_end1:
	.size	_Z22__device_stub__reduce1PfS_j, .Lfunc_end1-_Z22__device_stub__reduce1PfS_j
	.cfi_endproc
                                        # -- End function
	.globl	_Z22__device_stub__reduce2PfS_j # -- Begin function _Z22__device_stub__reduce2PfS_j
	.type	_Z22__device_stub__reduce2PfS_j,@function
_Z22__device_stub__reduce2PfS_j:        # @_Z22__device_stub__reduce2PfS_j
	.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	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%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	$_Z7reduce2PfS_j, %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	_Z22__device_stub__reduce2PfS_j, .Lfunc_end2-_Z22__device_stub__reduce2PfS_j
	.cfi_endproc
                                        # -- End function
	.globl	_Z22__device_stub__reduce3PfS_j # -- Begin function _Z22__device_stub__reduce3PfS_j
	.type	_Z22__device_stub__reduce3PfS_j,@function
_Z22__device_stub__reduce3PfS_j:        # @_Z22__device_stub__reduce3PfS_j
	.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	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%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	$_Z7reduce3PfS_j, %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_end3:
	.size	_Z22__device_stub__reduce3PfS_j, .Lfunc_end3-_Z22__device_stub__reduce3PfS_j
	.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	.LBB4_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB4_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z7reduce0PfS_j, %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	$_Z7reduce1PfS_j, %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	$_Z7reduce2PfS_j, %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	$_Z7reduce3PfS_j, %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_end4:
	.size	__hip_module_ctor, .Lfunc_end4-__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	.LBB5_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
.LBB5_2:
	retq
.Lfunc_end5:
	.size	__hip_module_dtor, .Lfunc_end5-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z7reduce0PfS_j,@object         # @_Z7reduce0PfS_j
	.section	.rodata,"a",@progbits
	.globl	_Z7reduce0PfS_j
	.p2align	3, 0x0
_Z7reduce0PfS_j:
	.quad	_Z22__device_stub__reduce0PfS_j
	.size	_Z7reduce0PfS_j, 8

	.type	_Z7reduce1PfS_j,@object         # @_Z7reduce1PfS_j
	.globl	_Z7reduce1PfS_j
	.p2align	3, 0x0
_Z7reduce1PfS_j:
	.quad	_Z22__device_stub__reduce1PfS_j
	.size	_Z7reduce1PfS_j, 8

	.type	_Z7reduce2PfS_j,@object         # @_Z7reduce2PfS_j
	.globl	_Z7reduce2PfS_j
	.p2align	3, 0x0
_Z7reduce2PfS_j:
	.quad	_Z22__device_stub__reduce2PfS_j
	.size	_Z7reduce2PfS_j, 8

	.type	_Z7reduce3PfS_j,@object         # @_Z7reduce3PfS_j
	.globl	_Z7reduce3PfS_j
	.p2align	3, 0x0
_Z7reduce3PfS_j:
	.quad	_Z22__device_stub__reduce3PfS_j
	.size	_Z7reduce3PfS_j, 8

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

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z7reduce1PfS_j"
	.size	.L__unnamed_2, 16

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"_Z7reduce2PfS_j"
	.size	.L__unnamed_3, 16

	.type	.L__unnamed_4,@object           # @3
.L__unnamed_4:
	.asciz	"_Z7reduce3PfS_j"
	.size	.L__unnamed_4, 16

	.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 _Z22__device_stub__reduce0PfS_j
	.addrsig_sym _Z22__device_stub__reduce1PfS_j
	.addrsig_sym _Z22__device_stub__reduce2PfS_j
	.addrsig_sym _Z22__device_stub__reduce3PfS_j
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z7reduce0PfS_j
	.addrsig_sym _Z7reduce1PfS_j
	.addrsig_sym _Z7reduce2PfS_j
	.addrsig_sym _Z7reduce3PfS_j
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z7reduce0PfS_j         ; -- Begin function _Z7reduce0PfS_j
	.globl	_Z7reduce0PfS_j
	.p2align	8
	.type	_Z7reduce0PfS_j,@function
_Z7reduce0PfS_j:                        ; @_Z7reduce0PfS_j
; %bb.0:
	s_clause 0x2
	s_load_b32 s3, s[0:1], 0x24
	s_load_b32 s12, s[0:1], 0x10
	s_load_b128 s[8:11], s[0:1], 0x0
	s_mov_b32 s2, s15
	v_lshl_add_u32 v5, v0, 2, 0
	v_cmp_eq_u32_e64 s0, 0, v0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s6, s3, 0xffff
	s_mov_b32 s3, 0
	v_mad_u64_u32 v[1:2], null, s2, s6, v[0:1]
	v_mov_b32_e32 v2, 0
	s_cmp_gt_u32 s6, 1
	s_cselect_b32 s7, -1, 0
	s_lshl_b64 s[4:5], s[2:3], 2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	s_add_u32 s4, s10, s4
	v_lshlrev_b64 v[3:4], 2, v[1:2]
	v_cmp_gt_u32_e64 s1, s12, v1
	s_addc_u32 s5, s11, s5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v3, vcc_lo, s8, v3
	v_add_co_ci_u32_e32 v4, vcc_lo, s9, v4, vcc_lo
.LBB0_1:                                ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_5 Depth 2
	v_mov_b32_e32 v1, 0
	s_and_saveexec_b32 s2, s1
	s_cbranch_execz .LBB0_3
; %bb.2:                                ;   in Loop: Header=BB0_1 Depth=1
	global_load_b32 v1, v[3:4], off
.LBB0_3:                                ;   in Loop: Header=BB0_1 Depth=1
	s_or_b32 exec_lo, exec_lo, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s7
	s_waitcnt vmcnt(0)
	ds_store_b32 v5, v1
	s_waitcnt lgkmcnt(0)
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	s_cbranch_vccnz .LBB0_8
; %bb.4:                                ; %.lr.ph.preheader
                                        ;   in Loop: Header=BB0_1 Depth=1
	s_mov_b32 s8, 1
.LBB0_5:                                ; %.lr.ph
                                        ;   Parent Loop BB0_1 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_lshl_b32 s2, s8, 1
	s_add_i32 s9, s2, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_and_b32_e32 v1, s9, v0
	s_mov_b32 s9, exec_lo
	v_cmpx_eq_u32_e32 0, v1
	s_cbranch_execz .LBB0_7
; %bb.6:                                ;   in Loop: Header=BB0_5 Depth=2
	v_lshl_add_u32 v1, s8, 2, v5
	ds_load_b32 v1, v1
	ds_load_b32 v6, v5
	s_waitcnt lgkmcnt(0)
	v_add_f32_e32 v1, v1, v6
	ds_store_b32 v5, v1
.LBB0_7:                                ;   in Loop: Header=BB0_5 Depth=2
	s_or_b32 exec_lo, exec_lo, s9
	s_cmp_ge_u32 s2, s6
	s_mov_b32 s8, s2
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_cbranch_scc0 .LBB0_5
.LBB0_8:                                ; %._crit_edge
                                        ;   in Loop: Header=BB0_1 Depth=1
	s_and_saveexec_b32 s2, s0
	s_cbranch_execz .LBB0_10
; %bb.9:                                ;   in Loop: Header=BB0_1 Depth=1
	ds_load_b32 v1, v2
	s_waitcnt lgkmcnt(0)
	global_store_b32 v2, v1, s[4:5]
.LBB0_10:                               ;   in Loop: Header=BB0_1 Depth=1
	s_or_b32 exec_lo, exec_lo, s2
	s_add_i32 s3, s3, 1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_eq_u32 s3, 0x186a0
	s_cbranch_scc0 .LBB0_1
; %bb.11:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z7reduce0PfS_j
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 280
		.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_end0:
	.size	_Z7reduce0PfS_j, .Lfunc_end0-_Z7reduce0PfS_j
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 372
; 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 : 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	_Z7reduce1PfS_j         ; -- Begin function _Z7reduce1PfS_j
	.globl	_Z7reduce1PfS_j
	.p2align	8
	.type	_Z7reduce1PfS_j,@function
_Z7reduce1PfS_j:                        ; @_Z7reduce1PfS_j
; %bb.0:
	s_clause 0x2
	s_load_b32 s3, s[0:1], 0x24
	s_load_b32 s12, s[0:1], 0x10
	s_load_b128 s[8:11], s[0:1], 0x0
	s_mov_b32 s2, s15
	v_lshl_add_u32 v5, v0, 2, 0
	v_cmp_eq_u32_e64 s0, 0, v0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s6, s3, 0xffff
	s_mov_b32 s3, 0
	v_mad_u64_u32 v[1:2], null, s2, s6, v[0:1]
	v_mov_b32_e32 v2, 0
	s_cmp_gt_u32 s6, 1
	s_cselect_b32 s7, -1, 0
	s_lshl_b64 s[4:5], s[2:3], 2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	s_add_u32 s4, s10, s4
	v_lshlrev_b64 v[3:4], 2, v[1:2]
	v_cmp_gt_u32_e64 s1, s12, v1
	s_addc_u32 s5, s11, s5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v3, vcc_lo, s8, v3
	v_add_co_ci_u32_e32 v4, vcc_lo, s9, v4, vcc_lo
.LBB1_1:                                ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB1_5 Depth 2
	v_mov_b32_e32 v1, 0
	s_and_saveexec_b32 s2, s1
	s_cbranch_execz .LBB1_3
; %bb.2:                                ;   in Loop: Header=BB1_1 Depth=1
	global_load_b32 v1, v[3:4], off
.LBB1_3:                                ;   in Loop: Header=BB1_1 Depth=1
	s_or_b32 exec_lo, exec_lo, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s7
	s_waitcnt vmcnt(0)
	ds_store_b32 v5, v1
	s_waitcnt lgkmcnt(0)
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	s_cbranch_vccnz .LBB1_8
; %bb.4:                                ; %.lr.ph.preheader
                                        ;   in Loop: Header=BB1_1 Depth=1
	s_mov_b32 s8, 1
.LBB1_5:                                ; %.lr.ph
                                        ;   Parent Loop BB1_1 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	s_lshl_b32 s2, s8, 1
	s_mov_b32 s9, exec_lo
	v_mul_lo_u32 v1, s2, v0
	v_cmpx_gt_u32_e64 s6, v1
	s_cbranch_execz .LBB1_7
; %bb.6:                                ;   in Loop: Header=BB1_5 Depth=2
	v_lshl_add_u32 v1, v1, 2, 0
	s_delay_alu instid0(VALU_DEP_1)
	v_lshl_add_u32 v6, s8, 2, v1
	ds_load_b32 v6, v6
	ds_load_b32 v7, v1
	s_waitcnt lgkmcnt(0)
	v_add_f32_e32 v6, v6, v7
	ds_store_b32 v1, v6
.LBB1_7:                                ;   in Loop: Header=BB1_5 Depth=2
	s_or_b32 exec_lo, exec_lo, s9
	s_cmp_ge_u32 s2, s6
	s_mov_b32 s8, s2
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_cbranch_scc0 .LBB1_5
.LBB1_8:                                ; %._crit_edge
                                        ;   in Loop: Header=BB1_1 Depth=1
	s_and_saveexec_b32 s2, s0
	s_cbranch_execz .LBB1_10
; %bb.9:                                ;   in Loop: Header=BB1_1 Depth=1
	ds_load_b32 v1, v2
	s_waitcnt lgkmcnt(0)
	global_store_b32 v2, v1, s[4:5]
.LBB1_10:                               ;   in Loop: Header=BB1_1 Depth=1
	s_or_b32 exec_lo, exec_lo, s2
	s_add_i32 s3, s3, 1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_eq_u32 s3, 0x186a0
	s_cbranch_scc0 .LBB1_1
; %bb.11:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z7reduce1PfS_j
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 280
		.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 8
		.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	_Z7reduce1PfS_j, .Lfunc_end1-_Z7reduce1PfS_j
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 384
; NumSgprs: 18
; NumVgprs: 8
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 8
; 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	_Z7reduce2PfS_j         ; -- Begin function _Z7reduce2PfS_j
	.globl	_Z7reduce2PfS_j
	.p2align	8
	.type	_Z7reduce2PfS_j,@function
_Z7reduce2PfS_j:                        ; @_Z7reduce2PfS_j
; %bb.0:
	s_clause 0x2
	s_load_b32 s3, s[0:1], 0x24
	s_load_b32 s12, s[0:1], 0x10
	s_load_b128 s[8:11], s[0:1], 0x0
	s_mov_b32 s2, s15
	v_lshl_add_u32 v5, v0, 2, 0
	v_cmp_eq_u32_e64 s0, 0, v0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s6, s3, 0xffff
	s_mov_b32 s3, 0
	v_mad_u64_u32 v[1:2], null, s2, s6, v[0:1]
	v_mov_b32_e32 v2, 0
	s_cmp_gt_u32 s6, 1
	s_cselect_b32 s7, -1, 0
	s_lshl_b64 s[4:5], s[2:3], 2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	s_add_u32 s4, s10, s4
	v_lshlrev_b64 v[3:4], 2, v[1:2]
	v_cmp_gt_u32_e64 s1, s12, v1
	s_addc_u32 s5, s11, s5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v3, vcc_lo, s8, v3
	v_add_co_ci_u32_e32 v4, vcc_lo, s9, v4, vcc_lo
.LBB2_1:                                ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB2_4 Depth 2
	v_mov_b32_e32 v1, 0
	s_and_saveexec_b32 s2, s1
	s_cbranch_execz .LBB2_3
; %bb.2:                                ;   in Loop: Header=BB2_1 Depth=1
	global_load_b32 v1, v[3:4], off
.LBB2_3:                                ;   in Loop: Header=BB2_1 Depth=1
	s_or_b32 exec_lo, exec_lo, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s7
	s_mov_b32 s2, s6
	s_waitcnt vmcnt(0)
	ds_store_b32 v5, v1
	s_waitcnt lgkmcnt(0)
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	s_cbranch_vccnz .LBB2_7
.LBB2_4:                                ; %.lr.ph
                                        ;   Parent Loop BB2_1 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_lshr_b32 s8, s2, 1
	s_mov_b32 s9, exec_lo
	v_cmpx_gt_u32_e64 s8, v0
	s_cbranch_execz .LBB2_6
; %bb.5:                                ;   in Loop: Header=BB2_4 Depth=2
	v_lshl_add_u32 v1, s8, 2, v5
	ds_load_b32 v1, v1
	ds_load_b32 v6, v5
	s_waitcnt lgkmcnt(0)
	v_add_f32_e32 v1, v1, v6
	ds_store_b32 v5, v1
.LBB2_6:                                ;   in Loop: Header=BB2_4 Depth=2
	s_or_b32 exec_lo, exec_lo, s9
	s_cmp_lt_u32 s2, 4
	s_mov_b32 s2, s8
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_cbranch_scc0 .LBB2_4
.LBB2_7:                                ; %._crit_edge
                                        ;   in Loop: Header=BB2_1 Depth=1
	s_and_saveexec_b32 s2, s0
	s_cbranch_execz .LBB2_9
; %bb.8:                                ;   in Loop: Header=BB2_1 Depth=1
	ds_load_b32 v1, v2
	s_waitcnt lgkmcnt(0)
	global_store_b32 v2, v1, s[4:5]
.LBB2_9:                                ;   in Loop: Header=BB2_1 Depth=1
	s_or_b32 exec_lo, exec_lo, s2
	s_add_i32 s3, s3, 1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_eq_u32 s3, 0x186a0
	s_cbranch_scc0 .LBB2_1
; %bb.10:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z7reduce2PfS_j
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 280
		.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	_Z7reduce2PfS_j, .Lfunc_end2-_Z7reduce2PfS_j
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 360
; 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 : 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	_Z7reduce3PfS_j         ; -- Begin function _Z7reduce3PfS_j
	.globl	_Z7reduce3PfS_j
	.p2align	8
	.type	_Z7reduce3PfS_j,@function
_Z7reduce3PfS_j:                        ; @_Z7reduce3PfS_j
; %bb.0:
	s_clause 0x2
	s_load_b32 s3, s[0:1], 0x24
	s_load_b32 s12, s[0:1], 0x10
	s_load_b128 s[4:7], s[0:1], 0x0
	s_mov_b32 s2, s15
	v_lshl_add_u32 v7, v0, 2, 0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s8, s3, 0xffff
	s_mov_b32 s3, 0
	s_mul_i32 s0, s15, s8
	s_cmp_gt_u32 s8, 1
	v_lshl_add_u32 v1, s0, 1, v0
	v_mov_b32_e32 v2, 0
	s_cselect_b32 s9, -1, 0
	s_lshl_b64 s[10:11], s[2:3], 2
	v_cmp_eq_u32_e64 s0, 0, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_dual_mov_b32 v4, v2 :: v_dual_add_nc_u32 v3, s8, v1
	v_lshlrev_b64 v[5:6], 2, v[1:2]
	v_cmp_gt_u32_e64 s1, s12, v1
	v_cmp_gt_u32_e64 s2, s12, v3
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_lshlrev_b64 v[8:9], 2, v[3:4]
	v_add_co_u32 v3, vcc_lo, s4, v5
	v_add_co_ci_u32_e32 v4, vcc_lo, s5, v6, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v5, vcc_lo, s4, v8
	v_add_co_ci_u32_e32 v6, vcc_lo, s5, v9, vcc_lo
	s_add_u32 s4, s6, s10
	s_addc_u32 s5, s7, s11
.LBB3_1:                                ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB3_6 Depth 2
	v_mov_b32_e32 v1, 0
	s_and_saveexec_b32 s6, s1
	s_cbranch_execz .LBB3_3
; %bb.2:                                ;   in Loop: Header=BB3_1 Depth=1
	global_load_b32 v1, v[3:4], off
.LBB3_3:                                ;   in Loop: Header=BB3_1 Depth=1
	s_or_b32 exec_lo, exec_lo, s6
	s_waitcnt vmcnt(0)
	v_cvt_i32_f32_e32 v1, v1
	s_and_saveexec_b32 s6, s2
	s_cbranch_execz .LBB3_5
; %bb.4:                                ;   in Loop: Header=BB3_1 Depth=1
	global_load_b32 v8, v[5:6], off
	v_cvt_f32_i32_e32 v1, v1
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v1, v8, v1
	v_cvt_i32_f32_e32 v1, v1
.LBB3_5:                                ;   in Loop: Header=BB3_1 Depth=1
	s_or_b32 exec_lo, exec_lo, s6
	s_delay_alu instid0(VALU_DEP_1)
	v_cvt_f32_i32_e32 v8, v1
	s_and_not1_b32 vcc_lo, exec_lo, s9
	s_mov_b32 s6, s8
	ds_store_b32 v7, v8
	s_waitcnt lgkmcnt(0)
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	s_cbranch_vccnz .LBB3_10
.LBB3_6:                                ; %.lr.ph
                                        ;   Parent Loop BB3_1 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_lshr_b32 s7, s6, 1
	s_mov_b32 s10, exec_lo
	v_cmpx_gt_u32_e64 s7, v0
	s_cbranch_execz .LBB3_8
; %bb.7:                                ;   in Loop: Header=BB3_6 Depth=2
	v_lshl_add_u32 v8, s7, 2, v7
	v_cvt_f32_i32_e32 v1, v1
	ds_load_b32 v8, v8
	s_waitcnt lgkmcnt(0)
	v_add_f32_e32 v1, v8, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_i32_f32_e32 v1, v1
	v_cvt_f32_i32_e32 v8, v1
	ds_store_b32 v7, v8
.LBB3_8:                                ;   in Loop: Header=BB3_6 Depth=2
	s_or_b32 exec_lo, exec_lo, s10
	s_cmp_lt_u32 s6, 4
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_cbranch_scc1 .LBB3_10
; %bb.9:                                ;   in Loop: Header=BB3_6 Depth=2
	s_mov_b32 s6, s7
	s_branch .LBB3_6
.LBB3_10:                               ; %Flow44
                                        ;   in Loop: Header=BB3_1 Depth=1
	s_and_saveexec_b32 s6, s0
	s_cbranch_execz .LBB3_12
; %bb.11:                               ;   in Loop: Header=BB3_1 Depth=1
	v_cvt_f32_i32_e32 v1, v1
	global_store_b32 v2, v1, s[4:5]
.LBB3_12:                               ;   in Loop: Header=BB3_1 Depth=1
	s_or_b32 exec_lo, exec_lo, s6
	s_add_i32 s3, s3, 1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_eq_u32 s3, 0x186a0
	s_cbranch_scc0 .LBB3_1
; %bb.13:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z7reduce3PfS_j
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 280
		.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 10
		.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	_Z7reduce3PfS_j, .Lfunc_end3-_Z7reduce3PfS_j
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 456
; NumSgprs: 18
; NumVgprs: 10
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 10
; 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:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         28
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         36
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         38
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         40
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         42
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         44
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         46
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         88
        .size:           2
        .value_kind:     hidden_grid_dims
      - .offset:         144
        .size:           4
        .value_kind:     hidden_dynamic_lds_size
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 280
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z7reduce0PfS_j
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z7reduce0PfS_j.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:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         28
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         36
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         38
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         40
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         42
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         44
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         46
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         88
        .size:           2
        .value_kind:     hidden_grid_dims
      - .offset:         144
        .size:           4
        .value_kind:     hidden_dynamic_lds_size
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 280
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z7reduce1PfS_j
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z7reduce1PfS_j.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     8
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         28
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         36
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         38
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         40
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         42
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         44
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         46
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         88
        .size:           2
        .value_kind:     hidden_grid_dims
      - .offset:         144
        .size:           4
        .value_kind:     hidden_dynamic_lds_size
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 280
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z7reduce2PfS_j
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z7reduce2PfS_j.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:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         28
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         36
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         38
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         40
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         42
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         44
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         46
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         88
        .size:           2
        .value_kind:     hidden_grid_dims
      - .offset:         144
        .size:           4
        .value_kind:     hidden_dynamic_lds_size
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 280
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z7reduce3PfS_j
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z7reduce3PfS_j.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     10
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
