/* Howto run ?
If executable name is 'reduce' then type the following in command terminal
./reduce 0 for global reduce
./reduce for shared memory reduce
 */

/* For global access, reduce run time - 0.911 ms approx
For shared memory access, reduce run time - 0.771 ms approx
*/


#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
/*
cuda.h -  defines the public host functions and types for the CUDA driver API.
cuda_runtime_api.h - defines the public host functions and types for the CUDA
runtime API
cuda_runtime.h - defines everything cuda_runtime_api.h does, as well as built-in
type definitions and function overlays for the CUDA language extensions and
device intrinsic functions.*/

/*All info about the .h files is from stach overflow*/

/*If you were writing host code to be compiled with the host compiler which
 * includes API calls, you would include either cuda.h or cuda_runtime_api.h.
 * If you needed other CUDA language built-ins, like types, and were using the
 * runtime API and compiling with the host compiler, you would include
 * cuda_runtime.h
*/

__global__ void global_reduce_kernel(float * d_out, float * d_in){
   int myId = threadIdx.x + blockDim.x * blockIdx.x;
    int tid  = threadIdx.x;

    // do reduction in global mem
    for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1)
    {
        if (tid < s)
        {
            d_in[myId] += d_in[myId + s];
        }
        __syncthreads();        // make sure all adds at one stage are
//done!
    }

    // only thread 0 writes result for this block back to global mem
    if (tid == 0)
    {
        d_out[blockIdx.x] = d_in[myId];
    }
}

__global__ void shmem_reduce_kernel(float * d_out, const float * d_in)
{
    // sdata is allocated in the kernel call: 3rd arg to <<<b, t,
    // shmem>>>
    extern __shared__ float sdata[];

    int myId = threadIdx.x + blockDim.x * blockIdx.x;
    int tid  = threadIdx.x;

    // load shared mem from global mem
    sdata[tid] = d_in[myId];
    __syncthreads();            // make sure entire block is loaded!

    // 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();        // make sure all adds at one stage are
//done!
    }

    // only thread 0 writes result for this block back to global mem
    if (tid == 0)
    {
        d_out[blockIdx.x] = sdata[0];
    }
}

void reduce(float * d_out, float * d_intermediate, float * d_in, int size, bool usesSharedMemory)
{
    // assumes that size is not greater than maxThreadsPerBlock^2
    // and that size is a multiple of maxThreadsPerBlock
    const int maxThreadsPerBlock = 1024;
    int threads = maxThreadsPerBlock;
    int blocks = size / maxThreadsPerBlock;
    if (usesSharedMemory)
    {
        shmem_reduce_kernel<<<blocks, threads, threads *sizeof(float)>>>(d_intermediate, d_in);
    }
    else
    {
        global_reduce_kernel<<<blocks, threads>>>
            (d_intermediate, d_in);
    }
    // now we're down to one block left, so reduce it
    threads = blocks; // launch one thread for each block in prev step
    blocks = 1;
    if (usesSharedMemory)
    {
      //  shmem_reduce_kernel<<<blocks, threads, threads *sizeof(float)>>>(d_out, d_intermediate);
    	global_reduce_kernel<<<blocks, threads>>>
            (d_out, d_intermediate);
    }
    else
    {
        global_reduce_kernel<<<blocks, threads>>>
            (d_out, d_intermediate);
    }
}

int main(int argc, char **argv)
{
    int deviceCount;
    cudaGetDeviceCount(&deviceCount);
/*cudaGetDeviceCount ( int* count )

PARAMETERS
count

RETURNS
cudaSuccess, cudaErrorNoDevice, cudaErrorInsufficientDriver

DESCRIPTION
Returns in *count the number of devices with compute capability greater or equal to 2.0 that are available for execution. If there is no such device then
cudaGetDeviceCount() will return cudaErrorNoDevice. If no driver can be loaded to determine if any such devices exist then cudaGetDeviceCount() will return
cudaErrorInsufficientDriver.*/


    if (deviceCount == 0) {
        fprintf(stderr, "error: no devices supporting CUDA.\n");
        exit(EXIT_FAILURE);
    }
    int dev = 0;
    cudaSetDevice(dev);
/*cudaError_t cudaSetDevice ( int  device )
PARAMETERS
device
- Device on which the active host thread should execute the device code.

RETURNS
cudaSuccess, cudaErrorInvalidDevice, cudaErrorDeviceAlreadyInUse

DESCRIPTION
Sets device as the current device for the calling host thread. Valid device id's are 0 to (cudaGetDeviceCount() - 1).

Any device memory subsequently allocated from this host thread using cudaMalloc(), cudaMallocPitch() or cudaMallocArray() will be physically resident on
device. Any host memory allocated from this host thread using cudaMallocHost() or cudaHostAlloc() or cudaHostRegister() will have its lifetime associated with
device. Any streams or events created from this host thread will be associated with device. Any kernels launched from this host thread using the <<<>>>
operator or cudaLaunchKernel() will be executed on device.

This call may be made from any host thread, to any device, and at any time. This function will do no synchronization with the previous or new device, and
should be considered a very low overhead call.
*/

    cudaDeviceProp devProps;
    if (cudaGetDeviceProperties(&devProps, dev) == 0)
    {
        printf("Using device %d:\n", dev);
        printf("%s; global mem: %dB; compute v%d.%d; clock: %d kHz\n",
               devProps.name, (int)devProps.totalGlobalMem, 
               (int)devProps.major, (int)devProps.minor, 
               (int)devProps.clockRate);
    }

    const int ARRAY_SIZE = 1 << 20;
    const int ARRAY_BYTES = ARRAY_SIZE * sizeof(float);

    // generate the input array on the host
    float h_in[ARRAY_SIZE];
    float sum = 0.0f;
   /* for(int i = 0; i < ARRAY_SIZE; i++) {
        // generate random float in [-1.0f, 1.0f]
        h_in[i] = -1.0f + (float)random()/((float)RAND_MAX/2.0f);
        sum += h_in[i];
    }*/

//Additional part for input array on host - used for teesting correctness added by - Kshitij
	for(int i = 0; i < ARRAY_SIZE; i++) {
        // generate non random float
        h_in[i] = 1.0f;
        sum += h_in[i];
    }



    // declare GPU memory pointers
    float * d_in, * d_intermediate, * d_out;

    // allocate GPU memory
    cudaMalloc((void **) &d_in, ARRAY_BYTES);
    cudaMalloc((void **) &d_intermediate, ARRAY_BYTES); //overallocated
    cudaMalloc((void **) &d_out, sizeof(float));

    // transfer the input array to the GPU
    cudaMemcpy(d_in, h_in, ARRAY_BYTES, cudaMemcpyHostToDevice); 

    int whichKernel = 0;
    if (argc == 2) {
        whichKernel = atoi(argv[1]);
	//whichKernel = argv[1]; -  incorrect. Explanation below.    
}

/*atoi example in this directory only
Run and see for understanding it 
Hence, when running reduce.cu we use ./reduce 0 or ./reduce 1
*/  
/*Why use the atoi ?
    When we receive argumnets from command line and stored in argv, they are stored as strings. Basically, argv is an array of strings.
	Now, while checking for shared memory or global memory usage, we use switch vase statements with number (int) - whichKernel.
	So,whichKernel = argv[1] won't work here as RHS is a string and LHS is an int.
	We need whichKernel to be int as used in the switch case statements.
	Hence, convert the received number - 0 or 1 which is stored as a string to an int and then use it for comparison ini switch case.
*/
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    // launch the kernel
    switch(whichKernel) {
    case 0:
        printf("Running global reduce\n");
        cudaEventRecord(start, 0);
        /*for (int i = 0; i < 100; i++)
        {
           reduce(d_out, d_intermediate, d_in, ARRAY_SIZE, false);
        }*/

	 //Addition by Kshitij - single call instead of 100 calls of the kernel
        reduce(d_out, d_intermediate,d_in,ARRAY_SIZE,false);
	cudaEventRecord(stop, 0);
        
	
	break;
    case 1:
        printf("Running reduce with shared mem\n");
        cudaEventRecord(start, 0);
        /*for (int i = 0; i < 100; i++)
        {
            reduce(d_out, d_intermediate, d_in, ARRAY_SIZE, true);
        }*/
	reduce(d_out, d_intermediate, d_in, ARRAY_SIZE, true);
        cudaEventRecord(stop, 0);
        break;
    default:
        fprintf(stderr, "error: ran no kernel\n");
        exit(EXIT_FAILURE);
    }
    cudaEventSynchronize(stop);
    float elapsedTime;
    cudaEventElapsedTime(&elapsedTime, start, stop);    
    elapsedTime /= 100.0f;      // 100 trials

    // copy back the sum from GPU
    float h_out;
    cudaMemcpy(&h_out, d_out, sizeof(float), cudaMemcpyDeviceToHost);

    printf("average time elapsed: %f\n", elapsedTime);
	
//Additional part - Kshitij
printf("Reduce answer :%f",h_out);
printf("Sum answer on host :%f",sum);



    // free GPU memory allocation
    cudaFree(d_in);
    cudaFree(d_intermediate);
    cudaFree(d_out);
        
    return 0;
}

	.file	"tmpxft_0031bd60_00000000-6_reduce.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	_Z42__device_stub__Z20global_reduce_kernelPfS_PfS_
	.type	_Z42__device_stub__Z20global_reduce_kernelPfS_PfS_, @function
_Z42__device_stub__Z20global_reduce_kernelPfS_PfS_:
.LFB2053:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z20global_reduce_kernelPfS_(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L2:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2053:
	.size	_Z42__device_stub__Z20global_reduce_kernelPfS_PfS_, .-_Z42__device_stub__Z20global_reduce_kernelPfS_PfS_
	.globl	_Z20global_reduce_kernelPfS_
	.type	_Z20global_reduce_kernelPfS_, @function
_Z20global_reduce_kernelPfS_:
.LFB2054:
	.cfi_startproc
	endbr64
	jmp	_Z42__device_stub__Z20global_reduce_kernelPfS_PfS_
	.cfi_endproc
.LFE2054:
	.size	_Z20global_reduce_kernelPfS_, .-_Z20global_reduce_kernelPfS_
	.globl	_Z42__device_stub__Z19shmem_reduce_kernelPfPKfPfPKf
	.type	_Z42__device_stub__Z19shmem_reduce_kernelPfPKfPfPKf, @function
_Z42__device_stub__Z19shmem_reduce_kernelPfPKfPfPKf:
.LFB2055:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L8
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z19shmem_reduce_kernelPfPKf(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L8:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2055:
	.size	_Z42__device_stub__Z19shmem_reduce_kernelPfPKfPfPKf, .-_Z42__device_stub__Z19shmem_reduce_kernelPfPKfPfPKf
	.globl	_Z19shmem_reduce_kernelPfPKf
	.type	_Z19shmem_reduce_kernelPfPKf, @function
_Z19shmem_reduce_kernelPfPKf:
.LFB2056:
	.cfi_startproc
	endbr64
	jmp	_Z42__device_stub__Z19shmem_reduce_kernelPfPKfPfPKf
	.cfi_endproc
.LFE2056:
	.size	_Z19shmem_reduce_kernelPfPKf, .-_Z19shmem_reduce_kernelPfPKf
	.globl	_Z6reducePfS_S_ib
	.type	_Z6reducePfS_S_ib, @function
_Z6reducePfS_S_ib:
.LFB2027:
	.cfi_startproc
	endbr64
	movl	%ecx, %eax
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movq	%rdx, %r13
	movl	$0, %r9d
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	cltd
	movq	%rdi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movq	%rsi, %rbp
	movl	$1024, %esi
	idivl	%esi
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movabsq	$4294968320, %rdx
	subq	$40, %rsp
	.cfi_def_cfa_offset 80
	testb	%r8b, %r8b
	movq	%rdx, 20(%rsp)
	movl	$1, 28(%rsp)
	movl	%eax, 8(%rsp)
	movl	%eax, %ebx
	movabsq	$4294967297, %rax
	movq	%rax, 12(%rsp)
	je	.L14
	movq	8(%rsp), %rdi
	movl	16(%rsp), %esi
	movl	$4096, %r8d
	movl	$1, %ecx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L18
	movq	%r13, %rsi
	movq	%rbp, %rdi
	call	_Z42__device_stub__Z19shmem_reduce_kernelPfPKfPfPKf
	jmp	.L18
.L14:
	movq	8(%rsp), %rdi
	movl	16(%rsp), %esi
	xorl	%r8d, %r8d
	movl	$1, %ecx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L18
	movq	%r13, %rsi
	movq	%rbp, %rdi
	call	_Z42__device_stub__Z20global_reduce_kernelPfS_PfS_
.L18:
	movl	%ebx, 20(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %esi
	movabsq	$4294967297, %rdi
	movl	$1, 16(%rsp)
	movq	%rdi, 24(%rsp)
	movl	28(%rsp), %ecx
	movq	20(%rsp), %rdx
	movq	%rdi, 8(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L13
	addq	$40, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 40
	movq	%rbp, %rsi
	movq	%r12, %rdi
	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	_Z42__device_stub__Z20global_reduce_kernelPfS_PfS_
.L13:
	.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
.LFE2027:
	.size	_Z6reducePfS_S_ib, .-_Z6reducePfS_S_ib
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"error: no devices supporting CUDA.\n"
.LC2:
	.string	"Using device %d:\n"
.LC3:
	.string	"%s; global mem: %dB; compute v%d.%d; clock: %d kHz\n"
.LC5:
	.string	"Running global reduce\n"
.LC6:
	.string	"Running reduce with shared mem\n"
.LC7:
	.string	"error: ran no kernel\n"
.LC9:
	.string	"average time elapsed: %f\n"
.LC10:
	.string	"Reduce answer :%f"
.LC11:
	.string	"Sum answer on host :%f"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	leaq	-4194304(%rsp), %r11
	.cfi_def_cfa 11, 4194344
.LPSRL0:
	subq	$4096, %rsp
	orq	$0, (%rsp)
	cmpq	%r11, %rsp
	jne	.LPSRL0
	.cfi_def_cfa_register 7
	subq	$1112, %rsp
	.cfi_def_cfa_offset 4195456
	movq	%fs:40, %rax
	movq	%rax, 4195400(%rsp)
	xorl	%eax, %eax
	movl	%edi, %r12d
	leaq	12(%rsp), %rdi
	movq	%rsi, %rbp
	call	cudaGetDeviceCount@PLT
	cmpl	$0, 12(%rsp)
	leaq	.LC1(%rip), %rdx
	je	.L41
	xorl	%edi, %edi
	leaq	64(%rsp), %rbx
	call	cudaSetDevice@PLT
	xorl	%esi, %esi
	movq	%rbx, %rdi
	call	cudaGetDeviceProperties_v2@PLT
	testl	%eax, %eax
	jne	.L28
	xorl	%edx, %edx
	leaq	.LC2(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	%rbx, %rdx
	movl	$2, %edi
	leaq	.LC3(%rip), %rsi
	pushq	%rax
	.cfi_def_cfa_offset 4195464
	movl	420(%rsp), %eax
	pushq	%rax
	.cfi_def_cfa_offset 4195472
	movl	444(%rsp), %r9d
	xorl	%eax, %eax
	movl	440(%rsp), %r8d
	movl	368(%rsp), %ecx
	call	__printf_chk@PLT
	popq	%rdx
	.cfi_def_cfa_offset 4195464
	popq	%rcx
	.cfi_def_cfa_offset 4195456
.L28:
	movss	.LC4(%rip), %xmm0
	xorl	%eax, %eax
	xorl	%r14d, %r14d
.L29:
	movd	%r14d, %xmm1
	leaq	1096(%rsp), %rbx
	addss	%xmm0, %xmm1
	movss	%xmm0, (%rbx,%rax,4)
	incq	%rax
	movd	%xmm1, %r14d
	cmpq	$1048576, %rax
	jne	.L29
	leaq	24(%rsp), %rdi
	movl	$4194304, %esi
	call	cudaMalloc@PLT
	leaq	32(%rsp), %rdi
	movl	$4194304, %esi
	call	cudaMalloc@PLT
	leaq	40(%rsp), %rdi
	movl	$4, %esi
	call	cudaMalloc@PLT
	movq	24(%rsp), %rdi
	movq	%rbx, %rsi
	movl	$1, %ecx
	movl	$4194304, %edx
	xorl	%ebx, %ebx
	call	cudaMemcpy@PLT
	cmpl	$2, %r12d
	jne	.L30
	movq	8(%rbp), %rdi
	call	atoi@PLT
	movl	%eax, %ebx
.L30:
	leaq	48(%rsp), %rdi
	call	cudaEventCreate@PLT
	leaq	56(%rsp), %rdi
	call	cudaEventCreate@PLT
	testl	%ebx, %ebx
	je	.L31
	decl	%ebx
	jne	.L42
	leaq	.LC6(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	48(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	$1, %r8d
	jmp	.L40
.L31:
	leaq	.LC5(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	48(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	xorl	%r8d, %r8d
.L40:
	movq	24(%rsp), %rdx
	movq	32(%rsp), %rsi
	movl	$1048576, %ecx
	movq	40(%rsp), %rdi
	call	_Z6reducePfS_S_ib
	movq	56(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	56(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	56(%rsp), %rdx
	movq	48(%rsp), %rsi
	leaq	16(%rsp), %rdi
	call	cudaEventElapsedTime@PLT
	movq	40(%rsp), %rsi
	movl	$2, %ecx
	movl	$4, %edx
	leaq	20(%rsp), %rdi
	movss	16(%rsp), %xmm0
	divss	.LC8(%rip), %xmm0
	movss	%xmm0, 16(%rsp)
	call	cudaMemcpy@PLT
	leaq	.LC9(%rip), %rsi
	movl	$2, %edi
	movb	$1, %al
	cvtss2sd	16(%rsp), %xmm0
	call	__printf_chk@PLT
	movl	$2, %edi
	movb	$1, %al
	leaq	.LC10(%rip), %rsi
	cvtss2sd	20(%rsp), %xmm0
	call	__printf_chk@PLT
	movd	%r14d, %xmm2
	leaq	.LC11(%rip), %rsi
	movb	$1, %al
	movl	$2, %edi
	cvtss2sd	%xmm2, %xmm0
	call	__printf_chk@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	32(%rsp), %rdi
	call	cudaFree@PLT
	movq	40(%rsp), %rdi
	call	cudaFree@PLT
	movq	4195400(%rsp), %rax
	subq	%fs:40, %rax
	je	.L35
	call	__stack_chk_fail@PLT
.L42:
	leaq	.LC7(%rip), %rdx
.L41:
	movq	stderr(%rip), %rdi
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	movl	$1, %edi
	call	exit@PLT
.L35:
	addq	$4195416, %rsp
	.cfi_def_cfa_offset 40
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2028:
	.size	main, .-main
	.section	.rodata.str1.1
.LC12:
	.string	"_Z19shmem_reduce_kernelPfPKf"
.LC13:
	.string	"_Z20global_reduce_kernelPfS_"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2058:
	.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	.LC12(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z19shmem_reduce_kernelPfPKf(%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	.LC13(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z20global_reduce_kernelPfS_(%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
.LFE2058:
	.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
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC4:
	.long	1065353216
	.align 4
.LC8:
	.long	1120403456
	.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:

/* Howto run ?
If executable name is 'reduce' then type the following in command terminal
./reduce 0 for global reduce
./reduce for shared memory reduce
 */

/* For global access, reduce run time - 0.911 ms approx
For shared memory access, reduce run time - 0.771 ms approx
*/


#include <stdio.h>
#include <stdlib.h>
#include <hip/hip_runtime.h>
/*
cuda.h -  defines the public host functions and types for the CUDA driver API.
cuda_runtime_api.h - defines the public host functions and types for the CUDA
runtime API
cuda_runtime.h - defines everything cuda_runtime_api.h does, as well as built-in
type definitions and function overlays for the CUDA language extensions and
device intrinsic functions.*/

/*All info about the .h files is from stach overflow*/

/*If you were writing host code to be compiled with the host compiler which
 * includes API calls, you would include either cuda.h or cuda_runtime_api.h.
 * If you needed other CUDA language built-ins, like types, and were using the
 * runtime API and compiling with the host compiler, you would include
 * cuda_runtime.h
*/

__global__ void global_reduce_kernel(float * d_out, float * d_in){
   int myId = threadIdx.x + blockDim.x * blockIdx.x;
    int tid  = threadIdx.x;

    // do reduction in global mem
    for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1)
    {
        if (tid < s)
        {
            d_in[myId] += d_in[myId + s];
        }
        __syncthreads();        // make sure all adds at one stage are
//done!
    }

    // only thread 0 writes result for this block back to global mem
    if (tid == 0)
    {
        d_out[blockIdx.x] = d_in[myId];
    }
}

__global__ void shmem_reduce_kernel(float * d_out, const float * d_in)
{
    // sdata is allocated in the kernel call: 3rd arg to <<<b, t,
    // shmem>>>
    extern __shared__ float sdata[];

    int myId = threadIdx.x + blockDim.x * blockIdx.x;
    int tid  = threadIdx.x;

    // load shared mem from global mem
    sdata[tid] = d_in[myId];
    __syncthreads();            // make sure entire block is loaded!

    // 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();        // make sure all adds at one stage are
//done!
    }

    // only thread 0 writes result for this block back to global mem
    if (tid == 0)
    {
        d_out[blockIdx.x] = sdata[0];
    }
}

void reduce(float * d_out, float * d_intermediate, float * d_in, int size, bool usesSharedMemory)
{
    // assumes that size is not greater than maxThreadsPerBlock^2
    // and that size is a multiple of maxThreadsPerBlock
    const int maxThreadsPerBlock = 1024;
    int threads = maxThreadsPerBlock;
    int blocks = size / maxThreadsPerBlock;
    if (usesSharedMemory)
    {
        shmem_reduce_kernel<<<blocks, threads, threads *sizeof(float)>>>(d_intermediate, d_in);
    }
    else
    {
        global_reduce_kernel<<<blocks, threads>>>
            (d_intermediate, d_in);
    }
    // now we're down to one block left, so reduce it
    threads = blocks; // launch one thread for each block in prev step
    blocks = 1;
    if (usesSharedMemory)
    {
      //  shmem_reduce_kernel<<<blocks, threads, threads *sizeof(float)>>>(d_out, d_intermediate);
    	global_reduce_kernel<<<blocks, threads>>>
            (d_out, d_intermediate);
    }
    else
    {
        global_reduce_kernel<<<blocks, threads>>>
            (d_out, d_intermediate);
    }
}

int main(int argc, char **argv)
{
    int deviceCount;
    hipGetDeviceCount(&deviceCount);
/*cudaGetDeviceCount ( int* count )

PARAMETERS
count

RETURNS
cudaSuccess, cudaErrorNoDevice, cudaErrorInsufficientDriver

DESCRIPTION
Returns in *count the number of devices with compute capability greater or equal to 2.0 that are available for execution. If there is no such device then
cudaGetDeviceCount() will return cudaErrorNoDevice. If no driver can be loaded to determine if any such devices exist then cudaGetDeviceCount() will return
cudaErrorInsufficientDriver.*/


    if (deviceCount == 0) {
        fprintf(stderr, "error: no devices supporting CUDA.\n");
        exit(EXIT_FAILURE);
    }
    int dev = 0;
    hipSetDevice(dev);
/*cudaError_t cudaSetDevice ( int  device )
PARAMETERS
device
- Device on which the active host thread should execute the device code.

RETURNS
cudaSuccess, cudaErrorInvalidDevice, cudaErrorDeviceAlreadyInUse

DESCRIPTION
Sets device as the current device for the calling host thread. Valid device id's are 0 to (cudaGetDeviceCount() - 1).

Any device memory subsequently allocated from this host thread using cudaMalloc(), cudaMallocPitch() or cudaMallocArray() will be physically resident on
device. Any host memory allocated from this host thread using cudaMallocHost() or cudaHostAlloc() or cudaHostRegister() will have its lifetime associated with
device. Any streams or events created from this host thread will be associated with device. Any kernels launched from this host thread using the <<<>>>
operator or cudaLaunchKernel() will be executed on device.

This call may be made from any host thread, to any device, and at any time. This function will do no synchronization with the previous or new device, and
should be considered a very low overhead call.
*/

    hipDeviceProp_t devProps;
    if (hipGetDeviceProperties(&devProps, dev) == 0)
    {
        printf("Using device %d:\n", dev);
        printf("%s; global mem: %dB; compute v%d.%d; clock: %d kHz\n",
               devProps.name, (int)devProps.totalGlobalMem, 
               (int)devProps.major, (int)devProps.minor, 
               (int)devProps.clockRate);
    }

    const int ARRAY_SIZE = 1 << 20;
    const int ARRAY_BYTES = ARRAY_SIZE * sizeof(float);

    // generate the input array on the host
    float h_in[ARRAY_SIZE];
    float sum = 0.0f;
   /* for(int i = 0; i < ARRAY_SIZE; i++) {
        // generate random float in [-1.0f, 1.0f]
        h_in[i] = -1.0f + (float)random()/((float)RAND_MAX/2.0f);
        sum += h_in[i];
    }*/

//Additional part for input array on host - used for teesting correctness added by - Kshitij
	for(int i = 0; i < ARRAY_SIZE; i++) {
        // generate non random float
        h_in[i] = 1.0f;
        sum += h_in[i];
    }



    // declare GPU memory pointers
    float * d_in, * d_intermediate, * d_out;

    // allocate GPU memory
    hipMalloc((void **) &d_in, ARRAY_BYTES);
    hipMalloc((void **) &d_intermediate, ARRAY_BYTES); //overallocated
    hipMalloc((void **) &d_out, sizeof(float));

    // transfer the input array to the GPU
    hipMemcpy(d_in, h_in, ARRAY_BYTES, hipMemcpyHostToDevice); 

    int whichKernel = 0;
    if (argc == 2) {
        whichKernel = atoi(argv[1]);
	//whichKernel = argv[1]; -  incorrect. Explanation below.    
}

/*atoi example in this directory only
Run and see for understanding it 
Hence, when running reduce.cu we use ./reduce 0 or ./reduce 1
*/  
/*Why use the atoi ?
    When we receive argumnets from command line and stored in argv, they are stored as strings. Basically, argv is an array of strings.
	Now, while checking for shared memory or global memory usage, we use switch vase statements with number (int) - whichKernel.
	So,whichKernel = argv[1] won't work here as RHS is a string and LHS is an int.
	We need whichKernel to be int as used in the switch case statements.
	Hence, convert the received number - 0 or 1 which is stored as a string to an int and then use it for comparison ini switch case.
*/
    hipEvent_t start, stop;
    hipEventCreate(&start);
    hipEventCreate(&stop);
    // launch the kernel
    switch(whichKernel) {
    case 0:
        printf("Running global reduce\n");
        hipEventRecord(start, 0);
        /*for (int i = 0; i < 100; i++)
        {
           reduce(d_out, d_intermediate, d_in, ARRAY_SIZE, false);
        }*/

	 //Addition by Kshitij - single call instead of 100 calls of the kernel
        reduce(d_out, d_intermediate,d_in,ARRAY_SIZE,false);
	hipEventRecord(stop, 0);
        
	
	break;
    case 1:
        printf("Running reduce with shared mem\n");
        hipEventRecord(start, 0);
        /*for (int i = 0; i < 100; i++)
        {
            reduce(d_out, d_intermediate, d_in, ARRAY_SIZE, true);
        }*/
	reduce(d_out, d_intermediate, d_in, ARRAY_SIZE, true);
        hipEventRecord(stop, 0);
        break;
    default:
        fprintf(stderr, "error: ran no kernel\n");
        exit(EXIT_FAILURE);
    }
    hipEventSynchronize(stop);
    float elapsedTime;
    hipEventElapsedTime(&elapsedTime, start, stop);    
    elapsedTime /= 100.0f;      // 100 trials

    // copy back the sum from GPU
    float h_out;
    hipMemcpy(&h_out, d_out, sizeof(float), hipMemcpyDeviceToHost);

    printf("average time elapsed: %f\n", elapsedTime);
	
//Additional part - Kshitij
printf("Reduce answer :%f",h_out);
printf("Sum answer on host :%f",sum);



    // free GPU memory allocation
    hipFree(d_in);
    hipFree(d_intermediate);
    hipFree(d_out);
        
    return 0;
}

	.text
	.file	"reduce.hip"
	.globl	_Z35__device_stub__global_reduce_kernelPfS_ # -- Begin function _Z35__device_stub__global_reduce_kernelPfS_
	.type	_Z35__device_stub__global_reduce_kernelPfS_,@function
_Z35__device_stub__global_reduce_kernelPfS_: # @_Z35__device_stub__global_reduce_kernelPfS_
	.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	$80, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	16(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%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	$_Z20global_reduce_kernelPfS_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$96, %rsp
	.cfi_adjust_cfa_offset -96
	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	_Z35__device_stub__global_reduce_kernelPfS_, .Lfunc_end0-_Z35__device_stub__global_reduce_kernelPfS_
	.cfi_endproc
                                        # -- End function
	.globl	_Z34__device_stub__shmem_reduce_kernelPfPKf # -- Begin function _Z34__device_stub__shmem_reduce_kernelPfPKf
	.type	_Z34__device_stub__shmem_reduce_kernelPfPKf,@function
_Z34__device_stub__shmem_reduce_kernelPfPKf: # @_Z34__device_stub__shmem_reduce_kernelPfPKf
	.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	$80, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	16(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%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	$_Z19shmem_reduce_kernelPfPKf, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$96, %rsp
	.cfi_adjust_cfa_offset -96
	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	_Z34__device_stub__shmem_reduce_kernelPfPKf, .Lfunc_end1-_Z34__device_stub__shmem_reduce_kernelPfPKf
	.cfi_endproc
                                        # -- End function
	.globl	_Z6reducePfS_S_ib               # -- Begin function _Z6reducePfS_S_ib
	.type	_Z6reducePfS_S_ib,@function
_Z6reducePfS_S_ib:                      # @_Z6reducePfS_S_ib
	.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
                                        # kill: def $ecx killed $ecx def $rcx
	movq	%rdx, %r13
	movq	%rsi, %rbx
	movq	%rdi, %r14
	leal	1023(%rcx), %eax
	testl	%ecx, %ecx
	cmovnsl	%ecx, %eax
	movabsq	$4294967297, %r15               # imm = 0x100000001
	sarl	$10, %eax
	leaq	(%rax,%r15), %r12
	decq	%r12
	leaq	1023(%r15), %rdx
	testl	%r8d, %r8d
	je	.LBB2_3
# %bb.1:
	movl	$4096, %r8d                     # imm = 0x1000
	movq	%r12, %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_5
# %bb.2:
	movq	%rbx, %rdi
	movq	%r13, %rsi
	callq	_Z34__device_stub__shmem_reduce_kernelPfPKf
	jmp	.LBB2_5
.LBB2_3:
	movq	%r12, %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_5
# %bb.4:
	movq	%rbx, %rdi
	movq	%r13, %rsi
	callq	_Z35__device_stub__global_reduce_kernelPfS_
.LBB2_5:
	movq	%r15, %rdi
	movl	$1, %esi
	movq	%r12, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	je	.LBB2_7
# %bb.6:
	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
.LBB2_7:                                # %.sink.split
	.cfi_def_cfa_offset 48
	movq	%r14, %rdi
	movq	%rbx, %rsi
	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	_Z35__device_stub__global_reduce_kernelPfS_ # TAILCALL
.Lfunc_end2:
	.size	_Z6reducePfS_S_ib, .Lfunc_end2-_Z6reducePfS_S_ib
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function main
.LCPI3_0:
	.long	0x3f800000                      # float 1
.LCPI3_1:
	.long	0x42c80000                      # float 100
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	subq	$4195840, %rsp                  # imm = 0x400600
	.cfi_def_cfa_offset 4195872
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %rbp, -16
	movq	%rsi, %rbx
	movl	%edi, %ebp
	leaq	60(%rsp), %r14
	movq	%r14, %rdi
	callq	hipGetDeviceCount
	cmpl	$0, (%r14)
	je	.LBB3_1
# %bb.3:
	xorl	%edi, %edi
	callq	hipSetDevice
	leaq	64(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipGetDevicePropertiesR0600
	testl	%eax, %eax
	jne	.LBB3_5
# %bb.4:
	movl	$.L.str.1, %edi
	xorl	%esi, %esi
	xorl	%eax, %eax
	callq	printf
	leaq	64(%rsp), %rsi
	movl	360(%rsi), %ecx
	movl	364(%rsi), %r8d
	movl	288(%rsi), %edx
	movl	348(%rsi), %r9d
	movl	$.L.str.2, %edi
	xorl	%eax, %eax
	callq	printf
.LBB3_5:
	xorps	%xmm1, %xmm1
	xorl	%eax, %eax
	movss	.LCPI3_0(%rip), %xmm0           # xmm0 = mem[0],zero,zero,zero
.LBB3_6:                                # =>This Inner Loop Header: Depth=1
	movl	$1065353216, 1536(%rsp,%rax,4)  # imm = 0x3F800000
	addss	%xmm0, %xmm1
	incq	%rax
	cmpq	$1048576, %rax                  # imm = 0x100000
	jne	.LBB3_6
# %bb.7:
	movss	%xmm1, 48(%rsp)                 # 4-byte Spill
	leaq	40(%rsp), %r14
	movl	$4194304, %esi                  # imm = 0x400000
	movq	%r14, %rdi
	callq	hipMalloc
	leaq	32(%rsp), %rdi
	movl	$4194304, %esi                  # imm = 0x400000
	callq	hipMalloc
	leaq	24(%rsp), %rdi
	movl	$4, %esi
	callq	hipMalloc
	movq	(%r14), %rdi
	leaq	1536(%rsp), %rsi
	movl	$4194304, %edx                  # imm = 0x400000
	movl	$1, %ecx
	callq	hipMemcpy
	cmpl	$2, %ebp
	jne	.LBB3_8
# %bb.9:
	movq	8(%rbx), %rdi
	callq	atoi
	movl	%eax, %ebx
	leaq	16(%rsp), %rdi
	callq	hipEventCreate
	leaq	8(%rsp), %rdi
	callq	hipEventCreate
	cmpl	$1, %ebx
	je	.LBB3_13
# %bb.10:
	testl	%ebx, %ebx
	je	.LBB3_11
# %bb.12:
	movq	stderr(%rip), %rcx
	movl	$.L.str.5, %edi
	movl	$21, %esi
	jmp	.LBB3_2
.LBB3_8:                                # %.thread
	leaq	16(%rsp), %rdi
	callq	hipEventCreate
	leaq	8(%rsp), %rdi
	callq	hipEventCreate
.LBB3_11:
	movl	$.Lstr.1, %edi
	callq	puts@PLT
	movq	16(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	24(%rsp), %rdi
	movq	32(%rsp), %rsi
	movq	40(%rsp), %rdx
	movl	$1048576, %ecx                  # imm = 0x100000
	xorl	%r8d, %r8d
	jmp	.LBB3_14
.LBB3_13:
	movl	$.Lstr, %edi
	callq	puts@PLT
	movq	16(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	24(%rsp), %rdi
	movq	32(%rsp), %rsi
	movq	40(%rsp), %rdx
	movl	$1048576, %ecx                  # imm = 0x100000
	movl	$1, %r8d
.LBB3_14:
	callq	_Z6reducePfS_S_ib
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	8(%rsp), %rdi
	callq	hipEventSynchronize
	movq	16(%rsp), %rsi
	movq	8(%rsp), %rdx
	leaq	56(%rsp), %rbx
	movq	%rbx, %rdi
	callq	hipEventElapsedTime
	movss	(%rbx), %xmm0                   # xmm0 = mem[0],zero,zero,zero
	divss	.LCPI3_1(%rip), %xmm0
	movss	%xmm0, (%rbx)
	movq	24(%rsp), %rsi
	leaq	52(%rsp), %r14
	movl	$4, %edx
	movq	%r14, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	xorps	%xmm0, %xmm0
	cvtss2sd	(%rbx), %xmm0
	movl	$.L.str.6, %edi
	movb	$1, %al
	callq	printf
	xorps	%xmm0, %xmm0
	cvtss2sd	(%r14), %xmm0
	movl	$.L.str.7, %edi
	movb	$1, %al
	callq	printf
	xorps	%xmm0, %xmm0
	cvtss2sd	48(%rsp), %xmm0                 # 4-byte Folded Reload
	movl	$.L.str.8, %edi
	movb	$1, %al
	callq	printf
	movq	40(%rsp), %rdi
	callq	hipFree
	movq	32(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	xorl	%eax, %eax
	addq	$4195840, %rsp                  # imm = 0x400600
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.LBB3_1:
	.cfi_def_cfa_offset 4195872
	movq	stderr(%rip), %rcx
	movl	$.L.str, %edi
	movl	$35, %esi
.LBB3_2:
	movl	$1, %edx
	callq	fwrite@PLT
	movl	$1, %edi
	callq	exit
.Lfunc_end3:
	.size	main, .Lfunc_end3-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	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	$_Z20global_reduce_kernelPfS_, %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	$_Z19shmem_reduce_kernelPfPKf, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %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	_Z20global_reduce_kernelPfS_,@object # @_Z20global_reduce_kernelPfS_
	.section	.rodata,"a",@progbits
	.globl	_Z20global_reduce_kernelPfS_
	.p2align	3, 0x0
_Z20global_reduce_kernelPfS_:
	.quad	_Z35__device_stub__global_reduce_kernelPfS_
	.size	_Z20global_reduce_kernelPfS_, 8

	.type	_Z19shmem_reduce_kernelPfPKf,@object # @_Z19shmem_reduce_kernelPfPKf
	.globl	_Z19shmem_reduce_kernelPfPKf
	.p2align	3, 0x0
_Z19shmem_reduce_kernelPfPKf:
	.quad	_Z34__device_stub__shmem_reduce_kernelPfPKf
	.size	_Z19shmem_reduce_kernelPfPKf, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"error: no devices supporting CUDA.\n"
	.size	.L.str, 36

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"Using device %d:\n"
	.size	.L.str.1, 18

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"%s; global mem: %dB; compute v%d.%d; clock: %d kHz\n"
	.size	.L.str.2, 52

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"error: ran no kernel\n"
	.size	.L.str.5, 22

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"average time elapsed: %f\n"
	.size	.L.str.6, 26

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"Reduce answer :%f"
	.size	.L.str.7, 18

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"Sum answer on host :%f"
	.size	.L.str.8, 23

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z20global_reduce_kernelPfS_"
	.size	.L__unnamed_1, 29

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z19shmem_reduce_kernelPfPKf"
	.size	.L__unnamed_2, 29

	.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

	.type	.Lstr,@object                   # @str
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr:
	.asciz	"Running reduce with shared mem"
	.size	.Lstr, 31

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"Running global reduce"
	.size	.Lstr.1, 22

	.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 _Z35__device_stub__global_reduce_kernelPfS_
	.addrsig_sym _Z34__device_stub__shmem_reduce_kernelPfPKf
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z20global_reduce_kernelPfS_
	.addrsig_sym _Z19shmem_reduce_kernelPfPKf
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z20global_reduce_kernelPfS_ ; -- Begin function _Z20global_reduce_kernelPfS_
	.globl	_Z20global_reduce_kernelPfS_
	.p2align	8
	.type	_Z20global_reduce_kernelPfS_,@function
_Z20global_reduce_kernelPfS_:           ; @_Z20global_reduce_kernelPfS_
; %bb.0:
	s_clause 0x1
	s_load_b32 s5, s[0:1], 0x1c
	s_load_b128 s[0:3], s[0:1], 0x0
	s_mov_b32 s4, s15
	s_waitcnt lgkmcnt(0)
	s_and_b32 s5, s5, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s4, s5, v[0:1]
	s_cmp_lt_u32 s5, 2
	v_ashrrev_i32_e32 v2, 31, v1
	s_cbranch_scc1 .LBB0_5
; %bb.1:                                ; %.lr.ph
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[3:4], 2, v[1:2]
	v_mov_b32_e32 v6, 0
	v_add_co_u32 v3, vcc_lo, s2, v3
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v4, vcc_lo, s3, v4, vcc_lo
.LBB0_2:                                ; =>This Inner Loop Header: Depth=1
	s_lshr_b32 s6, s5, 1
	s_mov_b32 s7, exec_lo
	v_cmpx_gt_u32_e64 s6, v0
	s_cbranch_execz .LBB0_4
; %bb.3:                                ;   in Loop: Header=BB0_2 Depth=1
	v_add_nc_u32_e32 v5, s6, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[7:8], 2, v[5:6]
	v_add_co_u32 v7, vcc_lo, s2, v7
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v8, vcc_lo, s3, v8, vcc_lo
	s_clause 0x1
	global_load_b32 v5, v[7:8], off
	global_load_b32 v7, v[3:4], off
	s_waitcnt vmcnt(0)
	v_add_f32_e32 v5, v5, v7
	global_store_b32 v[3:4], v5, off
.LBB0_4:                                ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s7
	s_cmp_lt_u32 s5, 4
	s_mov_b32 s5, s6
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	s_cbranch_scc0 .LBB0_2
.LBB0_5:                                ; %._crit_edge
	s_mov_b32 s5, 0
	s_mov_b32 s6, exec_lo
	v_cmpx_eq_u32_e32 0, v0
	s_cbranch_execz .LBB0_7
; %bb.6:
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s2, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s3, v1, vcc_lo
	s_lshl_b64 s[2:3], s[4:5], 2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s0, s0, s2
	global_load_b32 v0, v[0:1], off
	v_mov_b32_e32 v1, 0
	s_addc_u32 s1, s1, s3
	s_waitcnt vmcnt(0)
	global_store_b32 v1, v0, s[0:1]
.LBB0_7:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z20global_reduce_kernelPfS_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 272
		.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 9
		.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	_Z20global_reduce_kernelPfS_, .Lfunc_end0-_Z20global_reduce_kernelPfS_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 304
; NumSgprs: 18
; NumVgprs: 9
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 9
; 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	_Z19shmem_reduce_kernelPfPKf ; -- Begin function _Z19shmem_reduce_kernelPfPKf
	.globl	_Z19shmem_reduce_kernelPfPKf
	.p2align	8
	.type	_Z19shmem_reduce_kernelPfPKf,@function
_Z19shmem_reduce_kernelPfPKf:           ; @_Z19shmem_reduce_kernelPfPKf
; %bb.0:
	s_clause 0x1
	s_load_b32 s5, s[0:1], 0x1c
	s_load_b128 s[0:3], s[0:1], 0x0
	s_mov_b32 s4, s15
	s_waitcnt lgkmcnt(0)
	s_and_b32 s5, s5, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s4, s5, v[0:1]
	s_cmp_lt_u32 s5, 2
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[1:2], 2, v[1:2]
	v_add_co_u32 v1, vcc_lo, s2, v1
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v2, vcc_lo, s3, v2, vcc_lo
	global_load_b32 v2, v[1:2], off
	v_lshl_add_u32 v1, v0, 2, 0
	s_waitcnt vmcnt(0)
	ds_store_b32 v1, v2
	s_waitcnt lgkmcnt(0)
	s_barrier
.LBB1_1:                                ; =>This Inner Loop Header: Depth=1
	buffer_gl0_inv
	s_cbranch_scc1 .LBB1_5
; %bb.2:                                ; %.lr.ph
                                        ;   in Loop: Header=BB1_1 Depth=1
	s_lshr_b32 s2, s5, 1
	s_mov_b32 s3, exec_lo
	v_cmpx_gt_u32_e64 s2, v0
	s_cbranch_execz .LBB1_4
; %bb.3:                                ;   in Loop: Header=BB1_1 Depth=1
	v_lshl_add_u32 v2, s2, 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_4:                                ;   in Loop: Header=BB1_1 Depth=1
	s_or_b32 exec_lo, exec_lo, s3
	s_waitcnt lgkmcnt(0)
	s_barrier
	s_cmp_lt_u32 s5, 4
	s_mov_b32 s5, s2
	s_branch .LBB1_1
.LBB1_5:                                ; %._crit_edge
	s_mov_b32 s2, exec_lo
	v_cmpx_eq_u32_e32 0, v0
	s_cbranch_execz .LBB1_7
; %bb.6:
	v_mov_b32_e32 v0, 0
	s_mov_b32 s5, 0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_lshl_b64 s[2:3], s[4:5], 2
	s_add_u32 s0, s0, s2
	ds_load_b32 v1, v0
	s_addc_u32 s1, s1, s3
	s_waitcnt lgkmcnt(0)
	global_store_b32 v0, v1, s[0:1]
.LBB1_7:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z19shmem_reduce_kernelPfPKf
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 272
		.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	_Z19shmem_reduce_kernelPfPKf, .Lfunc_end1-_Z19shmem_reduce_kernelPfPKf
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 284
; NumSgprs: 18
; NumVgprs: 4
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 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
	.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:     hidden_block_count_x
      - .offset:         20
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         28
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         30
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         32
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         34
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         36
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         38
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         56
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         80
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 272
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z20global_reduce_kernelPfS_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z20global_reduce_kernelPfS_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     9
    .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:     hidden_block_count_x
      - .offset:         20
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         28
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         30
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         32
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         34
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         36
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         38
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         56
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         80
        .size:           2
        .value_kind:     hidden_grid_dims
      - .offset:         136
        .size:           4
        .value_kind:     hidden_dynamic_lds_size
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 272
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z19shmem_reduce_kernelPfPKf
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z19shmem_reduce_kernelPfPKf.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     4
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
